diff options
Diffstat (limited to 'tools/cuda-ab/transformer_mt.cu')
| -rw-r--r-- | tools/cuda-ab/transformer_mt.cu | 8 |
1 files changed, 8 insertions, 0 deletions
diff --git a/tools/cuda-ab/transformer_mt.cu b/tools/cuda-ab/transformer_mt.cu index ffd23aa..5cfe55e 100644 --- a/tools/cuda-ab/transformer_mt.cu +++ b/tools/cuda-ab/transformer_mt.cu @@ -142,6 +142,14 @@ static void launch_chain(cudaStream_t st, const float *x, struct attn_bufs b, dim3 g1(DH / TILE, SEQ / TILE); dim3 g2(SEQ / TILE, SEQ / TILE); dim3 tb(TILE, TILE); + static int chain_n; + + /* r226 GRID-WIT: the exact geometry every chain launch uses, so the + * host-side QMD geometry capture (GRID-WIT in vgsp) can be compared + * against what the client actually launched. Chains 0..2 are the + * multi-stream class, 3..7 the 5-launch loop. */ + printf("GEO chain=%d gemm_nn grid=(%u,%u,1) block=(%u,%u,1) gemm_nt grid=(%u,%u,1) softmax grid=(64,1,1) block=(64,1,1)\n", + chain_n++, g1.x, g1.y, tb.x, tb.y, g2.x, g2.y); gemm_nn<<<g1, tb, 0, st>>>(x, wq, b.q, SEQ, DH, DH); gemm_nn<<<g1, tb, 0, st>>>(x, wk, b.k, SEQ, DH, DH); |