Verify each finding against the current code and only fix it if needed.

Inline comments:
In
`@cpp/tensorrt_llm/cutlass_extensions/include/cutlass_extensions/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input_.hpp`:
- Around line 1493-1527: The zero-mode branch for
ConversionMode::ConvertAndScaleWithZero updates only the zero tensormap but
omits updating the scale tensormap, causing tma_load_scale to read stale or
pointer-table data; in the functions that call
cute::tma_descriptor_replace_addr_in_shared_mem /
cute::tma_descriptor_replace_addr_in_global_mem (look for
tensormaps_replace_global_address/static_assert text and the branches using
KernelConversionMode), add the same scale update as in the ConvertAndScale
branch: replace the shared/global smem_tensormap_scale/get<2>(input_tensormaps)
with mainloop_params.ptr_S[next_batch] wherever the ConvertAndScaleWithZero
branch touches A/zero metadata, and also ensure smem_tensormap_scale is
initialized/staged inside tensormaps_init() so per-batch scale updates exist for
grouped GEMM; repeat the same fix in the equivalent locations (the other similar
blocks noted in the review).
- Around line 1015-1016: The LDSM thread-slice is using the global thread_idx
instead of the warpgroup-local index, causing incorrect lane mapping when
multiple consumer warpgroups participate; change the call to
smem_tiled_copy_A_LDSM.get_thread_slice(thread_idx) to use
get_thread_slice(warp_group_thread_idx) (consistent with the earlier A-copy path
and other SM90 collectives) so SmemCopyAtomA_LDSM / make_tiled_copy_A with
tiled_mma uses the warpgroup-local index for the slice.

In `@cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp`:
- Around line 313-314: The single-line if statement checking "if (tile_config ==
CutlassTileConfigSM90::CtaShape128x128x128B) continue;" should be rewritten to
use brace-delimited body per style guidelines: locate the check in
cutlass_heuristic.cpp (the loop handling tile_config values) and replace the
single-statement form with an explicit block using braces around the continue;
so the condition and its action are enclosed in { } to ensure consistent
formatting and avoid style warnings.

In `@cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.cu`:
- Around line 104-106: The two wrapper functions that currently launch
interleave_fp4_weights_for_Hopper_mixed_gemm_kernel with <<<1024, block>>> must
accept a cudaStream_t stream parameter, use a computed grid (e.g., dim3
grid(1024,1,1) or existing grid variable) and launch with <<<grid, block, 0,
stream>>> instead of the implicit default stream, and then call
cudaGetLastError() (or cudaPeekAtLastError()) immediately after the launch to
catch failures; update both function signatures and all call sites (e.g., where
weightOnlyQuantOp.cpp invokes these wrappers) to pass the caller's stream, and
follow the same error-checking pattern used in unfusedAttentionKernels.cu /
cudaCoreGemmNVFP4.cu.

In `@cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_mixed_utils.h`:
- Line 17: Replace the lone `#pragma once` with the repository-required include
guard: add a preprocessor guard using the macro name
TRTLLM_MOE_GEMM_MIXED_UTILS_H (open with `#ifndef` TRTLLM_MOE_GEMM_MIXED_UTILS_H /
`#define` TRTLLM_MOE_GEMM_MIXED_UTILS_H and close with `#endif`), ensuring the macro
matches the header filename moe_gemm_mixed_utils.h in all caps and without
directory names or trailing underscores; update the top and bottom of the file
(moe_gemm_mixed_utils.h) accordingly so other symbols in this header (e.g.,
declarations used by moe_gemm_mixed_utils) remain unchanged.

In `@cpp/tensorrt_llm/thop/weightOnlyQuantOp.cpp`:
- Around line 412-435: Add explicit shape validation before computing n/k and
before dispatching to the Hopper kernels: ensure weight.dim() == 2,
weight.size(0) % 16 == 0, and weight.size(1) % 32 == 0, and emit TORCH_CHECK
errors if any condition fails. Place these checks at the start of the
function/block handling the public entrypoint (before the lines that compute n =
weight.size(0) and k = weight.size(1) * 2) so that inputs violating the kernel
tile assumptions are rejected prior to calling
interleave_int4_weights_for_Hopper_mixed_gemm or
interleave_fp4_weights_for_Hopper_mixed_gemm.
- Around line 416-427: The output tensor is allocated on the current CUDA device
instead of the input tensor's device, causing mismatched allocations on
multi-GPU; add `#include` <c10/cuda/CUDAGuard.h>, create a c10::cuda::CUDAGuard
guard{weight.device()}; (or at::cuda::CUDAGuard guard(weight.device().index());)
immediately before allocating weight_interleaved so the torch::empty call uses
the same device as weight, and ensure subsequent get_ptr<uint8_t>(weight) /
get_ptr<uint8_t>(weight_interleaved) dereference the correct device pointers.

