From 7bca9dcd9062fd5d3f22aa886e4b146cb88630af Mon Sep 17 00:00:00 2001 From: Basil Milanich Date: Mon, 6 Apr 2026 11:03:58 -0500 Subject: [PATCH] Fix CUDA Graph GPU zones with proper cuGraphLaunch correlation MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Replace the synthetic APICallInfo hack with proper correlation via CUPTI_ACTIVITY_KIND_GRAPH_TRACE. When cuGraphLaunch fires an API callback, its correlationId is stored in cudaCallSiteInfo. The GRAPH_TRACE activity record carries the same correlationId plus the graphId, which lets us build a graphId→APICallInfo map. Kernel/memcpy/ memset activities then look up this map via their graphId field. Key changes: - Add cuGraphLaunch/cuGraphLaunch_ptsz to cbidDriverTrackers so the API callback machinery captures the CPU call site - Enable CUPTI_ACTIVITY_KIND_GRAPH_TRACE and handle it in DoProcessDeviceEvent to populate cudaGraphCurrentLaunch[graphId] - Add cudaGraphCurrentLaunch map to PersistentState - Two-pass buffer processing in OnBufferCompleted so GRAPH_TRACE records (which complete last on GPU) are processed before the kernel/memcpy/memset records that depend on them - Replace graphId=0 fallback in kernel/memcpy/memset with proper cudaGraphCurrentLaunch lookup; fall through to matchError if the graphId is not found - Update repro to include TracyCUDA headers and properly test GPU zone correlation Co-Authored-By: Claude Sonnet 4.6 --- examples/CUDAGraphRepro/Makefile | 37 +++++++++++++++---- examples/CUDAGraphRepro/repro.cu | 50 +++++++++++++------------- public/tracy/TracyCUDA.hpp | 61 ++++++++++++++++++++++---------- 3 files changed, 98 insertions(+), 50 deletions(-) diff --git a/examples/CUDAGraphRepro/Makefile b/examples/CUDAGraphRepro/Makefile index e0393ab3..665f2e89 100644 --- a/examples/CUDAGraphRepro/Makefile +++ b/examples/CUDAGraphRepro/Makefile @@ -1,12 +1,37 @@ -NVCC := nvcc -CFLAGS := -O2 +TRACY_PUBLIC := ../../public +NVCC := nvcc +CXX := g++ +CUPTI_INC := /usr/local/cuda/include +CUPTI_LIB := /usr/local/cuda/lib64 -.PHONY: all clean +TRACY_SRCS := $(TRACY_PUBLIC)/TracyClient.cpp +INCLUDES := -I$(TRACY_PUBLIC) -I$(CUPTI_INC) +LIBS := -L$(CUPTI_LIB) -lcuda -lcupti -lpthread -ldl + +CXXFLAGS_REL := -O2 -DTRACY_ENABLE +CXXFLAGS_DBG := -g -O0 -DTRACY_ENABLE +NVCCFLAGS_REL := -O2 -DTRACY_ENABLE +NVCCFLAGS_DBG := -g -O0 -DTRACY_ENABLE + +.PHONY: all debug clean all: repro -repro: repro.cu - $(NVCC) $(CFLAGS) -o $@ $< +debug: repro_debug + +# Release build +repro: repro.cu tracy_client.o + $(NVCC) $(NVCCFLAGS_REL) $(INCLUDES) -o $@ $< tracy_client.o $(LIBS) + +tracy_client.o: $(TRACY_SRCS) + $(CXX) $(CXXFLAGS_REL) $(INCLUDES) -c -o $@ $< + +# Debug build (asserts enabled, no NDEBUG) +repro_debug: repro.cu tracy_client_debug.o + $(NVCC) $(NVCCFLAGS_DBG) $(INCLUDES) -o $@ $< tracy_client_debug.o $(LIBS) + +tracy_client_debug.o: $(TRACY_SRCS) + $(CXX) $(CXXFLAGS_DBG) $(INCLUDES) -c -o $@ $< clean: - rm -f repro + rm -f repro repro_debug tracy_client.o tracy_client_debug.o diff --git a/examples/CUDAGraphRepro/repro.cu b/examples/CUDAGraphRepro/repro.cu index 5428ca6e..71d5c14d 100644 --- a/examples/CUDAGraphRepro/repro.cu +++ b/examples/CUDAGraphRepro/repro.cu @@ -1,23 +1,24 @@ // Tracy CUDA Graph GPU Zone Repro // -// Demonstrates that Tracy (unpatched) fails to show GPU zones for kernels -// launched via CUDA Graphs. The CUPTI activity records arrive but have no -// matching API callback correlation, so matchActivityToAPICall() fails and -// matchError() silently drops every GPU zone. +// 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. // // Build: -// nvcc -o repro repro.cu -lcuda -lcupti -I/path/to/tracy/public \ -// -DTRACY_ENABLE -DTRACY_ON_DEMAND +// make # release build +// make debug # debug build (asserts enabled) // -// Run with Tracy profiler connected to see: -// - Unpatched: 0 GPU zones from the graph-launched kernels -// - Patched: GPU zones appear for each kernel invocation +// Run (start tracy-capture first, then run repro): +// tracy-capture -o out.tracy & +// ./repro #include #include #include -// A trivial kernel — just increments each element. +#include "tracy/Tracy.hpp" +#include "tracy/TracyCUDA.hpp" + __global__ void vector_add(float* a, float* b, float* c, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) { @@ -36,16 +37,19 @@ __global__ void vector_add(float* a, float* b, float* c, int n) { } while (0) int main() { - const int N = 1 << 20; // 1M elements + ZoneScoped; + + auto ctx = TracyCUDAContext(); + TracyCUDAStartProfiling(ctx); + + const int N = 1 << 20; const size_t bytes = N * sizeof(float); - // Allocate device memory 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)); - // Initialize with some data float* h_a = (float*)malloc(bytes); float* h_b = (float*)malloc(bytes); for (int i = 0; i < N; i++) { @@ -59,44 +63,37 @@ int main() { cudaStream_t stream; CHECK_CUDA(cudaStreamCreate(&stream)); - // Begin capture CHECK_CUDA(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal)); - // Record operations into the graph 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); - // End capture cudaGraph_t graph; CHECK_CUDA(cudaStreamEndCapture(stream, &graph)); - // Instantiate the graph 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"); - // --- Launch the graph multiple times --- - // With unpatched Tracy, these produce 0 GPU zones. - // With patched Tracy, each launch produces 3 GPU zones (2 kernels + 1 memcpy). + // 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)); } CHECK_CUDA(cudaStreamSynchronize(stream)); - printf("Done. Expected ~30 GPU zones in Tracy (10 launches x 3 ops).\n"); - printf("Unpatched Tracy will show 0 GPU zones.\n"); + printf("Done. Expected 30 GPU zones in Tracy (10 launches x 3 ops).\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 after two additions)\n", h_c[0]); + printf("Result check: c[0] = %.1f (expected 4.0)\n", h_c[0]); - // Cleanup CHECK_CUDA(cudaGraphExecDestroy(graphExec)); CHECK_CUDA(cudaGraphDestroy(graph)); CHECK_CUDA(cudaStreamDestroy(stream)); @@ -107,5 +104,8 @@ int main() { free(h_b); free(h_c); + TracyCUDAStopProfiling(ctx); + TracyCUDAContextDestroy(ctx); + return 0; } diff --git a/public/tracy/TracyCUDA.hpp b/public/tracy/TracyCUDA.hpp index b941e676..1fdb3158 100644 --- a/public/tracy/TracyCUDA.hpp +++ b/public/tracy/TracyCUDA.hpp @@ -661,9 +661,18 @@ namespace tracy ZoneScoped; tracy::SetThreadName("NVIDIA CUPTI Worker"); CUptiResult status; + // Two-pass processing: GRAPH_TRACE records complete last on the GPU (after + // all their child kernels), so they appear at the end of the buffer. Process + // them first to populate the graphId->APICallInfo map before kernels look it up. CUpti_Activity* record = nullptr; + while (cuptiActivityGetNextRecord(buffer, validSize, &record) == CUPTI_SUCCESS) { + if (record->kind == CUPTI_ACTIVITY_KIND_GRAPH_TRACE) + DoProcessDeviceEvent(record); + } + record = nullptr; while ((status = cuptiActivityGetNextRecord(buffer, validSize, &record)) == CUPTI_SUCCESS) { - DoProcessDeviceEvent(record); + if (record->kind != CUPTI_ACTIVITY_KIND_GRAPH_TRACE) + DoProcessDeviceEvent(record); } if (status != CUPTI_ERROR_MAX_LIMIT_REACHED) { CUptiCallChecked(status, "cuptiActivityGetNextRecord", TracyFile, TracyLine); @@ -833,6 +842,10 @@ namespace tracy { CUPTI_DRIVER_TRACE_CBID_cuEventSynchronize, NON_STREAM_FUNC() }, { CUPTI_DRIVER_TRACE_CBID_cuCtxSynchronize, NON_STREAM_FUNC() }, { CUPTI_DRIVER_TRACE_CBID_cuStreamWaitEvent, GET_STREAM_FUNC(cuStreamWaitEvent_params, hStream) }, + // Graph launch: tracked so we can correlate GPU activities back to + // the cuGraphLaunch call site via CUPTI_ACTIVITY_KIND_GRAPH_TRACE + { CUPTI_DRIVER_TRACE_CBID_cuGraphLaunch, GET_STREAM_FUNC(cuGraphLaunch_params, hStream) }, + { CUPTI_DRIVER_TRACE_CBID_cuGraphLaunch_ptsz, GET_STREAM_FUNC(cuGraphLaunch_params, hStream) }, }; #undef NON_STREAM_FUNC #undef GET_STREAM_FUNC @@ -998,12 +1011,12 @@ namespace tracy CUpti_ActivityKernel9* kernel9 = (CUpti_ActivityKernel9*) record; APICallInfo apiCall; if (!matchActivityToAPICall(kernel9->correlationId, apiCall)) { - // Fallback for CUDA Graph-launched kernels: create GPU zone - // using kernel timestamps when no API callback correlation exists. - auto* host = PersistentState::Get().profilerHost; - if (!host) return; - TracyTimestamp cpuTime = tracyFromCUpti(kernel9->start); - apiCall = APICallInfo{ cpuTime, cpuTime, kernel9->start, host }; + uint32_t graphId = kernel9->graphId; + if (graphId == 0 || !PersistentState::Get().cudaGraphCurrentLaunch.fetch(graphId, apiCall)) { + return matchError(kernel9->correlationId, "KERNEL"); + } + // Don't erase cudaGraphCurrentLaunch: multiple kernels in the same + // graph launch share this APICallInfo entry. } apiCall.host->EmitGpuZone(apiCall.start, apiCall.end, kernel9->start, kernel9->end, getKernelSourceLocation(kernel9->name), kernel9->contextId, kernel9->streamId); auto latency_ms = (kernel9->start - apiCall.cupti) / 1'000'000.0; @@ -1017,12 +1030,10 @@ namespace tracy CUpti_ActivityMemcpy5* memcpy5 = (CUpti_ActivityMemcpy5*) record; APICallInfo apiCall; if (!matchActivityToAPICall(memcpy5->correlationId, apiCall)) { - // Fallback for CUDA Graph memcpy: create GPU zone using - // activity timestamps when no API callback correlation exists. - auto* host = PersistentState::Get().profilerHost; - if (!host) return; - TracyTimestamp cpuTime = tracyFromCUpti(memcpy5->start); - apiCall = APICallInfo{ cpuTime, cpuTime, memcpy5->start, host }; + uint32_t graphId = memcpy5->graphId; + if (graphId == 0 || !PersistentState::Get().cudaGraphCurrentLaunch.fetch(graphId, apiCall)) { + return matchError(memcpy5->correlationId, "MEMCPY"); + } } static constexpr tracy::SourceLocationData TracyCUPTISrcLocDeviceMemcpy { "CUDA::memcpy", TracyFunction, TracyFile, (uint32_t)TracyLine, tracy::Color::Blue }; apiCall.host->EmitGpuZone(apiCall.start, apiCall.end, memcpy5->start, memcpy5->end, &TracyCUPTISrcLocDeviceMemcpy, memcpy5->contextId, memcpy5->streamId); @@ -1038,12 +1049,10 @@ namespace tracy CUpti_ActivityMemset4* memset4 = (CUpti_ActivityMemset4*) record; APICallInfo apiCall; if (!matchActivityToAPICall(memset4->correlationId, apiCall)) { - // Fallback for CUDA Graph memset: create GPU zone using - // activity timestamps when no API callback correlation exists. - auto* host = PersistentState::Get().profilerHost; - if (!host) return; - TracyTimestamp cpuTime = tracyFromCUpti(memset4->start); - apiCall = APICallInfo{ cpuTime, cpuTime, memset4->start, host }; + uint32_t graphId = memset4->graphId; + if (graphId == 0 || !PersistentState::Get().cudaGraphCurrentLaunch.fetch(graphId, apiCall)) { + return matchError(memset4->correlationId, "MEMSET"); + } } static constexpr tracy::SourceLocationData TracyCUPTISrcLocDeviceMemset { "CUDA::memset", TracyFunction, TracyFile, (uint32_t)TracyLine, tracy::Color::Blue }; apiCall.host->EmitGpuZone(apiCall.start, apiCall.end, memset4->start, memset4->end, &TracyCUPTISrcLocDeviceMemset, memset4->contextId, memset4->streamId); @@ -1117,6 +1126,18 @@ namespace tracy } break; } + case CUPTI_ACTIVITY_KIND_GRAPH_TRACE: + { + // Correlate the cuGraphLaunch API call to the graph's graphId so that + // kernel/memcpy/memset activities can look it up via their graphId field. + // cuGraphLaunch's correlationId == graphTrace->correlationId (per CUPTI docs). + auto* graphTrace = (CUpti_ActivityGraphTrace2*)record; + APICallInfo apiCall; + if (matchActivityToAPICall(graphTrace->correlationId, apiCall)) { + PersistentState::Get().cudaGraphCurrentLaunch.emplace(graphTrace->graphId, apiCall); + } + break; + } case CUPTI_ACTIVITY_KIND_CUDA_EVENT : { // NOTE(marcos): a byproduct of CUPTI_ACTIVITY_KIND_SYNCHRONIZATION @@ -1151,6 +1172,7 @@ namespace tracy CUPTI_ACTIVITY_KIND_MEMSET, CUPTI_ACTIVITY_KIND_SYNCHRONIZATION, CUPTI_ACTIVITY_KIND_MEMORY2, + CUPTI_ACTIVITY_KIND_GRAPH_TRACE, //CUPTI_ACTIVITY_KIND_MEMCPY2, //CUPTI_ACTIVITY_KIND_OVERHEAD, //CUPTI_ACTIVITY_KIND_INTERNAL_LAUNCH_API, @@ -1279,6 +1301,7 @@ namespace tracy // NOTE(marcos): these objects do not need to persist, but their relative // footprint is trivial enough that we don't care if we let them leak ConcurrentHashMap cudaCallSiteInfo; + ConcurrentHashMap cudaGraphCurrentLaunch; ConcurrentHashMap memAllocAddress; CUpti_SubscriberHandle subscriber = {}; CUDACtx* profilerHost = nullptr;