Hi,
In the debugging environment, compute-sanitizer can effectively capture GPU memory out-of-bounds issues. However, it is difficult to apply compute-sanitizer in actual production environments because it severely impacts performance.
I am trying to find an alternative tool that can capture information when CUDA memory out-of-bounds occurs. But when I attempted to capture core dumps, I discovered that a core dump is only generated if the memory out-of-bounds access goes far beyond the allocated range.
In my environment (Windows, RTX 3060, CUDA 12.9), I allocated 1 byte of memory and tried writing to addresses with larger offsets. My actual tests showed that a core dump is only generated when the write offset exceeds approximately 80MB relative to the start of the allocated memory address.
My questions are:
Why is the threshold around 80MB?
Is there an effective method to automatically detect GPU memory out-of-bounds with minimal performance impact?
Thanks for your reply.
MyDemo：
#include <cuda_runtime.h> #include <stdio.h> #include <windows.h> __global__ void OutOfBoundsKernel(unsigned char* ptr, size_t offset) { unsigned long long write_addr = (unsigned long long)(ptr + offset); unsigned long long base_addr = (unsigned long long)ptr; if (threadIdx.x == 0 && blockIdx.x == 0) { printf("[Kernel] Base address: 0x%016llx\n", base_addr); printf("[Kernel] Write offset: %zu bytes (%.2f MB)\n", offset, offset / (1024.0 * 1024.0)); printf("[Kernel] Write address: 0x%016llx\n", write_addr); } // Out-of-bounds write ptr[offset] = 0xFF; } int main() { cudaError_t err; unsigned char* d_ptr = NULL; printf("========================================\n"); printf("GPU Memory Protection Range Test\n"); printf("========================================\n\n"); // 1. Get device information int deviceCount; cudaGetDeviceCount(&deviceCount); if (deviceCount == 0) { printf("No CUDA device found!\n"); return -1; } cudaDeviceProp prop; cudaGetDeviceProperties(&prop, 0); printf("Device: %s\n", prop.name); printf("Total Memory: %.2f GB\n", prop.totalGlobalMem / (1024.0 * 1024.0 * 1024.0)); // 2. Allocate memory (only 1 byte) size_t alloc_size = 1; err = cudaMalloc(&d_ptr, alloc_size); if (err != cudaSuccess) { printf("cudaMalloc failed: %s\n", cudaGetErrorString(err)); return -1; } printf("\nAllocated %zu bytes\n", alloc_size); printf("Allocated address: 0x%016llx\n", (unsigned long long)d_ptr); // 3. Get memory range info cudaPointerAttributes attrs; cudaPointerGetAttributes(&attrs, d_ptr); printf("Memory type: %d\n", attrs.type); // 4. Test different offsets size_t test_offsets[] = { 1 * 1024 * 1024, // 1MB 2 * 1024 * 1024, // 2MB 4 * 1024 * 1024, // 4MB 8 * 1024 * 1024, // 8MB 16 * 1024 * 1024, // 16MB 32 * 1024 * 1024, // 32MB 64 * 1024 * 1024, // 64MB 79 * 1024 * 1024 - 1, // 79MB 80 * 1024 * 1024, // 80MB 100 * 1024 * 1024, // 100MB }; printf("\n========================================\n"); printf("Testing Different Offsets\n"); printf("========================================\n"); for (int i = 0; i < sizeof(test_offsets) / sizeof(test_offsets[0]); i++) { size_t offset = test_offsets[i]; unsigned long long write_addr = (unsigned long long)d_ptr + offset; printf("\n[Test %d] Offset %zu bytes (%.2f MB) -> Address 0x%016llx\n", i + 1, offset, offset / (1024.0 * 1024.0), write_addr); // Launch kernel OutOfBoundsKernel << <1, 1 >> > (d_ptr, offset); // Synchronize err = cudaDeviceSynchronize(); if (err != cudaSuccess) { printf(" CRASHED! %s (code=%d)\n", cudaGetErrorString(err), err); // Reallocate memory after crash cudaFree(d_ptr); err = cudaMalloc(&d_ptr, alloc_size); if (err != cudaSuccess) { printf("Failed to reallocate memory, test terminated\n"); break; } printf(" Reallocated address: 0x%016llx\n", (unsigned long long)d_ptr); } else { printf(" SUCCESS\n"); } Sleep(500); } return 0; }
Result：
GPU Memory Protection Range Test
Device: NVIDIA GeForce RTX 3060
Total Memory: 12.00 GBAllocated 1 bytes
Allocated address: 0x0000000b06200000
Memory type: 2========================================
Testing Different Offsets
[Test 1] Offset 1048576 bytes (1.00 MB) → Address 0x0000000b06300000
[Kernel] Base address: 0x0000000b06200000
[Kernel] Write offset: 1048576 bytes (0.00 MB)
[Kernel] Write address: 0x0000000b06300000
SUCCESS[Test 2] Offset 2097152 bytes (2.00 MB) → Address 0x0000000b06400000
[Kernel] Base address: 0x0000000b06200000
[Kernel] Write offset: 2097152 bytes (0.00 MB)
[Kernel] Write address: 0x0000000b06400000
SUCCESS[Test 3] Offset 4194304 bytes (4.00 MB) → Address 0x0000000b06600000
Kernel] Base address: 0x0000000b06200000
[Kernel] Write offset: 4194304 bytes (0.00 MB)
[Kernel] Write address: 0x0000000b06600000
SUCCESS[Test 4] Offset 8388608 bytes (8.00 MB) → Address 0x0000000b06a00000
Kernel] Base address: 0x0000000b06200000
[Kernel] Write offset: 8388608 bytes (0.00 MB)
[Kernel] Write address: 0x0000000b06a00000
SUCCESS[Test 5] Offset 16777216 bytes (16.00 MB) → Address 0x0000000b07200000
Kernel] Base address: 0x0000000b06200000
[Kernel] Write offset: 16777216 bytes (0.00 MB)
[Kernel] Write address: 0x0000000b07200000
SUCCESS[Test 6] Offset 33554432 bytes (32.00 MB) → Address 0x0000000b08200000
Kernel] Base address: 0x0000000b06200000
[Kernel] Write offset: 33554432 bytes (0.00 MB)
[Kernel] Write address: 0x0000000b08200000
SUCCESS[Test 7] Offset 67108864 bytes (64.00 MB) → Address 0x0000000b0a200000
Kernel] Base address: 0x0000000b06200000
[Kernel] Write offset: 67108864 bytes (0.00 MB)
[Kernel] Write address: 0x0000000b0a200000
SUCCESS[Test 8] Offset 82837503 bytes (79.00 MB) → Address 0x0000000b0b0fffff
Kernel] Base address: 0x0000000b06200000
[Kernel] Write offset: 82837503 bytes (0.00 MB)
[Kernel] Write address: 0x0000000b0b0fffff
SUCCESS[Test 9] Offset 83886080 bytes (80.00 MB) → Address 0x0000000b0b200000
[13:47:48.876775] coredump: Starting GPU coredump generation
[13:47:48.877101] coredump: Detected an exception of type CUDBG_EXCEPTION_WARP_ILLEGAL_ADDRESS (14)
[13:47:48.877214] coredump: - Device: 0
[13:47:48.877298] coredump: - SM: 0
[13:47:48.877407] coredump: - Warp: 0
[13:47:48.877487] coredump: - PC 0xb00e7bf90
[13:47:48.877920] coredump: Stack trace (lane masks: active 0x00000001, valid 0x00000001):
[13:47:48.878024] coredump: #0 0xb00e7bfa0 _Z17OutOfBoundsKernelPhy
[13:47:48.878922] coredump: SM 1/28 has finished state collection
[13:47:48.878981] coredump: SM 2/28 has finished state collection
[13:47:48.878998] coredump: SM 3/28 has finished state collection
[13:47:48.879014] coredump: SM 4/28 has finished state collection
[13:47:48.879030] coredump: SM 5/28 has finished state collection
[13:47:48.879046] coredump: SM 6/28 has finished state collection
[13:47:48.879062] coredump: SM 7/28 has finished state collection
[13:47:48.879080] coredump: SM 8/28 has finished state collection
[13:47:48.879097] coredump: SM 9/28 has finished state collection
[13:47:48.879115] coredump: SM 10/28 has finished state collection
[13:47:48.879132] coredump: SM 11/28 has finished state collection
[13:47:48.879148] coredump: SM 12/28 has finished state collection
[13:47:48.879164] coredump: SM 13/28 has finished state collection
[13:47:48.879181] coredump: SM 14/28 has finished state collection
[13:47:48.879199] coredump: SM 15/28 has finished state collection
[13:47:48.879217] coredump: SM 16/28 has finished state collection
[13:47:48.879235] coredump: SM 17/28 has finished state collection
[13:47:48.879252] coredump: SM 18/28 has finished state collection
[13:47:48.879270] coredump: SM 19/28 has finished state collection
[13:47:48.879288] coredump: SM 20/28 has finished state collection
[13:47:48.879305] coredump: SM 21/28 has finished state collection
[13:47:48.879323] coredump: SM 22/28 has finished state collection
[13:47:48.879340] coredump: SM 23/28 has finished state collection
[13:47:48.879357] coredump: SM 24/28 has finished state collection
[13:47:48.879375] coredump: SM 25/28 has finished state collection
[13:47:48.879393] coredump: SM 26/28 has finished state collection
[13:47:48.879410] coredump: SM 27/28 has finished state collection
[13:47:48.879428] coredump: SM 28/28 has finished state collection
[13:47:48.879450] coredump: Device 1/1 has finished state collection
[13:47:48.879471] coredump: Calculating ELF file layout
[13:47:48.879513] coredump: ELF file layout calculated
[13:47:48.879530] coredump: Writing ELF file to D:\Dumps\coreDump\core.dmp
[13:47:48.879551] coredump: Current working directory is D:\code\CompressTool\Project1
[13:47:48.880590] coredump: Writing out global memory (12582912 bytes)
[13:47:48.881128] coredump: 5%…
[13:47:48.881193] coredump: 10%…
[13:47:48.881212] coredump: 15%…
[13:47:48.933808] coredump: 20%…
[13:47:48.933912] coredump: 25%…
[13:47:48.933938] coredump: 30%…
[13:47:48.984957] coredump: 35%…
[13:47:48.985029] coredump: 40%…
[13:47:48.985076] coredump: 45%…
[13:47:48.985099] coredump: 50%…
[13:47:49.040064] coredump: 55%…
[13:47:49.040124] coredump: 60%…
[13:47:49.040146] coredump: 65%…
[13:47:49.090164] coredump: 70%…
[13:47:49.090309] coredump: 75%…
[13:47:49.090346] coredump: 80%…
[13:47:49.141109] coredump: 85%…
[13:47:49.141196] coredump: 90%…
[13:47:49.141225] coredump: 95%…
[13:47:49.141248] coredump: 100%…
[13:47:49.141367] coredump: Writing out device table
[13:47:49.142994] coredump: Writing out metadata
[13:47:49.143027] coredump: Finalizing
[13:47:49.143554] coredump: Writing done
[13:47:49.143586] coredump: All done (took 01s)