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

A developer released gemmcheck3, a self-checking rocBLAS GEMM stress test for AMD GPUs targeting RDNA3/gfx11 hardware such as the RX 7900 XTX. The tool runs five rocblas_gemm_ex calls in LLM layer shapes (fp16-in/fp32-out and int8-in/int32-out) back to back on a single stream, comparing each output bit-for-bit against a startup reference and recording launch number, word index, and values on any mismatch. It supports fault injection for self-testing and exits with distinct codes for clean runs, mismatches, HIP/rocBLAS errors, and failed determinism checks.

| | // 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; | | | } |