vllm - ✅(Solved) Fix [Bug]: CUDA illegal memory access on GPTQ Marlin [1 pull requests, 17 comments, 3 participants]

vllm2026-03-11 18:34:02

ON THIS PAGE

Recommended Tools

×6

Utilities matched from this issue’s tags and category — try them while you read without losing context.

GitHub issue graph ai analysis

Paste a GitHub issue URL. We fetch that issue, discover linked issues from bodies/comments/timeline, collect linked pull requests, and produce a structured English report.

The report is written in English Markdown for sharing and archival.

GitHub issue URL

Helpful · Quick feedback

GitHub stats

vllm-project/vllm#36811•Fetched 2026-04-08 00:34:30

View on GitHub

Comments

Participants

Timeline

Reactions

Author

Participants

Timeline (top)

commented ×17subscribed ×12mentioned ×11referenced ×2

Error Message

(EngineCore_DP0 pid=121)     model_output = self._model_forward(
(EngineCore_DP0 pid=121)                    ^^^^^^^^^^^^^^^^^^^^
(EngineCore_DP0 pid=121)   File "/usr/local/lib/python3.12/dist-packages/vllm/v1/worker/gpu_model_runner.py", line 3277, in _model_forward
(EngineCore_DP0 pid=121)     return self.model(
(EngineCore_DP0 pid=121)            ^^^^^^^^^^^
(EngineCore_DP0 pid=121)   File "/usr/local/lib/python3.12/dist-packages/vllm/compilation/cuda_graph.py", line 342, in __call__
(EngineCore_DP0 pid=121)     entry.cudagraph.replay()
(EngineCore_DP0 pid=121)   File "/usr/local/lib/python3.12/dist-packages/torch/cuda/graphs.py", line 143, in replay
(EngineCore_DP0 pid=121)     super().replay()
(EngineCore_DP0 pid=121) torch.AcceleratorError: CUDA error: an illegal memory access was encountered

https://gist.githubusercontent.com/tonibofarull/da64693c9f94d566cae20761a58f910a/raw/3831d39700cdf8e28a9b7431594b5983ddb671de/gistfile1.txt

Fix Action

Fix / Workaround

============================== CPU Info

Architecture: x86_64 CPU op-mode(s): 32-bit, 64-bit Address sizes: 40 bits physical, 57 bits virtual Byte Order: Little Endian CPU(s): 8 On-line CPU(s) list: 0-7 Vendor ID: AuthenticAMD Model name: AMD EPYC-Genoa Processor CPU family: 25 Model: 17 Thread(s) per core: 1 Core(s) per socket: 1 Socket(s): 8 Stepping: 0 BogoMIPS: 4999.71 Flags: fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush mmx fxsr sse sse2 syscall nx mmxext fxsr_opt pdpe1gb rdtscp lm rep_good nopl xtopology cpuid extd_apicid pni pclmulqdq ssse3 fma cx16 pcid sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand hypervisor lahf_lm cmp_legacy svm cr8_legacy abm sse4a misalignsse 3dnowprefetch osvw topoext perfctr_core ssbd ibrs ibpb stibp ibrs_enhanced vmmcall fsgsbase tsc_adjust bmi1 avx2 smep bmi2 erms invpcid avx512f avx512dq rdseed adx smap avx512ifma clflushopt clwb avx512cd sha_ni avx512bw avx512vl xsaveopt xsavec xgetbv1 xsaves avx512_bf16 clzero xsaveerptr wbnoinvd arat npt lbrv nrip_save tsc_scale vmcb_clean flushbyasid pausefilter pfthreshold v_vmsave_vmload vgif vnmi avx512vbmi umip pku ospke avx512_vbmi2 gfni vaes vpclmulqdq avx512_vnni avx512_bitalg avx512_vpopcntdq la57 rdpid fsrm flush_l1d arch_capabilities Virtualization: AMD-V L1d cache: 256 KiB (8 instances) L1i cache: 256 KiB (8 instances) L2 cache: 8 MiB (8 instances) L3 cache: 256 MiB (8 instances) NUMA node(s): 1 NUMA node0 CPU(s): 0-7 Vulnerability Gather data sampling: Not affected Vulnerability Ghostwrite: Not affected Vulnerability Indirect target selection: Not affected Vulnerability Itlb multihit: Not affected Vulnerability L1tf: Not affected Vulnerability Mds: Not affected Vulnerability Meltdown: Not affected Vulnerability Mmio stale data: Not affected Vulnerability Old microcode: Not affected Vulnerability Reg file data sampling: Not affected Vulnerability Retbleed: Not affected Vulnerability Spec rstack overflow: Vulnerable: Safe RET, no microcode Vulnerability Spec store bypass: Mitigation; Speculative Store Bypass disabled via prctl Vulnerability Spectre v1: Mitigation; usercopy/swapgs barriers and __user pointer sanitization Vulnerability Spectre v2: Mitigation; Enhanced / Automatic IBRS; IBPB conditional; STIBP disabled; PBRSB-eIBRS Not affected; BHI Not affected Vulnerability Srbds: Not affected Vulnerability Tsa: Vulnerable: No microcode Vulnerability Tsx async abort: Not affected Vulnerability Vmscape: Not affected

PR fix notes

PR #36889: [WIP][Bugfix] Fix CUDA illegal memory access for GPTQ Marlin MoE with CUDA graphs

Repository: vllm-project/vllm
Author: haosdent
State: open | merged: False
Link: https://github.com/vllm-project/vllm/pull/36889

Description (problem / solution / changelog)

Purpose

Close #36811

Fix CUDA illegal memory access when running GPTQ Marlin quantized MoE models (e.g., Qwen/Qwen3.5-35B-A3B-GPTQ-Int4) with CUDA graphs enabled.

Root Cause: `c_tmp` buffer out-of-bounds access

In csrc/moe/marlin_moe_wna16/ops.cu, the FP32 reduction buffer c_tmp is sized using:

long max_c_tmp_size = min(
    (long)size_n * sorted_token_ids.size(0),
    (long)sms * 4 * moe_block_size * max_thread_n);

The first term (size_n * sorted_token_ids.size(0)) underestimates the required buffer when:

Small batch size (M=1): sorted_token_ids.size(0) is small (e.g., 64 for E=64, topk=8, block=8)
Small size_n (w2/down-projection call): uses hidden_size (e.g., 2560 for Qwen3.5-35B-A3B)

The kernel accesses C_tmp[locks_off * c_size + ...] where c_size = tb_m * tb_n / 4 (int4 units). The first term doesn't account for the tb_m factor in the stride, causing the buffer to be too small.

Concrete example on RTX A6000 (84 SMs), M=1, w2 call (size_n=2560):

term1 = 2560 * 64 = 163,840, after *2 for block_size=8 → 327,680 FP32 = 81,920 int4
Kernel launches 168 blocks, 160 active tiles, max locks_off = 159
Max access index: 159 * 512 + 256 * 6 + 255 = 83,199 int4
OOB by 1,280 int4 elements (20,480 bytes)

In eager mode, PyTorch's caching allocator over-allocates memory blocks, so the OOB silently reads/writes adjacent valid memory. With CUDA graphs, the graph memory pool has tight allocation boundaries, causing a crash.

Fix (single file: `csrc/moe/marlin_moe_wna16/ops.cu`)

c_tmp sizing: Remove the min() to always use sms * 4 * moe_block_size * max_thread_n, matching the non-MoE Marlin pattern (csrc/quantization/marlin/marlin.cu:685-687). Memory cost increase is ~5 MB (negligible).
cudaFuncSetAttribute hardening (defensive): Always set the attribute to the device maximum shared memory, and use sh_cache_size (the exact amount needed) for the kernel launch parameter. This prevents potential issues on GPUs where different batch sizes might select different blocks_per_sm values for the same kernel function.

Test Plan

Added test_fused_marlin_moe_cuda_graph regression test that captures CUDA graphs at 7 different batch sizes (1, 2, 4, 8, 16, 24, 32) with 64 experts and verifies replay produces correct results matching eager mode.

Test Result

All 7 batch sizes pass:

Capturing CUDA graph for M=1...
Capturing CUDA graph for M=2...
Capturing CUDA graph for M=4...
...
  M=1: PASSED
  M=2: PASSED
  M=4: PASSED
  M=8: PASSED
  M=16: PASSED
  M=24: PASSED
  M=32: PASSED

Changed files

csrc/moe/marlin_moe_wna16/ops.cu (modified, +5/-5)
tests/kernels/moe/test_moe.py (modified, +113/-0)

Code Example

Collecting environment information...
==============================
        System Info
==============================
OS                           : Ubuntu 22.04.5 LTS (x86_64)
GCC version                  : (Ubuntu 11.4.0-1ubuntu1~22.04.3) 11.4.0
Clang version                : Could not collect
CMake version                : Could not collect
Libc version                 : glibc-2.35

==============================
       PyTorch Info
==============================
PyTorch version              : 2.10.0+cu129
Is debug build               : False
CUDA used to build PyTorch   : 12.9
ROCM used to build PyTorch   : N/A

==============================
      Python Environment
==============================
Python version               : 3.12.13 (main, Mar  4 2026, 09:23:07) [GCC 11.4.0] (64-bit runtime)
Python platform              : Linux-6.16.11-200.fc42.x86_64-x86_64-with-glibc2.35

==============================
       CUDA / GPU Info
==============================
Is CUDA available            : True
CUDA runtime version         : 12.9.86
CUDA_MODULE_LOADING set to   : 
GPU models and configuration : GPU 0: NVIDIA RTX A6000
Nvidia driver version        : 580.126.18
cuDNN version                : Could not collect
HIP runtime version          : N/A
MIOpen runtime version       : N/A
Is XNNPACK available         : True

==============================
          CPU Info
==============================
Architecture:                            x86_64
CPU op-mode(s):                          32-bit, 64-bit
Address sizes:                           40 bits physical, 57 bits virtual
Byte Order:                              Little Endian
CPU(s):                                  8
On-line CPU(s) list:                     0-7
Vendor ID:                               AuthenticAMD
Model name:                              AMD EPYC-Genoa Processor
CPU family:                              25
Model:                                   17
Thread(s) per core:                      1
Core(s) per socket:                      1
Socket(s):                               8
Stepping:                                0
BogoMIPS:                                4999.71
Flags:                                   fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush mmx fxsr sse sse2 syscall nx mmxext fxsr_opt pdpe1gb rdtscp lm rep_good nopl xtopology cpuid extd_apicid pni pclmulqdq ssse3 fma cx16 pcid sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand hypervisor lahf_lm cmp_legacy svm cr8_legacy abm sse4a misalignsse 3dnowprefetch osvw topoext perfctr_core ssbd ibrs ibpb stibp ibrs_enhanced vmmcall fsgsbase tsc_adjust bmi1 avx2 smep bmi2 erms invpcid avx512f avx512dq rdseed adx smap avx512ifma clflushopt clwb avx512cd sha_ni avx512bw avx512vl xsaveopt xsavec xgetbv1 xsaves avx512_bf16 clzero xsaveerptr wbnoinvd arat npt lbrv nrip_save tsc_scale vmcb_clean flushbyasid pausefilter pfthreshold v_vmsave_vmload vgif vnmi avx512vbmi umip pku ospke avx512_vbmi2 gfni vaes vpclmulqdq avx512_vnni avx512_bitalg avx512_vpopcntdq la57 rdpid fsrm flush_l1d arch_capabilities
Virtualization:                          AMD-V
L1d cache:                               256 KiB (8 instances)
L1i cache:                               256 KiB (8 instances)
L2 cache:                                8 MiB (8 instances)
L3 cache:                                256 MiB (8 instances)
NUMA node(s):                            1
NUMA node0 CPU(s):                       0-7
Vulnerability Gather data sampling:      Not affected
Vulnerability Ghostwrite:                Not affected
Vulnerability Indirect target selection: Not affected
Vulnerability Itlb multihit:             Not affected
Vulnerability L1tf:                      Not affected
Vulnerability Mds:                       Not affected
Vulnerability Meltdown:                  Not affected
Vulnerability Mmio stale data:           Not affected
Vulnerability Old microcode:             Not affected
Vulnerability Reg file data sampling:    Not affected
Vulnerability Retbleed:                  Not affected
Vulnerability Spec rstack overflow:      Vulnerable: Safe RET, no microcode
Vulnerability Spec store bypass:         Mitigation; Speculative Store Bypass disabled via prctl
Vulnerability Spectre v1:                Mitigation; usercopy/swapgs barriers and __user pointer sanitization
Vulnerability Spectre v2:                Mitigation; Enhanced / Automatic IBRS; IBPB conditional; STIBP disabled; PBRSB-eIBRS Not affected; BHI Not affected
Vulnerability Srbds:                     Not affected
Vulnerability Tsa:                       Vulnerable: No microcode
Vulnerability Tsx async abort:           Not affected
Vulnerability Vmscape:                   Not affected

==============================
Versions of relevant libraries
==============================
[pip3] flashinfer-python==0.6.4
[pip3] numpy==2.2.6
[pip3] nvidia-cublas-cu12==12.9.1.4
[pip3] nvidia-cuda-cupti-cu12==12.9.79
[pip3] nvidia-cuda-nvrtc-cu12==12.9.86
[pip3] nvidia-cuda-runtime-cu12==12.9.79
[pip3] nvidia-cudnn-cu12==9.10.2.21
[pip3] nvidia-cudnn-frontend==1.18.0
[pip3] nvidia-cufft-cu12==11.4.1.4
[pip3] nvidia-cufile-cu12==1.14.1.1
[pip3] nvidia-curand-cu12==10.3.10.19
[pip3] nvidia-cusolver-cu12==11.7.5.82
[pip3] nvidia-cusparse-cu12==12.5.10.65
[pip3] nvidia-cusparselt-cu12==0.7.1
[pip3] nvidia-cutlass-dsl==4.4.1
[pip3] nvidia-cutlass-dsl-libs-base==4.4.1
[pip3] nvidia-ml-py==13.590.48
[pip3] nvidia-nccl-cu12==2.27.5
[pip3] nvidia-nvjitlink-cu12==12.9.86
[pip3] nvidia-nvshmem-cu12==3.4.5
[pip3] nvidia-nvtx-cu12==12.9.79
[pip3] pyzmq==27.1.0
[pip3] torch==2.10.0+cu129
[pip3] torch_c_dlpack_ext==0.1.5
[pip3] torchaudio==2.10.0+cu129
[pip3] torchvision==0.25.0+cu129
[pip3] transformers==4.57.6
[pip3] triton==3.6.0
[conda] Could not collect

==============================
         vLLM Info
==============================
ROCM Version                 : Could not collect
vLLM Version                 : 0.17.1rc1.dev23+g76c6e6da0 (git sha: 76c6e6da0)
vLLM Build Flags:
  CUDA Archs: 7.0 7.5 8.0 8.9 9.0 10.0 12.0; ROCm: Disabled
GPU Topology:
        GPU0    CPU Affinity    NUMA Affinity   GPU NUMA ID
GPU0     X      0-7     0               N/A

Legend:

  X    = Self
  SYS  = Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI)
  NODE = Connection traversing PCIe as well as the interconnect between PCIe Host Bridges within a NUMA node
  PHB  = Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU)
  PXB  = Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge)
  PIX  = Connection traversing at most a single PCIe bridge
  NV#  = Connection traversing a bonded set of # NVLinks

