docs(csrc): condense verbose comments (comments-only, no code change)

csrc/ck_batched_gemm_a8w8/batched_gemm_a8w8.cu

@@ -15,8 +15,7 @@ using BatchedRowwiseKernel = torch::Tensor (*)(torch::Tensor&,
                                                std::optional<torch::Tensor>,
                                                int);
 
-// For certain high priority shapes, we directly use the best kernel rather
-// than use heuristics.
+// For certain high priority shapes, we directly use the best kernel rather than use heuristics.
 using BatchedRowwiseKernelMap = BatchedGemmDispatchMap<BatchedRowwiseKernel>;
 
 template <typename DDataType, typename EDataType = DDataType>

csrc/ck_batched_gemm_a8w8/batched_gemm_a8w8_tune.cu

@@ -12,8 +12,7 @@ using BatchedRowwiseKernel = std::function<
                   torch::Tensor &, std::optional<torch::Tensor>,
                   int)>;
 
-// For certain high priority shapes, we directly use the best kernel rather
-// than use heuristics.
+// For certain high priority shapes, we directly use the best kernel rather than use heuristics.
 using BatchedRowwiseKernelMap = std::unordered_map<
     int,
     BatchedRowwiseKernel>;

csrc/ck_batched_gemm_bf16/batched_gemm_bf16.cu

@@ -10,8 +10,7 @@
 using BatchedKernel = torch::Tensor (*)(
     torch::Tensor&, torch::Tensor&, torch::Tensor&, std::optional<torch::Tensor>, int);
 
-// For certain high priority shapes, we directly use the best kernel rather
-// than use heuristics.
+// For certain high priority shapes, we directly use the best kernel rather than use heuristics.
 using BatchedKernelMap = BatchedGemmDispatchMap<BatchedKernel>;
 
 BatchedKernel batched_heuristic_dispatch(int B, int M, int N, int K)

csrc/ck_batched_gemm_bf16/batched_gemm_bf16_tune.cu

@@ -11,8 +11,7 @@ using BatchedKernel = std::function<
                   torch::Tensor &, std::optional<torch::Tensor>,
                   int)>;
 
-// For certain high priority shapes, we directly use the best kernel rather
-// than use heuristics.
+// For certain high priority shapes, we directly use the best kernel rather than use heuristics.
 using BatchedKernelMap = std::unordered_map<
     int,
     BatchedKernel>;

csrc/ck_deepgemm/deepgemm.cu

@@ -15,8 +15,7 @@ using RowwiseKernel = torch::Tensor (*)(torch::Tensor&,
                                         std::optional<torch::Tensor>,
                                         std::optional<torch::Tensor>);
 
-// For certain high priority shapes, we directly use the best kernel rather
-// than use heuristics.
+// For certain high priority shapes, we directly use the best kernel rather than use heuristics.
 using RowwiseKernelMap = GemmDispatchMap<RowwiseKernel>;
 
 template <typename ABDataType, typename AccDataType, typename CDataType>

csrc/ck_gemm_a4w4_blockscale/gemm_a4w4_blockscale.cu

@@ -14,20 +14,16 @@
 using BlockwiseKernel = torch::Tensor (*)(
     torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, int);
 
-// Name-keyed dispatch table; see gemm_a8w8_blockscale.cu for the rationale
-// behind std::string_view keys + raw fn-ptr values (constant-init into
-// .data.rel.ro, matching PR #3255's GemmDispatchMap style).
+// Name-keyed dispatch table (std::string_view keys + raw fn-ptr values); see gemm_a8w8_blockscale.cu for the rationale.
 //
-// Note: the a4w4 tuned CSV mixes CK and ASM kernels; the ASM rows are
-// filtered out at codegen (see csrc/ck_gemm_a4w4_blockscale/gen_instances.py)
-// and dispatched separately by aiter/ops/gemm_op_a4w4.py based on the
-// kernelName prefix (_ZN... = mangled ASM symbol).  This map only contains
-// CK kernels.
+// Note: the a4w4 tuned CSV mixes CK and ASM kernels; the ASM rows are filtered
+// out at codegen (gen_instances.py) and dispatched separately by
+// aiter/ops/gemm_op_a4w4.py by kernelName prefix (_ZN... = mangled ASM symbol).
+// This map only contains CK kernels.
 using BlockwiseKernelMap = std::unordered_map<std::string_view, BlockwiseKernel>;
 
