vLLM 使用 CUDA Core Dump 调试非法内存访问问题

Introduction; What is a Core Dump?; How to Enable CUDA Core Dump; Example of Using CUDA Core Dump; Debugging Improper Kernel Launch; Debugging Kernel Exceptions in CUDA Graphs; Limitations and Considerations; Conclusion; Acknowledgement

TL;DR: If you hit an illegal memory access was encountered error, you can enable CUDA core dump to debug the issue. Simply set the following environment variables and run your program again to collect the coredump file, then you can use cuda-gdb to debug the issue.

Have you ever felt you are developing cuda kernels and your tests often run into illegal memory access (IMA for short) and you have no idea how to debug? We definitely felt this pain again and again while working on vLLM, a high-performance inference engine for LLM models.

If you are one of the developers who have faced this issue, this blog is for you! We will uncover some of advanced debugging techniques we use that can help users debug complicated issues in vLLM, such as IMA.

For example, here’s an error from PyTorch:

The challenging bit here is: CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect. In our experience the python stack traces for these types of exceptions are basically always incorrect and pretty worthless. To resolve this the error message suggests adding CUDA_LAUNCH_BLOCKING=1 when running the code. However, there are still two problems:

Many people launch CUDA kernels using the kernel >> syntax without adding error checking for the kernel launch status, for example, this code. In such cases, even with CUDA_LAUNCH_BLOCKING=1, it’s still impossible to locate the faulty kernel; If the illegal memory access occurs inside a kernel within a CUDA graph, then even with CUDA_LAUNCH_BLOCKING=1, we can only see that there’s an issue when launching the CUDA graph, but still cannot pinpoint the exact kernel that failed

To accurately pinpoint this kind of problem, we need to react immediately when an illegal memory access occurs. Of course, this isn’t something users can do directly — it must be supported by the CUDA driver itself.

The CUDA core dump functionality, is exactly designed for this purpose. It allows the CUDA driver to dump the GPU state when an illegal memory access occurs, so that users can analyze the GPU state later to find out which kernel caused the issue and what the illegal memory access was.

A GPU is essentially a massively parallel processor, and many of its concepts can find counterparts in CPUs.

A core dump is a feature jointly provided by the CPU and the operating system. When a program crashes during execution, the operating system can record the program's memory data, runtime state, and other information for subsequent analysis and debugging. A program crash is a hardware-level concept. When the CPU encounters an error while executing certain instructions, it enters a trap state. At this point, the operating system takes over the program and executes the corresponding exception handling procedure (by default, this simply terminates the program, but options can be configured to generate a core dump for analysis. For example, ulimit -c 1 can enable core dump generation, and echo "core.%e.%p" > /proc/sys/kernel/core_pattern can specify the path for the core dump file).

By analogy, the core dump functionality on GPUs requires collaboration between GPU hardware and GPU drivers. When a thread on the GPU crashes during execution, the GPU hardware needs to trigger an exception and pass it to the GPU driver, which then immediately handles the exception. However, according to forum discussions, the default behavior of the GPU driver when handling exceptions is to mark the current CUDA context as unusable, rather than terminating the program.

Enabling CUDA core dump is very straightforward; you just need to set the CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1 environment variable. However, for a smoother experience, you should also set a few additional environment variables:

By default, the CUDA core dump saves the coredump file in the current directory without printing the file path. You can enable the CUDA_COREDUMP_SHOW_PROGRESS=1 environment variable to display the progress and details of the coredump procedure. Most importantly, it shows the path of the coredump file after the procedure is complete, making it easier for subsequent debugging and analysis; Many tasks run inside containers, and when a task fails, the container is destroyed, making it impossible to retain the coredump file. In such cases, you can use the CUDA_COREDUMP_FILE environment variable to specify a file path template for the coredump file. For example, you can store the coredump file in a persistent storage directory: CUDA_COREDUMP_FILE="/persistent_dir/cuda_coredump_%h.%p.%t", where %h is the hostname, %p is the process ID, and %t is the timestamp of the coredump; By default, the coredump procedure saves the entire GPU context. For programs like large model inference that occupy almost all GPU memory, a full coredump is impractical (hundreds of GiB of data). You can use the CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory' environment variable to skip saving GPU memory, shared memory, and local memory, thereby reducing the size of the coredump file. The skip_constbank_memory flag is missing in the documentation, but it is actually supported by the CUDA core dump feature, and would be necessary sometimes when we have many GPU threads hitting errors at the same time

