Repository navigation
[https://nvbugs/6939562][fix] Bound module_out_sizes read in LoRA param-fill kernel - #20477
Conversation
|
/bot run --disable-fail-fast |
|
PR_Github #78706 [ run ] triggered by Bot. Commit: |
|
@NVIDIA/trt-llm-kernels-devs: requesting a code-owner review. This is a one-hunk fix for NVBug 6939562 (post-merge
|
|
/bot run --disable-fail-fast |
|
PR_Github #78706 [ run ] completed with state
|
|
PR_Github #78713 [ run ] triggered by Bot. Commit: |
|
PR_Github #78713 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #78729 [ run ] triggered by Bot. Commit: |
|
PR_Github #78729 [ run ] completed with state
|
|
Pre-merge CI on head
That perf test passed in other PRs' builds 64839, 64840 and 64841. The LoRA tests that ran in #64848 passed. Could someone from CI infra look at the H100_PCIe perf pool, and could a code owner (@NVIDIA/trt-llm-kernels-devs) review the change? |
|
/bot run --disable-fail-fast |
|
PR_Github #78737 [ run ] triggered by Bot. Commit: |
|
PR_Github #78737 [ run ] completed with state
|
…am-fill kernel The OUT_SIZES block of loraGroupGEMMParamFillRowReorderFusionKernel runs with at least 256 threads, but module_out_sizes holds only MODULE_COUNT entries. Every thread past max_lora_count * MODULE_COUNT read module_out_sizes[module_id] out of bounds; the rows it fed are never stored, so the read was harmless only while the bytes after that small tensor happened to be mapped. Guard the read by module_id < MODULE_COUNT. Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>
739cdb2 to
df4e531
Compare
|
/bot run --disable-fail-fast |
|
PR_Github #78832 [ run ] triggered by Bot. Commit: |
|
@NVIDIA/trt-llm-kernels-devs: following up on the review request. This PR is a single bounds check in |
|
Automatically added "ci: full pre-merge approved" because this PR has satisfied the required GitHub review approvals. Unresolved review conversations and other required checks remain independent merge requirements. |
|
PR_Github #78832 [ run ] completed with state |
Description
examples/test_gpt.py::test_gpt_oss_20b_lora_torch[gpt-oss-20b-lora-adapter_NIM_r8-gpt-oss-20b]fails on H100 (L0_PostMerge 3021, DGX_H100-PyTorch-Post-Merge-1) withCUDA error: an illegal memory access was encounteredat the first CUDA-graph LoRA decode step.Root cause.
loraGroupGEMMParamFillRowReorderFusionKernelruns with at least 256 threads per block. In theOUT_SIZES_INDEXblock, every thread readsmodule_out_sizes[module_id], wheremodule_id = threadIdx.x / max_lora_count. That tensor has onlyMODULE_COUNTentries (3 int32 = 12 bytes for fused QKV), so threads pastmax_lora_count * MODULE_COUNTread out of bounds.BlockStore3drops their rows, so the result was never wrong, but the read faults whenever the allocation sits at the end of mapped memory. A GPU coredump names this kernel and block (Warp MMU Fault on themodule_out_sizesload). With the caching allocator disabled, compute-sanitizer reports 10 invalid 4-byte reads past the 12-byte allocation. #20188 changed the allocation layout enough that this test now hits it every time; the bug itself dates back to #8279.Fix. Read
module_out_sizesonly formodule_id < MODULE_COUNT. The output for real problems is unchanged.Test Coverage
On computelab H100-80GB with the CI image
pytorch-26.09-...-202610130652-20175:PYTORCH_NO_CUDA_MEMORY_CACHING=1): 10 errors before, 0 after;out_sizesis correct.tests/unittest/_torch/peft/{test_lora,test_lora_autotuner,test_fp8_lora_grouped_gemm_regressions,test_moe_lora_cuda_graph_params}.pyandthop/parallel_hw_agnostic/test_custom_ops.py: 37 passed.unittest/_torch/modules/tests_lora_modulesandtest_qwen_moe_routed_expert_multi_lora_varying_ranks[eager|cudagraph]: passed.No waiver exists for this bug.
Fixes https://nvbugs/6939562
🤖 Generated with Claude Code
Dev Engineer Review
module_id < MODULE_COUNTbefore readingmodule_out_sizes. This prevents out-of-range reads from threads beyond the valid module range.b_ptris non-null and the module ID is valid. The change does not alter the public API.QA Engineer Review
No test changes.
Per-File QA Perspective
cpp/tensorrt_llm/kernels/lora/loraGroupGEMMParamFillRowReorderFusion.cu: Verify output-size rows for valid modules and nullb_ptrvalues, plus threads beyondMODULE_COUNT. Confirm CUDA-graph LoRA decode no longer triggers invalid reads.