==============================
     Environment Variables
==============================
NVIDIA_VISIBLE_DEVICES=GPU-4ce9082e-c19f-c259-3101-1a12778ffe97
NVIDIA_REQUIRE_CUDA=cuda>=12.9 brand=unknown,driver>=535,driver<536 brand=grid,driver>=535,driver<536 brand=tesla,driver>=535,driver<536 brand=nvidia,driver>=535,driver<536 brand=quadro,driver>=535,driver<536 brand=quadrortx,driver>=535,driver<536 brand=nvidiartx,driver>=535,driver<536 brand=vapps,driver>=535,driver<536 brand=vpc,driver>=535,driver<536 brand=vcs,driver>=535,driver<536 brand=vws,driver>=535,driver<536 brand=cloudgaming,driver>=535,driver<536 brand=unknown,driver>=550,driver<551 brand=grid,driver>=550,driver<551 brand=tesla,driver>=550,driver<551 brand=nvidia,driver>=550,driver<551 brand=quadro,driver>=550,driver<551 brand=quadrortx,driver>=550,driver<551 brand=nvidiartx,driver>=550,driver<551 brand=vapps,driver>=550,driver<551 brand=vpc,driver>=550,driver<551 brand=vcs,driver>=550,driver<551 brand=vws,driver>=550,driver<551 brand=cloudgaming,driver>=550,driver<551 brand=unknown,driver>=560,driver<561 brand=grid,driver>=560,driver<561 brand=tesla,driver>=560,driver<561 brand=nvidia,driver>=560,driver<561 brand=quadro,driver>=560,driver<561 brand=quadrortx,driver>=560,driver<561 brand=nvidiartx,driver>=560,driver<561 brand=vapps,driver>=560,driver<561 brand=vpc,driver>=560,driver<561 brand=vcs,driver>=560,driver<561 brand=vws,driver>=560,driver<561 brand=cloudgaming,driver>=560,driver<561 brand=unknown,driver>=565,driver<566 brand=grid,driver>=565,driver<566 brand=tesla,driver>=565,driver<566 brand=nvidia,driver>=565,driver<566 brand=quadro,driver>=565,driver<566 brand=quadrortx,driver>=565,driver<566 brand=nvidiartx,driver>=565,driver<566 brand=vapps,driver>=565,driver<566 brand=vpc,driver>=565,driver<566 brand=vcs,driver>=565,driver<566 brand=vws,driver>=565,driver<566 brand=cloudgaming,driver>=565,driver<566 brand=unknown,driver>=570,driver<571 brand=grid,driver>=570,driver<571 brand=tesla,driver>=570,driver<571 brand=nvidia,driver>=570,driver<571 brand=quadro,driver>=570,driver<571 brand=quadrortx,driver>=570,driver<571 brand=nvidiartx,driver>=570,driver<571 brand=vapps,driver>=570,driver<571 brand=vpc,driver>=570,driver<571 brand=vcs,driver>=570,driver<571 brand=vws,driver>=570,driver<571 brand=cloudgaming,driver>=570,driver<571
TORCH_CUDA_ARCH_LIST=7.0 7.5 8.0 8.9 9.0 10.0 12.0
CUDA_DEVICE_SM_LIMIT=0
NVIDIA_DRIVER_CAPABILITIES=compute,utility
VLLM_USAGE_SOURCE=production-docker-image
CUDA_VERSION=12.9.1
CUDA_LAUNCH_BLOCKING=1
VLLM_ENABLE_CUDA_COMPATIBILITY=0
TORCH_USE_CUDA_DSA=1
CUDA_DEVICE_MEMORY_LIMIT_0=38000m
CUDA_DEVICE_MEMORY_SHARED_CACHE=/usr/local/vgpu/fcbf0e86-47f8-4241-80e7-4979f1f524b9.cache
LD_LIBRARY_PATH=/usr/local/nvidia/lib64:/usr/local/cuda/lib64:/usr/local/cuda/lib64
VLLM_LOGGING_LEVEL=DEBUG
PYTORCH_NVML_BASED_CUDA_CHECK=1
TORCHINDUCTOR_COMPILE_THREADS=1
TORCHINDUCTOR_CACHE_DIR=/tmp/torchinductor_root

