Martas/Towards
1
0
Fork 0
Code Issues 7 Pull Requests Actions Packages Projects 1 Releases Wiki Activity

initial

Browse Source
This commit is contained in:
Martin Slachta
2026-07-18 14:31:15 +02:00
commit a04f0dc262
3343 changed files with 1140208 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
external/tracy/examples/CUDAGraphRepro/Makefile
Vendored
+49
View File
@@ -0,0 +1,49 @@
TRACY_PUBLIC := ../../public
NVCC := nvcc
CXX := g++
CUPTI_INC := /usr/local/cuda/include
CUPTI_LIB := /usr/local/cuda/lib64
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 := -arch=native -O2 -DTRACY_ENABLE
NVCCFLAGS_DBG := -arch=native -g -O0 -DTRACY_ENABLE
.PHONY: all debug investigate investigate2 clean
all: repro
debug: repro_debug
investigate: test_corr_reuse
investigate2: test_graphid_recycle
# 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 $@ $<
# Investigation: correlationId uniqueness per graph launch (no Tracy dependency)
test_corr_reuse: test_corr_reuse.cu
$(NVCC) $(NVCCFLAGS_REL) $(INCLUDES) -o $@ $< $(LIBS)
# Investigation: does CUPTI recycle graphId values after cudaGraphExecDestroy?
test_graphid_recycle: test_graphid_recycle.cu
$(NVCC) $(NVCCFLAGS_REL) $(INCLUDES) -o $@ $< $(LIBS)
clean:
rm -f repro repro_debug test_corr_reuse test_graphid_recycle tracy_client.o tracy_client_debug.o
external/tracy/examples/CUDAGraphRepro/README.md
Vendored
+35
View File
@@ -0,0 +1,35 @@
# Tracy CUDA Graph GPU Zone Repro
Demonstrates that unpatched Tracy fails to show GPU zones for kernels
launched via CUDA Graphs (`cudaGraphLaunch`).
## Root cause
When kernels are launched through CUDA Graphs, CUPTI delivers
`CONCURRENT_KERNEL` and `MEMCPY` activity records but no corresponding
API callback fires for the individual kernel launches. Tracy's
`matchActivityToAPICall()` always fails, and `matchError()` silently
drops every GPU zone.
## Build and run
```bash
make
./repro
```
## What to expect
| Tracy version | GPU zones shown |
|---|---|
| Unpatched | 0 |
| Patched (cuda-graph-gpu-zones.patch) | ~30 (10 launches x 3 ops) |
## The graph structure
Each graph launch contains:
1. `vector_add` kernel (c = a + b)
2. Device-to-device memcpy
3. `vector_add` kernel (c = a + c)
The graph is launched 10 times, so 30 GPU operations total.
external/tracy/examples/CUDAGraphRepro/repro.cu
Vendored
+143
View File
@@ -0,0 +1,143 @@
// Tracy CUDA Graph GPU Zone Repro
//
// 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:
// tracy-capture -o out.tracy -f & sleep 1 && ./repro
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#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) 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) \
do { \
cudaError_t err = (call); \
if (err != cudaSuccess) { \
fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__, \
cudaGetErrorString(err)); \
exit(1); \
} \
} while (0)
int main() {
ZoneScoped;
auto ctx = TracyCUDAContext();
TracyCUDAStartProfiling(ctx);
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, *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; }
CHECK_CUDA(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
CHECK_CUDA(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));
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));
// --- 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));
printf("Graph A: kernel + memcpy + kernel (3 nodes)\n");
printf("Graph B: scale + add + scale (3 nodes)\n");
printf("Interleaving 5 launches each...\n");
// 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.\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 6.0: (a+b)*2 after last graphB)\n", h_c[0]);
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);
TracyCUDAStopProfiling(ctx);
TracyCUDAContextDestroy(ctx);
return 0;
}
external/tracy/examples/CUDAGraphRepro/test_corr_reuse.cu
Vendored
+106
View File
@@ -0,0 +1,106 @@
// Investigate: does relaunching the same cudaGraphExec produce a new correlationId
// each time, or is the correlationId reused/fixed per exec handle?
//
// Also checks: do different graphExec handles for the same graph share graphId?
#include <cstdio>
#include <cuda_runtime.h>
#include <cupti.h>
#define CHECK(x) do { cudaError_t e=(x); if(e!=cudaSuccess){fprintf(stderr,"CUDA %s:%d: %s\n",__FILE__,__LINE__,cudaGetErrorString(e));exit(1);} } while(0)
#define CHECK_CUPTI(x) do { CUptiResult e=(x); if(e!=CUPTI_SUCCESS){const char*s;cuptiGetResultString(e,&s);fprintf(stderr,"CUPTI %s:%d: %s\n",__FILE__,__LINE__,s);exit(1);} } while(0)
struct Record { uint32_t corr; uint32_t graphId; };
static Record records[64];
static int nrecords = 0;
static void CUPTIAPI bufferRequested(uint8_t** buf, size_t* size, size_t* maxNumRecords) {
*size = 1 << 20; *buf = (uint8_t*)malloc(*size); *maxNumRecords = 0;
}
static void CUPTIAPI bufferCompleted(CUcontext ctx, uint32_t streamId,
uint8_t* buf, size_t size, size_t validSize) {
CUpti_Activity* record = nullptr;
while (cuptiActivityGetNextRecord(buf, validSize, &record) == CUPTI_SUCCESS) {
if (record->kind == CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL) {
auto* k = (CUpti_ActivityKernel9*)record;
if (nrecords < 64)
records[nrecords++] = { k->correlationId, k->graphId };
}
}
free(buf);
}
__global__ void dummy(int* x) { atomicAdd(x, 1); }
static uint32_t launchCorrId[32];
static int nlaunch = 0;
// Intercept cudaGraphLaunch via CUPTI callback to capture the correlationId
// assigned to each launch on the CPU side
static void CUPTIAPI onCallback(void* userdata, CUpti_CallbackDomain domain,
CUpti_CallbackId cbid, const void* cbdata) {
if (domain != CUPTI_CB_DOMAIN_RUNTIME_API) return;
if (cbid != CUPTI_RUNTIME_TRACE_CBID_cudaGraphLaunch_v10000) return;
auto* api = (CUpti_CallbackData*)cbdata;
if (api->callbackSite == CUPTI_API_ENTER && nlaunch < 32)
launchCorrId[nlaunch++] = api->correlationId;
}
int main() {
CHECK_CUPTI(cuptiActivityRegisterCallbacks(bufferRequested, bufferCompleted));
CHECK_CUPTI(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL));
CUpti_SubscriberHandle sub;
CHECK_CUPTI(cuptiSubscribe(&sub, onCallback, nullptr));
CHECK_CUPTI(cuptiEnableCallback(1, sub, CUPTI_CB_DOMAIN_RUNTIME_API,
CUPTI_RUNTIME_TRACE_CBID_cudaGraphLaunch_v10000));
cudaStream_t stream; CHECK(cudaStreamCreate(&stream));
int* d_x; CHECK(cudaMalloc(&d_x, sizeof(int)));
// Build a simple graph with one kernel
cudaGraph_t graph; cudaGraphExec_t exec;
CHECK(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal));
dummy<<<1,1,0,stream>>>(d_x);
CHECK(cudaStreamEndCapture(stream, &graph));
CHECK(cudaGraphInstantiate(&exec, graph, nullptr, nullptr, 0));
// Launch the SAME exec handle 5 times
printf("=== Same exec, 5 launches ===\n");
for (int i = 0; i < 5; i++)
CHECK(cudaGraphLaunch(exec, stream));
CHECK(cudaStreamSynchronize(stream));
CHECK_CUPTI(cuptiActivityFlushAll(1));
printf("CPU-side launch correlationIds:\n");
for (int i = 0; i < nlaunch; i++)
printf(" launch[%d] corr=%u\n", i, launchCorrId[i]);
printf("GPU activity records (CONCURRENT_KERNEL):\n");
for (int i = 0; i < nrecords; i++)
printf(" kernel[%d] corr=%-4u graphId=%u\n", i, records[i].corr, records[i].graphId);
// Check: do CPU launch corrIds match GPU activity corrIds?
bool allMatch = (nlaunch == nrecords);
for (int i = 0; allMatch && i < nlaunch; i++)
allMatch = (launchCorrId[i] == records[i].corr);
printf("All launch corrIds match kernel corrIds: %s\n", allMatch ? "YES" : "NO (order may differ)");
// Now test: two different exec handles from the same graph — same graphId?
printf("\n=== Two exec handles from same graph ===\n");
nlaunch = 0; nrecords = 0;
cudaGraphExec_t exec2;
CHECK(cudaGraphInstantiate(&exec2, graph, nullptr, nullptr, 0));
CHECK(cudaGraphLaunch(exec, stream));
CHECK(cudaGraphLaunch(exec2, stream));
CHECK(cudaStreamSynchronize(stream));
CHECK_CUPTI(cuptiActivityFlushAll(1));
for (int i = 0; i < nrecords; i++)
printf(" kernel[%d] corr=%-4u graphId=%u\n", i, records[i].corr, records[i].graphId);
CHECK(cudaGraphExecDestroy(exec));
CHECK(cudaGraphExecDestroy(exec2));
CHECK(cudaGraphDestroy(graph));
CHECK(cudaFree(d_x));
CHECK(cudaStreamDestroy(stream));
CHECK_CUPTI(cuptiUnsubscribe(sub));
return 0;
}
external/tracy/examples/CUDAGraphRepro/test_graphid_recycle.cu
Vendored
+139
View File
@@ -0,0 +1,139 @@
// test_graphid_recycle.cu
//
// Investigates whether CUPTI recycles graphId values after cudaGraphExecDestroy.
//
// Question: Can two *different* graph exec handles ever produce the same graphId?
// (If yes, the graphLaunchCache in TracyCUDA.hpp could serve stale entries.)
//
// Rounds:
// Round 1: create + instantiate + launch + destroy graph A → record graphId
// Round 2: create + instantiate + launch + destroy graph B → does it reuse A's graphId?
// Round 3: 20 rapid create/launch/destroy cycles → attempt to exhaust any counter
#include <cstdio>
#include <cstdlib>
#include <cassert>
#include <vector>
#include <set>
#include <cuda_runtime.h>
#include <cupti.h>
#define CHECK_CUDA(call) \
do { \
cudaError_t err = (call); \
if (err != cudaSuccess) { \
fprintf(stderr, "CUDA error at %s:%d: %s\n", \
__FILE__, __LINE__, cudaGetErrorString(err)); \
exit(1); \
} \
} while (0)
#define CHECK_CUPTI(call) \
do { \
CUptiResult err = (call); \
if (err != CUPTI_SUCCESS) { \
const char* msg; \
cuptiGetResultString(err, &msg); \
fprintf(stderr, "CUPTI error at %s:%d: %s\n", \
__FILE__, __LINE__, msg); \
exit(1); \
} \
} while (0)
static std::vector<uint32_t> g_observed_graphIds;
static void CUPTIAPI bufferRequested(uint8_t** buffer, size_t* size, size_t* maxNumRecords) {
*size = 1 << 20;
*buffer = (uint8_t*)malloc(*size);
*maxNumRecords = 0;
}
static void CUPTIAPI bufferCompleted(CUcontext ctx, uint32_t streamId,
uint8_t* buffer, size_t size, size_t validSize) {
CUpti_Activity* record = nullptr;
while (cuptiActivityGetNextRecord(buffer, validSize, &record) == CUPTI_SUCCESS) {
if (record->kind == CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL) {
auto* kernel = (CUpti_ActivityKernel9*)record;
g_observed_graphIds.push_back(kernel->graphId);
}
}
free(buffer);
}
__global__ void dummy_kernel() {}
static uint32_t launchAndGetGraphId(cudaStream_t stream) {
size_t before = g_observed_graphIds.size();
cudaGraph_t graph;
cudaGraphExec_t exec;
CHECK_CUDA(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal));
dummy_kernel<<<1, 1, 0, stream>>>();
CHECK_CUDA(cudaStreamEndCapture(stream, &graph));
CHECK_CUDA(cudaGraphInstantiate(&exec, graph, nullptr, nullptr, 0));
CHECK_CUDA(cudaGraphLaunch(exec, stream));
CHECK_CUDA(cudaStreamSynchronize(stream));
CHECK_CUPTI(cuptiActivityFlushAll(CUPTI_ACTIVITY_FLAG_FLUSH_FORCED));
CHECK_CUDA(cudaGraphExecDestroy(exec));
CHECK_CUDA(cudaGraphDestroy(graph));
if (g_observed_graphIds.size() <= before) {
fprintf(stderr, "ERROR: No CONCURRENT_KERNEL record received\n");
return 0;
}
return g_observed_graphIds.back();
}
int main() {
// Initialize CUDA context
cudaFree(0);
CHECK_CUPTI(cuptiActivityRegisterCallbacks(bufferRequested, bufferCompleted));
CHECK_CUPTI(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL));
cudaStream_t stream;
CHECK_CUDA(cudaStreamCreate(&stream));
// Round 1
uint32_t id1 = launchAndGetGraphId(stream);
printf("Round 1 graphId: %u\n", id1);
// Round 2: new exec after round 1 destroyed
uint32_t id2 = launchAndGetGraphId(stream);
printf("Round 2 graphId: %u\n", id2);
if (id1 == id2) {
printf("*** RECYCLE DETECTED: id1 == id2 == %u ***\n", id1);
} else {
printf("Round 2 got different graphId (no recycle after 1 destroy)\n");
}
// Round 3: 20 rapid cycles — try to exhaust monotonic counter
printf("\nRound 3: 20 rapid create/launch/destroy cycles\n");
std::set<uint32_t> seen;
seen.insert(id1);
seen.insert(id2);
bool recycle_seen = false;
for (int i = 0; i < 20; i++) {
uint32_t id = launchAndGetGraphId(stream);
printf(" cycle %2d: graphId = %u", i + 1, id);
if (seen.count(id)) {
printf(" *** RECYCLED (seen before) ***");
recycle_seen = true;
}
printf("\n");
seen.insert(id);
}
if (!recycle_seen) {
printf("\nNo graphId recycling observed across %zu total launches.\n", seen.size());
printf("graphId appears to be a monotonically increasing counter.\n");
printf("Min=%u Max=%u Count=%zu\n",
*seen.begin(), *seen.rbegin(), seen.size());
}
CHECK_CUDA(cudaStreamDestroy(stream));
CHECK_CUPTI(cuptiActivityDisable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL));
return 0;
}