summaryrefslogtreecommitdiffstats
path: root/tools/cuda-ab
diff options
context:
space:
mode:
Diffstat (limited to 'tools/cuda-ab')
-rw-r--r--tools/cuda-ab/frame_producer.cu172
-rwxr-xr-xtools/cuda-ab/run/frame_producerbin0 -> 841544 bytes
-rwxr-xr-xtools/cuda-ab/run/transformerbin0 -> 846232 bytes
-rwxr-xr-xtools/cuda-ab/run/wedge_abbin0 -> 825144 bytes
-rwxr-xr-xtools/cuda-ab/run/wedge_bisectbin0 -> 821000 bytes
-rw-r--r--tools/cuda-ab/transformer.cu472
-rw-r--r--tools/cuda-ab/wedge_ab.cu100
-rw-r--r--tools/cuda-ab/wedge_bisect.cu104
8 files changed, 848 insertions, 0 deletions
diff --git a/tools/cuda-ab/frame_producer.cu b/tools/cuda-ab/frame_producer.cu
new file mode 100644
index 0000000..609114c
--- /dev/null
+++ b/tools/cuda-ab/frame_producer.cu
@@ -0,0 +1,172 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * frame_producer.cu - render an animated frame on the guest GPU and publish
+ * it into the ivshmem (kvmfr) BAR so the host can see a real GPU-rendered
+ * frame from the guest.
+ *
+ * Path: CUDA kernel -> device buffer -> host buffer -> ivshmem BAR
+ * (/sys/bus/pci/devices/0000:00:03.0/resource2, the 128 MiB kvmfr
+ * device QEMU backs with /dev/shm/kvmfr-<uuid>).
+ *
+ * Layout in the BAR (offset 0x100000, clear of the kvmfr protocol header):
+ * struct frame_hdr { magic, version, width, height, stride, format, seq }
+ * then width*height*4 bytes of RGBA8.
+ *
+ * Usage: sudo ./frame_producer [frames] [width] [height] [fps]
+ * Prints: FRAME seq=<n> ms=<render+copy ms> (one line per frame, first 5 and
+ * every 30th) and FRAME_DONE frames=<n>.
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+#include <cstring>
+#include <ctime>
+#include <fcntl.h>
+#include <sys/mman.h>
+#include <unistd.h>
+#include <cuda_runtime.h>
+
+#define CHECK(x) \
+ do { \
+ cudaError_t e_ = (x); \
+ if (e_ != cudaSuccess) { \
+ printf("CUDA_ERROR %s at line %d\n", \
+ cudaGetErrorString(e_), __LINE__); \
+ return 1; \
+ } \
+ } while (0)
+
+#define FRAME_OFFSET 0x100000UL
+#define HDR_MAGIC 0x4b564d46u /* "FMVK" little-endian-ish */
+
+struct frame_hdr {
+ uint32_t magic;
+ uint32_t version;
+ uint32_t width;
+ uint32_t height;
+ uint32_t stride;
+ uint32_t format; /* 0 = RGBA8 */
+ uint32_t seq;
+ uint32_t reserved;
+ uint64_t timestamp_ns;
+};
+
+__global__ void render_julia(uchar4 *fb, int w, int h, float t)
+{
+ int x = blockIdx.x * blockDim.x + threadIdx.x;
+ int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+ if (x >= w || y >= h)
+ return;
+
+ float fx = (x - w * 0.5f) / (h * 0.45f);
+ float fy = (y - h * 0.5f) / (h * 0.45f);
+ float cr = 0.7885f * cosf(t);
+ float ci = 0.7885f * sinf(t);
+ float zr = fx, zi = fy;
+ int it = 0;
+ const int maxit = 96;
+
+ while (zr * zr + zi * zi < 4.0f && it < maxit) {
+ float nzr = zr * zr - zi * zi + cr;
+ float nzi = 2.0f * zr * zi + ci;
+
+ zr = nzr;
+ zi = nzi;
+ it++;
+ }
+ {
+ float v = (float)it / (float)maxit;
+ unsigned char r = (unsigned char)(255.0f * v);
+ unsigned char g = (unsigned char)(255.0f * (1.0f - v) * v * 4.0f);
+ unsigned char b = (unsigned char)(255.0f * (1.0f - v));
+ unsigned char a = 255;
+
+ fb[y * w + x] = make_uchar4(r, g, b, a);
+ }
+}
+
+static uint64_t now_ns(void)
+{
+ struct timespec ts;
+
+ clock_gettime(CLOCK_MONOTONIC, &ts);
+ return (uint64_t)ts.tv_sec * 1000000000ull + ts.tv_nsec;
+}
+
+int main(int argc, char **argv)
+{
+ int frames = argc > 1 ? atoi(argv[1]) : 60;
+ int w = argc > 2 ? atoi(argv[2]) : 800;
+ int h = argc > 3 ? atoi(argv[3]) : 600;
+ int fps = argc > 4 ? atoi(argv[4]) : 10;
+ const char *bar_path =
+ "/sys/bus/pci/devices/0000:00:03.0/resource2";
+ size_t fb_bytes = (size_t)w * h * 4;
+ uchar4 *d_fb = NULL, *h_fb = NULL;
+ unsigned char *bar;
+ struct frame_hdr *hdr;
+ volatile uint32_t *seqp;
+ int fd, i;
+
+ setvbuf(stdout, NULL, _IONBF, 0);
+
+ fd = open(bar_path, O_RDWR | O_SYNC);
+ if (fd < 0) {
+ printf("OPEN_FAIL %s\n", bar_path);
+ return 1;
+ }
+ bar = (unsigned char *)mmap(NULL, 128UL << 20, PROT_READ | PROT_WRITE,
+ MAP_SHARED, fd, 0);
+ if (bar == MAP_FAILED) {
+ printf("MMAP_FAIL\n");
+ return 1;
+ }
+ hdr = (struct frame_hdr *)(bar + FRAME_OFFSET);
+ seqp = &hdr->seq;
+
+ CHECK(cudaMalloc(&d_fb, fb_bytes));
+ h_fb = (uchar4 *)malloc(fb_bytes);
+ if (!h_fb) {
+ printf("HOST_ALLOC_FAIL\n");
+ return 1;
+ }
+
+ hdr->magic = HDR_MAGIC;
+ hdr->version = 1;
+ hdr->width = (uint32_t)w;
+ hdr->height = (uint32_t)h;
+ hdr->stride = (uint32_t)w * 4;
+ hdr->format = 0;
+
+ for (i = 0; i < frames; i++) {
+ uint64_t t0 = now_ns();
+ float phase = (float)i * 0.15f;
+ dim3 block(16, 16);
+ dim3 grid((w + 15) / 16, (h + 15) / 16);
+
+ render_julia<<<grid, block>>>(d_fb, w, h, phase);
+ CHECK(cudaGetLastError());
+ CHECK(cudaDeviceSynchronize());
+ CHECK(cudaMemcpy(h_fb, d_fb, fb_bytes,
+ cudaMemcpyDeviceToHost));
+ memcpy(bar + FRAME_OFFSET + sizeof(struct frame_hdr), h_fb,
+ fb_bytes);
+ hdr->timestamp_ns = now_ns();
+ __sync_synchronize();
+ *seqp = (uint32_t)(i + 1);
+ __sync_synchronize();
+ if (i < 5 || (i % 30) == 29)
+ printf("FRAME seq=%d ms=%.1f\n", i + 1,
+ (double)(now_ns() - t0) / 1e6);
+ if (fps > 0)
+ usleep((useconds_t)(1000000 / fps));
+ }
+ printf("FRAME_DONE frames=%d %dx%d\n", frames, w, h);
+ CHECK(cudaFree(d_fb));
+ free(h_fb);
+ munmap(bar, 128UL << 20);
+ close(fd);
+ return 0;
+}
diff --git a/tools/cuda-ab/run/frame_producer b/tools/cuda-ab/run/frame_producer
new file mode 100755
index 0000000..a202ac7
--- /dev/null
+++ b/tools/cuda-ab/run/frame_producer
Binary files differ
diff --git a/tools/cuda-ab/run/transformer b/tools/cuda-ab/run/transformer
new file mode 100755
index 0000000..bcde39e
--- /dev/null
+++ b/tools/cuda-ab/run/transformer
Binary files differ
diff --git a/tools/cuda-ab/run/wedge_ab b/tools/cuda-ab/run/wedge_ab
new file mode 100755
index 0000000..23a90aa
--- /dev/null
+++ b/tools/cuda-ab/run/wedge_ab
Binary files differ
diff --git a/tools/cuda-ab/run/wedge_bisect b/tools/cuda-ab/run/wedge_bisect
new file mode 100755
index 0000000..5898fcf
--- /dev/null
+++ b/tools/cuda-ab/run/wedge_bisect
Binary files differ
diff --git a/tools/cuda-ab/transformer.cu b/tools/cuda-ab/transformer.cu
new file mode 100644
index 0000000..3044522
--- /dev/null
+++ b/tools/cuda-ab/transformer.cu
@@ -0,0 +1,472 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * transformer.cu - single-head attention forward + 4-layer FFN mini-transformer.
+ *
+ * Two independent checks in one binary, both verified against an internal
+ * double-precision CPU reference (tolerance 1e-3 relative, 1e-6 absolute floor):
+ *
+ * 1. Attention: X[64x256], Wq/Wk/Wv/Wo[256x256] (head dim 256)
+ * Q = X*Wq, K = X*Wk, V = X*Wv
+ * S = Q*K^T / 16 (scale = 1/sqrt(head_dim))
+ * A = softmax(S) (row-wise)
+ * Y = (A*V)*Wo
+ * 2. Mini-transformer FFN head: h <- ReLU(h*Wf[l]) for 4 layers, then
+ * logits = h*Wh, P = softmax(logits) (row-wise over 256)
+ *
+ * Inputs and weights are multiples of 0.25/0.125, so every GEMM is exact in
+ * fp32; only the softmax exponentials differ from the CPU reference, hence
+ * the relative tolerance.
+ *
+ * Output lines are stable and greppable:
+ *
+ * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
+ * TRANSFORMER dh=256 seq=64 bad=<mismatches> ysum=<sum>
+ * FFN4 layers=4 bad=<mismatches> psum=<sum>
+ * STATUS OK|FAIL
+ *
+ * Build with a static CUDA runtime (--cudart=static) so the guest only needs
+ * the driver's libcuda.so, not a CUDA toolkit installation:
+ *
+ * nvcc -O2 --cudart=static -arch=sm_75 -o transformer transformer.cu
+ *
+ * Then, from the guest: /tmp/transformer
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+#include <cmath>
+#include <cuda_runtime.h>
+
+#define CHECK(x) \
+ do { \
+ cudaError_t e_ = (x); \
+ if (e_ != cudaSuccess) { \
+ printf("CUDA_ERROR %s at line %d\n", \
+ cudaGetErrorString(e_), __LINE__); \
+ printf("STATUS FAIL\n"); \
+ return 1; \
+ } \
+ } while (0)
+
+#define SEQ 64
+#define DH 256
+#define TILE 16
+#define TOL 1e-3f
+
+__global__ void gemm_nn(const float *A, const float *B, float *C,
+ int M, int N, int K)
+{
+ __shared__ float As[TILE][TILE];
+ __shared__ float Bs[TILE][TILE];
+ int row = blockIdx.y * TILE + threadIdx.y;
+ int col = blockIdx.x * TILE + threadIdx.x;
+ float acc = 0.0f;
+
+ for (int t = 0; t < K / TILE; t++) {
+ As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x];
+ Bs[threadIdx.y][threadIdx.x] =
+ B[(t * TILE + threadIdx.y) * N + col];
+ __syncthreads();
+ for (int k = 0; k < TILE; k++)
+ acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
+ __syncthreads();
+ }
+ C[row * N + col] = acc;
+}
+
+__global__ void gemm_nt(const float *A, const float *B, float *C,
+ int M, int N, int K)
+{
+ /* C[i][j] = sum_k A[i][k] * B[j][k] (B is [N][K], row-major). */
+ __shared__ float As[TILE][TILE];
+ __shared__ float Bs[TILE][TILE];
+ int row = blockIdx.y * TILE + threadIdx.y;
+ int col = blockIdx.x * TILE + threadIdx.x;
+ float acc = 0.0f;
+
+ for (int t = 0; t < K / TILE; t++) {
+ As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x];
+ Bs[threadIdx.y][threadIdx.x] =
+ B[(blockIdx.x * TILE + threadIdx.y) * K +
+ t * TILE + threadIdx.x];
+ __syncthreads();
+ for (int k = 0; k < TILE; k++)
+ acc += As[threadIdx.y][k] * Bs[threadIdx.x][k];
+ __syncthreads();
+ }
+ C[row * N + col] = acc;
+}
+
+__global__ void softmax_rows(const float *S, float *P, int rows, int cols,
+ float scale)
+{
+ extern __shared__ float sm[];
+ int row = blockIdx.x;
+ int t = threadIdx.x;
+ float v = S[row * cols + t] * scale;
+ int stride;
+
+ sm[t] = v;
+ __syncthreads();
+ for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
+ if (t < stride)
+ sm[t] = fmaxf(sm[t], sm[t + stride]);
+ __syncthreads();
+ }
+ v = expf(v - sm[0]);
+ sm[t] = v;
+ __syncthreads();
+ for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
+ if (t < stride)
+ sm[t] += sm[t + stride];
+ __syncthreads();
+ }
+ P[row * cols + t] = v / sm[0];
+}
+
+__global__ void relu_inplace(float *h, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n && h[i] < 0.0f)
+ h[i] = 0.0f;
+}
+
+/* ------------------------------------------------------------------ host */
+
+static void launch_gemm_nn(const float *A, const float *B, float *C,
+ int M, int N, int K)
+{
+ gemm_nn<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K);
+}
+
+static void launch_gemm_nt(const float *A, const float *B, float *C,
+ int M, int N, int K)
+{
+ gemm_nt<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K);
+}
+
+static void launch_softmax(const float *S, float *P, int rows, int cols,
+ float scale)
+{
+ softmax_rows<<<rows, cols, (size_t)cols * sizeof(float)>>>(S, P, rows,
+ cols,
+ scale);
+}
+
+/* Deterministic inputs: multiples of 0.25 (X) and 0.125 (weights). */
+static float xval(int i)
+{
+ return (float)(((i % 7) - 3)) * 0.25f;
+}
+
+static float wval(int i, int seed)
+{
+ return (float)(((i * 13 + seed * 101 + 7) % 5) - 2) * 0.125f;
+}
+
+static void ref_gemm(const double *A, const double *B, double *C,
+ int M, int N, int K)
+{
+ for (int i = 0; i < M; i++)
+ for (int j = 0; j < N; j++) {
+ double acc = 0.0;
+
+ for (int k = 0; k < K; k++)
+ acc += A[i * K + k] * B[k * N + j];
+ C[i * N + j] = acc;
+ }
+}
+
+static void ref_gemm_nt(const double *A, const double *B, double *C,
+ int M, int N, int K)
+{
+ for (int i = 0; i < M; i++)
+ for (int j = 0; j < N; j++) {
+ double acc = 0.0;
+
+ for (int k = 0; k < K; k++)
+ acc += A[i * K + k] * B[j * K + k];
+ C[i * N + j] = acc;
+ }
+}
+
+static void ref_softmax(const double *S, double *P, int rows, int cols,
+ double scale)
+{
+ for (int i = 0; i < rows; i++) {
+ const double *s = S + i * cols;
+ double *p = P + i * cols;
+ double m = s[0] * scale, sum = 0.0;
+
+ for (int j = 1; j < cols; j++)
+ if (s[j] * scale > m)
+ m = s[j] * scale;
+ for (int j = 0; j < cols; j++) {
+ p[j] = exp(s[j] * scale - m);
+ sum += p[j];
+ }
+ for (int j = 0; j < cols; j++)
+ p[j] /= sum;
+ }
+}
+
+static void ref_relu(double *h, int n)
+{
+ for (int i = 0; i < n; i++)
+ if (h[i] < 0.0)
+ h[i] = 0.0;
+}
+
+static int count_bad(const float *gpu, const double *ref, int n)
+{
+ int bad = 0;
+
+ for (int i = 0; i < n; i++) {
+ double g = (double)gpu[i];
+ double r = ref[i];
+
+ if (fabs(g - r) > (double)TOL * (fabs(r) + 1e-6))
+ bad++;
+ }
+ return bad;
+}
+
+static double dsum(const float *h, int n)
+{
+ double s = 0.0;
+
+ for (int i = 0; i < n; i++)
+ s += (double)h[i];
+ return s;
+}
+
+#define NF 4 /* FFN layers */
+
+int main(void)
+{
+ cudaDeviceProp prop;
+ const int seqdh = SEQ * DH, dh2 = DH * DH, seq2 = SEQ * SEQ;
+ const size_t b_seqdh = (size_t)seqdh * sizeof(float);
+ const size_t b_dh2 = (size_t)dh2 * sizeof(float);
+ const size_t b_seq2 = (size_t)seq2 * sizeof(float);
+ float *h_x, *h_y, *h_p;
+ double *r_x, *r_q, *r_k, *r_v, *r_s, *r_a, *r_av, *r_y;
+ double *r_f0, *r_f1, *r_logits, *r_p;
+ float *d_x, *d_q, *d_k, *d_v, *d_s, *d_a, *d_av, *d_y;
+ float *d_wq, *d_wk, *d_wv, *d_wo;
+ float *d_wf[NF], *d_wh, *d_f0, *d_f1, *d_logits, *d_p;
+ float *d_ping, *d_pong, *tmpf;
+ double *r_ping, *r_pong, *tmpd;
+ int bad_y, bad_p, i, l;
+
+ setvbuf(stdout, NULL, _IONBF, 0);
+
+ CHECK(cudaGetDeviceProperties(&prop, 0));
+ printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name,
+ prop.major, prop.minor, prop.multiProcessorCount,
+ (unsigned long long)prop.totalGlobalMem);
+
+ /* host inputs + CPU reference workspace */
+ h_x = (float *)malloc(b_seqdh);
+ h_y = (float *)malloc(b_seqdh);
+ h_p = (float *)malloc(b_seqdh);
+ r_x = (double *)malloc((size_t)seqdh * sizeof(double));
+ r_q = (double *)malloc((size_t)seqdh * sizeof(double));
+ r_k = (double *)malloc((size_t)seqdh * sizeof(double));
+ r_v = (double *)malloc((size_t)seqdh * sizeof(double));
+ r_s = (double *)malloc((size_t)seq2 * sizeof(double));
+ r_a = (double *)malloc((size_t)seq2 * sizeof(double));
+ r_av = (double *)malloc((size_t)seqdh * sizeof(double));
+ r_y = (double *)malloc((size_t)seqdh * sizeof(double));
+ r_f0 = (double *)malloc((size_t)seqdh * sizeof(double));
+ r_f1 = (double *)malloc((size_t)seqdh * sizeof(double));
+ r_logits = (double *)malloc((size_t)seqdh * sizeof(double));
+ r_p = (double *)malloc((size_t)seqdh * sizeof(double));
+ if (!h_x || !h_y || !h_p || !r_x || !r_q || !r_k || !r_v || !r_s ||
+ !r_a || !r_av || !r_y || !r_f0 || !r_f1 || !r_logits || !r_p)
+ return 1;
+
+ for (i = 0; i < seqdh; i++) {
+ h_x[i] = xval(i);
+ r_x[i] = (double)xval(i);
+ }
+
+ /* device buffers */
+ CHECK(cudaMalloc(&d_x, b_seqdh));
+ CHECK(cudaMalloc(&d_q, b_seqdh));
+ CHECK(cudaMalloc(&d_k, b_seqdh));
+ CHECK(cudaMalloc(&d_v, b_seqdh));
+ CHECK(cudaMalloc(&d_s, b_seq2));
+ CHECK(cudaMalloc(&d_a, b_seq2));
+ CHECK(cudaMalloc(&d_av, b_seqdh));
+ CHECK(cudaMalloc(&d_y, b_seqdh));
+ CHECK(cudaMalloc(&d_wq, b_dh2));
+ CHECK(cudaMalloc(&d_wk, b_dh2));
+ CHECK(cudaMalloc(&d_wv, b_dh2));
+ CHECK(cudaMalloc(&d_wo, b_dh2));
+ CHECK(cudaMalloc(&d_wh, b_dh2));
+ CHECK(cudaMalloc(&d_f0, b_seqdh));
+ CHECK(cudaMalloc(&d_f1, b_seqdh));
+ CHECK(cudaMalloc(&d_logits, b_seqdh));
+ CHECK(cudaMalloc(&d_p, b_seqdh));
+ for (l = 0; l < NF; l++)
+ CHECK(cudaMalloc(&d_wf[l], b_dh2));
+
+ /* weights: deterministic, uploaded once */
+ {
+ float *w = (float *)malloc(b_dh2);
+
+ if (!w)
+ return 1;
+ for (i = 0; i < dh2; i++)
+ w[i] = wval(i, 1);
+ CHECK(cudaMemcpy(d_wq, w, b_dh2, cudaMemcpyHostToDevice));
+ for (i = 0; i < dh2; i++)
+ w[i] = wval(i, 2);
+ CHECK(cudaMemcpy(d_wk, w, b_dh2, cudaMemcpyHostToDevice));
+ for (i = 0; i < dh2; i++)
+ w[i] = wval(i, 3);
+ CHECK(cudaMemcpy(d_wv, w, b_dh2, cudaMemcpyHostToDevice));
+ for (i = 0; i < dh2; i++)
+ w[i] = wval(i, 4);
+ CHECK(cudaMemcpy(d_wo, w, b_dh2, cudaMemcpyHostToDevice));
+ for (i = 0; i < dh2; i++)
+ w[i] = wval(i, 20);
+ CHECK(cudaMemcpy(d_wh, w, b_dh2, cudaMemcpyHostToDevice));
+ for (l = 0; l < NF; l++) {
+ for (i = 0; i < dh2; i++)
+ w[i] = wval(i, 10 + l);
+ CHECK(cudaMemcpy(d_wf[l], w, b_dh2,
+ cudaMemcpyHostToDevice));
+ }
+ free(w);
+ }
+ CHECK(cudaMemcpy(d_x, h_x, b_seqdh, cudaMemcpyHostToDevice));
+
+ /* ---- check 1: attention forward ---- */
+ launch_gemm_nn(d_x, d_wq, d_q, SEQ, DH, DH);
+ launch_gemm_nn(d_x, d_wk, d_k, SEQ, DH, DH);
+ launch_gemm_nn(d_x, d_wv, d_v, SEQ, DH, DH);
+ launch_gemm_nt(d_q, d_k, d_s, SEQ, SEQ, DH);
+ launch_softmax(d_s, d_a, SEQ, SEQ, 1.0f / 16.0f);
+ launch_gemm_nn(d_a, d_v, d_av, SEQ, DH, SEQ);
+ launch_gemm_nn(d_av, d_wo, d_y, SEQ, DH, DH);
+ CHECK(cudaGetLastError());
+ CHECK(cudaDeviceSynchronize());
+ CHECK(cudaMemcpy(h_y, d_y, b_seqdh, cudaMemcpyDeviceToHost));
+
+ {
+ double *r_w = (double *)malloc((size_t)dh2 * sizeof(double));
+
+ if (!r_w)
+ return 1;
+ for (i = 0; i < dh2; i++)
+ r_w[i] = (double)wval(i, 1);
+ ref_gemm(r_x, r_w, r_q, SEQ, DH, DH);
+ for (i = 0; i < dh2; i++)
+ r_w[i] = (double)wval(i, 2);
+ ref_gemm(r_x, r_w, r_k, SEQ, DH, DH);
+ for (i = 0; i < dh2; i++)
+ r_w[i] = (double)wval(i, 3);
+ ref_gemm(r_x, r_w, r_v, SEQ, DH, DH);
+ ref_gemm_nt(r_q, r_k, r_s, SEQ, SEQ, DH);
+ ref_softmax(r_s, r_a, SEQ, SEQ, 1.0 / 16.0);
+ ref_gemm(r_a, r_v, r_av, SEQ, DH, SEQ);
+ for (i = 0; i < dh2; i++)
+ r_w[i] = (double)wval(i, 4);
+ ref_gemm(r_av, r_w, r_y, SEQ, DH, DH);
+ free(r_w);
+ }
+
+ bad_y = count_bad(h_y, r_y, seqdh);
+ printf("TRANSFORMER dh=256 seq=64 bad=%d ysum=%.6f\n", bad_y,
+ dsum(h_y, seqdh));
+
+ /* ---- check 2: 4-layer FFN stack + softmax head ---- */
+ CHECK(cudaMemcpy(d_f0, h_x, b_seqdh, cudaMemcpyHostToDevice));
+ d_ping = d_f0;
+ d_pong = d_f1;
+ for (l = 0; l < NF; l++) {
+ launch_gemm_nn(d_ping, d_wf[l], d_pong, SEQ, DH, DH);
+ relu_inplace<<<(seqdh + 255) / 256, 256>>>(d_pong, seqdh);
+ tmpf = d_ping;
+ d_ping = d_pong;
+ d_pong = tmpf;
+ }
+ launch_gemm_nn(d_ping, d_wh, d_logits, SEQ, DH, DH);
+ launch_softmax(d_logits, d_p, SEQ, DH, 1.0f);
+ CHECK(cudaGetLastError());
+ CHECK(cudaDeviceSynchronize());
+ CHECK(cudaMemcpy(h_p, d_p, b_seqdh, cudaMemcpyDeviceToHost));
+
+ {
+ double *r_w = (double *)malloc((size_t)dh2 * sizeof(double));
+
+ if (!r_w)
+ return 1;
+ memcpy(r_f0, r_x, (size_t)seqdh * sizeof(double));
+ r_ping = r_f0;
+ r_pong = r_f1;
+ for (l = 0; l < NF; l++) {
+ for (i = 0; i < dh2; i++)
+ r_w[i] = (double)wval(i, 10 + l);
+ ref_gemm(r_ping, r_w, r_pong, SEQ, DH, DH);
+ ref_relu(r_pong, seqdh);
+ tmpd = r_ping;
+ r_ping = r_pong;
+ r_pong = tmpd;
+ }
+ for (i = 0; i < dh2; i++)
+ r_w[i] = (double)wval(i, 20);
+ ref_gemm(r_ping, r_w, r_logits, SEQ, DH, DH);
+ ref_softmax(r_logits, r_p, SEQ, DH, 1.0);
+ free(r_w);
+ }
+
+ bad_p = count_bad(h_p, r_p, seqdh);
+ printf("FFN4 layers=4 bad=%d psum=%.6f\n", bad_p, dsum(h_p, seqdh));
+
+ if (bad_y != 0 || bad_p != 0) {
+ printf("STATUS FAIL\n");
+ return 1;
+ }
+
+ printf("STATUS OK\n");
+ CHECK(cudaFree(d_x));
+ CHECK(cudaFree(d_q));
+ CHECK(cudaFree(d_k));
+ CHECK(cudaFree(d_v));
+ CHECK(cudaFree(d_s));
+ CHECK(cudaFree(d_a));
+ CHECK(cudaFree(d_av));
+ CHECK(cudaFree(d_y));
+ CHECK(cudaFree(d_wq));
+ CHECK(cudaFree(d_wk));
+ CHECK(cudaFree(d_wv));
+ CHECK(cudaFree(d_wo));
+ CHECK(cudaFree(d_wh));
+ for (l = 0; l < NF; l++)
+ CHECK(cudaFree(d_wf[l]));
+ CHECK(cudaFree(d_f0));
+ CHECK(cudaFree(d_f1));
+ CHECK(cudaFree(d_logits));
+ CHECK(cudaFree(d_p));
+ free(h_x);
+ free(h_y);
+ free(h_p);
+ free(r_x);
+ free(r_q);
+ free(r_k);
+ free(r_v);
+ free(r_s);
+ free(r_a);
+ free(r_av);
+ free(r_y);
+ free(r_f0);
+ free(r_f1);
+ free(r_logits);
+ free(r_p);
+ return 0;
+}
diff --git a/tools/cuda-ab/wedge_ab.cu b/tools/cuda-ab/wedge_ab.cu
new file mode 100644
index 0000000..2d88f92
--- /dev/null
+++ b/tools/cuda-ab/wedge_ab.cu
@@ -0,0 +1,100 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * wedge_ab.cu - discriminate the wedge trigger: is it the second launch of
+ * the SAME program, or the second launch of the context regardless of
+ * program?
+ *
+ * Sequence (each launch verified against mutated inputs):
+ * 1. vadd (program A)
+ * 2. vsub (program B, different code)
+ * 3. vadd (program A again)
+ * 4. vsub (program B again)
+ *
+ * If launch 2 (B) executes -> the trigger is same-program reuse.
+ * If launch 2 (B) wedges -> the trigger is the second launch in the context.
+ *
+ * Build: nvcc -O2 --cudart=static -o wedge_ab wedge_ab.cu
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cuda_runtime.h>
+
+#define CHECK(x) \
+ do { \
+ cudaError_t e_ = (x); \
+ if (e_ != cudaSuccess) { \
+ printf("CUDA_ERROR %s at line %d\n", \
+ cudaGetErrorString(e_), __LINE__); \
+ printf("STATUS FAIL\n"); \
+ return 1; \
+ } \
+ } while (0)
+
+__global__ void vadd(const float *a, const float *b, float *c, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n)
+ c[i] = a[i] + b[i];
+}
+
+__global__ void vsub(const float *a, const float *b, float *c, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n)
+ c[i] = a[i] - b[i];
+}
+
+int main(void)
+{
+ const int n = 1 << 15;
+ size_t bytes = (size_t)n * sizeof(float);
+ float *ha, *hb, *hc;
+ float *da, *db, *dc;
+ int i, bad, step;
+
+ setvbuf(stdout, NULL, _IONBF, 0);
+ ha = (float *)malloc(bytes);
+ hb = (float *)malloc(bytes);
+ hc = (float *)malloc(bytes);
+ if (!ha || !hb || !hc)
+ return 1;
+ for (i = 0; i < n; i++) {
+ ha[i] = (float)(i % 977) * 0.25f;
+ hb[i] = (float)(i % 251) * 0.5f;
+ }
+ CHECK(cudaMalloc(&da, bytes));
+ CHECK(cudaMalloc(&db, bytes));
+ CHECK(cudaMalloc(&dc, bytes));
+
+ for (step = 0; step < 4; step++) {
+ const char *which = (step & 1) ? "vsub" : "vadd";
+ int is_sub = step & 1;
+
+ if (step > 0) {
+ ha[step] += 1.0f;
+ hb[step * 7 + 3] += 2.0f;
+ }
+ CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice));
+ CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice));
+ if (is_sub)
+ vsub<<<(n + 255) / 256, 256>>>(da, db, dc, n);
+ else
+ vadd<<<(n + 255) / 256, 256>>>(da, db, dc, n);
+ CHECK(cudaGetLastError());
+ CHECK(cudaDeviceSynchronize());
+ CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
+ bad = 0;
+ for (i = 0; i < n; i++) {
+ float expect = is_sub ? ha[i] - hb[i] : ha[i] + hb[i];
+
+ if (hc[i] != expect)
+ bad++;
+ }
+ printf("STEP %d %s bad=%d/%d\n", step, which, bad, n);
+ }
+ printf("STATUS OK\n");
+ return 0;
+}
diff --git a/tools/cuda-ab/wedge_bisect.cu b/tools/cuda-ab/wedge_bisect.cu
new file mode 100644
index 0000000..e062a2b
--- /dev/null
+++ b/tools/cuda-ab/wedge_bisect.cu
@@ -0,0 +1,104 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * wedge_bisect.cu - bisect the wedge trigger: which inter-launch activity
+ * makes the SECOND vadd launch no-execute?
+ *
+ * Every iteration is: mutate host inputs -> H2D da/db -> launch -> sync ->
+ * D2H dc -> verify (the wedge_recover shape). Modes add activity between
+ * the verify and the next launch:
+ *
+ * mode A: none (pure second-launch class)
+ * mode B: an extra H2D of the same buffers (CE write before the launch)
+ * mode C: an extra D2H readback (CE read before the launch)
+ * mode D: a 500 ms sleep (no CE traffic)
+ *
+ * A no-execute launch is reported (bad>0) not masked: the inputs mutate every
+ * iteration.
+ *
+ * Build: nvcc -O2 --cudart=static -o wedge_bisect wedge_bisect.cu
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <cuda_runtime.h>
+#include <unistd.h>
+
+#define CHECK(x) \
+ do { \
+ cudaError_t e_ = (x); \
+ if (e_ != cudaSuccess) { \
+ printf("CUDA_ERROR %s at line %d\n", \
+ cudaGetErrorString(e_), __LINE__); \
+ printf("STATUS FAIL\n"); \
+ return 1; \
+ } \
+ } while (0)
+
+__global__ void vadd(const float *a, const float *b, float *c, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n)
+ c[i] = a[i] + b[i];
+}
+
+int main(int argc, char **argv)
+{
+ const int n = 1 << 15;
+ size_t bytes = (size_t)n * sizeof(float);
+ float *ha, *hb, *hc;
+ float *da, *db, *dc;
+ const char *mode = argc > 1 ? argv[1] : "A";
+ int iters = argc > 2 ? atoi(argv[2]) : 4;
+ int i, bad = 0;
+
+ setvbuf(stdout, NULL, _IONBF, 0);
+ ha = (float *)malloc(bytes);
+ hb = (float *)malloc(bytes);
+ hc = (float *)malloc(bytes);
+ if (!ha || !hb || !hc)
+ return 1;
+ for (i = 0; i < n; i++) {
+ ha[i] = (float)(i % 977) * 0.25f;
+ hb[i] = (float)(i % 251) * 0.5f;
+ }
+ CHECK(cudaMalloc(&da, bytes));
+ CHECK(cudaMalloc(&db, bytes));
+ CHECK(cudaMalloc(&dc, bytes));
+ CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice));
+ CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice));
+
+ for (i = 0; i < iters; i++) {
+ int j, itbad = 0;
+
+ if (i > 0) {
+ ha[i % n] += 1.0f;
+ hb[(i * 7 + 3) % n] += 2.0f;
+ }
+ CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice));
+ CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice));
+ vadd<<<(n + 255) / 256, 256>>>(da, db, dc, n);
+ CHECK(cudaGetLastError());
+ CHECK(cudaDeviceSynchronize());
+ CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost));
+ for (j = 0; j < n; j++)
+ if (hc[j] != ha[j] + hb[j])
+ itbad++;
+ printf("MODE %s iter=%d bad=%d/%d\n", mode, i, itbad, n);
+ bad += itbad;
+ if (i + 1 < iters) {
+ if (mode[0] == 'B') {
+ CHECK(cudaMemcpy(da, ha, bytes,
+ cudaMemcpyHostToDevice));
+ } else if (mode[0] == 'C') {
+ CHECK(cudaMemcpy(hc, dc, bytes,
+ cudaMemcpyDeviceToHost));
+ } else if (mode[0] == 'D') {
+ usleep(500000);
+ }
+ }
+ }
+ printf("STATUS %s\n", bad ? "FAIL" : "OK");
+ return bad ? 1 : 0;
+}
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-21 20:37:50 +0000