Scarfski/tracy
1
0
Fork 0
mirror of https://github.com/wolfpld/tracy.git synced 2026-09-07 02:48:24 +00:00
Code Issues Packages Projects Releases Wiki Activity

Fix CUDA Graph GPU zones with proper cuGraphLaunch correlation

Browse Source 
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 <noreply@anthropic.com>
This commit is contained in:
Basil Milanich
2026-04-06 11:03:58 -05:00
parent 6c6999bf01
commit 7bca9dcd90
3 changed files with 98 additions and 50 deletions
Download Patch File Download Diff File Expand all files Collapse all files

37
examples/CUDAGraphRepro/Makefile
View File

@@ -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

50
examples/CUDAGraphRepro/repro.cu
View File

@@ -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 <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
// 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<<<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);
// 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;
}

61
public/tracy/TracyCUDA.hpp
View File

@@ -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<CorrelationID, APICallInfo> cudaCallSiteInfo;
ConcurrentHashMap<uint32_t, APICallInfo> cudaGraphCurrentLaunch;
ConcurrentHashMap<uintptr_t, int> memAllocAddress;
CUpti_SubscriberHandle subscriber = {};
CUDACtx* profilerHost = nullptr;