debug-cuda-crash

# Tutorial: Debugging CUDA Crashes with API Logging This tutorial shows you how to debug CUDA crashes and errors in FlashInfer using the `@flashinfer_api` logging decorator. ## Goal When your code crashes with CUDA errors (illegal memory access, out-of-bounds, NaN/Inf), use API logging to: - Capture input tensors BEFORE the crash occurs - Understand what data caused the problem - Track tensor shapes, dtypes, and values through your pipeline - Detect numerical issues (NaN, Inf, wrong shapes) ## Why Use API Logging? **Problem**: CUDA errors often crash the program, leaving no debugging information. **Solution**: FlashInfer's `@flashinfer_api` decorator logs inputs BEFORE execution, so you can see what caused the crash even after the program terminates. ## Step 1: Enable API Logging ### Basic Logging (Function Names Only) ```bash export FLASHINFER_LOGLEVEL=1 # Log function names export FLASHINFER_LOGDEST=stdout # Log to console python my_script.py ``` Output: ``` [2025-12-18 10:30:45] FlashInfer API Call: batch_decode_with_padded_kv_cache ``` ### Detailed Logging (Inputs/Outputs with Metadata) ```bash export FLASHINFER_LOGLEVEL=3 # Log inputs/outputs with metadata export FLASHINFER_LOGDEST=debug.log # Save to file python my_script.py ``` Output in `debug.log`: ``` ================================================================================ [2025-12-18 10:30:45] FlashInfer API Logging - System Information ================================================================================ FlashInfer version: 0.6.0 CUDA toolkit version: 12.1 GPU 0: NVIDIA H100 PCIe Compute capability: 9.0 (SM90) PyTorch version: 2.1.0 ================================================================================ ================================================================================ [2025-12-18 10:30:46] FlashInfer API Call: batch_decode_with_padded_kv_cache -------------------------------------------------------------------------------- Positional input arguments: arg[0]: Tensor( shape=(32, 8, 128) dtype=torch.bfloat16 device=cuda:0 requires_grad=False is_contiguous=True ) Keyword input arguments: kv_cache= Tensor( shape=(1024, 2, 8, 128) dtype=torch.bfloat16 device=cuda:0 requires_grad=False is_contiguous=True ) ``` ### Full Logging (With Tensor Statistics) ```bash export FLASHINFER_LOGLEVEL=5 # Log with min/max/mean/nan/inf export FLASHINFER_LOGDEST=debug.log python my_script.py ``` Additional output: ``` Tensor( shape=(32, 8, 128) dtype=torch.bfloat16 device=cuda:0 requires_grad=False is_contiguous=True min=-3.125000 max=4.250000 mean=0.015625 nan_count=0 inf_count=0 ) ``` ## Step 2: Reproduce the Crash ### Example: Shape Mismatch Your code crashes with: ``` RuntimeError: CUDA error: an illegal memory access was encountered ``` Enable logging and run again: ```bash export FLASHINFER_LOGLEVEL=3 export FLASHINFER_LOGDEST=crash_log.txt python my_script.py ``` The log shows inputs before the crash: ``` [2025-12-18 10:32:15] FlashInfer API Call: batch_decode_with_padded_kv_cache Positional input arguments: arg[0]: Tensor( shape=(32, 8, 128) # Query tensor ... ) Keyword input arguments: kv_cache= Tensor( shape=(1024, 2, 8, 64) # ❌ Wrong! Should be (..., 128) not (..., 64) ... ) ``` **Found the bug**: `head_dim` mismatch (64 vs 128) ## Step 3: Common CUDA Errors and How to Debug ### Error 1: Illegal Memory Access **Error Message**: ``` RuntimeError: CUDA error: an illegal memory access was encountered ``` **Enable logging**: ```bash export FLASHINFER_LOGLEVEL=3 python my_script.py ``` **What to check in logs**: - ✅ Tensor shapes match expected dimensions - ✅ All tensors are on CUDA (not CPU) - ✅ Tensor strides are reasonable - ✅ `is_contiguous=True` (if required) **Common causes**: - Wrong tensor dimensions - CPU tensor passed to GPU kernel - Incorrect stride patterns ### Error 2: NaN or Inf Values **Error Message**: ``` RuntimeError: Function ... returned nan or inf ``` **Enable statistics logging**: ```bash export FLASHINFER_LOGLEVEL=5 # Level 5 shows nan_count, inf_count python my_script.py ``` **What to check in logs**: ``` Tensor( ... min=-1234567.000000 # ❌ Suspiciously large max=9876543.000000 # ❌ Suspiciously large mean=nan # ❌ NaN detected nan_count=128 # ❌ 128 NaN values! inf_count=0 ) ``` **Common causes**: - Division by zero in previous operation - Numerical overflow/underflow - Uninitialized memory ### Error 3: Out of Memory **Error Message**: ``` RuntimeError: CUDA out of memory ``` **Enable logging**: ```bash export FLASHINFER_LOGLEVEL=3 python my_script.py ``` **What to check in logs**: - ✅ Tensor shapes (are they unexpectedly large?) - ✅ Batch size - ✅ Sequence length Example: ``` Tensor( shape=(1024, 8192, 128, 128) # ❌ Way too large! Should be (1024, 128, 128)? ... ) ``` ### Error 4: Wrong Dtype **Error Message**: ``` RuntimeError: expected scalar type BFloat16 but found Float16 ``` **Enable logging**: ```bash export FLASHINFER_LOGLEVEL=3 python my_script.py ``` **What to check in logs**: ``` Tensor( dtype=torch.float16 # ❌ Should be torch.bfloat16 ... ) ``` ## Step 4: Multi-Process Debugging When running with multiple GPUs/processes, use `%i` pattern: ```bash export FLASHINFER_LOGLEVEL=3 export FLASHINFER_LOGDEST=debug_rank_%i.txt # %i = process ID torchrun --nproc_per_node=4 my_script.py ``` This creates separate logs: - `debug_rank_12345.txt` (process 12345) - `debug_rank_12346.txt` (process 12346) - `debug_rank_12347.txt` (process 12347) - `debug_rank_12348.txt` (process 12348) Now you can debug each rank independently. ## Step 5: Advanced Debugging with compute-sanitizer For harder bugs, combine API logging with CUDA tools: ### Use compute-sanitizer (Memory Checker) ```bash export FLASHINFER_LOGLEVEL=3 export FLASHINFER_LOGDEST=debug.log compute-sanitizer --tool memcheck python my_script.py ``` Output shows exact memory errors: ``` ========= COMPUTE-SANITIZER ========= Invalid __global__ write of size 4 bytes ========= at 0x1234 in ScaleKernel<float> ========= by thread (256,0,0) in block (10,0,0) ========= Address 0x7f1234567890 is out of bounds ``` Check `debug.log` to see what inputs caused this kernel to fail. ### Use cuda-gdb (Debugger) ```bash export FLASHINFER_LOGLEVEL=3 export FLASHINFER_LOGDEST=debug.log cuda-gdb --args python my_script.py ``` In gdb: ``` (cuda-gdb) run (cuda-gdb) where # Show stack trace when it crashes ``` Check `debug.log` for the inputs that led to the crash. ## Step 6: Kernel-Level Debugging with printf() You can use `printf()` inside CUDA kernels for debugging: ### Basic Usage ```cpp __global__ void MyKernel(const float* input, float* output, int n) { int idx = blockIdx.x * blockDim.x + threadIdx.x; // Print from one thread to avoid spam if (threadIdx.x == 0 && blockIdx.x == 0) { printf("n=%d, input[0]=%f\n", n, input[0]); } if (idx < n) { output[idx] = input[idx] * 2.0f; } } ``` **Important**: Flush printf buffer after kernel: ```python my_kernel(input, output) torch.cuda.synchronize() # ← Flushes printf output ``` ### ⚠️ Warp-Specialized Kernels: Choosing the Right Print Thread **Problem**: `threadIdx.x == 0` doesn't work for all warps (warp starting at thread 32 won't have thread 0). **Solution**: Choose one representative thread per specialization group. ```cpp __global__ void WarpSpecializedKernel(...) { // Define your group's representative thread // e.g., first thread of each warp: threadIdx.x % 32 == 0 // e.g., first thread of each 4-warp group: threadIdx.x % 128 == 0 if (is_group_representative) { printf("Group %d processing\n", group_id); } } ``` **Common mistake** ❌: ```cpp // ❌ Only warp 0 will print! if (threadIdx.x == 0) { printf("Warp %d processing\n", threadIdx.x / 32); } ``` ### Quick Reference | Kernel Type | Print Condition | Notes | |-------------|-----------------|-------| | Simple kernel | `threadIdx.x == 0` | One thread per block | | Warp-specialized | One thread per group | Depends on kernel design | ### Other Kernel Debugging Tools ```cpp // Assert for invariants assert(value >= 0.0f && "Value must be non-negative"); // Compile-time checks static_assert(BLOCK_SIZE % 32 == 0, "BLOCK_SIZE must be multiple of warp size"); ``` ## Environment Variables Reference | Variable | Values | Description | |----------|--------|-------------| | `FLASHINFER_LOGLEVEL` | `0` | No logging (default) | | | `1` | Function names only | | | `3` | Inputs/outputs with metadata | | | `5` | + Tensor statistics (min/max/mean/nan/inf) |