The documentation also mentions that adding skip_abort to CUDA_COREDUMP_GENERATION_FLAGS prevents the CPU process from aborting after the coredump is complete. This allows the CPU process to add its own error trace, providing more debugging information. However, experiments have shown that this feature has a significant bug, which may cause illegal memory access errors on the GPU to be ignored. In such cases, subsequent code may continue to run normally, but the program's memory data might already be corrupted. This is unacceptable for training tasks and undesirable for inference tasks. Therefore, this feature is generally unreliable and not recommended.

Additionally, the documentation states that enabling CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1 not only enables CUDA core dump but also generates a CPU coredump by default. However, in practice, we find that the CPU coredump contains little useful information and is difficult to analyze.

If you want live data for debugging, you can also enable CUDA_DEVICE_WAITS_ON_EXCEPTION=1 environment variable, which does not use CUDA core dump, but stops GPU execution immediately when an exception occurs, and hangs there, waiting for users to attach a debugger (like cuda-gdb) to inspect the GPU state, where the full GPU memory is still intact. However, this approach is less automatic and requires more manual intervention.

In summary, when using the CUDA core dump feature, it is recommended to use the following combination of environment variables:

CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1 CUDA_COREDUMP_SHOW_PROGRESS=1 CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory' CUDA_COREDUMP_FILE="/persistent_dir/cuda_coredump_%h.%p.%t"

Let's use some code to verify the effectiveness of CUDA core dump.

Debugging Improper Kernel Launch

This code launches two kernels consecutively (illegalMemoryAccessKernel and normalKernel). During execution, you would encounter an error message: CUDA Error at test.cu:62 - cudaMemcpy(h_data, d_data, size * sizeof(int), cudaMemcpyDeviceToHost): an illegal memory access was encountered, and the error would only be detected in the return value of cudaMemcpy. Even with CUDA_LAUNCH_BLOCKING=1, it is still impossible to identify the specific kernel that caused the error.

By adding the CUDA core dump-related environment variables, we can observe:

After a GPU thread triggers an illegal memory access, the CPU immediately generates a coredump file and then triggers a CPU exception, directly terminating the program. At this point, we obtain a coredump file /tmp/cuda_coredump_xxx.1799919.1754898045. We can open it using cuda-gdb (command: target cudacore /path/to/coredump_file, where cudacore refers to the coredump on CUDA):

We can clearly see that the exception is caused by illegalMemoryAccessKernel at kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 124, warp 0, lane 0.

Here’s a more complex example where an illegal memory access kernel is inserted into a CUDA graph:

Direct execution results in the following error:

The error is not printed until the output.sum() triggers a device synchronization and reveals the illegal memory access. However, we don't know which kernel caused the illegal memory access since cuda kernels are executed asynchronously.

After adding CUDA_LAUNCH_BLOCKING=1, the error message changed to:

It can be inferred that an exception occurs in a kernel within the CUDA graph. However, conventional methods can only provide information up to this point.

By adding the environment variables CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1 CUDA_COREDUMP_SHOW_PROGRESS=1 CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory' CUDA_COREDUMP_FILE="/tmp/cuda_coredump_%h.%p.%t", we can clearly identify the kernel that caused the error:

Clearly, this is a fill function, and the grid size of 40960 is very large. With this information, we can easily pinpoint that the lines y = from_buffer(x.data_ptr(), x.numel() * 1024 * 1024); y.fill_(1); forcibly expand the length of x by a million times and then fill it entirely with 1s, thereby triggering the illegal memory access exception.

On some GPUs, this line might cause invalid argument error instead of illegal memory access, because the grid size exceeds the maximum limit. In such cases, the CUDA core dump feature cannot be triggered, and you need to turn down the expansion factor 1024 * 1024 a little bit to avoid exceeding the grid size limit.