---

--model
Qwen/Qwen3.5-35B-A3B-GPTQ-Int4
--served-model-name
main
--max-model-len
8192
--host
0.0.0.0
--port
8080
--language-model-only
--enable-auto-tool-choice
--tool-call-parser
qwen3_coder
--reasoning-parser
qwen3
--default-chat-template-kwargs
{"enable_thinking": false}
--max-cudagraph-capture-size
32
--gpu-memory-utilization
0.8

---

VLLM_LOGGING_LEVEL=DEBUG
CUDA_LAUNCH_BLOCKING=1
TORCH_USE_CUDA_DSA=1

---

(EngineCore_DP0 pid=121)     model_output = self._model_forward(
(EngineCore_DP0 pid=121)                    ^^^^^^^^^^^^^^^^^^^^
(EngineCore_DP0 pid=121)   File "/usr/local/lib/python3.12/dist-packages/vllm/v1/worker/gpu_model_runner.py", line 3277, in _model_forward
(EngineCore_DP0 pid=121)     return self.model(
(EngineCore_DP0 pid=121)            ^^^^^^^^^^^
(EngineCore_DP0 pid=121)   File "/usr/local/lib/python3.12/dist-packages/vllm/compilation/cuda_graph.py", line 342, in __call__
(EngineCore_DP0 pid=121)     entry.cudagraph.replay()
(EngineCore_DP0 pid=121)   File "/usr/local/lib/python3.12/dist-packages/torch/cuda/graphs.py", line 143, in replay
(EngineCore_DP0 pid=121)     super().replay()
(EngineCore_DP0 pid=121) torch.AcceleratorError: CUDA error: an illegal memory access was encountered

RAW_BUFFERClick to expand / collapse

Your current environment

<details> <summary>The output of <code>python collect_env.py</code></summary>

Collecting environment information...
==============================
        System Info
==============================
OS                           : Ubuntu 22.04.5 LTS (x86_64)
GCC version                  : (Ubuntu 11.4.0-1ubuntu1~22.04.3) 11.4.0
Clang version                : Could not collect
CMake version                : Could not collect
Libc version                 : glibc-2.35

==============================
       PyTorch Info
==============================
PyTorch version              : 2.10.0+cu129
Is debug build               : False
CUDA used to build PyTorch   : 12.9
ROCM used to build PyTorch   : N/A

==============================
      Python Environment
==============================
Python version               : 3.12.13 (main, Mar  4 2026, 09:23:07) [GCC 11.4.0] (64-bit runtime)
Python platform              : Linux-6.16.11-200.fc42.x86_64-x86_64-with-glibc2.35

==============================
       CUDA / GPU Info
==============================
Is CUDA available            : True
CUDA runtime version         : 12.9.86
CUDA_MODULE_LOADING set to   : 
GPU models and configuration : GPU 0: NVIDIA RTX A6000
Nvidia driver version        : 580.126.18
cuDNN version                : Could not collect
HIP runtime version          : N/A
MIOpen runtime version       : N/A
Is XNNPACK available         : True

==============================
          CPU Info
==============================
Architecture:                            x86_64
CPU op-mode(s):                          32-bit, 64-bit
Address sizes:                           40 bits physical, 57 bits virtual
Byte Order:                              Little Endian
CPU(s):                                  8
On-line CPU(s) list:                     0-7
Vendor ID:                               AuthenticAMD
Model name:                              AMD EPYC-Genoa Processor
CPU family:                              25
Model:                                   17
Thread(s) per core:                      1
Core(s) per socket:                      1
Socket(s):                               8
Stepping:                                0
BogoMIPS:                                4999.71
Flags:                                   fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush mmx fxsr sse sse2 syscall nx mmxext fxsr_opt pdpe1gb rdtscp lm rep_good nopl xtopology cpuid extd_apicid pni pclmulqdq ssse3 fma cx16 pcid sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand hypervisor lahf_lm cmp_legacy svm cr8_legacy abm sse4a misalignsse 3dnowprefetch osvw topoext perfctr_core ssbd ibrs ibpb stibp ibrs_enhanced vmmcall fsgsbase tsc_adjust bmi1 avx2 smep bmi2 erms invpcid avx512f avx512dq rdseed adx smap avx512ifma clflushopt clwb avx512cd sha_ni avx512bw avx512vl xsaveopt xsavec xgetbv1 xsaves avx512_bf16 clzero xsaveerptr wbnoinvd arat npt lbrv nrip_save tsc_scale vmcb_clean flushbyasid pausefilter pfthreshold v_vmsave_vmload vgif vnmi avx512vbmi umip pku ospke avx512_vbmi2 gfni vaes vpclmulqdq avx512_vnni avx512_bitalg avx512_vpopcntdq la57 rdpid fsrm flush_l1d arch_capabilities
Virtualization:                          AMD-V
L1d cache:                               256 KiB (8 instances)
L1i cache:                               256 KiB (8 instances)
L2 cache:                                8 MiB (8 instances)
L3 cache:                                256 MiB (8 instances)
NUMA node(s):                            1
NUMA node0 CPU(s):                       0-7
Vulnerability Gather data sampling:      Not affected
Vulnerability Ghostwrite:                Not affected
Vulnerability Indirect target selection: Not affected
Vulnerability Itlb multihit:             Not affected
Vulnerability L1tf:                      Not affected
Vulnerability Mds:                       Not affected
Vulnerability Meltdown:                  Not affected
Vulnerability Mmio stale data:           Not affected
Vulnerability Old microcode:             Not affected
Vulnerability Reg file data sampling:    Not affected
Vulnerability Retbleed:                  Not affected
Vulnerability Spec rstack overflow:      Vulnerable: Safe RET, no microcode
Vulnerability Spec store bypass:         Mitigation; Speculative Store Bypass disabled via prctl
Vulnerability Spectre v1:                Mitigation; usercopy/swapgs barriers and __user pointer sanitization
Vulnerability Spectre v2:                Mitigation; Enhanced / Automatic IBRS; IBPB conditional; STIBP disabled; PBRSB-eIBRS Not affected; BHI Not affected
Vulnerability Srbds:                     Not affected
Vulnerability Tsa:                       Vulnerable: No microcode
Vulnerability Tsx async abort:           Not affected
Vulnerability Vmscape:                   Not affected

==============================
Versions of relevant libraries
==============================
[pip3] flashinfer-python==0.6.4
[pip3] numpy==2.2.6
[pip3] nvidia-cublas-cu12==12.9.1.4
[pip3] nvidia-cuda-cupti-cu12==12.9.79
[pip3] nvidia-cuda-nvrtc-cu12==12.9.86
[pip3] nvidia-cuda-runtime-cu12==12.9.79
[pip3] nvidia-cudnn-cu12==9.10.2.21
[pip3] nvidia-cudnn-frontend==1.18.0
[pip3] nvidia-cufft-cu12==11.4.1.4
[pip3] nvidia-cufile-cu12==1.14.1.1
[pip3] nvidia-curand-cu12==10.3.10.19
[pip3] nvidia-cusolver-cu12==11.7.5.82
[pip3] nvidia-cusparse-cu12==12.5.10.65
[pip3] nvidia-cusparselt-cu12==0.7.1
[pip3] nvidia-cutlass-dsl==4.4.1
[pip3] nvidia-cutlass-dsl-libs-base==4.4.1
[pip3] nvidia-ml-py==13.590.48
[pip3] nvidia-nccl-cu12==2.27.5
[pip3] nvidia-nvjitlink-cu12==12.9.86
[pip3] nvidia-nvshmem-cu12==3.4.5
[pip3] nvidia-nvtx-cu12==12.9.79
[pip3] pyzmq==27.1.0
[pip3] torch==2.10.0+cu129
[pip3] torch_c_dlpack_ext==0.1.5
[pip3] torchaudio==2.10.0+cu129
[pip3] torchvision==0.25.0+cu129
[pip3] transformers==4.57.6
[pip3] triton==3.6.0
[conda] Could not collect

==============================
         vLLM Info
==============================
ROCM Version                 : Could not collect
vLLM Version                 : 0.17.1rc1.dev23+g76c6e6da0 (git sha: 76c6e6da0)
vLLM Build Flags:
  CUDA Archs: 7.0 7.5 8.0 8.9 9.0 10.0 12.0; ROCm: Disabled
GPU Topology:
        GPU0    CPU Affinity    NUMA Affinity   GPU NUMA ID
GPU0     X      0-7     0               N/A

Legend:

  X    = Self
  SYS  = Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI)
  NODE = Connection traversing PCIe as well as the interconnect between PCIe Host Bridges within a NUMA node
  PHB  = Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU)
  PXB  = Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge)
  PIX  = Connection traversing at most a single PCIe bridge
  NV#  = Connection traversing a bonded set of # NVLinks

