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