Skip to content

[https://nvbugs/6604925][fix] Swap the grid axes so N maps to the unbounded grid.x and M to grid.y… #17660

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
JunyiXu-nv merged 1 commit into from
Aug 18, 2026
+7 −5
Merged

[https://nvbugs/6604925][fix] Swap the grid axes so N maps to the unbounded grid.x and M to grid.y… #17660

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
11 changes: 7 additions & 4 deletions cpp/tensorrt_llm/kernels/weightOnlyBatchedGemv/cudaCoreGemmNVFP4.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2025, NVIDIA CORPORATION. All rights reserved.
* Copyright (c) 2025-2026, NVIDIA CORPORATION. All rights reserved.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -40,8 +40,8 @@ __device__ void cudaCoreGemmImpl(InputType const* __restrict__ act, InputType co
static constexpr SizeType32 nvfp4_scale_granularity = 16;
static constexpr SizeType32 step_k_scale = step_k / nvfp4_scale_granularity;
static constexpr SizeType32 tile_k = step_k * BLOCK_SIZE;
auto tile_id_m = static_cast<SizeType32>(blockIdx.x * TILE_M);
auto tile_id_n = static_cast<SizeType32>(blockIdx.y * TILE_N);
auto tile_id_m = static_cast<SizeType32>(blockIdx.y * TILE_M);
auto tile_id_n = static_cast<SizeType32>(blockIdx.x * TILE_N);
auto tid = static_cast<SizeType32>(threadIdx.x);
float tile_a[step_k];
float tile_w[TILE_N * step_k];
Expand Down Expand Up @@ -185,7 +185,10 @@ template <typename InputType, typename OutputType, typename ScaleType, SizeType3
void cudaCoreGemmKernel(Params const& params, cudaStream_t stream)
{
dim3 block(BLOCK_SIZE);
dim3 grid(params.m / TILE_M, params.n / TILE_N);
// N rides grid.x: its tile count is unbounded (a quantized LM head can be hundreds of
// thousands wide) and only grid.x allows more than 65535 blocks. M is safe on grid.y
// because it never exceeds cudaCoreGemmTemplateMaxM.
dim3 grid(params.n / TILE_N, params.m / TILE_M);

if (tensorrt_llm::common::getEnvEnablePDL())
{
Expand Down
1 change: 0 additions & 1 deletion tests/integration/test_lists/waives.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -90,7 +90,6 @@ accuracy/test_llm_api_pytorch.py::TestQwen3_5_397B_A17B::test_nvfp4[adp4_trtllm]
accuracy/test_llm_api_pytorch.py::TestQwen3_5_397B_A17B::test_nvfp4[tep4_trtllm] SKIP (https://nvbugs/6535779)
accuracy/test_llm_api_pytorch.py::TestQwen3_5_397B_A17B::test_nvfp4_4gpus_static_eplb[moe_backend=TRTLLM] SKIP (https://nvbugs/6418830)
accuracy/test_llm_api_pytorch.py::TestQwen3_5_397B_A17B::test_nvfp4_mtp3_gdn_replay_tep4 SKIP (https://nvbugs/6535779)
accuracy/test_llm_api_pytorch.py::TestQwen3_6_35B_A3B::test_nvfp4[TRTLLM] SKIP (https://nvbugs/6604925)
accuracy/test_llm_api_pytorch.py::TestQwen3_8B::test_bf16[latency] SKIP (https://nvbugs/6412098)
accuracy/test_llm_api_pytorch.py::TestQwen3_8B::test_w4a8_mxfp4[fp8-latency] SKIP (https://nvbugs/6566768)
accuracy/test_llm_api_pytorch_multimodal.py::TestMistralLarge3_675B::test_nvfp4_4gpus[latency_moe_trtllm] SKIP (https://nvbugs/6248827)
Expand Down
Loading