In `@tensorrt_llm/_torch/modules/fused_moe/quantization.py`:
- Around line 1735-1737: The interleave op
torch.ops.trtllm.interleave_4bit_weights_for_Hopper_mixed_gemm is CUDA-only, so
before calling it move the FP4 shard tensors to CUDA (e.g., w31_weight_shard =
w31_weight_shard.to(dst_device or torch.device('cuda'), non_blocking=True) or
use .cuda()) and ensure the same for the other shard(s) used in the nearby block
(the similar call around the w13/w31 handling at the later location). Update the
code that calls interleave_4bit_weights_for_Hopper_mixed_gemm to stage the
shard(s) onto CUDA first (preserving dtype/contiguity) so the CUDA-only op runs
on GPU-hosted tensors.
- Around line 1341-1344: The bare .cuda() calls moving w31_weight_shard to the
current CUDA device break multi-GPU placement; replace those .cuda() uses (the
ones before calling
torch.ops.trtllm.interleave_4bit_weights_for_Hopper_mixed_gemm) with an explicit
.to(device) (using the previously captured device variable `device =
dst_w3_w1_weight.device`) so the tensor is moved to the intended device before
invoking the Hopper interleave operator.

In `@tests/unittest/_torch/modules/test_fused_moe_jiangs.py`:
- Around line 229-234: The reference dequant path loads q2 and q3 swapped: when
weight_loading_mode == MoEWeightLoadingMode.VANILLA set q2 and q3 from the
correct weight keys (q2 should come from f"{e_idx}.w2.{lut['weight_scale_2']}"
and q3 from f"{e_idx}.w3.{lut['weight_scale_2']}") so fc2 receives the proper
tensor and q3_q1 combines w3 with w1 for fc13; update both occurrences (the
block around q1/q2/q3/q3_q1 and the similar block at lines ~263-275) to swap the
assignments accordingly while keeping q1 loaded from w1.
- Around line 466-469: The test currently forces a runner lookup with
next(iter(MoERunner.runner_dict.values())) which can raise StopIteration when
the backend doesn't populate MoERunner.runner_dict; change this to a safe
retrieval (e.g., use next(iter(...), None) or guard with a try/except) and
ensure any use of cpp_runner (the commented diagnostic block) only runs when
cpp_runner is not None so the test won't error when the registry is empty.
- Around line 1-3: This test module is missing the required NVIDIA Apache 2.0
header; add the repository-mandated NVIDIA copyright/license block at the very
top of the file (above the import lines) using the current year of the latest
meaningful modification and the standard Apache-2.0 header format used across
the repo so the file (test_fused_moe_jiangs.py) includes the exact NVIDIA
copyright and Apache 2.0 license block.

---

Outside diff comments:
In `@cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp`:
- Line 2: Update the file header in cutlass_heuristic.cpp to reflect the current
modified-year range and include the NVIDIA copyright header; specifically change
the existing "2020-2023" year span to include the current year (e.g.,
"2020-2026") and ensure the full NVIDIA copyright header text used across the
repo is present and matches formatting of other modified files.

In
`@cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_mixed_input_launcher.inl`:
- Line 2: Update the copyright header in the file
moe_gemm_tma_ws_mixed_input_launcher.inl to reflect the current modification
year (change the existing "2020-2023" to include the current year), i.e., update
the top-of-file NVIDIA copyright comment block in
cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_mixed_input_launcher.inl
so the range covers through the current year; keep the existing header text and
formatting intact except for the year range.

In `@cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu`:
- Line 2: Update the copyright header in the top-of-file comment in
moe_kernels.cu to include 2026 (e.g., change "2020-2025" to "2020-2026") so the
modified source file reflects the current year; locate the header comment near
the top of cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu (the
file-level copyright comment) and update the year range accordingly.

---

Nitpick comments:
In
`@cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_mixed_input_launcher.inl`:
- Around line 262-263: Replace the magic literal 2 used for
arguments.scheduler.max_swizzle_size with a named constant (e.g.,
kDefaultMaxSwizzleSize or MAX_SWIZZLE_SIZE) and use that constant in the
assignment to arguments.scheduler.max_swizzle_size; define the constant as a
constexpr integer with an appropriate scope (file-scope in this implementation
file or in the module header alongside other scheduler constants) so calls like
arguments.scheduler.max_swizzle_size = kDefaultMaxSwizzleSize; replace the
hard-coded 2 and keep the adjacent use of RasterOrderOptions::Heuristic
unchanged.

In `@cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_kernels.cu`:
- Line 4616: Remove the dead local variable declaration [[maybe_unused]] bool
use_w4_groupwise = use_w4afp8 || use_wfp4a16; in moe_kernels.cu (the computed
value is never read); simply delete this unused variable (or if the intention
was to use it later, replace unused declaration with the appropriate logic that
actually references use_w4_groupwise in the relevant control flow) so there is
no unused local symbol left.
- Around line 2157-2158: Declare the local variables that are not modified after
initialization as const: change the declarations of global_scale_val and
num_tokens_before_expert to use const (e.g., const float global_scale_val = ...;
const int64_t num_tokens_before_expert = ...;). Apply the same change to the
other immutable locals noted in the diff region around 2402-2407 so any
variables that are initialized once and never reassigned become const (use the
exact symbol names from those lines when updating).

In `@tests/unittest/_torch/modules/test_fused_moe_jiangs.py`:
- Around line 667-680: The constructor currently uses a shared ModelConfig
default (model_config: ModelConfig = ModelConfig()), which can be mutated across
tests; change the signature to accept model_config: Optional[ModelConfig] = None
and inside __init__ set self.model_config = model_config if model_config is not
None else ModelConfig() so each instance gets a fresh ModelConfig; update any
references to model_config in the constructor/body accordingly (look for the
__init__ method and the symbol model_config in this class).