-// Python-driven name-keyed dispatch (see gemm_a8w8_blockscale.cu for the
-// rationale).  Empty kernelName -> default heuristic; non-empty but unknown
-// kernelName -> hard error.
+// Python-driven name-keyed dispatch (see gemm_a8w8_blockscale.cu for the rationale).
+// Empty kernelName -> default heuristic; non-empty but unknown kernelName -> hard error.
 template <typename CDataType>
 BlockwiseKernel blockscale_dispatch(const std::string& kernelName)
 {

csrc/ck_gemm_a4w4_blockscale/gemm_a4w4_blockscale_tune.cu

@@ -11,8 +11,7 @@ using BlockwiseKernel = std::function<
                   torch::Tensor &, torch::Tensor &,
                   torch::Tensor &, int)>;
 
-// For certain high priority shapes, we directly use the best kernel rather
-// than use heuristics.
+// For certain high priority shapes, we directly use the best kernel rather than use heuristics.
 using BlockwiseKernelMap = std::unordered_map<
     int,
     BlockwiseKernel>;

csrc/ck_gemm_a8w8/gemm_a8w8.cu

@@ -16,8 +16,7 @@ using RowwiseKernel = torch::Tensor (*)(torch::Tensor&,
                                         std::optional<torch::Tensor>,
                                         int);
 
-// For certain high priority shapes, we directly use the best kernel rather
-// than use heuristics.
+// For certain high priority shapes, we directly use the best kernel rather than use heuristics.
 using RowwiseKernelMap = GemmDispatchMap<RowwiseKernel>;
 
 template <typename ABDataType, typename DDataType, typename EDataType>

csrc/ck_gemm_a8w8/gemm_a8w8_tune.cu

@@ -13,8 +13,7 @@ using RowwiseKernel = std::function<
                   torch::Tensor &, std::optional<torch::Tensor>,
                   int)>;
 
-// For certain high priority shapes, we directly use the best kernel rather
-// than use heuristics.
+// For certain high priority shapes, we directly use the best kernel rather than use heuristics.
 using RowwiseKernelMap = std::unordered_map<
     int,
     RowwiseKernel>;

csrc/ck_gemm_a8w8_blockscale/gemm_a8w8_blockscale.cu

@@ -8,38 +8,28 @@
 #include <torch/extension.h>
 
 #include "gemm_a8w8_blockscale_common.cuh"
-// lookup.h is included by gemm_a8w8_blockscale_lookup_fp16.cu and
-// gemm_a8w8_blockscale_lookup_bf16.cu, not here, so ninja can expand
-// GENERATE_LOOKUP_TABLE for FP16 and BF16 in parallel.
+// lookup.h is included by gemm_a8w8_blockscale_lookup_fp16.cu and gemm_a8w8_blockscale_lookup_bf16.cu, not here, so
+// ninja can expand GENERATE_LOOKUP_TABLE for FP16 and BF16 in parallel.
 #include "gemm_a8w8_blockscale_manifest.h"
 
 using BlockwiseKernel = torch::Tensor (*)(
     torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, int);
 
-// Name-keyed dispatch table.  Keys are std::string_view onto the kernel-name
-// string literals embedded in the generated *_lookup.h (static storage,
-// permanently valid).  Values are raw function pointers so the table is
-// trivially destructible and gets constant-initialized into .data.rel.ro,
-// matching the style of GemmDispatchMap introduced for the tuple-keyed
-// modules in PR #3255 (no per-pair std::function ctor / landing pad cost).
+// Name-keyed dispatch table.  Keys are std::string_view onto the kernel-name string literals in the generated
+// *_lookup.h (static storage, permanently valid).  Values are raw fn ptrs so the table is trivially destructible and
+// constant-initialized into .data.rel.ro (no per-pair std::function cost).
 using BlockwiseKernelMap = std::unordered_map<std::string_view, BlockwiseKernel>;
 
 // Defined in gemm_a8w8_blockscale_lookup_{fp16,bf16}.cu.
 extern const BlockwiseKernelMap& get_blockscale_lookup_fp16();
 extern const BlockwiseKernelMap& get_blockscale_lookup_bf16();
 
