From 7f31204c680383e5dbaf148046e841c2e3959446 Mon Sep 17 00:00:00 2001 From: trtllm-agent <296075020+trtllm-agent@users.noreply.github.com> Date: Thu, 13 Aug 2026 13:12:51 -0700 Subject: [PATCH] [https://nvbugs/6604925][fix] Map NVFP4 cuda-core GEMM N tiles onto grid.x cudaCoreGemmKernel built its grid as (m / TILE_M, n / TILE_N), placing the N tile count on grid.y. CUDA caps grid.y at 65535 while grid.x allows 2^31-1, so a wide output overflowed the axis and every launch returned cudaErrorInvalidArgument. Qwen3.6-35B-A3B has an NVFP4 LM head of N=248320, which with TILE_N=2 needs 124160 blocks. Swap the axes so N rides grid.x and M rides grid.y; M is bounded by cudaCoreGemmTemplateMaxM (16) and can never overflow. The tile-id reads are swapped to match, which are the only blockIdx uses in the kernel. The autotuner caught the failing tactic and fell back to a working backend, so this usually surfaced only as a profiling warning -- hence the ~1% flake rate rather than a hard failure. Verified bit-exact (max|diff| = 0) against nvfp4_gemm_cutlass for m in {1,4,8,16} and n up to 248320, and the 65536-block boundary that previously failed to launch now succeeds. Removes the waiver this bug added. Signed-off-by: trtllm-agent <296075020+trtllm-agent@users.noreply.github.com> --- .../weightOnlyBatchedGemv/cudaCoreGemmNVFP4.cu | 11 +++++++---- tests/integration/test_lists/waives.txt | 1 - 2 files changed, 7 insertions(+), 5 deletions(-) diff --git a/cpp/tensorrt_llm/kernels/weightOnlyBatchedGemv/cudaCoreGemmNVFP4.cu b/cpp/tensorrt_llm/kernels/weightOnlyBatchedGemv/cudaCoreGemmNVFP4.cu index 8df4bf96b616..5c8f6615d8b2 100644 --- a/cpp/tensorrt_llm/kernels/weightOnlyBatchedGemv/cudaCoreGemmNVFP4.cu +++ b/cpp/tensorrt_llm/kernels/weightOnlyBatchedGemv/cudaCoreGemmNVFP4.cu @@ -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. @@ -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(blockIdx.x * TILE_M); - auto tile_id_n = static_cast(blockIdx.y * TILE_N); + auto tile_id_m = static_cast(blockIdx.y * TILE_M); + auto tile_id_n = static_cast(blockIdx.x * TILE_N); auto tid = static_cast(threadIdx.x); float tile_a[step_k]; float tile_w[TILE_N * step_k]; @@ -185,7 +185,10 @@ template