I have a minimal CUDA repro where a pure global-memory copy kernel
out[i] = in[i];
shows nondeterministic mismatches across iterations.
Key isolation result: I compare dev_out vs dev_in on device and copy back only a tiny result struct (no D2H copy of the large output buffer).
I still get mismatches, so corruption appears to be on the GPU side (not just D2H/host RAM path).
System under test:
- GPU: NVIDIA GeForce RTX 3060 Ti
- Driver: 535.274.02
nvcc: CUDA 12.1 (V12.1.105)
Useful additional info:
- Exact board partner SKU + VBIOS version
- OS + motherboard + PCIe generation/link width
- Whether behavior changes with different driver versions
- Whether behavior changes with memory underclock / stock vs OC settings
Repro code (repro.cu)
#include <cuda_runtime.h>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
namespace {
constexpr int kN = 500000;
constexpr int kIters = 500;
constexpr int kBlock = 256;
constexpr int kGrid = (kN + kBlock - 1) / kBlock;
void check_cuda(cudaError_t status, const char *ctx) {
if (status != cudaSuccess) {
std::fprintf(stderr, "%s failed: %s\n", ctx, cudaGetErrorString(status));
std::exit(EXIT_FAILURE);
}
}
__global__ void copy64_kernel(const uint64_t *in, uint64_t *out) {
const int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < kN) out[i] = in[i];
}
struct DeviceCompareResult {
int found;
unsigned int idx;
uint64_t in_val;
uint64_t out_val;
uint64_t xor_mask;
};
__global__ void find_mismatch_kernel(const uint64_t *in,
const uint64_t *out,
DeviceCompareResult *result) {
const int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < kN) {
const uint64_t a = in[i];
const uint64_t b = out[i];
if (a != b) {
if (atomicCAS(&result->found, 0, 1) == 0) {
result->idx = static_cast<unsigned int>(i);
result->in_val = a;
result->out_val = b;
result->xor_mask = a ^ b;
}
}
}
}
} // namespace
int main() {
uint64_t *host_in = nullptr;
check_cuda(cudaMallocHost(&host_in, static_cast<size_t>(kN) * sizeof(uint64_t)),
"cudaMallocHost(host_in)");
for (int i = 0; i < kN; ++i) {
host_in[i] = static_cast<uint64_t>(i) * 0x9e3779b97f4a7c15ULL + 0x123456789ULL;
}
uint64_t *dev_in = nullptr;
uint64_t *dev_out = nullptr;
DeviceCompareResult *dev_result = nullptr;
check_cuda(cudaMalloc(&dev_in, static_cast<size_t>(kN) * sizeof(uint64_t)), "cudaMalloc(dev_in)");
check_cuda(cudaMalloc(&dev_out, static_cast<size_t>(kN) * sizeof(uint64_t)), "cudaMalloc(dev_out)");
check_cuda(cudaMalloc(&dev_result, sizeof(DeviceCompareResult)), "cudaMalloc(dev_result)");
check_cuda(cudaMemcpy(dev_in,
host_in,
static_cast<size_t>(kN) * sizeof(uint64_t),
cudaMemcpyHostToDevice),
"cudaMemcpy(dev_in <- host_in)");
check_cuda(cudaDeviceSynchronize(), "cudaDeviceSynchronize(after H2D)");
DeviceCompareResult host_result{};
DeviceCompareResult reset_result{};
reset_result.found = 0;
reset_result.idx = 0xFFFFFFFFu;
reset_result.in_val = 0;
reset_result.out_val = 0;
reset_result.xor_mask = 0;
for (int iter = 0; iter < kIters; ++iter) {
check_cuda(cudaMemcpy(dev_result,
&reset_result,
sizeof(DeviceCompareResult),
cudaMemcpyHostToDevice),
"reset dev_result");
copy64_kernel<<<kGrid, kBlock>>>(dev_in, dev_out);
check_cuda(cudaGetLastError(), "copy64_kernel launch");
find_mismatch_kernel<<<kGrid, kBlock>>>(dev_in, dev_out, dev_result);
check_cuda(cudaGetLastError(), "find_mismatch_kernel launch");
check_cuda(cudaDeviceSynchronize(), "compare sync");
check_cuda(cudaMemcpy(&host_result,
dev_result,
sizeof(DeviceCompareResult),
cudaMemcpyDeviceToHost),
"cudaMemcpy(host_result <- dev_result)");
if (host_result.found != 0) {
const uint64_t delta = host_result.out_val - host_result.in_val;
std::printf("DEVICE_MISMATCH iter=%d first_diff=%u\n", iter, host_result.idx);
std::printf("in =%016llx out=%016llx\n",
static_cast<unsigned long long>(host_result.in_val),
static_cast<unsigned long long>(host_result.out_val));
std::printf("xor=%016llx delta=%016llx\n",
static_cast<unsigned long long>(host_result.xor_mask),
static_cast<unsigned long long>(delta));
return 0;
}
}
std::printf("No on-device mismatch in %d iterations.\n", kIters);
return 0;
}
Build and run
nvcc -arch=sm_86 -O3 -lineinfo repro.cu -o repro
./repro
Typical output:
DEVICE_MISMATCH iter=2 first_diff=484151
in =c6105e3eaed52f0c out=c6105e1eaed52f0c
xor=0000002000000000 delta=ffffffe000000000
Additional checks
compute-sanitizer --tool memcheck ./repro
compute-sanitizer --tool racecheck ./repro
Both report clean summaries on my side (0 errors, 0 hazards), while mismatch still reproduces.
Question
Given this is now isolated as an on-device mismatch in a minimal deterministic copy path, is there any known driver/runtime issue here, or should this be treated as likely hardware instability/defect first?