-// Python-driven name-keyed dispatch.  The Python frontend
-// (aiter/ops/gemm_op_a8w8.py) reads aiter/configs/a8w8_blockscale_tuned_gemm.csv
-// (cached via lru_cache) and passes the resolved `kernelName` here; we only
-// look it up.  This makes the CSV the single source of truth — editing it no
-// longer requires rebuilding a tuple-keyed C++ table.
+// Python-driven name-keyed dispatch.  The Python frontend (aiter/ops/gemm_op_a8w8.py) reads
+// a8w8_blockscale_tuned_gemm.csv and passes the resolved `kernelName` here; the CSV is the single source of truth.
 //
-// Empty kernelName  : Python had no tuned row (or AITER_BYPASS_TUNE_CONFIG=1)
-//                     -> use the heuristic default kernel.
-// Non-empty kernelName not in registry: hard error — the CSV references a
-//                     kernel that was not compiled into this .so, almost
-//                     certainly because the user updated the CSV without
-//                     rebuilding aiter.  Surface, don't hide.
+// Empty kernelName -> Python had no tuned row (or AITER_BYPASS_TUNE_CONFIG=1); use the heuristic default kernel.
+// Non-empty kernelName not in registry -> hard error: the CSV references a
+//   kernel not compiled into this .so (CSV updated without rebuilding aiter).
 template <typename DDataType, typename EDataType = DDataType>
 static BlockwiseKernel blockscale_dispatch(const std::string& kernelName)
 {

csrc/ck_gemm_a8w8_blockscale/gemm_a8w8_blockscale_cktile.cu

@@ -14,14 +14,12 @@
 using BlockwiseKernel = torch::Tensor (*)(
     torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, bool, int);
 
-// Name-keyed dispatch table; see gemm_a8w8_blockscale.cu for the rationale
-// behind std::string_view keys + raw fn-ptr values (constant-init into
-// .data.rel.ro, matching PR #3255's GemmDispatchMap style).
+// Name-keyed dispatch table; see gemm_a8w8_blockscale.cu for the rationale behind std::string_view keys + raw fn-ptr values
+// (constant-init into .data.rel.ro, matching PR #3255's GemmDispatchMap style).
 using BlockwiseKernelMap = std::unordered_map<std::string_view, BlockwiseKernel>;
 
-// Python-driven name-keyed dispatch (see gemm_a8w8_blockscale.cu for the
-// rationale).  Empty kernelName -> default heuristic; non-empty but unknown
-// kernelName -> hard error.
+// Python-driven name-keyed dispatch (see gemm_a8w8_blockscale.cu for the rationale).
+// Empty kernelName -> default heuristic; non-empty but unknown kernelName -> hard error.
 template <typename DDataType, typename EDataType = DDataType>
 static BlockwiseKernel blockscale_dispatch(const std::string& kernelName)
 {

csrc/ck_gemm_a8w8_blockscale/gemm_a8w8_blockscale_cktile_tune.cu

@@ -14,8 +14,7 @@
 using BlockwiseKernel = std::function<torch::Tensor(
     torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, bool, int)>;
 
-// For certain high priority shapes, we directly use the best kernel rather
-// than use heuristics.
+// For certain high priority shapes, we directly use the best kernel rather than use heuristics.
 using BlockwiseKernelMap = std::unordered_map<int, BlockwiseKernel>;
 
 template <typename DDataType, typename EDataType = DDataType>

csrc/ck_gemm_a8w8_blockscale/gemm_a8w8_blockscale_tune.cu

@@ -14,8 +14,7 @@
 using BlockwiseKernel = std::function<torch::Tensor(
     torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, int)>;
 
-// For certain high priority shapes, we directly use the best kernel rather
-// than use heuristics.
+// For certain high priority shapes, we directly use the best kernel rather than use heuristics.
 using BlockwiseKernelMap = std::unordered_map<int, BlockwiseKernel>;
 
 template <typename DDataType, typename EDataType = DDataType>

csrc/ck_gemm_a8w8_blockscale/include/gemm_a8w8_blockscale_cktile_common.cuh

@@ -312,8 +312,7 @@ __forceinline__ torch::Tensor gemm_a8w8_blockscale_cktile_impl(torch::Tensor& XQ
     const int N = WQ.size(0);
     const int K = XQ.size(1);
 
-    // Whether this kernel configuration uses column-major AQ layout,
-    // requiring a host-side transpose of x_scale.
+    // Whether this kernel configuration uses column-major AQ layout, requiring a host-side transpose of x_scale.
     constexpr bool aq_col_major =
         BQuantGroupSize::kN == 128 &&
         (GemmInstance::M_Warp_v * GemmInstance::N_Warp_v * GemmInstance::K_Warp_v == 8) &&
@@ -329,8 +328,7 @@ __forceinline__ torch::Tensor gemm_a8w8_blockscale_cktile_impl(torch::Tensor& XQ
     ck_tile::QuantGemmHostArgs args;
     args.a_ptr = XQ.data_ptr();
 
-    // Declared at function scope so the transposed tensor stays alive
-    // through the async kernel launch.
+    // Declared at function scope so the transposed tensor stays alive through the async kernel launch.
     torch::Tensor x_scale_t;
 
     if constexpr(aq_col_major)
@@ -377,10 +375,9 @@ __forceinline__ torch::Tensor gemm_a8w8_blockscale_cktile_impl(torch::Tensor& XQ
     const int BQK = ck_tile::integer_divide_ceil(K, BQuantGroupSize::kK);
     const int BQN = ck_tile::integer_divide_ceil(N, BQuantGroupSize::kN);
 
-    // Read leading-dimension strides from tensor metadata instead of
-    // assuming dense layout.  vLLM's _maybe_pad_fp8_weight can produce
-    // row-major tensors whose leading-dimension stride exceeds the
-    // logical column count (e.g. shape [N,K] with stride [K+pad, 1]).
+    // Read leading-dimension strides from tensor metadata instead of assuming dense layout.  vLLM's _maybe_pad_fp8_weight
+    // can produce row-major tensors whose leading-dimension stride exceeds the logical column count
+    // (e.g. shape [N,K] with stride [K+pad, 1]).
     const int stride_A  = XQ.stride(0);
     const int stride_B  = WQ.stride(0);
     const int stride_C  = Y.stride(0);

csrc/ck_gemm_a8w8_blockscale_bpreshuffle/gemm_a8w8_blockscale_bpreshuffle.cu

@@ -14,14 +14,12 @@
 using BlockwiseKernel = torch::Tensor (*)(
     torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&);
 
-// Name-keyed dispatch table; see gemm_a8w8_blockscale.cu for the rationale
-// behind std::string_view keys + raw fn-ptr values (constant-init into
-// .data.rel.ro, matching PR #3255's GemmDispatchMap style).
+// Name-keyed dispatch table; see gemm_a8w8_blockscale.cu for the rationale behind std::string_view keys + raw fn-ptr values
+// (constant-init into .data.rel.ro, matching PR #3255's GemmDispatchMap style).
 using BlockwiseKernelMap = std::unordered_map<std::string_view, BlockwiseKernel>;
 
-// Python-driven name-keyed dispatch (see gemm_a8w8_blockscale.cu for the
-// rationale).  Empty kernelName -> default heuristic; non-empty but unknown
-// kernelName -> hard error.
+// Python-driven name-keyed dispatch (see gemm_a8w8_blockscale.cu for the rationale).
+// Empty kernelName -> default heuristic; non-empty but unknown kernelName -> hard error.
 template <typename DDataType, typename EDataType = DDataType>
 BlockwiseKernel blockscale_bpreshuffle_dispatch(const std::string& kernelName)
 {

csrc/ck_gemm_a8w8_blockscale_bpreshuffle/gemm_a8w8_blockscale_bpreshuffle_tune.cu

@@ -9,8 +9,7 @@
 using BlockwiseKernel = std::function<torch::Tensor(
     torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&)>;
 
-// For certain high priority shapes, we directly use the best kernel rather
-// than use heuristics.
+// For certain high priority shapes, we directly use the best kernel rather than use heuristics.
 using BlockwiseKernelMap = std::unordered_map<int, BlockwiseKernel>;
 
 // Helper function to return the next largest power of 2

csrc/ck_gemm_a8w8_bpreshuffle/gemm_a8w8_bpreshuffle_tune.cu

@@ -10,8 +10,7 @@
 using BlockwiseKernel = std::function<torch::Tensor(
     torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, int)>;
 
-// For certain high priority shapes, we directly use the best kernel rather
-// than use heuristics.
+// For certain high priority shapes, we directly use the best kernel rather than use heuristics.
 using BlockwiseKernelMap = std::unordered_map<int, BlockwiseKernel>;
 
 // Helper function to return the next largest power of 2

csrc/ck_gemm_moe_2stages_codegen/gemm_moe_ck2stages_common.cuh

@@ -41,7 +41,7 @@ void ck_moe_stage1_gemm(const hipStream_t& stream,
                         std::optional<int> splitk,     // splitk
                         std::optional<bool> nt)
 {
-    // ~~~~~~~~~~~~~~~~~~~~~~~~following start with ck things
+    // CK setup below
     ck::index_t StrideA = K;
     ck::index_t StrideB = K;
     ck::index_t StrideD = 0;
@@ -230,7 +230,7 @@ void ck_moe_stage2_gemm(const hipStream_t& stream,
                         std::optional<int> splitk,     // splitk
                         std::optional<bool> nt)
 {
-    // ~~~~~~~~~~~~~~~~~~~~~~~~following start with ck things
+    // CK setup below
     ck::index_t StrideA = K;
     ck::index_t StrideB = K;
     ck::index_t StrideD = 0;

csrc/ck_gemm_moe_2stages_codegen/gemm_moe_ck2stages_common_mxfp4.cuh

@@ -42,7 +42,7 @@ void ck_moe_stage1_gemm(const hipStream_t& stream,
                         std::optional<int> splitk,     // splitk
                         std::optional<bool> nt)
 {
-    // ~~~~~~~~~~~~~~~~~~~~~~~~following start with ck things
+    // CK setup below
     using A1DataType                            = E8M0;
     using B1DataType                            = E8M0;
     static constexpr ck::index_t ScaleBlockSize = 32;
@@ -235,7 +235,7 @@ void ck_moe_stage2_gemm(const hipStream_t& stream,
                         std::optional<int> splitk,     // splitk
                         std::optional<bool> nt)
 {
-    // ~~~~~~~~~~~~~~~~~~~~~~~~following start with ck things
+    // CK setup below
     using A1DataType                            = E8M0;
     using B1DataType                            = E8M0;
     static constexpr ck::index_t ScaleBlockSize = 32;

csrc/ck_gemm_moe_2stages_codegen/gemm_moe_ck2stages_common_mxfp4_bns.cuh

@@ -41,7 +41,7 @@ void ck_moe_stage1_gemm(const hipStream_t& stream,
                         std::optional<int> splitk,     // splitk
                         std::optional<bool> nt)
 {
-    // ~~~~~~~~~~~~~~~~~~~~~~~~following start with ck things
+    // CK setup below
     using A1DataType                            = E8M0;
     using B1DataType                            = E8M0;
     static constexpr ck::index_t ScaleBlockSize = 32;
@@ -233,7 +233,7 @@ void ck_moe_stage2_gemm(const hipStream_t& stream,
                         std::optional<int> splitk,     // splitk
                         std::optional<bool> nt)
 {
-    // ~~~~~~~~~~~~~~~~~~~~~~~~following start with ck things
+    // CK setup below
     using A1DataType                            = E8M0;
     using B1DataType                            = E8M0;
     static constexpr ck::index_t ScaleBlockSize = 32;

csrc/cktile_gemm_a8w8_bpreshuffle/gemm_a8w8_bpreshuffle_cktile_tune.cu

@@ -10,8 +10,7 @@
 using RowwiseKernel = std::function<torch::Tensor(
     torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, torch::Tensor&, int)>;
 
-// For certain high priority shapes, we directly use the best kernel rather
-// than use heuristics.
+// For certain high priority shapes, we directly use the best kernel rather than use heuristics.
 using RowwiseKernelMap = std::unordered_map<int, RowwiseKernel>;
 
 // Helper function to return the next largest power of 2

csrc/cpp_itfs/mha_bwd.cu

@@ -460,10 +460,8 @@ float fmha_v3_bwd(mha_bwd_args a, const ck_tile::stream_config& s)
         AITER_LOG_WARNING("fmha_v3_bwd: unsupported mask type for asm kernels.");
         return -1;
     }
-    // On gfx942, a16 (atomic32=0) has no mask_type=2 (bottom-right causal) kernels,
-    // only mask_type=1 (top-left causal). When seqlen_q == seqlen_k the two masks
-    // are mathematically equivalent, so we can safely convert 2 → 1 to hit the
-    // existing a16 causal kernels.
+    // gfx942 a16 (atomic32=0) has no mask_type=2 (bottom-right) kernels, only mask_type=1
+    // (top-left); the two are equivalent when seqlen_q == seqlen_k, so convert 2 -> 1.
     // See config: hsa/gfx942/fmha_v3_bwd/fmha_bwd_dqdkdv.csv
     if(arch_id == "gfx942" && !a.v3_atomic_fp32 && mt == 2 && a.seqlen_q == a.seqlen_k)
     {