==============================
     Environment Variables
==============================
NVIDIA_VISIBLE_DEVICES=GPU-4ce9082e-c19f-c259-3101-1a12778ffe97
NVIDIA_REQUIRE_CUDA=cuda>=12.9 brand=unknown,driver>=535,driver<536 brand=grid,driver>=535,driver<536 brand=tesla,driver>=535,driver<536 brand=nvidia,driver>=535,driver<536 brand=quadro,driver>=535,driver<536 brand=quadrortx,driver>=535,driver<536 brand=nvidiartx,driver>=535,driver<536 brand=vapps,driver>=535,driver<536 brand=vpc,driver>=535,driver<536 brand=vcs,driver>=535,driver<536 brand=vws,driver>=535,driver<536 brand=cloudgaming,driver>=535,driver<536 brand=unknown,driver>=550,driver<551 brand=grid,driver>=550,driver<551 brand=tesla,driver>=550,driver<551 brand=nvidia,driver>=550,driver<551 brand=quadro,driver>=550,driver<551 brand=quadrortx,driver>=550,driver<551 brand=nvidiartx,driver>=550,driver<551 brand=vapps,driver>=550,driver<551 brand=vpc,driver>=550,driver<551 brand=vcs,driver>=550,driver<551 brand=vws,driver>=550,driver<551 brand=cloudgaming,driver>=550,driver<551 brand=unknown,driver>=560,driver<561 brand=grid,driver>=560,driver<561 brand=tesla,driver>=560,driver<561 brand=nvidia,driver>=560,driver<561 brand=quadro,driver>=560,driver<561 brand=quadrortx,driver>=560,driver<561 brand=nvidiartx,driver>=560,driver<561 brand=vapps,driver>=560,driver<561 brand=vpc,driver>=560,driver<561 brand=vcs,driver>=560,driver<561 brand=vws,driver>=560,driver<561 brand=cloudgaming,driver>=560,driver<561 brand=unknown,driver>=565,driver<566 brand=grid,driver>=565,driver<566 brand=tesla,driver>=565,driver<566 brand=nvidia,driver>=565,driver<566 brand=quadro,driver>=565,driver<566 brand=quadrortx,driver>=565,driver<566 brand=nvidiartx,driver>=565,driver<566 brand=vapps,driver>=565,driver<566 brand=vpc,driver>=565,driver<566 brand=vcs,driver>=565,driver<566 brand=vws,driver>=565,driver<566 brand=cloudgaming,driver>=565,driver<566 brand=unknown,driver>=570,driver<571 brand=grid,driver>=570,driver<571 brand=tesla,driver>=570,driver<571 brand=nvidia,driver>=570,driver<571 brand=quadro,driver>=570,driver<571 brand=quadrortx,driver>=570,driver<571 brand=nvidiartx,driver>=570,driver<571 brand=vapps,driver>=570,driver<571 brand=vpc,driver>=570,driver<571 brand=vcs,driver>=570,driver<571 brand=vws,driver>=570,driver<571 brand=cloudgaming,driver>=570,driver<571
TORCH_CUDA_ARCH_LIST=7.0 7.5 8.0 8.9 9.0 10.0 12.0
CUDA_DEVICE_SM_LIMIT=0
NVIDIA_DRIVER_CAPABILITIES=compute,utility
VLLM_USAGE_SOURCE=production-docker-image
CUDA_VERSION=12.9.1
CUDA_LAUNCH_BLOCKING=1
VLLM_ENABLE_CUDA_COMPATIBILITY=0
TORCH_USE_CUDA_DSA=1
CUDA_DEVICE_MEMORY_LIMIT_0=38000m
CUDA_DEVICE_MEMORY_SHARED_CACHE=/usr/local/vgpu/fcbf0e86-47f8-4241-80e7-4979f1f524b9.cache
LD_LIBRARY_PATH=/usr/local/nvidia/lib64:/usr/local/cuda/lib64:/usr/local/cuda/lib64
VLLM_LOGGING_LEVEL=DEBUG
PYTORCH_NVML_BASED_CUDA_CHECK=1
TORCHINDUCTOR_COMPILE_THREADS=1
TORCHINDUCTOR_CACHE_DIR=/tmp/torchinductor_root

