# gemmcheck3: self-checking rocBLAS GEMM test (RX 7900 XTX corruption)

> Source: <https://gist.github.com/roverdrubber/7cd347a9340abdad2f707f2252aaaee0>
> Published: 2026-09-18 18:12:03+00:00

|  | // gemmcheck3: a self-checking rocBLAS GEMM stress test for AMD GPUs (RDNA3 / gfx11 tested). | 
|  | // | 
|  | // Five rocblas_gemm_ex calls in LLM layer shapes (fp16 in / fp32 out, and int8 in / int32 out), random inputs, | 
|  | // atomics disabled, launched back to back on one stream. A reference is computed once at startup and a | 
|  | // determinism check must pass. After every GEMM a compare kernel checks the output bit for bit; on a mismatch it | 
|  | // records the launch number, word index, the value read, the reference, and both re-read through an atomic | 
|  | // (served by L2). Every few seconds the host prints a status line and describes any bad launch: GEMM | 
|  | // coordinates (rows, columns), checker coordinates, magnitudes, and whether the re-read agrees. | 
|  | // | 
|  | // On a healthy GPU with a deterministic library, a mismatch should never happen. | 
|  | // | 
|  | //   hipcc -O2 --offload-arch=gfx1100 gemmcheck3.hip -o gemmcheck3 -L$ROCM_PATH/lib -lrocblas -Wl,-rpath,$ROCM_PATH/lib | 
|  | //   GEMMCHECK_STOP=1 ./gemmcheck3 600 5                run up to 600 s, report every 5 s, stop at the first mismatch | 
|  | //   GEMMCHECK_INJECT=7 ./gemmcheck3 10 5               self-test: plants a known 64x16 fault at launch 7 (exit 1) | 
|  | //   GEMMCHECK_GRID=512,128 ./gemmcheck3 ...            checker geometry (blocks,threads), default 1024,256 | 
|  | //   GEMMCHECK_TRACE=1 AMD_LOG_LEVEL=3 ./gemmcheck3 1 1  one launch per case between markers, then exit | 
|  | // Exit codes: 0 clean, 1 mismatches found, 2 HIP/rocBLAS error, 3 determinism check failed. | 
|  | #include <hip/hip_runtime.h> | 
|  | #include <hip/hip_fp16.h> | 
|  | #include <rocblas/rocblas.h> | 
|  | #include <algorithm> | 
|  | #include <chrono> | 
|  | #include <cmath> | 
|  | #include <cstdint> | 
|  | #include <cstdio> | 
|  | #include <cstdlib> | 
|  | #include <cstring> | 
|  | #include <map> | 
|  | #include <random> | 
|  | #include <set> | 
|  | #include <utility> | 
|  | #include <vector> | 
|  |  | 
|  | #define CK(x) do { hipError_t e_ = (x); if (e_ != hipSuccess) { fprintf(stderr, "HIP %s at %s:%d\n", hipGetErrorString(e_), __FILE__, __LINE__); exit(2); } } while (0) | 
|  | #define RB(x) do { rocblas_status s_ = (x); if (s_ != rocblas_status_success) { fprintf(stderr, "rocBLAS %s at %s:%d\n", rocblas_status_to_string(s_), __FILE__, __LINE__); exit(2); } } while (0) | 
|  |  | 
|  | struct Rec { | 
|  | unsigned long long launch; | 
|  | unsigned int ci, idx, got, want, got2, want2; // got2/want2: the same words read again through an atomic | 
|  | }; | 
|  |  | 
|  | // counts 32-bit words that differ and records each one (up to cap); first keeps one differing index. | 
|  | // got2/want2 are atomic read-modify-writes of 0: served by L2 past the per-CU caches, but an RMW, not a passive probe | 
|  | __global__ void compare_words(uint32_t * a, uint32_t * ref, size_t n, unsigned long long * count, unsigned long long * first, | 
|  | Rec * recs, unsigned long long * nrec, unsigned long long cap, unsigned long long launch, int ci) { | 
|  | size_t i = (size_t) blockIdx.x * blockDim.x + threadIdx.x; | 
|  | size_t stride = (size_t) gridDim.x * blockDim.x; | 
|  | unsigned long long local = 0; | 
|  | for (; i < n; i += stride) { | 
|  | const uint32_t got = a[i], want = ref[i]; | 
|  | if (got != want) { | 
|  | // re-read first, before any counter traffic can delay it | 
|  | const uint32_t got2 = atomicAdd(&a[i], 0u), want2 = atomicAdd(&ref[i], 0u); | 
|  | if (local == 0) { atomicCAS(first, ~0ull, (unsigned long long) i); } | 
|  | local++; | 
|  | const unsigned long long slot = atomicAdd(nrec, 1ull); | 
|  | if (slot < cap) { | 
|  | Rec r; | 
|  | r.launch = launch; r.ci = (unsigned int) ci; r.idx = (unsigned int) i; r.got = got; r.want = want; | 
|  | r.got2 = got2; r.want2 = want2; | 
|  | recs[slot] = r; | 
|  | } | 
|  | } | 
|  | } | 
|  | if (local) { atomicAdd(count, local); } | 
|  | } | 
|  |  | 
|  | // test hook (GEMMCHECK_INJECT=launch): flips bit 0 of rows 100-163, columns 5-20 of that launch's output | 
|  | __global__ void inject(uint32_t * c, int m) { | 
|  | const int row = 100 + threadIdx.x, col = 5 + blockIdx.x; | 
|  | c[(size_t) col * m + row] ^= 1u; | 
|  | } | 
|  |  | 
|  | struct Case { | 
|  | const char * name; | 
|  | int m, n, k; | 
|  | bool int8; | 
|  | void * a = nullptr; void * b = nullptr; void * c = nullptr; void * ref = nullptr; | 
|  | size_t cwords = 0; | 
|  | unsigned long long * d_count = nullptr; unsigned long long * d_first = nullptr; | 
|  | unsigned long long launches = 0; | 
|  | }; | 
|  |  | 
|  | static void run_gemm(rocblas_handle h, Case & cs, void * out) { | 
|  | if (cs.int8) { | 
|  | int32_t alpha = 1, beta = 0; | 
|  | RB(rocblas_gemm_ex(h, rocblas_operation_transpose, rocblas_operation_none, cs.m, cs.n, cs.k, &alpha, | 
|  | cs.a, rocblas_datatype_i8_r, cs.k, cs.b, rocblas_datatype_i8_r, cs.k, &beta, | 
|  | out, rocblas_datatype_i32_r, cs.m, out, rocblas_datatype_i32_r, cs.m, | 
|  | rocblas_datatype_i32_r, rocblas_gemm_algo_standard, 0, rocblas_gemm_flags_none)); | 
|  | } else { | 
|  | float alpha = 1.0f, beta = 0.0f; | 
|  | RB(rocblas_gemm_ex(h, rocblas_operation_transpose, rocblas_operation_none, cs.m, cs.n, cs.k, &alpha, | 
|  | cs.a, rocblas_datatype_f16_r, cs.k, cs.b, rocblas_datatype_f16_r, cs.k, &beta, | 
|  | out, rocblas_datatype_f32_r, cs.m, out, rocblas_datatype_f32_r, cs.m, | 
|  | rocblas_datatype_f32_r, rocblas_gemm_algo_standard, 0, rocblas_gemm_flags_none)); | 
|  | } | 
|  | } | 
|  |  | 
|  | static float as_float(uint32_t u) { float f; memcpy(&f, &u, 4); return f; } | 
|  |  | 
|  | // one bad launch of one case: where the wrong words sit and what they look like | 
|  | static void describe(const std::vector<Case> & cases, unsigned long long launch, unsigned int ci, std::vector<Rec> & rs, | 
|  | int blocks, int threads) { | 
|  | const Case & cs = cases[ci]; | 
|  | const size_t T = (size_t) blocks * threads; | 
|  | std::sort(rs.begin(), rs.end(), [](const Rec & x, const Rec & y) { return x.idx < y.idx; }); | 
|  | std::set<int> cols, blks, iters; | 
|  | int rmin = cs.m, rmax = -1, cmin = cs.n, cmax = -1; | 
|  | size_t reread_differs = 0, reread_eq_want = 0, reread_eq_want2 = 0, ref_reread_differs = 0, nonfinite_got = 0, nonfinite_want = 0, ulp_small = 0; | 
|  | size_t xor1 = 0; | 
|  | std::set<unsigned int> uniq; | 
|  | double maxdiff = 0, maxref = 0; | 
|  | for (const Rec & r : rs) { | 
|  | const int row = (int) (r.idx % cs.m), col = (int) (r.idx / cs.m); | 
|  | rmin = std::min(rmin, row); rmax = std::max(rmax, row); cmin = std::min(cmin, col); cmax = std::max(cmax, col); | 
|  | cols.insert(col); | 
|  | const size_t tid = r.idx % T; | 
|  | blks.insert((int) (tid / threads)); iters.insert((int) (r.idx / T)); | 
|  | uniq.insert(r.idx); | 
|  | if ((r.got ^ r.want) == 1u) { xor1++; } | 
|  | if (r.got2 != r.got) { reread_differs++; } | 
|  | if (r.got2 == r.want) { reread_eq_want++; } | 
|  | if (r.got2 == r.want2) { reread_eq_want2++; } | 
|  | if (r.want2 != r.want) { ref_reread_differs++; } | 
|  | if (cs.int8) { | 
|  | maxdiff = std::max(maxdiff, std::fabs((double) (int32_t) r.got - (double) (int32_t) r.want)); | 
|  | maxref = std::max(maxref, std::fabs((double) (int32_t) r.want)); | 
|  | } else { | 
|  | const float g = as_float(r.got), w = as_float(r.want); | 
|  | const bool gf = std::isfinite(g), wf = std::isfinite(w); | 
|  | if (!gf) { nonfinite_got++; } | 
|  | if (!wf) { nonfinite_want++; } | 
|  | if (gf && wf) { | 
|  | maxdiff = std::max(maxdiff, std::fabs((double) g - (double) w)); | 
|  | if (((r.got ^ r.want) & 0x80000000u) == 0 && (r.got > r.want ? r.got - r.want : r.want - r.got) <= 16) { ulp_small++; } | 
|  | } | 
|  | if (wf) { maxref = std::max(maxref, std::fabs((double) w)); } | 
|  | } | 
|  | } | 
|  | printf("\n  EVENT launch %llu, %s: %zu wrong words recorded\n", launch, cs.name, rs.size()); | 
|  | printf("    GEMM coordinates: rows %d-%d, columns %d-%d (%zu distinct columns)\n", rmin, rmax, cmin, cmax, cols.size()); | 
|  | printf("    checker coordinates (%d x %d): %zu distinct blocks (%d-%d), %zu distinct iterations (%d-%d)\n", blocks, threads, | 
|  | blks.size(), *blks.begin(), *blks.rbegin(), iters.size(), *iters.begin(), *iters.rbegin()); | 
|  | printf("    second read (atomic, L2): differs from first read %zu, equals want %zu, equals want2 %zu; want2 differs from want %zu\n", | 
|  | reread_differs, reread_eq_want, reread_eq_want2, ref_reread_differs); | 
|  | printf("    checks: unique indices %zu, got^want==1 %zu, got2==got %zu, want2==want %zu, records %zu\n", | 
|  | uniq.size(), xor1, rs.size() - reread_differs, rs.size() - ref_reread_differs, rs.size()); | 
|  | if (cs.int8) { | 
|  | printf("    largest \|got - want\| %.0f (largest \|want\| %.0f)\n", maxdiff, maxref); | 
|  | } else { | 
|  | printf("    NaN/Inf in got %zu, in want %zu; within 16 ulp (both finite) %zu; largest finite \|got - want\| %.4g (largest \|want\| %.4g)\n", | 
|  | nonfinite_got, nonfinite_want, ulp_small, maxdiff, maxref); | 
|  | } | 
|  | for (size_t j = 0; j < rs.size() && j < 6; j++) { | 
|  | const Rec & r = rs[j]; | 
|  | printf("    word %u (row %u col %u, checker block %zu iter %zu): got 0x%08x want 0x%08x got2 0x%08x want2 0x%08x\n", r.idx, | 
|  | r.idx % cs.m, r.idx / cs.m, (r.idx % T) / threads, (size_t) (r.idx / T), r.got, r.want, r.got2, r.want2); | 
|  | } | 
|  | fflush(stdout); | 
|  | } | 
|  |  | 
|  | int main(int argc, char ** argv) { | 
|  | const double seconds = argc > 1 ? atof(argv[1]) : 600; | 
|  | const double every   = argc > 2 ? atof(argv[2]) : 30; | 
|  | int blocks = 1024, threads = 256; | 
|  | if (const char * g = getenv("GEMMCHECK_GRID")) { | 
|  | if (sscanf(g, "%d,%d", &blocks, &threads) != 2 \|\| blocks < 1 \|\| blocks > 65535 \|\| threads < 32 \|\| threads > 1024 \|\| threads % 32) { | 
|  | fprintf(stderr, "bad GEMMCHECK_GRID '%s' (want blocks,threads with 1..65535 blocks and 32..1024 threads, a multiple of 32)\n", g); | 
|  | return 2; | 
|  | } | 
|  | } | 
|  | const unsigned long long inject_at = getenv("GEMMCHECK_INJECT") ? strtoull(getenv("GEMMCHECK_INJECT"), nullptr, 10) : 0; | 
|  | const bool trace = getenv("GEMMCHECK_TRACE") != nullptr; | 
|  | const bool stop_on_mismatch = getenv("GEMMCHECK_STOP") != nullptr; | 
|  | const unsigned long long cap = 1ull << 16; | 
|  |  | 
|  | // Qwen3.8-27B at ubatch 512: ffn gate/up, ffn down, fused QKV; hidden 5120, ffn 17408, qkv 12288 | 
|  | std::vector<Case> cases = { | 
|  | {"f16 ffn_up   17408x512x5120",  17408, 512, 5120,  false}, | 
|  | {"f16 ffn_down  5120x512x17408",  5120, 512, 17408, false}, | 
|  | {"i8  ffn_up   17408x512x5120",  17408, 512, 5120,  true}, | 
|  | {"i8  ffn_down  5120x512x17408",  5120, 512, 17408, true}, | 
|  | {"i8  qkv      12288x512x5120",  12288, 512, 5120,  true}, | 
|  | }; | 
|  |  | 
|  | rocblas_handle h; RB(rocblas_create_handle(&h)); | 
|  | RB(rocblas_set_atomics_mode(h, rocblas_atomics_not_allowed)); | 
|  | hipStream_t st; CK(hipStreamCreate(&st)); RB(rocblas_set_stream(h, st)); | 
|  |  | 
|  | // allocation order identical to gemmcheck.hip; the record buffers come after all of it | 
|  | std::mt19937 rng(12345); | 
|  | for (auto & cs : cases) { | 
|  | size_t na = (size_t) cs.m * cs.k, nb = (size_t) cs.k * cs.n, nc = (size_t) cs.m * cs.n; | 
|  | size_t es = cs.int8 ? 1 : 2; | 
|  | std::vector<uint8_t> ha(na * es), hb(nb * es); | 
|  | if (cs.int8) { | 
|  | std::uniform_int_distribution<int> d(-127, 127); | 
|  | for (auto & v : ha) v = (uint8_t) (int8_t) d(rng); | 
|  | for (auto & v : hb) v = (uint8_t) (int8_t) d(rng); | 
|  | } else { | 
|  | std::uniform_real_distribution<float> d(-1.0f, 1.0f); | 
|  | __half * pa = (__half *) ha.data(); __half * pb = (__half *) hb.data(); | 
|  | for (size_t i = 0; i < na; i++) pa[i] = __float2half(d(rng)); | 
|  | for (size_t i = 0; i < nb; i++) pb[i] = __float2half(d(rng)); | 
|  | } | 
|  | CK(hipMalloc(&cs.a, na * es)); CK(hipMalloc(&cs.b, nb * es)); | 
|  | CK(hipMalloc(&cs.c, nc * 4)); CK(hipMalloc(&cs.ref, nc * 4)); | 
|  | CK(hipMemcpy(cs.a, ha.data(), na * es, hipMemcpyHostToDevice)); | 
|  | CK(hipMemcpy(cs.b, hb.data(), nb * es, hipMemcpyHostToDevice)); | 
|  | cs.cwords = nc; | 
|  | CK(hipMalloc(&cs.d_count, sizeof(unsigned long long))); CK(hipMemset(cs.d_count, 0, sizeof(unsigned long long))); | 
|  | CK(hipMalloc(&cs.d_first, sizeof(unsigned long long))); CK(hipMemset(cs.d_first, 0xff, sizeof(unsigned long long))); | 
|  | } | 
|  | // reference results, then a determinism check: recompute each and require bit-identical output. The record | 
|  | // buffers are allocated only after rocBLAS's first GEMMs, to reduce perturbation of its lazily allocated workspace (identical | 
|  | // addresses are not guaranteed) | 
|  | for (auto & cs : cases) { run_gemm(h, cs, cs.ref); } | 
|  | CK(hipStreamSynchronize(st)); | 
|  | Rec * d_recs; unsigned long long * d_nrec; | 
|  | CK(hipMalloc(&d_recs, sizeof(Rec) * cap)); CK(hipMalloc(&d_nrec, sizeof(unsigned long long))); CK(hipMemset(d_nrec, 0, sizeof(unsigned long long))); | 
|  | for (size_t ci = 0; ci < cases.size(); ci++) { | 
|  | auto & cs = cases[ci]; | 
|  | run_gemm(h, cs, cs.c); | 
|  | compare_words<<<blocks, threads, 0, st>>>((uint32_t *) cs.c, (uint32_t *) cs.ref, cs.cwords, cs.d_count, cs.d_first, d_recs, d_nrec, cap, 0, (int) ci); | 
|  | } | 
|  | CK(hipStreamSynchronize(st)); | 
|  | for (auto & cs : cases) { | 
|  | unsigned long long c = 0; CK(hipMemcpy(&c, cs.d_count, sizeof c, hipMemcpyDeviceToHost)); | 
|  | if (c) { printf("NONDETERMINISTIC: %s differs from its own reference (%llu words) before the stress starts; results would be meaningless\n", cs.name, c); return 3; } | 
|  | } | 
|  | CK(hipMemset(d_nrec, 0, sizeof(unsigned long long))); | 
|  | printf("determinism check passed for %zu cases; checker %d x %d; stressing for %.0f s\n", cases.size(), blocks, threads, seconds); fflush(stdout); | 
|  |  | 
|  | if (trace) { | 
|  | for (size_t ci = 0; ci < cases.size(); ci++) { | 
|  | fprintf(stderr, "MARK begin case %zu %s\n", ci, cases[ci].name); | 
|  | run_gemm(h, cases[ci], cases[ci].c); | 
|  | CK(hipStreamSynchronize(st)); | 
|  | fprintf(stderr, "MARK end case %zu\n", ci); | 
|  | } | 
|  | return 0; | 
|  | } | 
|  |  | 
|  | const auto t0 = std::chrono::steady_clock::now(); | 
|  | auto last = t0; | 
|  | unsigned long long iters = 0, total_bad = 0, launch = 0; | 
|  | unsigned long long seen = 0; | 
|  | std::vector<unsigned long long> prev(cases.size(), 0); | 
|  | while (true) { | 
|  | for (size_t ci = 0; ci < cases.size(); ci++) { | 
|  | auto & cs = cases[ci]; | 
|  | run_gemm(h, cs, cs.c); | 
|  | if (++launch == inject_at) { inject<<<16, 64, 0, st>>>((uint32_t *) cs.c, cs.m); } | 
|  | compare_words<<<blocks, threads, 0, st>>>((uint32_t *) cs.c, (uint32_t *) cs.ref, cs.cwords, cs.d_count, cs.d_first, d_recs, d_nrec, cap, launch, (int) ci); | 
|  | cs.launches++; | 
|  | } | 
|  | iters++; | 
|  | if ((iters & 63) == 0) { | 
|  | CK(hipStreamSynchronize(st)); | 
|  | const auto now = std::chrono::steady_clock::now(); | 
|  | const double el = std::chrono::duration<double>(now - t0).count(); | 
|  | const bool done = el >= seconds; | 
|  | if (std::chrono::duration<double>(now - last).count() >= every \|\| done) { | 
|  | last = now; | 
|  | total_bad = 0; | 
|  | printf("t=%6.0fs iters=%llu", el, iters); | 
|  | for (size_t i = 0; i < cases.size(); i++) { | 
|  | unsigned long long c = 0; | 
|  | CK(hipMemcpy(&c, cases[i].d_count, sizeof c, hipMemcpyDeviceToHost)); | 
|  | total_bad += c; | 
|  | if (c != prev[i]) { printf("\n  MISMATCH %s: +%llu words (%llu so far)", cases[i].name, c - prev[i], c); prev[i] = c; } | 
|  | } | 
|  | printf(" \| total mismatched words %llu\n", total_bad); | 
|  | unsigned long long nr = 0; CK(hipMemcpy(&nr, d_nrec, sizeof nr, hipMemcpyDeviceToHost)); | 
|  | const unsigned long long upto = std::min(nr, cap); | 
|  | if (upto > seen) { | 
|  | std::vector<Rec> rs(upto - seen); | 
|  | CK(hipMemcpy(rs.data(), d_recs + seen, sizeof(Rec) * rs.size(), hipMemcpyDeviceToHost)); | 
|  | std::map<std::pair<unsigned long long, unsigned int>, std::vector<Rec>> groups; | 
|  | size_t invalid = 0; | 
|  | for (const Rec & r : rs) { | 
|  | if (r.ci >= cases.size() \|\| r.idx >= cases[r.ci].cwords) { invalid++; continue; } | 
|  | groups[{r.launch, r.ci}].push_back(r); | 
|  | } | 
|  | if (invalid) { printf("  %zu records with an invalid case or index skipped\n", invalid); } | 
|  | for (auto & g : groups) { describe(cases, g.first.first, g.first.second, g.second, blocks, threads); } | 
|  | seen = upto; | 
|  | } | 
|  | if (nr > cap) { printf("  record buffer full: %llu mismatches seen, %llu recorded; event geometry above is TRUNCATED\n", nr, cap); } | 
|  | if (stop_on_mismatch && total_bad) { printf("STOPPED at the first mismatch, t=%.0fs\n", el); fflush(stdout); break; } | 
|  | fflush(stdout); | 
|  | } | 
|  | if (done) { break; } | 
|  | } | 
|  | } | 
|  | unsigned long long launches = 0; for (auto & cs : cases) launches += cs.launches; | 
|  | printf("RESULT %s: %llu GEMM launches, %llu mismatched words\n", total_bad ? "MISMATCHES" : "CLEAN", launches, total_bad); | 
|  | return total_bad ? 1 : 0; | 
|  | } |