From 393f1c676d0fc9c7b37ab6a52efe7e2d288d7e90 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Wed, 21 Oct 2026 14:02:15 -0400 Subject: gr: r420b author the guest QMD release on every outcome when release0 is redirected With gsp_phys_rmcall_gr_rel_mod=1 (release0 -> module semaphore, the ordering fix for the engine-release-before-copy-out race) the engine never writes the guest's own QMD release word, so the host must author it on EVERY backend outcome. The hold's success/fail edges do; a launch with no deferred release records (no hold engaged) fell through and its client hung on a word nothing would ever write - the r414 saxpy f0 cur=5 stall. The backend now reports qmd_rel_redirected in the copy result whenever it points release0 at the module semaphore (rel_mod, or no targetable guest release page). svc_fwd_gr_work folds that into the existing gr_rel_unobserved latch, which makes the bounded outcome check author the word on done (success or failure) - the same forward-only write the hold edges perform. rel_mod=0 paths are unchanged (the engine still authors the word and the latch is only set on an unresolved release page). --- tools/cuda-ab/transformer.cu | 472 +++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 472 insertions(+) create mode 100644 tools/cuda-ab/transformer.cu (limited to 'tools/cuda-ab/transformer.cu') 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 cc=. sm= mem= + * TRANSFORMER dh=256 seq=64 bad= ysum= + * FFN4 layers=4 bad= psum= + * 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 +#include +#include +#include +#include + +#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<<>>(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<<>>(A, B, C, M, N, K); +} + +static void launch_softmax(const float *S, float *P, int rows, int cols, + float scale) +{ + softmax_rows<<>>(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; +} -- cgit v1.2.3