Higher Bug Detection: How Compile-Time Instrumentation for Compute Sanitizer Enhances Memory Safety

Let’s dive in and use Compute Sanitizer to uncover a bug in the next program, an adaptation of an example from a previous blog post on Compute Sanitizer. Can you notice the logical error within the code below? Hint: It’s the classic off-by-one error. We’ve all been there, more times than we’d care to confess. In our code example, we not only launch scaleArray with too many threads, but we also use the mistaken condition in scaleArray’s if-condition. The result’s that thread 512 will access array[512], which is out of bounds.  

#include 
#include   

__global__ void scaleArray(float* array, size_t N, float value) {
  int threadGlobalID    = threadIdx.x + blockIdx.x * blockDim.x;
  if (threadGlobalID <= N) {
    array[threadGlobalID] = array[threadGlobalID]*value;
  }
}

int important() {
  float* array = nullptr;
  float* buffer = nullptr;
  const size_t N = 512;

  // Allocate N float-size elements, visible to each CPU and GPU
  cudaMallocManaged(&array, N*sizeof(float));  
  //cudaMallocManaged(&buffer, N*sizeof(float)); 

  for (int i=0; i>>(array, N, 3.0);
  cudaDeviceSynchronize();

  printf("After : Array 0, 1 .. N-1: %f %f %fn", array[0], array[1], array[N-1]);
  assert(array[N/2] == 3.0); // Check that it has worked

  cudaFree(array);
  cudaFree(buffer);
  exit(0);
}

Let’s see if this bug is observable after we compile and run this system:

$ /usr/local/cuda-13.1/bin/nvcc -o example1 -arch sm_86 example1.cu
$ ./example1
Before: Array 0, 1 .. N-1: 1.000000 1.000000 1.000000
After : Array 0, 1 .. N-1: 3.000000 3.000000 3.000000

On our test machine, this program runs without issue. Bugs just like the one in our example can easily go unnoticed during development, but manifest later as data corruption, segmentation faults, and even code exploits.

Now, let’s run this system through Compute Sanitizer to see if it uncovers the problem.

$ compute-sanitizer ./example1

========= COMPUTE-SANITIZER
Before: Array 0, 1 .. N-1: 1.000000 1.000000 1.000000
========= Invalid __global__ read of size 4 bytes
=========     at scaleArray(float *, unsigned long, float)+0xb0
=========     by thread (512,0,0) in block (0,0,0)
=========     Access to 0x7f7754000800 is out of bounds
=========     and is 1 bytes after the closest allocation at 0x7f7754000000 of size 2,048 bytes
=========     Saved host backtrace as much as driver entry point at kernel launch time
=========         Host Frame: important [0x8deb] in example1
========= 
========= Program hit cudaErrorLaunchFailure (error 719) resulting from "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
=========     Saved host backtrace as much as driver entry point at error
=========         Host Frame: important [0x8df0] in example1
========= 
After : Array 0, 1 .. N-1: 1.000000 1.000000 1.000000
example1: example1.cu:27: int important(): Assertion `array[N/2] == 3.0' failed.
========= Error: process didn't terminate successfully
========= Goal application returned an error
========= ERROR SUMMARY: 2 errors

Hurray, it does! Nonetheless, Compute Sanitizer, like all commercially available memory sanitizers, can have false negatives. That’s, it could miss some legitimate memory safety errors. Let’s change the instance above by uncommenting the second call to cudaMallocManaged. On our test machine, after recompiling and running this system through compute-sanitizer, we now get this message:

$ compute-sanitizer ./example1

========= COMPUTE-SANITIZER
Before: Array 0, 1 .. N-1: 1.000000 1.000000 1.000000
After : Array 0, 1 .. N-1: 3.000000 3.000000 3.000000
========= ERROR SUMMARY: 0 errors

No errors? What happened? The buffer overrun still exists, however the tool not reports the error. How did that minor change mask the buffer overrun?

The short answer is that Compute Sanitizer’s memcheck tool ensures that memory accesses go to valid memory address ranges, i.e., addresses which were allocated and never freed. During our test, the CUDA runtime’s memory allocator placed array and buffer contiguously in memory. That’s, buffer immediately follows array, and thus an overrun from array to buffer still accesses valid memory. Note that this shortcoming also applies to GCC’s and Clang’s “address sanitizer” tools.   

Let’s understand why Compute Sanitizer operates the way in which it does. Until now, Compute Sanitizer has relied solely on binary instrumentation to inject instrumentation code around memory and synchronization instructions in an application, allowing the tool to find unsound behavior at run time. We’re simplifying the operation, but we will consider the tool inserting a elaborate assertion before each memory reference that checks whether the address is inside a currently-allocated object:

__global__ void scaleArray(float* array, size_t N, float value) {
  int threadGlobalID    = threadIdx.x + blockIdx.x * blockDim.x;
  if (threadGlobalID <= N) {
    assert(isAllocated(&array[threadGlobalID])); // Inserted instrumentation.
    float tmp = array[threadGlobalID]*value;
    assert(isAllocated(&array[threadGlobalID])); // Inserted instrumentation.
    array[threadGlobalID] = tmp;
  }
}

After Compute Sanitizer inserts those assertions, it runs the applying to see whether any of them fire at run time.

Compute Sanitizer didn’t have the luxurious of statically analyzing this system, partly since it only operated on raw binaries, that are notoriously difficult to investigate. Without evaluation, it could only instrument each memory reference independently, in a peephole manner. In our running example, despite the fact that some threads overflow from array to buffer, peephole instrumentation misses the illegal references because buffer is allocated. Subsequently, ‘isAllocated(&array[threadGlobalID]’ is true, and the assertion doesn’t trigger. 

There may be a Compute Sanitizer option, ‘--padding’, that adds invalid memory regions between allocations, allowing Compute Sanitizer to detect the overrun for this instance. But that option is just not on by default because it may possibly significantly increase memory usage. Moreover, the choice cannot pad all memory allocations, including shared memory and global variable allocations. 

We are able to do higher with compile-time instrumentation.

Improving coverage with a compiler evaluation

nvcc -o example1 -arch sm_86 -fdevice-sanitize=memcheck example1.cu

compute-sanitizer ./example1
========= COMPUTE-SANITIZER
Before: Array 0, 1 .. N-1: 1.000000 1.000000 1.000000
========= Invalid __global__ read of size 4 bytes
=========     at scaleArray(float *, unsigned long, float)+0x150
=========     by thread (512,0,0) in block (0,0,0)
=========     Access to 0x7b3566000800 is out of bounds
=========     and is contained in the nearest allocation at 0x7b3566000800 of size 2,048 bytes
=========     Saved host backtrace as much as driver entry point at kernel launch time
=========         Host Frame: important [0x8e01] in example1
========= 
========= Program hit cudaErrorLaunchFailure (error 719) resulting from "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
=========     Saved host backtrace as much as driver entry point at error
=========         Host Frame: important [0x8e06] in example1
========= 
After : Array 0, 1 .. N-1: 1.000000 1.000000 1.000000
example1: example1.cu:27: int important(): Assertion `array[N/2] == 3.0' failed.
========= Error: process didn't terminate successfully
========= Goal application returned an error
========= ERROR SUMMARY: 2 errors