Scarfski/tracy
1
0
Fork 0
mirror of https://github.com/wolfpld/tracy.git synced 2026-08-30 06:58:28 +00:00
Code Issues Packages Projects Releases Wiki Activity

Expand repro: multiple graphs, multiple kernels, interleaved launches

Browse Source 
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 <noreply@anthropic.com>
This commit is contained in:
Basil Milanich
2026-04-06 11:54:08 -05:00
parent d36ca27041
commit 4ccaea9f08
1 changed files with 72 additions and 40 deletions
Download Patch File Download Diff File Expand all files Collapse all files

112
examples/CUDAGraphRepro/repro.cu
View File

@@ -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 <cstdio>
#include <cstdlib>
@@ -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<<<blocks, threads, 0, stream>>>(d_a, d_b, d_c, N);
CHECK_CUDA(cudaMemcpyAsync(d_tmp, d_c, bytes, cudaMemcpyDeviceToDevice, stream));
vector_add<<<blocks, threads, 0, stream>>>(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<<<blocksPerGrid, threadsPerBlock, 0, stream>>>(d_a, d_b, d_c, N);
CHECK_CUDA(cudaMemcpyAsync(d_c, d_c, bytes, cudaMemcpyDeviceToDevice, stream));
vector_add<<<blocksPerGrid, threadsPerBlock, 0, stream>>>(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<<<blocks, threads, 0, stream>>>(d_c, 0.5f, N);
vector_add <<<blocks, threads, 0, stream>>>(d_a, d_b, d_c, N);
vector_scale<<<blocks, threads, 0, stream>>>(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);