We are seeing issues with compute-sanitizer when used with CUDA graphs that we think are false-positives.
Depending on the driver/CUDA version and GPU architecture, some tests that we run under compute-sanitizer sometimes or always fail, with diverse messages.
The below example always fails on our NVIDIA GeForce RTX 5070 Ti (NVIDIA-SMI 610.43.02, KMD Version: 610.43.02, CUDA UMD Version: 13.3), though we think it’s semantically valid. The example can be fixed by adding a stream synchronization before launching the graph, which is surprising as the graph launch should be stream-ordered.
========= COMPUTE-SANITIZER
Incrementing at index 0.
Incrementing at index 0.
Incrementing at index 0.
========= Invalid __global__ atomic of size 4 bytes
========= at __iAtomicAdd+0x80
========= by thread (0,0,0) in block (0,0,0)
========= Access to 0x13e0000000 is potentially made before memory is allocated
========= and is inside the nearest allocation at 0x13e0000000 of size 16 bytes
========= Device Frame: atomicAdd(int *, int)+0x180 in device_atomic_functions.hpp:107
========= Device Frame: MyFunctor<int *>::operator ()(unsigned int) const+0x3f0 in graph.conditional.cu:30
========= Device Frame: void launcher<MyFunctor<int *>>(T1)+0x200 in graph.conditional.cu:36
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame: cudaGraphLaunch [0x80191] in graph.conditional.exe
========= Host Frame: test_diamond() [0x9347] in graph.conditional.exe
========= Host Frame: main [0x95da] in graph.conditional.exe
=========
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaStreamSynchronize.
========= Saved host backtrace up to driver entry point at error
========= Host Frame: test_diamond() [0x93b0] in graph.conditional.exe
========= Host Frame: main [0x95da] in graph.conditional.exe
=========
/tmp/tmpbn9kn6xn/graph.conditional.cu:90: failure of statement cudaStreamSynchronize(stream): unspecified launch failure (719)
========= Error: process didn't terminate successfully
========= Target application returned an error
========= LEAK SUMMARY: 0 bytes leaked in 0 allocations
========= ERROR SUMMARY: 2 errors
Traceback (most recent call last):
File "//test.py", line 155, in <module>
raise RuntimeError(f"It failed after {trial} trial(s).")
RuntimeError: It failed after 0 trial(s).
import logging
import pathlib
import subprocess
import tempfile
CONTENT = \
"""
#include <array>
#include <cassert>
#include <cstdio>
#include <cuda_runtime.h>
//! Check the return code of an API call.
#define CHECK_CALL(call) \\
{ \\
const auto error_code = call; \\
if(error_code != cudaSuccess) \\
{ \\
printf("%s:%d: failure of statement %s: %s (%d)\\n", \\
__FILE__, __LINE__, \\
#call, \\
cudaGetErrorString(error_code), error_code); \\
std::abort(); \\
} \\
}
template <typename view_t>
struct MyFunctor
{
view_t data;
__device__
void operator()(const unsigned int index) const noexcept {
printf("Incrementing at index %d.\\n", index);
//++data[index];
atomicAdd(&data[index], 1);
}
};
template <typename Functor>
__global__ void launcher(const Functor functor) {
functor.operator()(threadIdx.x);
}
template <typename Functor>
auto get_launcher() { return launcher<Functor>; }
int test_diamond(void)
{
//! Working stream.
cudaStream_t stream;
CHECK_CALL(cudaStreamCreate(&stream));
//! Allocate device memory to use as input.
int *dPtr;
CHECK_CALL(cudaMallocAsync((void**)&dPtr, 4 * sizeof(int), stream));
//! Create graph (diamond pattern).
cudaGraph_t graph;
CHECK_CALL(cudaGraphCreate(&graph, 0));
cudaGraphNode_t node_A, node_B, node_C, node_D;
cudaKernelNodeParams params_A, params_B, params_C, params_D;
params_A.gridDim = params_B.gridDim = params_C.gridDim = params_D.gridDim = dim3(1, 1 ,1);
params_A.blockDim = params_B.blockDim = params_C.blockDim = params_D.blockDim = dim3(1, 1 ,1);
params_A.sharedMemBytes = params_B.sharedMemBytes = params_C.sharedMemBytes = params_D.sharedMemBytes = 0;
params_A.func = params_B.func = params_C.func = params_D.func =(void*)get_launcher<MyFunctor<int*>>();
MyFunctor<int*> functor_A{.data = dPtr}, functor_B{.data = dPtr}, functor_C{.data = dPtr}, functor_D{.data = dPtr};
std::array<void*, 1> inputs_A{(void*)&functor_A}, inputs_B{(void*)&functor_B}, inputs_C{(void*)&functor_C}, inputs_D{(void*)&functor_D};
params_A.kernelParams = inputs_A.data();
params_B.kernelParams = inputs_B.data();
params_C.kernelParams = inputs_C.data();
params_D.kernelParams = inputs_D.data();
params_A.extra = params_B.extra = params_C.extra = params_D.extra = nullptr;
CHECK_CALL(cudaGraphAddKernelNode(&node_A, graph, nullptr, 0, ¶ms_A));
CHECK_CALL(cudaGraphAddKernelNode(&node_B, graph, &node_A, 1, ¶ms_B));
CHECK_CALL(cudaGraphAddKernelNode(&node_C, graph, &node_A, 1, ¶ms_C));
CHECK_CALL(cudaGraphAddKernelNode(&node_D, graph, std::array<cudaGraphNode_t, 2>{node_B, node_C}.data(), 2, ¶ms_D));
//! Instantiate the graph and launch it.
cudaGraphExec_t graph_exec;
CHECK_CALL(cudaGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0));
CHECK_CALL(cudaGraphLaunch(graph_exec, stream));
CHECK_CALL(cudaStreamSynchronize(stream));
CHECK_CALL(cudaGraphExecDestroy(graph_exec));
CHECK_CALL(cudaGraphDestroy(graph));
CHECK_CALL(cudaFree(dPtr));
CHECK_CALL(cudaStreamDestroy(stream));
return EXIT_SUCCESS;
}
int main()
{
return test_diamond();
}
"""
def compile(source, target):
"""
Compile with device debug symbols.
"""
cmd = ['nvcc', '-G', source, '-o', target]
logging.info(f"Compiling {source} to {target} with {cmd}.")
subprocess.check_call(cmd)
def check(target):
"""
Run without `compute-sanitizer` to ensure that it runs fine.
Run it many times just to be sure.
"""
for _ in range(10):
subprocess.check_call(target)
logging.info("Checking was successful.")
def run(target):
return subprocess.run([
'compute-sanitizer', '--leak-check=full', target
])
if __name__ == "__main__":
with tempfile.TemporaryDirectory() as tmpdir:
output = pathlib.Path(tmpdir) / 'graph.conditional.tmp'
source = output.with_suffix('.cu')
target = output.with_suffix('.exe')
with open(source, 'w+') as fout:
fout.write(CONTENT)
compile(source = source, target = target)
check(target = target)
trial = 0
while True:
logging.info(f"Trial {trial}.")
if run(target = target).returncode != 0:
raise RuntimeError(f"It failed after {trial} trial(s).")
trial += 1