In our line of work, performing low-level performance optimization, the compiler is less of a language abstraction and more of a tool to directly interact with. Compilers are complex machines. Correctly converting C++-like code to assembly instructions is a gigantic challenge, producing optimized instructions even more so. It is therefore no surprise that the compiler itself may contain bugs. When interacting with the compiler at a low level you tend to run into such bugs.
In this article we’ll have a look what to do when running into a compiler bug: when to start suspecting the compiler, how to confirm your suspicions, and how to report the bug.
The Application
For the article, we’ll consider a situation in CUDA using the NVCC compiler. The application has threads that processes data one byte at a time, at different offsets from a large input array. A naive implementation has the threads load the bytes one at a time.
__global__ void kernel(uint8_t* data, int* offsets)
{
int index = offsets[threadIdx.x];
uint32_t buffer = 0; // 4-byte bit buffer in register
int bits_remaining = 0;
bool is_done = false;
while (!is_done) {
// load bits, assuming `process` needs at least four bits
if (bits_remaining < 4) {
uint8_t val = data[index++];
buffer |= (uint32_t{val} << (24 - bits_remaining));
bits_remaining += 8;
}
process(is_done, buffer, bits_remaining);
}
}
// placeholder function, consumes bits and updates the state variables
__device__ void process(bool& is_done, uint32_t& buffer, int& bits_remaining);Those familiar with performance-oriented programming will immediately spot that the above is inefficient. Depending on how the offsets are distributed, this pattern does not make efficient use of CUDA’s memory access model. If offsets are nicely aligned CUDA can use coalescing to take care of this issue, but that is not the case for this application. Instead, the chosen solution will use the shared memory to assign a “private cache” to each thread, and load a big block of memory at once. Using the assumption that the offsets are aligned to 16 bytes we can write:
template<int block_size>
__global__ void kernel(uint8_t* data, int* offsets)
{
int index = offsets[threadIdx.x];
assert(blockDim.x == block_size);
__shared__ ulonglong4 block_buffer[block_size];
ulonglong4& buffer = block_buffer[threadIdx.x]; // per-thread 32-byte bit buffer in shared memory
int bits_remaining = 0;
int sub_idx = 0; // indicate which 128 bits of the 256-bits buffer should be used
bool is_done = false;
while (!is_done) {
if (bits_remaining < 4) {
// load 16 bytes at once
uint4 val = *reinterpret_cast<uint4*>(data + index);
// swap the bytes per word
val.x = cuda::std::byteswap(val.x);
val.y = cuda::std::byteswap(val.y);
val.z = cuda::std::byteswap(val.z);
val.w = cuda::std::byteswap(val.w);
reinterpret_cast<uint4*>(buffer)[sub_idx] = val;
bits_remaining += 128;
}
process(is_done, buffer, bits_remaining, sub_idx);
}
}
// placeholder function, consumes bits and updates the state variables
__device__ void process(bool& is_done, ulonglong4& buffer, int& bits_remaining, int& sub_idx);This optimization works because CUDA accesses memory with 32-byte memory transactions. The limit for an individual load or store instruction is 16 bytes. Every thread makes its own large load, making sure only a minimal part of the memory transaction is wasted. Since the bytes are now not loaded individually but multiple at a time, an additional computation is needed. NVIDIA GPUs are little-endian architectures, so the loaded 4-byte integer is represented in a different order in memory than the individual bytes. A “byteswap” must be performed to change the order, for example converting 0xefbeadde to 0xdeadbeef. Note that the above code is not an exact representation of the code resulting in the bad behavior, just an illustration to explain the application.
The Investigation
After moving from a naive implementation to the optimized implementation, we start to notice test failures. The tests behave as if the data was wrongly loaded. Always the first thing to do is to double check: is the code correct? Is the buffer-updating code functioning? Have I not introduced any undefined behavior? The type casts here break the C++ strict aliasing rule. We are dereferencing a cast of a variable of type uint8_t to a uint4. However, relying on this behavior is well-documented and even encouraged when authoring CUDA kernels, as long as alignment requirements are met.
All in all, the code looks correct. Next thing to do is to debug. While CUDA offers a fancy kernel debugger, we’ll save that for another time. The simplest thing to do is to print the results from the kernel directly. We initialize the data array with bytes of incrementing values in order to distinguish them. Below are the results from the first thread starting from an offset of zero, after loading the vector types and swapping the bytes.
0x00010203
0x04050607
0x04050607 // expected 0x08090a0b
0x0c0d0e0fFrom this result it is immediately clear that the wrong data is printed. The third 4-byte value is the same as the second one. The expected value is 0x08090a0b, continuing the pattern. This is the exact moment the compiler itself becomes the topic of investigation.
One handy way to confirm such suspicions is to turn off optimization. With NVCC optimization is enabled by default and can be fully turned off with --device-debug (-G). Adding this flag when compiling results in the correct values being printed when running the application. While this method is not fool-proof as the compiler bug may not be in the optimization process, it is a simple way to confirm suspicions. There is also a caveat: if your code did contain undefined behavior to start with, there may also be difference in the results between optimized and non-optimized builds.
If we know with certainty that the CUDA source code is correct, to confirm the compiler is at fault we can have a look at the generated assembly. It is advisable to first reduce the code to a minimum working example, because assembly instructions aren’t exactly human-readable. Reducing your code to a minimal example usually is the hard part in narrowing down a compiler bug. Remove too much, and the compiler bug stops exhibiting or the intention of the code is lost. Remove too little, and there is too much assembly to go through. To instruct the compiler to output assembly, pass the --save-temps option, which instructs the compiler to store the results of the compilation phases to the current directory. While we use NVCC as our example here, other compilers like Clang and GCC also offer the option.
The minimal reproducing code is provided via Compiler Explorer. It is slightly different from the code above, which is only used to explain the intention of the code. The linked Compiler Explorer code has the actual minimal reproducer derived from the original application. CUDA compiles to PTX and SASS. PTX is a portable intermediate representation and SASS is the architecture-specific machine code. When saving the outputs of the compilation phases, the PTX is stored in .ptx files and the SASS must be extracted from the .cubin file, for example using nvdisasm. PTX is usually easier to look at because it is better documented. Below is the relevant PTX section from the reproducer.
.visible .entry kernel(unsigned char*)(
.param .u64 kernel(unsigned char*)_param_0
)
{
ld.param.u64 %rd1, [kernel(unsigned char*)_param_0];
cvta.to.global.u64 %rd2, %rd1;
mov.u32 %r1, %tid.x;
shl.b32 %r2, %r1, 5;
mov.u32 %r3, kernel(unsigned char*)::buffer;
add.s32 %r4, %r3, %r2;
ld.global.nc.v4.u32 {%r5, %r6, %r7, %r8}, [%rd2];
mov.u32 %r12, 291;
mov.u32 %r13, 0;
prmt.b32 %r14, %r8, %r13, %r12;
prmt.b32 %r15, %r6, %r13, %r12;
prmt.b32 %r16, %r5, %r13, %r12;
st.shared.v4.u32 [%r4], {%r16, %r15, %r15, %r14};
ret;
}Now, reading the PTX does take some getting used to though the mnemonics help. The meaning of most instructions can be guessed, such as mov (move), add (addition), and shl (shift left). One instruction that jumps out here is prmt.b32, for which we can look up the PTX documentation. The instruction is a byte permute that combines two 32-bit registers into an 8-byte value and produces a single 4-byte value based on a control register. In prmt.b32 %r14, %r8, %r13, %r12 register %r14 is the result, %r8 is one of the inputs, %r13 is the other (though it is always zero in the example), and %r12 is the control register. The control register determines which bytes get selected from the two input registers. The value of 291 (or more easily read as 0x0123) selects the bytes from %r8 from low to high, resulting in a byte swap. Hence, this permute instruction must be what the cuda::std::byteswap gets lowered into. The critical observation now is that our code uses the byteswaps in groups of four, and the generated PTX only has three. In fact, this seems to be the underlying problem. The instruction st.shared.v4.u32 stores the four byte-permuted integers to shared memory, but uses only three unique registers (%r15 is repeated).
While this is conclusive evidence that the compiler is indeed generating the wrong instructions, this is about as far we can take our investigation. If NVCC were open source or gave finer-grained control over optimization passes we could dig further. Unfortunately that is not the case, and our investigation ends here.
The Conclusion
Now that we have confirmed that the problem lies with the compiler and we’ve hunted down the conditions that cause it to appear, it is time to inform the compiler authors. This is done by filing a bug report with our minimal reproducer. Unfortunately for NVCC such reports are not publicly visible. After having filed the report, the bug was confirmed and their engineers indicated it’ll be fixed with CUDA 13.1. The final step is to confirm: open the linked Compiler Explorer instance, select the 13.1 compiler, and we can observe the generated PTX is correct.
In this article I hope to have given practical tips on how to debug compiler bugs: when to start suspecting the compiler, how to confirm these suspicions, and finally how to create an actionable bug report.
