debug-cuda-crash — Отладка сбоев CUDA для FlashInfer

--- name: debug-cuda-crash description: Tutorial for debugging CUDA crashes using API logging --- # 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]=%fn", 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 processingn", group_id); } } ``` **Common mistake** ❌: ```cpp // ❌ Only warp 0 will print! if (threadIdx.x == 0) { printf("Warp %d processingn", 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) | | `FLASHINFER_LOGDEST` | `stdout` | Log to console (default) | | | `stderr` | Log to stderr | | | `<path>` | Log to file | | | `log_%i.txt` | Multi-process: %i = process ID | ## Best Practices ### 1. Always Start with Level 3 ```bash export FLASHINFER_LOGLEVEL=3 ``` Level 3 provides tensor metadata (shape, dtype, device) without overwhelming output. ### 2. Use Level 5 for Numerical Issues ```bash export FLASHINFER_LOGLEVEL=5 ``` Only use level 5 when debugging NaN/Inf problems (adds statistics). ### 3. Log to File for Crashes ```bash export FLASHINFER_LOGDEST=crash_log.txt ``` Console output may be lost when program crashes. File logs persist. ### 4. Compare Before/After Enable logging and compare: - Last successful API call (inputs logged, outputs logged) ✅ - First failed API call (inputs logged, no outputs) ❌ ← This is where it crashed! ### 5. Disable Logging in Production ```bash unset FLASHINFER_LOGLEVEL # or export FLASHINFER_LOGLEVEL=0 ``` Logging has zero overhead when disabled (decorator returns original function). ## Troubleshooting ### No Logs Appearing **Problem**: Set `FLASHINFER_LOGLEVEL=3` but no logs appear **Solutions**: 1. **Check if API has the decorator**: Not all FlashInfer APIs have `@flashinfer_api` yet (work in progress) 2. **Verify environment variable**: ```bash echo $FLASHINFER_LOGLEVEL # Should print "3" ``` 3. **Check log destination**: ```bash echo $FLASHINFER_LOGDEST # Should print path or "stdout" ``` ### Too Much Output **Problem**: Level 5 produces too much output **Solution**: Use level 3 instead: ```bash export FLASHINFER_LOGLEVEL=3 # Skip tensor statistics ``` ### Statistics Skipped in CUDA Graph **Warning**: `[statistics skipped: CUDA graph capture in progress]` **What it means**: Level 5 statistics are automatically skipped during CUDA graph capture (to avoid synchronization) **This is normal**: The framework protects you from graph capture issues. ## Quick Examples ### Debug Shape Mismatch ```bash export FLASHINFER_LOGLEVEL=3 export FLASHINFER_LOGDEST=stdout python my_script.py # Check tensor shapes in output ``` ### Debug NaN/Inf ```bash export FLASHINFER_LOGLEVEL=5 # Show statistics export FLASHINFER_LOGDEST=debug.log python my_script.py # Check nan_count and inf_count in debug.log ``` ### Debug Multi-GPU Training ```bash export FLASHINFER_LOGLEVEL=3 export FLASHINFER_LOGDEST=rank_%i.log # Separate log per rank torchrun --nproc_per_node=8 train.py # Check rank_*.log files ``` ### Combine with Memory Checker ```bash export FLASHINFER_LOGLEVEL=3 export FLASHINFER_LOGDEST=inputs.log compute-sanitizer --tool memcheck python my_script.py # inputs.log shows what data caused the memory error ``` ## Example: Full Debug Session ### Your code crashes: ```python import torch from flashinfer import batch_decode_with_padded_kv_cache q = torch.randn(32, 8, 128, dtype=torch.bfloat16, device="cuda") kv = torch.randn(1024, 2, 8, 64, dtype=torch.bfloat16, device="cuda") # Wrong dim! output = batch_decode_with_padded_kv_cache(q, kv) # ❌ Crashes ``` ### Enable logging: ```bash export FLASHINFER_LOGLEVEL=3 export FLASHINFER_LOGDEST=debug.log python test.py ``` ### Check debug.log: ``` [2025-12-18 10:45:23] FlashInfer API Call: batch_decode_with_padded_kv_cache Positional input arguments: arg[0]: Tensor( shape=(32, 8, 128) dtype=torch.bfloat16 ... ) arg[1]: Tensor( shape=(1024, 2, 8, 64) # ❌ Found it! Last dim should be 128 dtype=torch.bfloat16 ... ) ``` ### Fix the bug: ```python kv = torch.randn(1024, 2, 8, 128, dtype=torch.bfloat16, device="cuda") # ✅ Fixed ``` ### Success! ```bash python test.py # No crash, outputs logged successfully ``` ## Summary 1. **Enable logging** before the crash: ```bash export FLASHINFER_LOGLEVEL=3 export FLASHINFER_LOGDEST=debug.log ``` 2. **Run your code** - inputs are logged BEFORE crash 3. **Check the log** - last API call shows what caused the crash 4. **Fix the issue** based on logged input metadata 5. **Disable logging** when done: ```bash export FLASHINFER_LOGLEVEL=0 ``` ## Related Documentation - See CLAUDE.md "API Logging with @flashinfer_api" for technical details - See `flashinfer/api_logging.py` for implementation - See CUDA documentation for compute-sanitizer and cuda-gdb