In theory, CUDA core dump should be able to capture various exceptions caused by a specific thread on the GPU. However, in practice, on certain GPU and driver versions, exceptions like operation not supported on global/shared address space may fail to trigger a CUDA core dump. Fortunately, illegal memory access can generally trigger CUDA core dumps reliably, which satisfies most debugging needs; For hardware-related errors, such as Invalid access of peer GPU memory over nvlink or a hardware error, these are not caused by a specific thread and cannot be attributed to a particular GPU thread. As a result, CUDA core dumps will not be triggered for such issues; Errors caused by improper use of the driver API are considered non-sticky errors and are unrelated to the GPU itself. These errors are reported at the driver API level and do not trigger CUDA core dumps. A common example is an out-of-memory error during cudaMalloc, which will not result in a CUDA core dump; For distributed programs involving multi-GPU communication, memory mapping is often used to map the memory of other GPUs to the current GPU. If the program on another GPU exits, the mapped memory becomes invalid, and accessing it will trigger an illegal memory access. However, this does not fall under the typical illegal memory access issues. Such problems are common during the shutdown process of distributed programs. If GPUs are communicating during shutdown, the order of shutdown may cause some GPUs to report illegal memory access. When using CUDA core dump for such programs, it is important to distinguish these false positives; Enabling CUDA core dump does have some performance impact on CUDA kernels (since it needs to check for errors and attribute them when GPU threads exit). Therefore, it is not advisable to enable CUDA core dump in production environments. It is recommended to enable CUDA core dump only after errors like illegal memory access can be reliably reproduced for debugging purposes; To get the maximum benefit from CUDA core dump, it is recommended to recompile vLLM with debug symbols, or at least embed line information during compilation. Unfortunately, the default build of vLLM does not contain such information due to the binary size limit. To enjoy the benefit, users have to compile vLLM from source with an envrionment variable export NVCC_PREPEND_FLAGS='-lineinfo' or export NVCC_PREPEND_FLAGS='-G'. It is recommended to start from -lineinfo, and only switch to -G when -lineinfo is not enough. With rich debug information, cuda core dump can trace back to the exact line of code that caused the exception

This blogpost analyzed the principles and use cases of CUDA core dump. This debugging method is effective for issues like improper kernel launches and kernel exceptions within CUDA graphs, making it a powerful tool for debugging illegal memory access issues and beyond.

As an example, we recently use this technique to debug a complex illegal memory access issue in vLLM, see this PR for more details. Basically, we add a triton kernel for MRope, but that kernel has an implicit assumption that head_size==rotary_dim (i.e. it's a full Rope). When head_size!=rotary_dim (i.e. it's a partial Rope), the kernel will trigger an illegal memory access, which is the case for the new GLM-4.5V model. Without CUDA core dump, the error is reported as Failed: Cuda error /workspace/csrc/custom_all_reduce.cuh:453 'an illegal memory access was encountered', which is very misleading. With CUDA core dump, we can easily pinpoint the error to the MRope kernel, and then fix it. Note that this example is caused by mis-configuration of the cuda kernel parameters, and finding the kernel that caused the issue is pretty enough for debugging. For more complicated illegal memory access issues, we still need to isolate the kernel and reproduce the issue in a minimal example instead of an end-to-end example, and then use more dedicated tools like Compute Sanitizer to further investigate the issue.

The vLLM project aims to provide easy, fast, and cheap LLM serving for everyone, and easy debugging is also an important aspect. We will continue to share more debugging tips and techniques in the future, to build a strong LLM inference ecosystem together. To share your story or usage with vLLM, please submit a PR at the blogpost repository.

We would like to thank Ze Long, Vikram Sharma Mailthody, Jeremy Iverson, and Sandarbh Jain from NVIDIA for their helpful discussions. Lucas Wilkinson from Red Hat helped polishing the draft.

Related Posts

How vLLM developers debug hanging and complex CUDA kernels by triggering GPU core dumps, identifying stuck kernels, and mapping failures back to source code lines for faster kernel debugging.