From 4ccaea9f0849fd21663b36e196de953b604ddce2 Mon Sep 17 00:00:00 2001 From: Basil Milanich Date: Mon, 6 Apr 2026 11:54:08 -0500 Subject: [PATCH] Expand repro: multiple graphs, multiple kernels, interleaved launches Tests: - Two distinct graphs (different graphIds) on the same stream - Graph A: kernel + memcpy + kernel (3 nodes) - Graph B: scale + add + scale (3 nodes) - 5 interleaved launches of each, stressing the graphId cache - Expected 30 graph GPU zones total Co-Authored-By: Claude Sonnet 4.6 --- examples/CUDAGraphRepro/repro.cu | 112 ++++++++++++++++++++----------- 1 file changed, 72 insertions(+), 40 deletions(-) diff --git a/examples/CUDAGraphRepro/repro.cu b/examples/CUDAGraphRepro/repro.cu index 71d5c14d..33d95e1c 100644 --- a/examples/CUDAGraphRepro/repro.cu +++ b/examples/CUDAGraphRepro/repro.cu @@ -1,16 +1,24 @@ // Tracy CUDA Graph GPU Zone Repro // -// Demonstrates that Tracy correctly shows GPU zones for kernels launched -// via CUDA Graphs (cuGraphLaunch). Uses TracyCUDA to create a GPU context -// and verifies that GPU zones appear with proper CPU-to-GPU correlation. +// Tests GPU zone correlation for CUDA Graph launches covering: +// - Multiple distinct graphs (different graphIds) +// - Multiple kernels per graph +// - Mixed kernel + memcpy nodes +// - Interleaved launches from different graphs on the same stream +// - Repeated launches of the same graph (cache overwrite path) +// +// Expected GPU zone counts: +// graphA (kernel + memcpy + kernel): 5 launches x 3 nodes = 15 zones +// graphB (kernel + kernel + kernel): 5 launches x 3 nodes = 15 zones +// Total graph zones: 30 +// Plus setup memcpys, syncs, etc. // // Build: // make # release build // make debug # debug build (asserts enabled) // -// Run (start tracy-capture first, then run repro): -// tracy-capture -o out.tracy & -// ./repro +// Run: +// tracy-capture -o out.tracy -f & sleep 1 && ./repro #include #include @@ -21,9 +29,12 @@ __global__ void vector_add(float* a, float* b, float* c, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) { - c[i] = a[i] + b[i]; - } + if (i < n) c[i] = a[i] + b[i]; +} + +__global__ void vector_scale(float* a, float scale, int n) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) a[i] *= scale; } #define CHECK_CUDA(call) \ @@ -44,62 +55,83 @@ int main() { const int N = 1 << 20; const size_t bytes = N * sizeof(float); + const int threads = 256; + const int blocks = (N + threads - 1) / threads; - float *d_a, *d_b, *d_c; - CHECK_CUDA(cudaMalloc(&d_a, bytes)); - CHECK_CUDA(cudaMalloc(&d_b, bytes)); - CHECK_CUDA(cudaMalloc(&d_c, bytes)); + float *d_a, *d_b, *d_c, *d_tmp; + CHECK_CUDA(cudaMalloc(&d_a, bytes)); + CHECK_CUDA(cudaMalloc(&d_b, bytes)); + CHECK_CUDA(cudaMalloc(&d_c, bytes)); + CHECK_CUDA(cudaMalloc(&d_tmp, bytes)); float* h_a = (float*)malloc(bytes); float* h_b = (float*)malloc(bytes); - for (int i = 0; i < N; i++) { - h_a[i] = 1.0f; - h_b[i] = 2.0f; - } + for (int i = 0; i < N; i++) { h_a[i] = 1.0f; h_b[i] = 2.0f; } CHECK_CUDA(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice)); CHECK_CUDA(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice)); - // --- Create a CUDA Graph via stream capture --- cudaStream_t stream; CHECK_CUDA(cudaStreamCreate(&stream)); + // --- Graph A: kernel(add) + memcpy + kernel(add) --- + // 3 nodes, graphId will be assigned by CUPTI CHECK_CUDA(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal)); + vector_add<<>>(d_a, d_b, d_c, N); + CHECK_CUDA(cudaMemcpyAsync(d_tmp, d_c, bytes, cudaMemcpyDeviceToDevice, stream)); + vector_add<<>>(d_a, d_tmp, d_c, N); + cudaGraph_t graphA; + cudaGraphExec_t execA; + CHECK_CUDA(cudaStreamEndCapture(stream, &graphA)); + CHECK_CUDA(cudaGraphInstantiate(&execA, graphA, nullptr, nullptr, 0)); - int threadsPerBlock = 256; - int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock; - vector_add<<>>(d_a, d_b, d_c, N); - CHECK_CUDA(cudaMemcpyAsync(d_c, d_c, bytes, cudaMemcpyDeviceToDevice, stream)); - vector_add<<>>(d_a, d_c, d_c, N); + // --- Graph B: kernel(scale) + kernel(add) + kernel(scale) --- + // 3 nodes, different graphId from A + CHECK_CUDA(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal)); + vector_scale<<>>(d_c, 0.5f, N); + vector_add <<>>(d_a, d_b, d_c, N); + vector_scale<<>>(d_c, 2.0f, N); + cudaGraph_t graphB; + cudaGraphExec_t execB; + CHECK_CUDA(cudaStreamEndCapture(stream, &graphB)); + CHECK_CUDA(cudaGraphInstantiate(&execB, graphB, nullptr, nullptr, 0)); - cudaGraph_t graph; - CHECK_CUDA(cudaStreamEndCapture(stream, &graph)); + printf("Graph A: kernel + memcpy + kernel (3 nodes)\n"); + printf("Graph B: scale + add + scale (3 nodes)\n"); + printf("Interleaving 5 launches each...\n"); - cudaGraphExec_t graphExec; - CHECK_CUDA(cudaGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - - printf("CUDA Graph created with 3 nodes (kernel + memcpy + kernel)\n"); - printf("Launching graph 10 times...\n"); - - // Each launch should produce 3 GPU zones (2 kernels + 1 memcpy), all - // correlated back to the cuGraphLaunch CPU call site. - for (int i = 0; i < 10; i++) { - ZoneScopedN("cuGraphLaunch iteration"); - CHECK_CUDA(cudaGraphLaunch(graphExec, stream)); + // Interleave launches: A, B, A, B, ... to stress graphId cache switching + for (int i = 0; i < 5; i++) { + { + ZoneScopedN("graphA launch"); + CHECK_CUDA(cudaGraphLaunch(execA, stream)); + } + { + ZoneScopedN("graphB launch"); + CHECK_CUDA(cudaGraphLaunch(execB, stream)); + } } CHECK_CUDA(cudaStreamSynchronize(stream)); - printf("Done. Expected 30 GPU zones in Tracy (10 launches x 3 ops).\n"); + printf("Done.\n"); + printf("Expected GPU zones:\n"); + printf(" graphA: 5 launches x 3 nodes = 15\n"); + printf(" graphB: 5 launches x 3 nodes = 15\n"); + printf(" Total graph zones: 30\n"); + // Verify correctness float* h_c = (float*)malloc(bytes); CHECK_CUDA(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost)); - printf("Result check: c[0] = %.1f (expected 4.0)\n", h_c[0]); + printf("Result check: c[0] = %.1f (expected 6.0: (a+b)*2 after last graphB)\n", h_c[0]); - CHECK_CUDA(cudaGraphExecDestroy(graphExec)); - CHECK_CUDA(cudaGraphDestroy(graph)); + CHECK_CUDA(cudaGraphExecDestroy(execA)); + CHECK_CUDA(cudaGraphExecDestroy(execB)); + CHECK_CUDA(cudaGraphDestroy(graphA)); + CHECK_CUDA(cudaGraphDestroy(graphB)); CHECK_CUDA(cudaStreamDestroy(stream)); CHECK_CUDA(cudaFree(d_a)); CHECK_CUDA(cudaFree(d_b)); CHECK_CUDA(cudaFree(d_c)); + CHECK_CUDA(cudaFree(d_tmp)); free(h_a); free(h_b); free(h_c);