Add graphId recycle investigation test

Tests whether CUPTI recycles graphId values after cudaGraphExecDestroy,
which would be the only scenario where the graphLaunchCache in TracyCUDA
could serve stale entries for a non-matching exec handle.

Result (H100, CUDA 12, CUPTI): graphId is a monotonically increasing
counter that is never recycled. 22 create/instantiate/launch/destroy
cycles produced unique IDs ranging from 2 to 65 (incrementing by 3 per
cycle — one unit per node created during graph construction).

This confirms that the stale-cache concern raised in code review is not
a real risk in practice: two distinct exec handles always have distinct
graphIds.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
This commit is contained in:
Basil Milanich
2026-04-09 09:34:51 -05:00
parent c807367099
commit 12dc23f67e
2 changed files with 147 additions and 2 deletions

View File

@@ -13,7 +13,7 @@ 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 clean
.PHONY: all debug investigate investigate2 clean
all: repro
@@ -21,6 +21,8 @@ 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)
@@ -39,5 +41,9 @@ tracy_client_debug.o: $(TRACY_SRCS)
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 tracy_client.o tracy_client_debug.o
rm -f repro repro_debug test_corr_reuse test_graphid_recycle tracy_client.o tracy_client_debug.o

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;
}