</details>

🐛 Describe the bug

Similar to https://github.com/vllm-project/vllm/issues/32834 but observed when running a GPTQ Marlin quantized model.

Image used: vllm/vllm-openai:nightly-76c6e6da08dbe73c2ee0d92dabe01786b44845d2.

Server Arguments

--model
Qwen/Qwen3.5-35B-A3B-GPTQ-Int4
--served-model-name
main
--max-model-len
8192
--host
0.0.0.0
--port
8080
--language-model-only
--enable-auto-tool-choice
--tool-call-parser
qwen3_coder
--reasoning-parser
qwen3
--default-chat-template-kwargs
{"enable_thinking": false}
--max-cudagraph-capture-size
32
--gpu-memory-utilization
0.8

Environment Variables

VLLM_LOGGING_LEVEL=DEBUG
CUDA_LAUNCH_BLOCKING=1
TORCH_USE_CUDA_DSA=1

Logs

(EngineCore_DP0 pid=121)     model_output = self._model_forward(
(EngineCore_DP0 pid=121)                    ^^^^^^^^^^^^^^^^^^^^
(EngineCore_DP0 pid=121)   File "/usr/local/lib/python3.12/dist-packages/vllm/v1/worker/gpu_model_runner.py", line 3277, in _model_forward
(EngineCore_DP0 pid=121)     return self.model(
(EngineCore_DP0 pid=121)            ^^^^^^^^^^^
(EngineCore_DP0 pid=121)   File "/usr/local/lib/python3.12/dist-packages/vllm/compilation/cuda_graph.py", line 342, in __call__
(EngineCore_DP0 pid=121)     entry.cudagraph.replay()
(EngineCore_DP0 pid=121)   File "/usr/local/lib/python3.12/dist-packages/torch/cuda/graphs.py", line 143, in replay
(EngineCore_DP0 pid=121)     super().replay()
(EngineCore_DP0 pid=121) torch.AcceleratorError: CUDA error: an illegal memory access was encountered

https://gist.githubusercontent.com/tonibofarull/da64693c9f94d566cae20761a58f910a/raw/3831d39700cdf8e28a9b7431594b5983ddb671de/gistfile1.txt

Before submitting a new issue...

Make sure you already searched for relevant issues, and asked the chatbot living at the bottom right corner of the documentation page, which can answer lots of frequently asked questions.

extent analysis

Fix Plan

The error message torch.AcceleratorError: CUDA error: an illegal memory access was encountered suggests a memory access issue on the GPU. To fix this, we can try the following steps:

Update CUDA and cuDNN versions: Ensure that the CUDA and cuDNN versions are compatible with the PyTorch version being used.
Check GPU memory utilization: Reduce the --gpu-memory-utilization parameter to a lower value (e.g., 0.5) to prevent GPU memory overflow.
Disable CUDA graph capture: Set --max-cudagraph-capture-size to 0 to disable CUDA graph capture, which can help reduce memory usage.
Update PyTorch and vLLM: Ensure that PyTorch and vLLM are updated to the latest versions.

Example code changes:

# Update CUDA and cuDNN versions
import torch
print(torch.cuda.get_device_name(0))  # Check CUDA device
print(torch.cuda.get_device_capability(0))  # Check CUDA capability

# Check GPU memory utilization
import gputil
gpu = gputil.getGPUs()[0]
print(gpu.memoryUtil * 100)  # Check GPU memory utilization

# Disable CUDA graph capture
# Set --max-cudagraph-capture-size to 0 in server arguments

# Update PyTorch and vLLM
# Run pip install --upgrade torch vllm

Verification

To verify that the fix worked, run the server with the updated parameters and check for any errors. You can also monitor GPU memory utilization using tools like nvidia-smi or gputil.

Extra Tips

Regularly update PyTorch, vLLM, and CUDA to ensure compatibility and fix any known issues.
Monitor GPU memory utilization and adjust the --gpu-memory-utilization parameter accordingly.
Consider using a lower --max-cudagraph-capture-size value to reduce memory usage.

Vote matrix · Quick signals

Works

Did the solution work? Tap to confirm.

Easy Fix

Was it a quick fix?

Time Saver

Did it save you time?

Blocking

Was it severely blocking?

Common Issue

Are others likely hitting this too?

Flaky / Intermittent

Is it intermittent?

Verified / Reproducible

Can you reproduce it reliably?

#api #ssr #installation #tensor shape #runtime error #dependency conflict #environment setup #environment variable

Still need to ship something?

×6

Another batch ranked right after the header list — different links, same matching logic.

Data

Security

Network

Code

UI/UX

Text

System

Multimedia

Protocol

API

Engineering

vllm - ✅(Solved) Fix [Bug]: CUDA illegal memory access on GPTQ Marlin [1 pull requests, 17 comments, 3 participants]

Recommended Tools

GitHub issue graph ai analysis

Error Message

Fix Action

Fix / Workaround

============================== CPU Info

PR fix notes

PR #36889: [WIP][Bugfix] Fix CUDA illegal memory access for GPTQ Marlin MoE with CUDA graphs

Description (problem / solution / changelog)

Purpose

Root Cause: `c_tmp` buffer out-of-bounds access

Fix (single file: `csrc/moe/marlin_moe_wna16/ops.cu`)

Test Plan

Test Result

Changed files

Code Example

Your current environment

🐛 Describe the bug

Server Arguments

Environment Variables

Logs

Before submitting a new issue...

extent analysis

Fix Plan

Verification

Extra Tips

Still need to ship something?

TRENDING

vllm - ✅(Solved) Fix [Bug]: CUDA illegal memory access on GPTQ Marlin [1 pull requests, 17 comments, 3 participants]

Recommended Tools

GitHub issue graph ai analysis

Error Message

Fix Action

Fix / Workaround

============================== CPU Info

PR fix notes

PR #36889: [WIP][Bugfix] Fix CUDA illegal memory access for GPTQ Marlin MoE with CUDA graphs

Description (problem / solution / changelog)

Purpose

Root Cause: c_tmp buffer out-of-bounds access

Fix (single file: csrc/moe/marlin_moe_wna16/ops.cu)

Test Plan

Test Result

Changed files

Code Example

Your current environment

🐛 Describe the bug

Server Arguments

Environment Variables

Logs

Before submitting a new issue...

extent analysis

Fix Plan

Verification

Extra Tips

Still need to ship something?

RELATED_DISCOVERY

TRENDING

Root Cause: `c_tmp` buffer out-of-bounds access

Fix (single file: `csrc/moe/marlin_moe_wna16/ops.cu`)