diff options
Diffstat (limited to 'tools')
28 files changed, 2833 insertions, 2 deletions
diff --git a/tools/cuda-ab/ip_probe.cu b/tools/cuda-ab/ip_probe.cu new file mode 100644 index 0000000..366fbc8 --- /dev/null +++ b/tools/cuda-ab/ip_probe.cu @@ -0,0 +1,128 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * ip_probe.cu - how many CTAs actually executed for an in-place kernel? + * + * scale_mark marks progress[blockIdx.x]=0xCAFEBABE after the c store; D2H of + * progress counts executed CTAs exactly. c (the in-place buffer) is read back + * and verified. seed (argv[1]) makes the run content-unique. + * + * Build: nvcc -O2 --cudart=static -o ip_probe ip_probe.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstring> +#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 scale_mark(float *c, float s, int n, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0xCAFEBABEu; +} + +int main(int argc, char **argv) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + const int grid = (n + 255) / 256; + size_t bytes = (size_t)n * sizeof(float); + float k = argc > 1 ? atof(argv[1]) : 0.0f; + float *ha, *hb, *hc; + float *da, *db, *dc; + unsigned *d_prog, h_prog[128]; + int i; + + setvbuf(stdout, NULL, _IONBF, 0); + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s seed=%g grid=%d\n", prop.name, (double)k, grid); + + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 1000) * 0.5f + k; + hb[i] = (float)(i % 997) * 0.25f + 2.0f * k; + } + memset(hc, 0xCD, bytes); + + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMalloc(&d_prog, 128 * sizeof(unsigned))); + printf("DPTRS da=%p db=%p dc=%p prog=%p\n", (void *)da, (void *)db, + (void *)dc, (void *)d_prog); + + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemset(d_prog, 0xEE, 128 * sizeof(unsigned))); + + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("VADD readback bad=%d/%d\n", b, n); + } + + /* in-place scale with CTA markers */ + scale_mark<<<grid, 256>>>(dc, 2.0f, n, d_prog); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_prog, d_prog, 128 * sizeof(unsigned), + cudaMemcpyDeviceToHost)); + { + int marked = 0, first_missing = -1; + + for (i = 0; i < grid; i++) { + if (h_prog[i] == 0xCAFEBABEu) + marked++; + else if (first_missing < 0) + first_missing = i; + } + printf("CTA-MARK marked=%d/%d first_missing=%d\n", marked, grid, + first_missing); + } + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != (ha[i] + hb[i]) * 2.0f) { + b++; + if (first < 0) + first = i; + } + printf("SCALE readback bad=%d/%d first=%d\n", b, n, first); + if (first >= 0) + printf("SCALE firstbad got=%g want=%g\n", hc[first], + (ha[first] + hb[first]) * 2.0f); + } + printf("STATUS DONE\n"); + return 0; +} diff --git a/tools/cuda-ab/ip_probe2.cu b/tools/cuda-ab/ip_probe2.cu new file mode 100644 index 0000000..a5fa2ec --- /dev/null +++ b/tools/cuda-ab/ip_probe2.cu @@ -0,0 +1,201 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * ip_probe2.cu - why do in-place c stores vanish while prog stores land? + * + * After a vadd that fills c, three variants run on the SAME c: + * 1. scale_if : if (i<n) c[i] *= s; (control, like ip_probe) + * 2. scale_u : c[i] *= s; (unconditional, grid exact) + * 3. fill_c : c[i] = 7.0f; (store only, no load) + * 4. add_c : c[i] = c[i] + 1.0f; (load + store) + * Each stage reads c back and compares against the expected value at that + * stage. A separate prog buffer marks CTA execution. + * + * Build: nvcc -O2 --cudart=static -o ip_probe2 ip_probe2.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstring> +#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 scale_if(float *c, float s, int n, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0xCAFEBABEu; +} + +__global__ void scale_u(float *c, float s, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + c[i] = c[i] * s; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0xDEADBEEFu; +} + +__global__ void fill_c(float *c, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + c[i] = 7.0f; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0xFEEDFACu; +} + +__global__ void add_c(float *c, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + c[i] = c[i] + 1.0f; + if (threadIdx.x == 0) + prog[blockIdx.x] = 0x0DDF00Du; +} + +int main(int argc, char **argv) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + const int grid = (n + 255) / 256; + size_t bytes = (size_t)n * sizeof(float); + float k = argc > 1 ? atof(argv[1]) : 0.0f; + float *ha, *hb, *hc; + float *da, *db, *dc; + unsigned *d_prog, h_prog[128]; + int i, stage; + float mult; + + setvbuf(stdout, NULL, _IONBF, 0); + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s seed=%g\n", prop.name, (double)k); + + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 1000) * 0.5f + k; + hb[i] = (float)(i % 997) * 0.25f + 2.0f * k; + } + memset(hc, 0xCD, bytes); + + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMalloc(&d_prog, 128 * sizeof(unsigned))); + printf("DPTRS da=%p db=%p dc=%p prog=%p\n", (void *)da, (void *)db, + (void *)dc, (void *)d_prog); + + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("VADD readback bad=%d/%d\n", b, n); + } + + mult = 2.0f; + for (stage = 0; stage < 4; stage++) { + unsigned mark = 0; + + CHECK(cudaMemset(d_prog, 0xEE, 128 * sizeof(unsigned))); + switch (stage) { + case 0: + scale_if<<<grid, 256>>>(dc, 2.0f, n, d_prog); + mult = 2.0f; + mark = 0xCAFEBABEu; + break; + case 1: + scale_u<<<grid, 256>>>(dc, 2.0f, d_prog); + mult = 4.0f; + mark = 0xDEADBEEFu; + break; + case 2: + fill_c<<<grid, 256>>>(dc, d_prog); + mark = 0xFEEDFACu; + break; + default: + add_c<<<grid, 256>>>(dc, d_prog); + mult = -1.0f; + mark = 0x0DDF00Du; + break; + } + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_prog, d_prog, 128 * sizeof(unsigned), + cudaMemcpyDeviceToHost)); + { + int marked = 0; + + for (i = 0; i < grid; i++) + if (h_prog[i] == mark) + marked++; + printf("S%d CTA-MARK marked=%d/%d\n", stage + 1, + marked, grid); + } + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + if (stage < 2) { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != (ha[i] + hb[i]) * mult) { + b++; + if (first < 0) + first = i; + } + printf("S%d readback bad=%d/%d first=%d\n", + stage + 1, b, n, first); + } else if (stage == 2) { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != 7.0f) { + b++; + if (first < 0) + first = i; + } + printf("S3 readback bad=%d/%d first=%d (fill 7.0)\n", + b, n, first); + } else { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != 8.0f) { + b++; + if (first < 0) + first = i; + } + printf("S4 readback bad=%d/%d first=%d (add 1.0 to fill)\n", + b, n, first); + } + } + printf("STATUS DONE\n"); + return 0; +} diff --git a/tools/cuda-ab/ip_probe3.cu b/tools/cuda-ab/ip_probe3.cu new file mode 100644 index 0000000..a50df0e --- /dev/null +++ b/tools/cuda-ab/ip_probe3.cu @@ -0,0 +1,142 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * ip_probe3.cu - what does the SM actually READ for the scale kernel? + * + * scale_report writes, per CTA: prog[2b] = bits of the float param s, + * prog[2b+1] = bits of c[b*256] (the load). + * The D2H of prog then shows exactly what the SM saw, independent of any + * store-visibility path for c. + * + * Build: nvcc -O2 --cudart=static -o ip_probe3 ip_probe3.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstring> +#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 scale_report(float *c, float s, int n, unsigned *prog) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; + if (threadIdx.x == 0) { + prog[2 * blockIdx.x] = __float_as_uint(s); + prog[2 * blockIdx.x + 1] = __float_as_uint(c[blockIdx.x * 256]); + } +} + +int main(int argc, char **argv) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + const int grid = (n + 255) / 256; + size_t bytes = (size_t)n * sizeof(float); + float k = argc > 1 ? atof(argv[1]) : 0.0f; + float *ha, *hb, *hc; + float *da, *db, *dc; + unsigned *d_prog, h_prog[256]; + int i; + + setvbuf(stdout, NULL, _IONBF, 0); + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s seed=%g\n", prop.name, (double)k); + + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 1000) * 0.5f + k; + hb[i] = (float)(i % 997) * 0.25f + 2.0f * k; + } + memset(hc, 0xCD, bytes); + + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMalloc(&d_prog, 256 * sizeof(unsigned))); + printf("DPTRS da=%p db=%p dc=%p prog=%p\n", (void *)da, (void *)db, + (void *)dc, (void *)d_prog); + + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemset(d_prog, 0xEE, 256 * sizeof(unsigned))); + + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("VADD readback bad=%d/%d\n", b, n); + } + + scale_report<<<grid, 256>>>(dc, 2.0f, n, d_prog); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_prog, d_prog, 256 * sizeof(unsigned), + cudaMemcpyDeviceToHost)); + { + int s_ok = 0, s_1 = 0, c_ok = 0, n_seen = 0; + unsigned first_s = 0, first_c = 0; + + for (i = 0; i < grid; i++) { + unsigned sb = h_prog[2 * i]; + unsigned cb = h_prog[2 * i + 1]; + + if (i == 0) { + first_s = sb; + first_c = cb; + } + if (sb == 0x40000000u) + s_ok++; + if (sb == 0x3F800000u) + s_1++; + { + float want = ha[i * 256] + hb[i * 256]; + + if (cb == *(unsigned *)&want) + c_ok++; + } + n_seen++; + } + printf("REPORT ctas=%d s_eq_2.0=%d s_eq_1.0=%d c_load_ok=%d first_s=0x%08x first_c=0x%08x\n", + n_seen, s_ok, s_1, c_ok, first_s, first_c); + } + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0, first = -1; + + for (i = 0; i < n; i++) + if (hc[i] != (ha[i] + hb[i]) * 2.0f) { + b++; + if (first < 0) + first = i; + } + printf("SCALE readback bad=%d/%d first=%d\n", b, n, first); + } + printf("STATUS DONE\n"); + return 0; +} diff --git a/tools/cuda-ab/memset2d_probe.cu b/tools/cuda-ab/memset2d_probe.cu new file mode 100644 index 0000000..acf43d4 --- /dev/null +++ b/tools/cuda-ab/memset2d_probe.cu @@ -0,0 +1,157 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * memset2d_probe.cu - per-phase diagnostics for victory_memset2d. + * + * victory_memset2d reports one aggregate bad count (~12.5M) that mixes all + * three phases. This probe runs the SAME three phases and reports each + * phase's bad count, the first mismatching byte (offset, got, want), and a + * per-phase 1MB-chunk histogram so the failing phase and the failing region + * are attributable in one run. + * + * Phases: + * P1 memset : cudaMemset(d_img, 0xA5, pitch*h); D2H 2D readback + * P2 roundtrip: H2D 2D pattern; D2H 2D readback + * P3 memset2d : cudaMemset2D quadrant 0x3C; D2H 2D readback + * + * Output lines: + * DEVICE ... + * PHASE p=1 bad=<n> first_off=<i> got=<g> want=<w> + * PHIST p=1 mb=<m> ok=<n> bad=<n> + * MEMSET2D2 bad_total=<n> + * STATUS OK|FAIL + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cstring> +#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 IMG_W 4096 +#define IMG_H 1024 +#define FILL_BYTE 0xA5 +#define QUAD_BYTE 0x3C + +static int phase_check(const char *tag, const unsigned char *back, + const unsigned char *want, size_t n, int p) +{ + size_t i, first = (size_t)-1; + int bad = 0, got = 0, wantv = 0; + size_t mb; + + for (i = 0; i < n; i++) { + if (back[i] != want[i]) { + if (first == (size_t)-1) { + first = i; + got = back[i]; + wantv = want[i]; + } + bad++; + } + } + printf("PHASE p=%d bad=%d first_off=%zu got=0x%02x want=0x%02x\n", + p, bad, first, got, wantv); + if (bad) { + size_t nb = (n + (1u << 20) - 1) >> 20; + + for (mb = 0; mb < nb && mb < 8; mb++) { + size_t lo = mb << 20, hi = lo + (1u << 20); + size_t j; + int bok = 0, bbad = 0; + + if (hi > n) + hi = n; + for (j = lo; j < hi; j++) { + if (back[j] == want[j]) + bok++; + else + bbad++; + } + printf("PHIST p=%d mb=%zu ok=%d bad=%d\n", + p, mb, bok, bbad); + } + } + (void)tag; + return bad; +} + +int main(void) +{ + cudaDeviceProp prop; + const size_t w = IMG_W, h = IMG_H; + unsigned char *h_img, *h_back, *h_want; + unsigned char *d_img; + size_t pitch = 0; + int bad = 0, row, col; + + 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); + + h_img = (unsigned char *)malloc(w * h); + h_back = (unsigned char *)malloc(w * h); + h_want = (unsigned char *)malloc(w * h); + if (!h_img || !h_back || !h_want) + return 1; + + CHECK(cudaMallocPitch(&d_img, &pitch, w, h)); + printf("PITCH pitch=%zu\n", pitch); + printf("DPTRS h_img=%p h_back=%p h_want=%p d_img=%p\n", (void *)h_img, + (void *)h_back, (void *)h_want, (void *)d_img); + + /* P1: full cudaMemset over the pitched allocation, 2D readback. */ + CHECK(cudaMemset(d_img, FILL_BYTE, pitch * h)); + memset(h_back, 0, w * h); + CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h, + cudaMemcpyDeviceToHost)); + memset(h_want, FILL_BYTE, w * h); + bad += phase_check("P1", h_back, h_want, w * h, 1); + + /* P2: pitched H2D pattern -> D2H round trip. */ + for (row = 0; row < (int)h; row++) + for (col = 0; col < (int)w; col++) + h_img[row * w + col] = (unsigned char)((row + col) & 0xFF); + CHECK(cudaMemcpy2D(d_img, pitch, h_img, w, w, h, + cudaMemcpyHostToDevice)); + memset(h_back, 0, w * h); + CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h, + cudaMemcpyDeviceToHost)); + bad += phase_check("P2", h_back, h_img, w * h, 2); + + /* P3: cudaMemset2D of the upper-left quadrant only. */ + CHECK(cudaMemset2D(d_img, pitch, QUAD_BYTE, w / 2, h / 2)); + for (row = 0; row < (int)h; row++) + for (col = 0; col < (int)w; col++) { + unsigned char expect = h_img[row * w + col]; + + if (row < (int)(h / 2) && col < (int)(w / 2)) + expect = QUAD_BYTE; + h_want[row * w + col] = expect; + } + memset(h_back, 0, w * h); + CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h, + cudaMemcpyDeviceToHost)); + bad += phase_check("P3", h_back, h_want, w * h, 3); + + printf("MEMSET2D2 bad_total=%d\n", bad); + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + printf("STATUS OK\n"); + return 0; +} diff --git a/tools/cuda-ab/rt_probe.cu b/tools/cuda-ab/rt_probe.cu new file mode 100644 index 0000000..025301d --- /dev/null +++ b/tools/cuda-ab/rt_probe.cu @@ -0,0 +1,165 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * rt_probe.cu - instrumented memcpy_roundtrip variant isolating the + * memset-phase D2H readback failure. + * + * Phases (argv[1]): + * "full" : H2D -> kernel -> D2H -> memset -> D2H (the original order) + * "sync" : same but cudaDeviceSynchronize() before the second D2H + * "memset" : memset -> D2H only (memset_probe shape, 3-buffer alloc) + * "rev" : memset -> D2H first, then H2D -> kernel -> D2H + * + * Forensics on the memset readback: first bad offset, byte histogram + * (0x5A / 0x00 / other), per-1MB chunk histogram, and the first 32 bytes. + * + * Build: nvcc -O2 --cudart=static -o rt_probe rt_probe.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cstring> +#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 MemsetByte 0x5A + +__global__ void bump(const float *in, float *out, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + out[i] = in[i] + 0.5f; +} + +static void memset_forensics(const unsigned char *h, size_t bytes, + const char *tag) +{ + size_t i, first_bad = (size_t)-1; + int bad = 0, n5a = 0, n00 = 0, nother = 0; + size_t chunk; + + for (i = 0; i < bytes; i++) { + if (h[i] == 0x5A) + n5a++; + else if (h[i] == 0x00) + n00++; + else + nother++; + if (h[i] != 0x5A) { + if (first_bad == (size_t)-1) + first_bad = i; + bad++; + } + } + printf("MEMSET-READ %s bad=%d first_bad=%zu n5a=%d n00=%d nother=%d\n", + tag, bad, first_bad, n5a, n00, nother); + printf("MEMSET-HEAD %s", tag); + for (i = 0; i < 32; i++) + printf(" %02x", h[i]); + printf("\n"); + if (first_bad != (size_t)-1) { + printf("MEMSET-ATBAD %s off=%zu:", tag, first_bad); + for (i = first_bad; i < first_bad + 16 && i < bytes; i++) + printf(" %02x", h[i]); + printf("\n"); + } + for (chunk = 0; chunk < bytes / (1u << 20); chunk++) { + size_t lo = chunk << 20, hi = lo + (1u << 20); + int c5a = 0, c00 = 0, cother = 0; + + for (i = lo; i < hi; i++) { + if (h[i] == 0x5A) + c5a++; + else if (h[i] == 0x00) + c00++; + else + cother++; + } + printf("MEMSET-MB %s mb=%zu 5a=%d 00=%d other=%d\n", + tag, chunk, c5a, c00, cother); + } +} + +int main(int argc, char **argv) +{ + const char *mode = argc > 1 ? argv[1] : "full"; + cudaDeviceProp prop; + const size_t bytes = 8u << 20; + const int n = (int)(bytes / sizeof(float)); + unsigned char *h_m; + float *h_in, *h_out; + float *d_in, *d_out; + unsigned char *d_m; + double csum = 0.0; + int bad = 0, i; + const int grid = (n + 255) / 256; + + setvbuf(stdout, NULL, _IONBF, 0); + + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s cc=%d.%d sm=%d mem=%llu mode=%s\n", prop.name, + prop.major, prop.minor, prop.multiProcessorCount, + (unsigned long long)prop.totalGlobalMem, mode); + + h_in = (float *)malloc(bytes); + h_out = (float *)malloc(bytes); + h_m = (unsigned char *)malloc(bytes); + if (!h_in || !h_out || !h_m) + return 1; + + for (i = 0; i < n; i++) + h_in[i] = (float)(i % 1024) * 0.25f; + + CHECK(cudaMalloc(&d_in, bytes)); + CHECK(cudaMalloc(&d_out, bytes)); + CHECK(cudaMalloc(&d_m, bytes)); + printf("DPTRS d_in=%p d_out=%p d_m=%p\n", (void *)d_in, + (void *)d_out, (void *)d_m); + + if (!strcmp(mode, "memset") || !strcmp(mode, "rev")) { + /* memset phase first */ + CHECK(cudaMemset(d_m, MemsetByte, bytes)); + if (!strcmp(mode, "sync")) + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_m, d_m, bytes, cudaMemcpyDeviceToHost)); + memset_forensics(h_m, bytes, "early"); + } + + if (!strcmp(mode, "full") || !strcmp(mode, "sync") || + !strcmp(mode, "rev")) { + CHECK(cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice)); + bump<<<grid, 256>>>(d_in, d_out, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost)); + + for (i = 0; i < n; i++) { + if (h_out[i] != h_in[i] + 0.5f) + bad++; + csum += (double)h_out[i]; + } + printf("BUMP bad=%d csum=%.6f\n", bad, csum); + } + + if (!strcmp(mode, "full") || !strcmp(mode, "sync")) { + CHECK(cudaMemset(d_m, MemsetByte, bytes)); + if (!strcmp(mode, "sync")) + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_m, d_m, bytes, cudaMemcpyDeviceToHost)); + memset_forensics(h_m, bytes, "late"); + } + + printf("STATUS OK\n"); + return 0; +} diff --git a/tools/cuda-ab/run/ip4 b/tools/cuda-ab/run/ip4 Binary files differnew file mode 100755 index 0000000..54937e6 --- /dev/null +++ b/tools/cuda-ab/run/ip4 diff --git a/tools/cuda-ab/run/ip5 b/tools/cuda-ab/run/ip5 Binary files differnew file mode 100755 index 0000000..c4b4a39 --- /dev/null +++ b/tools/cuda-ab/run/ip5 diff --git a/tools/cuda-ab/run/ip6 b/tools/cuda-ab/run/ip6 Binary files differnew file mode 100755 index 0000000..56815c2 --- /dev/null +++ b/tools/cuda-ab/run/ip6 diff --git a/tools/cuda-ab/run/ip_probe b/tools/cuda-ab/run/ip_probe Binary files differnew file mode 100755 index 0000000..08fb86e --- /dev/null +++ b/tools/cuda-ab/run/ip_probe diff --git a/tools/cuda-ab/run/ip_probe2 b/tools/cuda-ab/run/ip_probe2 Binary files differnew file mode 100755 index 0000000..dfe6ab9 --- /dev/null +++ b/tools/cuda-ab/run/ip_probe2 diff --git a/tools/cuda-ab/run/ip_probe3 b/tools/cuda-ab/run/ip_probe3 Binary files differnew file mode 100755 index 0000000..8a34c05 --- /dev/null +++ b/tools/cuda-ab/run/ip_probe3 diff --git a/tools/cuda-ab/run/memset2d_probe b/tools/cuda-ab/run/memset2d_probe Binary files differnew file mode 100755 index 0000000..f5bd071 --- /dev/null +++ b/tools/cuda-ab/run/memset2d_probe diff --git a/tools/cuda-ab/run/rt_probe b/tools/cuda-ab/run/rt_probe Binary files differnew file mode 100755 index 0000000..48d866f --- /dev/null +++ b/tools/cuda-ab/run/rt_probe diff --git a/tools/cuda-ab/run/scale_probe b/tools/cuda-ab/run/scale_probe Binary files differnew file mode 100755 index 0000000..ca3fc20 --- /dev/null +++ b/tools/cuda-ab/run/scale_probe diff --git a/tools/cuda-ab/run/streams2_probe b/tools/cuda-ab/run/streams2_probe Binary files differnew file mode 100755 index 0000000..206d6d6 --- /dev/null +++ b/tools/cuda-ab/run/streams2_probe diff --git a/tools/cuda-ab/run/streams_events_probe b/tools/cuda-ab/run/streams_events_probe Binary files differnew file mode 100755 index 0000000..00f3539 --- /dev/null +++ b/tools/cuda-ab/run/streams_events_probe diff --git a/tools/cuda-ab/run/tgemm_probe b/tools/cuda-ab/run/tgemm_probe Binary files differnew file mode 100755 index 0000000..5fcdec0 --- /dev/null +++ b/tools/cuda-ab/run/tgemm_probe diff --git a/tools/cuda-ab/run/transformer_async b/tools/cuda-ab/run/transformer_async Binary files differnew file mode 100755 index 0000000..2911b5c --- /dev/null +++ b/tools/cuda-ab/run/transformer_async diff --git a/tools/cuda-ab/run/transformer_stages b/tools/cuda-ab/run/transformer_stages Binary files differnew file mode 100755 index 0000000..1461a3e --- /dev/null +++ b/tools/cuda-ab/run/transformer_stages diff --git a/tools/cuda-ab/scale_probe.cu b/tools/cuda-ab/scale_probe.cu new file mode 100644 index 0000000..8a0df59 --- /dev/null +++ b/tools/cuda-ab/scale_probe.cu @@ -0,0 +1,245 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * scale_probe.cu - localize the deterministic "in-place scale writes only the + * first 4 KB page" failure. + * + * Modes (argv[1]): + * V : vadd(a,b->c); sync; D2H+verify; vadd again; D2H+verify + * I : vadd(a,b->c); sync; D2H+verify; scale(c,*2 in place); sync; + * D2H+verify (with page map); scale again; sync; D2H+verify; D2H again + * O : vadd(a,b->c); sync; D2H+verify; scale_out(c->d,*2); sync; + * D2H+verify (with page map) + * + * Every stage also runs a GPU-side verify kernel that counts matching + * elements reading the device buffers (engine's own view) and writes the + * counts to a small buffer, D2H'd afterwards: separates "device pages never + * written" (GPU count low) from "D2H readback stale" (GPU count high, host + * count low). + * + * Build: nvcc -O2 --cudart=static -o scale_probe scale_probe.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cstring> +#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 scale(float *c, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; +} + +__global__ void scale_out(const float *c, float *d, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + d[i] = c[i] * s; +} + +/* verify kernel: out[0]=count(c==a+b), out[1]=count(c==2*(a+b)) */ +__global__ void vdigest(const float *a, const float *b, const float *c, int n, + unsigned long long *out) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i >= n) + return; + if (c[i] == a[i] + b[i]) + atomicAdd(&out[0], 1ull); + if (c[i] == (a[i] + b[i]) * 2.0f) + atomicAdd(&out[1], 1ull); +} + +int main(int argc, char **argv) +{ + const char *mode = argc > 1 ? argv[1] : "V"; + cudaDeviceProp prop; + const int n = 1 << 15; + const int grid = (n + 255) / 256; + size_t bytes = (size_t)n * sizeof(float); + float *ha, *hb, *hc, *hd; + float *da, *db, *dc, *dd; + unsigned long long *d_out, h_out[2]; + int i; + float k = argc > 2 ? atof(argv[2]) : 0.0f; + + setvbuf(stdout, NULL, _IONBF, 0); + + CHECK(cudaGetDeviceProperties(&prop, 0)); + { + FILE *f = fopen("/proc/uptime", "r"); + double up = -1.0; + + if (f) { if (fscanf(f, "%lf", &up) != 1) up = -1.0; fclose(f); } + printf("UPTIME %.3f\n", up); + } + printf("DEVICE %s mode=%s seed=%g\n", prop.name, mode, (double)k); + + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + hd = (float *)malloc(bytes); + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 1000) * 0.5f + k; + hb[i] = (float)(i % 997) * 0.25f + 2.0f * k; + } + memset(hc, 0xCD, bytes); + memset(hd, 0xCD, bytes); + + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMalloc(&dd, bytes)); + CHECK(cudaMalloc(&d_out, 2 * sizeof(unsigned long long))); + printf("DPTRS da=%p db=%p dc=%p dd=%p\n", (void *)da, (void *)db, + (void *)dc, (void *)dd); + + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("S1 vadd readback bad=%d/%d\n", b, n); + CHECK(cudaMemset(d_out, 0, 2 * sizeof(unsigned long long))); + vdigest<<<grid, 256>>>(da, db, dc, n, d_out); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, + 2 * sizeof(unsigned long long), + cudaMemcpyDeviceToHost)); + printf("S1 GPU-DIGEST c_eq_ab=%llu c_eq_2ab=%llu\n", h_out[0], + h_out[1]); + } + + if (!strcmp(mode, "V")) { + vadd<<<grid, 256>>>(da, db, dc, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0; + + for (i = 0; i < n; i++) + if (hc[i] != ha[i] + hb[i]) + b++; + printf("S2 vadd-again readback bad=%d/%d\n", b, n); + } + printf("STATUS DONE\n"); + return 0; + } + + if (!strcmp(mode, "I")) { + int stage; + + for (stage = 0; stage < 2; stage++) { + float mult = stage == 0 ? 2.0f : 4.0f; + + scale<<<grid, 256>>>(dc, 2.0f, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0, page; + + for (i = 0; i < n; i++) + if (hc[i] != (ha[i] + hb[i]) * mult) + b++; + printf("S%d inplace-scale readback bad=%d/%d (mult=%g)\n", + 2 + stage, b, n, mult); + if (stage == 0) + for (page = 0; page < 32; page++) { + int pb = 0, k; + + for (k = 0; k < 256; k++) + if (hc[page * 256 + k] != + (ha[page * 256 + k] + + hb[page * 256 + k]) * mult) + pb++; + if (pb) + printf("S%d PAGE %d bad=%d/256 first=%g want=%g\n", + 2 + stage, page, + pb, hc[page * 256], + (ha[page * 256] + + hb[page * 256]) * mult); + } + } + if (stage == 0) { + CHECK(cudaMemset(d_out, 0, + 2 * sizeof(unsigned long long))); + vdigest<<<grid, 256>>>(da, db, dc, n, d_out); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, + 2 * sizeof(unsigned long long), + cudaMemcpyDeviceToHost)); + printf("S2 GPU-DIGEST c_eq_ab=%llu c_eq_2ab=%llu\n", + h_out[0], h_out[1]); + } + } + printf("STATUS DONE\n"); + return 0; + } + + if (!strcmp(mode, "O")) { + scale_out<<<grid, 256>>>(dc, dd, 2.0f, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hd, dd, bytes, cudaMemcpyDeviceToHost)); + { + int b = 0, page; + + for (i = 0; i < n; i++) + if (hd[i] != (ha[i] + hb[i]) * 2.0f) + b++; + printf("S2 outofplace-scale readback bad=%d/%d\n", b, n); + for (page = 0; page < 32; page++) { + int pb = 0, k; + + for (k = 0; k < 256; k++) + if (hd[page * 256 + k] != + (ha[page * 256 + k] + + hb[page * 256 + k]) * 2.0f) + pb++; + if (pb) + printf("S2 PAGE %d bad=%d/256\n", page, + pb); + } + } + CHECK(cudaMemset(d_out, 0, 2 * sizeof(unsigned long long))); + vdigest<<<grid, 256>>>(da, db, dd, n, d_out); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, 2 * sizeof(unsigned long long), + cudaMemcpyDeviceToHost)); + printf("S2 GPU-DIGEST d_eq_ab=%llu d_eq_2ab=%llu\n", h_out[0], + h_out[1]); + printf("STATUS DONE\n"); + return 0; + } + return 1; +} diff --git a/tools/cuda-ab/soak.sh b/tools/cuda-ab/soak.sh index c150692..4f1430f 100755 --- a/tools/cuda-ab/soak.sh +++ b/tools/cuda-ab/soak.sh @@ -192,7 +192,7 @@ log "SOAK START mode=$MODE seconds=$SECS guest_dir=$GUEST_DIR timeout=${TMO}s lo # cannot silently shrink the battery. found="" for i in 1 2 3; do - found="$(gssh "for b in $CLIENTS; do [ -x $GUEST_DIR/\$b ] && echo \$b; done")" || found="" + found="$(gssh "for b in $CLIENTS; do [ -x $GUEST_DIR/\$b ] && echo \$b; done; true")" || found="" [ -n "$found" ] && break sleep 3 done diff --git a/tools/cuda-ab/streams2_probe.cu b/tools/cuda-ab/streams2_probe.cu new file mode 100644 index 0000000..5bebab4 --- /dev/null +++ b/tools/cuda-ab/streams2_probe.cu @@ -0,0 +1,299 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * streams2_probe.cu - instrumented streams2 to localize the deterministic + * failure: per-array mismatch counts, first-bad indices with values, and a + * GPU-side verify-kernel digest of the device buffers (to separate "device + * pages never written" from "D2H readback wrong"). + * + * Modes (argv[1]): + * A : exactly streams2's shape (control) + * B : like A but the two D2H copy-backs use synchronous cudaMemcpy + * C : like A but no events cross the streams; both streams are joined + * with cudaStreamSynchronize before any copy-back + * D : single stream (everything on s1) + * + * Build: nvcc -O2 --cudart=static -o streams2_probe streams2_probe.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cstring> +#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 scale(float *c, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; +} + +/* GPU-side digest: counts of matches for c1==a1+b1 and c2==(a2+b2)*2 */ +__global__ void digest(const float *a1, const float *b1, const float *c1, + const float *a2, const float *b2, const float *c2, + int n, unsigned long long *out) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i >= n) + return; + if (c1[i] == a1[i] + b1[i]) + atomicAdd(&out[0], 1ull); + if (c2[i] == (a2[i] + b2[i]) * 2.0f) + atomicAdd(&out[1], 1ull); +} + +static int report(const char *tag, const float *h_c, const float *h_a, + const float *h_b, const float *h_a2, const float *h_b2, + int n, float scale2) +{ + int bad = 0, i, shown = 0; + + for (i = 0; i < n; i++) { + float want = h_a[i] + h_b[i]; + + if (scale2 != 1.0f) + want = (h_a2[i] + h_b2[i]) * scale2; + if (h_c[i] != want) { + bad++; + if (shown < 4) { + printf("FIRSTBAD %s i=%d got=%g want=%g\n", + tag, i, h_c[i], want); + shown++; + } + } + } + printf("BAD %s n=%d bad=%d\n", tag, n, bad); + return bad; +} + +int main(int argc, char **argv) +{ + const char *mode = argc > 1 ? argv[1] : "A"; + cudaDeviceProp prop; + const int n = 1 << 15; + const float s2 = 2.0f; + size_t bytes = (size_t)n * sizeof(float); + float *h_a1, *h_b1, *h_c1, *h_a2, *h_b2, *h_c2; + float *d_a1, *d_b1, *d_c1, *d_a2, *d_b2, *d_c2; + unsigned long long *d_out, h_out[2]; + cudaStream_t s1, s2s; + cudaEvent_t ev_add1, ev_scale2, t0, t1; + float ms = 0.0f; + int i; + const int grid = (n + 255) / 256; + + setvbuf(stdout, NULL, _IONBF, 0); + + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s mode=%s\n", prop.name, mode); + + h_a1 = (float *)malloc(bytes); + h_b1 = (float *)malloc(bytes); + h_c1 = (float *)malloc(bytes); + h_a2 = (float *)malloc(bytes); + h_b2 = (float *)malloc(bytes); + h_c2 = (float *)malloc(bytes); + + for (i = 0; i < n; i++) { + h_a1[i] = (float)(i % 1000) * 0.5f; + h_b1[i] = (float)(i % 997) * 0.25f; + h_a2[i] = (float)(i % 991) * 0.125f; + h_b2[i] = (float)(i % 983) * 0.0625f; + } + memset(h_c1, 0xCD, bytes); + memset(h_c2, 0xCD, bytes); + + CHECK(cudaMalloc(&d_a1, bytes)); + CHECK(cudaMalloc(&d_b1, bytes)); + CHECK(cudaMalloc(&d_c1, bytes)); + CHECK(cudaMalloc(&d_a2, bytes)); + CHECK(cudaMalloc(&d_b2, bytes)); + CHECK(cudaMalloc(&d_c2, bytes)); + CHECK(cudaMalloc(&d_out, 2 * sizeof(unsigned long long))); + printf("DPTRS da1=%p db1=%p dc1=%p da2=%p db2=%p dc2=%p\n", + (void *)d_a1, (void *)d_b1, (void *)d_c1, (void *)d_a2, + (void *)d_b2, (void *)d_c2); + + CHECK(cudaStreamCreate(&s1)); + CHECK(cudaStreamCreate(&s2s)); + CHECK(cudaEventCreate(&ev_add1)); + CHECK(cudaEventCreate(&ev_scale2)); + CHECK(cudaEventCreate(&t0)); + CHECK(cudaEventCreate(&t1)); + + CHECK(cudaEventRecord(t0)); + + if (!strcmp(mode, "D")) { + CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, + cudaMemcpyHostToDevice, s1)); + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + vadd<<<grid, 256, 0, s1>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaStreamSynchronize(s1)); + scale<<<grid, 256, 0, s1>>>(d_c2, s2, n); + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaMemcpy(h_c1, d_c1, bytes, cudaMemcpyDeviceToHost)); + CHECK(cudaMemcpy(h_c2, d_c2, bytes, cudaMemcpyDeviceToHost)); + goto verify; + } + + if (!strcmp(mode, "S")) { + /* staged verification: check device content after every step */ + float *hb1 = (float *)malloc(bytes); + float *hb2 = (float *)malloc(bytes); + float *hb3 = (float *)malloc(bytes); + + CHECK(cudaMemcpy(d_a1, h_a1, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_b1, h_b1, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_a2, h_a2, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_b2, h_b2, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(hb1, d_a1, bytes, cudaMemcpyDeviceToHost)); + { int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_a1[i]) b++; + printf("STAGE-IN da1 bad=%d\n", b); } + CHECK(cudaMemcpy(hb1, d_b1, bytes, cudaMemcpyDeviceToHost)); + { int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_b1[i]) b++; + printf("STAGE-IN db1 bad=%d\n", b); } + CHECK(cudaMemcpy(hb1, d_a2, bytes, cudaMemcpyDeviceToHost)); + { int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_a2[i]) b++; + printf("STAGE-IN da2 bad=%d\n", b); } + CHECK(cudaMemcpy(hb1, d_b2, bytes, cudaMemcpyDeviceToHost)); + { int b = 0; for (i = 0; i < n; i++) if (hb1[i] != h_b2[i]) b++; + printf("STAGE-IN db2 bad=%d\n", b); } + + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hb1, d_c1, bytes, cudaMemcpyDeviceToHost)); + { int b = 0, first = -1; for (i = 0; i < n; i++) + if (hb1[i] != h_a1[i] + h_b1[i]) { b++; + if (first < 0) { first = i; + printf("S1-FIRSTBAD i=%d got=%g want=%g a2b2=%g\n", + i, hb1[i], h_a1[i]+h_b1[i], h_a2[i]+h_b2[i]); } } + printf("STAGE vadd1->c1 readback bad=%d first=%d\n", b, first); } + + CHECK(cudaDeviceSynchronize()); + vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hb2, d_c2, bytes, cudaMemcpyDeviceToHost)); + { int b = 0, first = -1; for (i = 0; i < n; i++) + if (hb2[i] != h_a2[i] + h_b2[i]) { b++; if (first < 0) first = i; } + printf("STAGE vadd2->c2 readback bad=%d first=%d\n", b, first); } + + scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hb3, d_c2, bytes, cudaMemcpyDeviceToHost)); + { int b = 0, first = -1; for (i = 0; i < n; i++) + if (hb3[i] != (h_a2[i] + h_b2[i]) * 2.0f) { b++; if (first < 0) first = i; } + printf("STAGE scale->c2 readback bad=%d first=%d\n", b, first); } + /* second readback of c1 AFTER everything: does it change? */ + CHECK(cudaMemcpy(hb1, d_c1, bytes, cudaMemcpyDeviceToHost)); + { int b = 0, first = -1; for (i = 0; i < n; i++) + if (hb1[i] != h_a1[i] + h_b1[i]) { b++; if (first < 0) first = i; } + printf("STAGE c1-late readback bad=%d first=%d\n", b, first); } + printf("STATUS OK\n"); + return 0; + } + + if (!strcmp(mode, "C")) { + CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, + cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, + cudaMemcpyHostToDevice, s2s)); + CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, + cudaMemcpyHostToDevice, s2s)); + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaStreamSynchronize(s2s)); + scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n); + CHECK(cudaStreamSynchronize(s2s)); + CHECK(cudaMemcpyAsync(h_c1, d_c1, bytes, + cudaMemcpyDeviceToHost, s1)); + CHECK(cudaMemcpyAsync(h_c2, d_c2, bytes, + cudaMemcpyDeviceToHost, s2s)); + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaStreamSynchronize(s2s)); + goto verify; + } + + /* A and B: original shape */ + CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, cudaMemcpyHostToDevice, s2s)); + CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, cudaMemcpyHostToDevice, s2s)); + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaGetLastError()); + + CHECK(cudaEventRecord(ev_add1, s1)); + CHECK(cudaStreamWaitEvent(s2s, ev_add1, 0)); + scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n); + CHECK(cudaGetLastError()); + + CHECK(cudaEventRecord(ev_scale2, s2s)); + CHECK(cudaStreamWaitEvent(s1, ev_scale2, 0)); + if (!strcmp(mode, "B")) { + CHECK(cudaStreamSynchronize(s2s)); + CHECK(cudaMemcpy(h_c1, d_c1, bytes, cudaMemcpyDeviceToHost)); + CHECK(cudaMemcpy(h_c2, d_c2, bytes, cudaMemcpyDeviceToHost)); + } else { + CHECK(cudaMemcpyAsync(h_c1, d_c1, bytes, + cudaMemcpyDeviceToHost, s1)); + CHECK(cudaMemcpyAsync(h_c2, d_c2, bytes, + cudaMemcpyDeviceToHost, s2s)); + } + CHECK(cudaEventRecord(t1, s1)); + + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaStreamSynchronize(s2s)); + CHECK(cudaEventElapsedTime(&ms, t0, t1)); + +verify: + { + int bad1, bad2; + + bad1 = report("C1", h_c1, h_a1, h_b1, h_a2, h_b2, n, 1.0f); + bad2 = report("C2", h_c2, h_a2, h_b2, h_a2, h_b2, n, s2); + printf("STREAMS n=%d bad=%d\n", n, bad1 + bad2); + printf("TIMING ms=%.3f\n", ms); + + /* GPU-side digest of the device buffers */ + CHECK(cudaMemset(d_out, 0, 2 * sizeof(unsigned long long))); + digest<<<grid, 256>>>(d_a1, d_b1, d_c1, d_a2, d_b2, d_c2, n, + d_out); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, 2 * sizeof(unsigned long long), + cudaMemcpyDeviceToHost)); + printf("GPU-DIGEST c1_ok=%llu c2_ok=%llu\n", h_out[0], + h_out[1]); + printf("STATUS %s\n", (bad1 | bad2) ? "FAIL" : "OK"); + } + return 0; +} diff --git a/tools/cuda-ab/streams_events_probe.cu b/tools/cuda-ab/streams_events_probe.cu new file mode 100644 index 0000000..94aef69 --- /dev/null +++ b/tools/cuda-ab/streams_events_probe.cu @@ -0,0 +1,170 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * streams_events_probe.cu - instrumented victory_streams_events. + * + * Same 3-stream event chain (c0 = a0+b0; c1 = (a1+b1)+c0; c2 = (a2+b2)+c1) + * but with per-buffer diagnostics: device pointers, per-buffer bad counts, + * and the first 8 expected/actual values of every output so a stale-input + * vs stale-readback failure is distinguishable in one run. + * + * Output lines: + * DEVICE ... + * DPTRS da0..dc2 (device), ha0..hc2 (host) + * BUF s=<n> bad=<n> + * FIRST s=<n> exp=<v0..v7> got=<v0..v7> + * STREAMS3 n=<count> bad=<mismatches> + * STATUS OK|FAIL + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#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 NSTREAMS 3 + +__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 vadd_into(float *c, const float *s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] + s[i]; +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 16; + size_t bytes = (size_t)n * sizeof(float); + float *h_a[NSTREAMS], *h_b[NSTREAMS], *h_c[NSTREAMS]; + float *d_a[NSTREAMS], *d_b[NSTREAMS], *d_c[NSTREAMS]; + cudaStream_t st[NSTREAMS]; + cudaEvent_t ev[NSTREAMS], t0, t1; + float ms = 0.0f; + int bad = 0, i, s; + const int grid = (n + 255) / 256; + + 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); + + for (s = 0; s < NSTREAMS; s++) { + h_a[s] = (float *)malloc(bytes); + h_b[s] = (float *)malloc(bytes); + h_c[s] = (float *)malloc(bytes); + if (!h_a[s] || !h_b[s] || !h_c[s]) + return 1; + for (i = 0; i < n; i++) { + h_a[s][i] = (float)((i + s) % 1000) * 0.5f; + h_b[s][i] = (float)((i + 2 * s) % 997) * 0.25f; + } + CHECK(cudaMalloc(&d_a[s], bytes)); + CHECK(cudaMalloc(&d_b[s], bytes)); + CHECK(cudaMalloc(&d_c[s], bytes)); + CHECK(cudaStreamCreate(&st[s])); + CHECK(cudaEventCreate(&ev[s])); + } + CHECK(cudaEventCreate(&t0)); + CHECK(cudaEventCreate(&t1)); + + printf("DPTRS"); + for (s = 0; s < NSTREAMS; s++) + printf(" da%d=%p db%d=%p dc%d=%p", s, (void *)d_a[s], s, + (void *)d_b[s], s, (void *)d_c[s]); + printf("\n"); + printf("HPTRS"); + for (s = 0; s < NSTREAMS; s++) + printf(" ha%d=%p hb%d=%p hc%d=%p", s, (void *)h_a[s], s, + (void *)h_b[s], s, (void *)h_c[s]); + printf("\n"); + + CHECK(cudaEventRecord(t0)); + + for (s = 0; s < NSTREAMS; s++) { + CHECK(cudaMemcpyAsync(d_a[s], h_a[s], bytes, + cudaMemcpyHostToDevice, st[s])); + CHECK(cudaMemcpyAsync(d_b[s], h_b[s], bytes, + cudaMemcpyHostToDevice, st[s])); + vadd<<<grid, 256, 0, st[s]>>>(d_a[s], d_b[s], d_c[s], n); + CHECK(cudaGetLastError()); + CHECK(cudaEventRecord(ev[s], st[s])); + } + + CHECK(cudaStreamWaitEvent(st[1], ev[0], 0)); + vadd_into<<<grid, 256, 0, st[1]>>>(d_c[1], d_c[0], n); + CHECK(cudaGetLastError()); + CHECK(cudaEventRecord(ev[1], st[1])); + + CHECK(cudaStreamWaitEvent(st[2], ev[1], 0)); + vadd_into<<<grid, 256, 0, st[2]>>>(d_c[2], d_c[1], n); + CHECK(cudaGetLastError()); + + for (s = 0; s < NSTREAMS; s++) + CHECK(cudaMemcpyAsync(h_c[s], d_c[s], bytes, + cudaMemcpyDeviceToHost, st[s])); + CHECK(cudaEventRecord(t1, st[2])); + + for (s = 0; s < NSTREAMS; s++) + CHECK(cudaStreamSynchronize(st[s])); + CHECK(cudaEventElapsedTime(&ms, t0, t1)); + + for (s = 0; s < NSTREAMS; s++) { + int b = 0; + float exp0[8], got0[8]; + + for (i = 0; i < n; i++) { + float c0 = h_a[0][i] + h_b[0][i]; + float c1 = h_a[1][i] + h_b[1][i] + c0; + float c2 = h_a[2][i] + h_b[2][i] + c1; + float want = (s == 0) ? c0 : (s == 1) ? c1 : c2; + + if (i < 8) { + exp0[i] = want; + got0[i] = h_c[s][i]; + } + if (h_c[s][i] != want) + b++; + } + bad += b; + printf("BUF s=%d bad=%d\n", s, b); + printf("FIRST s=%d exp=", s); + for (i = 0; i < 8; i++) + printf("%.4f ", exp0[i]); + printf(" got="); + for (i = 0; i < 8; i++) + printf("%.4f ", got0[i]); + printf("\n"); + } + + printf("STREAMS3 n=%d bad=%d\n", n, bad); + printf("TIMING ms=%.3f\n", ms); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + printf("STATUS OK\n"); + return 0; +} diff --git a/tools/cuda-ab/tgemm_probe.cu b/tools/cuda-ab/tgemm_probe.cu new file mode 100644 index 0000000..f0a1084 --- /dev/null +++ b/tools/cuda-ab/tgemm_probe.cu @@ -0,0 +1,216 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * tgemm_probe.cu - minimal isolation for the transformer K/V divergence. + * + * Reproduces exactly the transformer's Q/K/V GEMM shape (X[64x256] x + * W[256x256], tile 16) for three weight matrices uploaded back-to-back + * (wq, wk, wv), then verifies each against a CPU reference and prints + * per-GEMM bad counts plus a bad-position histogram (first bad index per + * 16-row band and per 16-col band) so a copy-side hole is distinguishable + * from a kernel-side error. + * + * Output: + * DEVICE ... + * GEMM name=<q|k|v> bad=<n> first=(row,col) got=.. ref=.. + * BADPOS name=.. rows=<r0,r1,..> cols=<c0,c1,..> + * STATUS OK|FAIL + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cuda_runtime.h> + +#define SEQ 64 +#define DH 256 +#define TILE 16 + +__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; +} + +static float wval(int i, int seed) +{ + return (float)(((i * 13 + seed * 101 + 7) % 5) - 2) * 0.125f; +} + +static float xval(int i, int j) +{ + return (float)(((i * 7 + j * 3 + 1) % 9) - 4) * 0.25f; +} + +int main(void) +{ + cudaDeviceProp prop; + float *h_x, *h_w, *h_c, *h_ref; + float *d_x, *d_w[3], *d_c[3]; + const int seqdh = SEQ * DH, dh2 = DH * DH; + int g, i, j, k; + + setvbuf(stdout, NULL, _IONBF, 0); + if (cudaGetDeviceProperties(&prop, 0) != cudaSuccess) { + printf("STATUS FAIL\n"); + return 1; + } + printf("DEVICE %s cc=%d.%d\n", prop.name, prop.major, prop.minor); + + h_x = (float *)malloc(seqdh * 4); + h_w = (float *)malloc(dh2 * 4); + h_c = (float *)malloc(seqdh * 4); + h_ref = (float *)malloc(seqdh * 4); + + if (cudaMalloc(&d_x, seqdh * 4) != cudaSuccess) + return 1; + for (g = 0; g < 3; g++) { + if (cudaMalloc(&d_w[g], dh2 * 4) != cudaSuccess) + return 1; + if (cudaMalloc(&d_c[g], seqdh * 4) != cudaSuccess) + return 1; + } + + for (i = 0; i < seqdh; i++) + h_x[i] = xval(i / DH, i % DH); + + /* Upload X once (the transformer's shared input). */ + if (cudaMemcpy(d_x, h_x, seqdh * 4, cudaMemcpyHostToDevice) != + cudaSuccess) + return 1; + /* Verify the X upload itself (first row + full buffer). */ + { + float *h_xb = (float *)malloc(seqdh * 4); + int xbad = 0, xr0 = 0; + + if (!h_xb) + return 1; + if (cudaMemcpy(h_xb, d_x, seqdh * 4, + cudaMemcpyDeviceToHost) != cudaSuccess) + return 1; + for (i = 0; i < seqdh; i++) + if (h_xb[i] != h_x[i]) { + xbad++; + if (i < DH) + xr0++; + } + printf("XVERIFY bad=%d row0_bad=%d\n", xbad, xr0); + free(h_xb); + } + + for (g = 0; g < 3; g++) { + int bad = 0, first_r = -1, first_c = -1; + float got = 0, ref = 0; + int rows[64] = {0}, cols[64] = {0}; + int nr = 0, nc = 0; + int pass; + + for (i = 0; i < dh2; i++) + h_w[i] = wval(i, g + 1); + if (cudaMemcpy(d_w[g], h_w, dh2 * 4, + cudaMemcpyHostToDevice) != cudaSuccess) + return 1; + + dim3 blk(TILE, TILE); + dim3 grd(DH / TILE, SEQ / TILE); + + for (pass = 0; pass < 3; pass++) { + float *hc = h_c, *hxb = h_ref; + int xbad = 0; + + gemm_nn<<<grd, blk>>>(d_x, d_w[g], d_c[g], SEQ, DH, DH); + if (cudaDeviceSynchronize() != cudaSuccess) { + printf("SYNC FAIL g=%d\n", g); + return 1; + } + if (cudaMemcpy(hc, d_c[g], seqdh * 4, + cudaMemcpyDeviceToHost) != cudaSuccess) + return 1; + /* re-verify X after the launch */ + if (cudaMemcpy(hxb, d_x, seqdh * 4, + cudaMemcpyDeviceToHost) != cudaSuccess) + return 1; + for (i = 0; i < seqdh; i++) + if (hxb[i] != h_x[i]) { + printf("XDIFF pass=%d i=%d got=%.6f want=%.6f\n", + pass, i, hxb[i], h_x[i]); + xbad++; + } + bad = 0; + nr = 0; + nc = 0; + for (i = 0; i < SEQ; i++) + for (j = 0; j < DH; j++) { + double acc = 0.0; + float r, a, diff; + + for (k = 0; k < DH; k++) + acc += (double)h_x[i * DH + k] * + (double)h_w[k * DH + j]; + r = (float)acc; + a = hc[i * DH + j]; + diff = a - r; + if (diff < 0) + diff = -diff; + if (diff > 1e-4f) { + int rr = i, cc = j; + + if (bad == 0) { + first_r = rr; + first_c = cc; + got = a; + ref = r; + } + bad++; + if (nr < 64 && (bad == 1 || + rr != rows[nr - 1])) + rows[nr++] = rr; + if (nc < 64 && (bad == 1 || + cc != cols[nc - 1])) + cols[nc++] = cc; + } + } + printf("PASS name=%c pass=%d bad=%d xbad=%d first=(%d,%d) got=%.6f ref=%.6f\n", + "qkv"[g], pass, bad, xbad, first_r, first_c, + got, ref); + if (bad) { + printf("ROW0 name=%c got=", "qkv"[g]); + for (i = 0; i < 8; i++) + printf("%.6f ", hc[i]); + printf(" ref="); + for (i = 0; i < 8; i++) { + double acc = 0.0; + + for (k = 0; k < DH; k++) + acc += (double)h_x[k] * + (double)h_w[k * DH + i]; + printf("%.6f ", (float)acc); + } + printf("\n"); + } + } + (void)got; + (void)ref; + (void)first_r; + (void)first_c; + (void)nr; + (void)nc; + } + + printf("STATUS OK\n"); + return 0; +} diff --git a/tools/cuda-ab/transformer_async.cu b/tools/cuda-ab/transformer_async.cu new file mode 100644 index 0000000..213363a --- /dev/null +++ b/tools/cuda-ab/transformer_async.cu @@ -0,0 +1,522 @@ +// 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 (async chain, stage readbacks) ---- */ + 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)); + { + float *hs = (float *)malloc(b_seqdh); + float *hs2 = (float *)malloc(b_seq2); + double *rw = (double *)malloc((size_t)dh2 * sizeof(double)); + double *rq = (double *)malloc((size_t)seqdh * sizeof(double)); + double *rk = (double *)malloc((size_t)seqdh * sizeof(double)); + double *rv = (double *)malloc((size_t)seqdh * sizeof(double)); + double *rs = (double *)malloc((size_t)seq2 * sizeof(double)); + double *ra = (double *)malloc((size_t)seq2 * sizeof(double)); + double *rav = (double *)malloc((size_t)seqdh * sizeof(double)); + double *ry = (double *)malloc((size_t)seqdh * sizeof(double)); + int st; + + for (i = 0; i < dh2; i++) + rw[i] = (double)wval(i, 1); + ref_gemm(r_x, rw, rq, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + rw[i] = (double)wval(i, 2); + ref_gemm(r_x, rw, rk, SEQ, DH, DH); + for (i = 0; i < dh2; i++) + rw[i] = (double)wval(i, 3); + ref_gemm(r_x, rw, rv, SEQ, DH, DH); + ref_gemm_nt(rq, rk, rs, SEQ, SEQ, DH); + ref_softmax(rs, ra, SEQ, SEQ, 1.0 / 16.0); + ref_gemm(ra, rv, rav, SEQ, DH, SEQ); + for (i = 0; i < dh2; i++) + rw[i] = (double)wval(i, 4); + ref_gemm(rav, rw, ry, SEQ, DH, DH); + + CHECK(cudaMemcpy(hs, d_q, b_seqdh, cudaMemcpyDeviceToHost)); + st = count_bad(hs, rq, seqdh); + printf("ASYNC Q bad=%d\n", st); + CHECK(cudaMemcpy(hs, d_k, b_seqdh, cudaMemcpyDeviceToHost)); + st = count_bad(hs, rk, seqdh); + printf("ASYNC K bad=%d\n", st); + CHECK(cudaMemcpy(hs, d_v, b_seqdh, cudaMemcpyDeviceToHost)); + st = count_bad(hs, rv, seqdh); + printf("ASYNC V bad=%d\n", st); + CHECK(cudaMemcpy(hs2, d_s, b_seq2, cudaMemcpyDeviceToHost)); + st = count_bad(hs2, rs, seq2); + printf("ASYNC S bad=%d\n", st); + CHECK(cudaMemcpy(hs2, d_a, b_seq2, cudaMemcpyDeviceToHost)); + st = count_bad(hs2, ra, seq2); + printf("ASYNC A bad=%d\n", st); + CHECK(cudaMemcpy(hs, d_av, b_seqdh, cudaMemcpyDeviceToHost)); + st = count_bad(hs, rav, seqdh); + printf("ASYNC AV bad=%d\n", st); + st = count_bad(h_y, ry, seqdh); + printf("ASYNC Y bad=%d\n", st); + } + + { + 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/transformer_stages.cu b/tools/cuda-ab/transformer_stages.cu new file mode 100644 index 0000000..d7f67c0 --- /dev/null +++ b/tools/cuda-ab/transformer_stages.cu @@ -0,0 +1,553 @@ +// 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; +} + +static void dump_bad(const char *tag, const float *gpu, const double *ref, + int n) +{ + int i, shown = 0; + + for (i = 0; i < n && shown < 8; i++) { + double g = (double)gpu[i]; + double r = ref[i]; + + if (fabs(g - r) > (double)TOL * (fabs(r) + 1e-6)) { + printf("BAD %s i=%d gpu=%.6f ref=%.6f (tile %d,%d row %d col %d)\n", + tag, i, g, r, i / 16, (i % 256) / 16, + i / 256, i % 256); + shown++; + } + } +} + +#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, stage-by-stage verify ---- */ + { + float *h_stage = (float *)malloc(b_seqdh); + float *h_s2 = (float *)malloc(b_seq2); + double *r_stage = (double *)malloc((size_t)seqdh * sizeof(double)); + double *r_s2 = (double *)malloc((size_t)seq2 * sizeof(double)); + double *r_w2 = (double *)malloc((size_t)dh2 * sizeof(double)); + int st; + + if (!h_stage || !h_s2 || !r_stage || !r_s2 || !r_w2) + return 1; + + launch_gemm_nn(d_x, d_wq, d_q, SEQ, DH, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_stage, d_q, b_seqdh, cudaMemcpyDeviceToHost)); + for (i = 0; i < dh2; i++) + r_w2[i] = (double)wval(i, 1); + ref_gemm(r_x, r_w2, r_stage, SEQ, DH, DH); + st = count_bad(h_stage, r_stage, seqdh); + printf("STAGE Q bad=%d\n", st); + dump_bad("Q", h_stage, r_stage, seqdh); + + launch_gemm_nn(d_x, d_wk, d_k, SEQ, DH, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_stage, d_k, b_seqdh, cudaMemcpyDeviceToHost)); + for (i = 0; i < dh2; i++) + r_w2[i] = (double)wval(i, 2); + ref_gemm(r_x, r_w2, r_stage, SEQ, DH, DH); + st = count_bad(h_stage, r_stage, seqdh); + printf("STAGE K bad=%d\n", st); + dump_bad("K", h_stage, r_stage, seqdh); + + launch_gemm_nn(d_x, d_wv, d_v, SEQ, DH, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_stage, d_v, b_seqdh, cudaMemcpyDeviceToHost)); + for (i = 0; i < dh2; i++) + r_w2[i] = (double)wval(i, 3); + ref_gemm(r_x, r_w2, r_stage, SEQ, DH, DH); + st = count_bad(h_stage, r_stage, seqdh); + printf("STAGE V bad=%d\n", st); + dump_bad("V", h_stage, r_stage, seqdh); + + launch_gemm_nt(d_q, d_k, d_s, SEQ, SEQ, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_s2, d_s, b_seq2, cudaMemcpyDeviceToHost)); + ref_gemm_nt(r_q, r_k, r_s, SEQ, SEQ, DH); + st = count_bad(h_s2, r_s, seq2); + printf("STAGE S bad=%d\n", st); + dump_bad("S", h_s2, r_s, seq2); + + launch_softmax(d_s, d_a, SEQ, SEQ, 1.0f / 16.0f); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_s2, d_a, b_seq2, cudaMemcpyDeviceToHost)); + ref_softmax(r_s, r_a, SEQ, SEQ, 1.0 / 16.0); + st = count_bad(h_s2, r_a, seq2); + printf("STAGE A bad=%d\n", st); + dump_bad("A", h_s2, r_a, seq2); + + launch_gemm_nn(d_a, d_v, d_av, SEQ, DH, SEQ); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_stage, d_av, b_seqdh, cudaMemcpyDeviceToHost)); + ref_gemm(r_a, r_v, r_av, SEQ, DH, SEQ); + st = count_bad(h_stage, r_av, seqdh); + printf("STAGE AV bad=%d\n", st); + + launch_gemm_nn(d_av, d_wo, d_y, SEQ, DH, DH); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_y, d_y, b_seqdh, cudaMemcpyDeviceToHost)); + for (i = 0; i < dh2; i++) + r_w2[i] = (double)wval(i, 4); + ref_gemm(r_av, r_w2, r_y, SEQ, DH, DH); + st = count_bad(h_y, r_y, seqdh); + printf("STAGE Y bad=%d\n", st); + } + + { + 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/guest-diag/pin_gr_stream.py b/tools/guest-diag/pin_gr_stream.py index 6ba51ee..6eb5be7 100644 --- a/tools/guest-diag/pin_gr_stream.py +++ b/tools/guest-diag/pin_gr_stream.py @@ -211,7 +211,7 @@ def _demand_loop(): except Exception: continue now = time.time() - for m in re.finditer(r"(?:stalled VA #\d+|pending release starved va)=(0x[0-9a-f]+)", + for m in re.finditer(r"(?:stalled VA #\d+[:=]\s*|pending release starved va=)(0x[0-9a-f]+)", (r.stdout or "")[-200000:]): va = int(m.group(1), 16) track.setdefault(va, [now, 0.0]) diff --git a/tools/lg/frame_sequence.sh b/tools/lg/frame_sequence.sh new file mode 100755 index 0000000..958a34a --- /dev/null +++ b/tools/lg/frame_sequence.sh @@ -0,0 +1,33 @@ +#!/bin/bash +# SPDX-License-Identifier: BSD-3-Clause +# frame_sequence.sh - capture an animated GPU-rendered frame sequence end to +# end: run the guest's CUDA frame producer into the kvmfr ivshmem BAR and +# save every published frame on the host as PPM/PNG (the proven +# CUDA -> ivshmem -> host path; see docs/05-operations/06). +# +# Usage: tools/lg/frame_sequence.sh [frames] [width] [height] [fps] [outdir] +set -u +HERE="$(cd "$(dirname "$0")" && pwd)" +REPO="$(cd "$HERE/../.." && pwd)" +FRAMES="${1:-40}" +W="${2:-800}" +H="${3:-600}" +FPS="${4:-10}" +OUT="${5:-/tmp/opencode/frames-seq}" + +mkdir -p "$OUT" +rm -f "$OUT"/frame-*.ppm "$OUT"/frame-*.png + +# Start the host reader first so no published frame is missed. +python3 "$REPO/tools/lg/frame_reader.py" --frames "$FRAMES" --out "$OUT" \ + --timeout 240 > "$OUT/reader.log" 2>&1 & +READER=$! + +sleep 1 +"$REPO/scripts/ssh_guest.sh" \ + "sudo /tmp/frame_producer $FRAMES $W $H $FPS" 2>&1 | grep -v "Permanently added" \ + | tail -5 + +wait "$READER" +echo "reader rc=$?" +ls -1 "$OUT"/frame-*.ppm 2>/dev/null | wc -l |