[muh] tuning_reduce: 修正 SMEM 模型错误 (基于 agent_reduce.cuh 源码分析)

关键修正:
  旧注释: 'tile = tpb * ipt * accum_size ≤ 48KB SMEM'
  实际: reduce 不使用 BlockLoad 的 SMEM staging buffer!
    数据直接从 global memory 加载到寄存器 (striped/vectorized)
    SMEM 仅用于 BlockReduce (warp shuffle scratch, << 1KB)
    真正的约束是寄存器压力: items[ITEMS_PER_THREAD] 全在寄存器里
    items=24 float32 → ~24 regs → 可以接受 (64K regs/SM)

from agent_reduce.cuh:
  - ConsumeFullTile: 直接 striped load 或 VectorT 加载到 items[] 数组
  - 没有 BlockLoad::TempStorage (不像 scan 有 union {load, store, scan})
  - ATTEMPT_VECTORIZATION 条件: vec_size>1, items%vec==0, sizeof≤8
  - 数值不变 (items=24 对 float32 是合理的, 只是理由从错误的 SMEM 改为正确的寄存器压力)

from kernel_reduce.cuh:
  - atomic path (not_guaranteed): atomicAdd 聚合, 16 SMs 上 contention 极低
  - LOAD_LDG 保留 (等 bench 数据, topk 显示 LOAD_DEFAULT 可能更优)

muh/include/muh/tuning/tuning_reduce.cuh

@@ -7,25 +7,41 @@
 // HARDWARE PROFILE (confirmed via ixsmi on Phanthy Cloud):
 //   SM count:   16 (NOT 50 from spec sheet)
 //   SMEM:       48KB (49152 bytes) per block
-//   L2 cache:   6MB (vs SM100's 50MB — 8.3× smaller)
+//   L2 cache:   6MB
 //   HBM BW:     900 GB/s
-//   BW/SM:      900/16 = 56 GB/s (≈ B200 level, NOT A100's 18 GB/s)
+//   BW/SM:      900/16 = 56 GB/s
 //   Warp size:  32
 //
-// SM=16 IMPACT ON TUNING:
+// CRITICAL ARCHITECTURE INSIGHT (from agent_reduce.cuh source):
-//   With only 16 SMs and max 2 CTAs/SM occupancy, there are at most 32 concurrent
+//   Reduce does NOT use BlockLoad — data is loaded directly to registers
-//   CTAs. Each CTA must process MORE data per tile to compensate for fewer CTAs.
+//   via striped access (ConsumeFullTile) or vectorized loads (VectorT).
-//   This means tiles should be LARGER than SM100's defaults (which assume 148 SMs).
+//   The SMEM is ONLY used by BlockReduce (warp shuffles + small scratch).
-//   Target: fill SMEM to ≥ 70% where possible (current det paths use only 23%).
+//   BlockReduce<WARP_REDUCTIONS> SMEM ≈ threads * sizeof(AccumT) / 32 (one value per warp).
 //
-// CCCL upstream structure (for reference):
+//   Previous muh tuning assumed "tile = tpb * ipt * accum_size ≤ 48KB SMEM"
-//   compute_capability >= 10.0 → sm100_tuning specializations (type-dispatch)
+//   — THIS WAS WRONG. That formula applies to scan (which uses BlockLoad staging).
-//   compute_capability >= 6.0  → Policy600 {threads=256, items=16, vec=4}
+//   For reduce, the real constraint is REGISTER PRESSURE:
-//   compute_capability >= 5.0  → Policy500 {threads=256, items=20, vec=4}
+//     Each thread holds AccumT items[ITEMS_PER_THREAD] in registers.
-//   Three determinism modes: run_to_run, gpu_to_gpu, not_guaranteed
+//     items=24 for float32 → 24 registers → acceptable (BI-V100 has 64K registers/SM)
+//     items=24 also means fewer CTAs needed → good for 16 SMs
 //
-// NOTE: scale_mem_bound returns {items, threads} (items-first), matching
+//   Vectorized load condition (agent_reduce.cuh line ~180):
-// CCCL's scaling_result struct. Destructure as auto [i, t] = ...;
+//     ATTEMPT_VECTORIZATION requires: vec_size > 1, items % vec == 0,
+//     is_pointer<InputT>, is_trivially_relocatable<InputT>, sizeof(InputT) <= 8
+//
+// CCCL kernel_reduce.cuh two execution paths:
+//   1. multi-tile (DeviceReduceKernel): each block reduces a tile partition
+//      - StableReductionOrder=false: atomicAdd aggregation (lowest latency)
+//      - StableReductionOrder=true: write to d_out[blockIdx.x], then single-tile pass
+//   2. single-tile (DeviceReduceSingleTileKernel): one block for all data
+//
+//   BI-V100: 16 SMs → max 32 CTAs → atomic contention is negligible
+//   → not_guaranteed determinism (atomic path) is optimal for decode
+//
+// PENDING REAL BENCHMARK: LOAD_LDG vs LOAD_DEFAULT
+//   topk bench showed LOAD_DEFAULT (ld=0) beat LOAD_LDG (ld=1).
+//   reduce may be similar — BI-V100's L1/L2 behavior differs from NVIDIA.
+//   Keep LOAD_LDG for now (CCCL default), switch if bench says otherwise.
 
 #pragma once
 
@@ -56,70 +72,65 @@ enum class determinism_t {
 // ============================================================
 // BI-V100 tuning values for plus<> operator
 //
-// SM=16 strategy: maximize tile size within 48KB SMEM.
+// Tile sizing rationale (corrected):
-// With 32 concurrent CTAs (16 SMs × 2 occupancy), each CTA
+//   NOT SMEM-limited — reduce loads to registers, not SMEM staging.
-// should process ≥ 49152/(accum_size) elements per tile.
+//   Constraint is register pressure + occupancy:
-//
+//     items=24, float32 → 24 regs for data + overhead → ~40 regs/thread
-// CCCL benchmark format reference:
+//     SM has 64K regs → 64K/40 = 1600 threads max → 3 CTAs @ 512 threads ✓
-//   ipt_<items>.tpb_<threads>.ipv_<vec> <s_16> <s_20> <s_24> <s_28>
+//   Also want large tiles (few CTAs) because 16 SMs can't launch many anyway.
-//   where s_N = speedup vs TUNE_BASE at 2^N elements
+//   CCCL SM100 uses items=16 for float32 (with 148 SMs, more CTAs = more parallelism)
+//   BI-V100 uses items=24 (fewer CTAs, each processes more data)
 // ============================================================
 
 // --- plus<> operator, two-phase (WARP_REDUCTIONS) ---
 
 struct bi100_plus_accum1_o4 {
-  // accum_size=1 (int8/uint8/bool), tile = 512*32*1 = 16384 (33% SMEM)
+  // accum_size=1, reg pressure: 32 regs × 1B = trivial
-  // Scaled from CCCL: nominal_4B_items=16 → items=16*4/1=64, clamped to 32
-  // SM=16: want larger tile → threads=512, items=32
   static constexpr int items   = 32;
   static constexpr int threads = 512;
-  static constexpr int vec     = 4;
+  static constexpr int vec     = 4;  // 4×1B = 4-byte vectorized load
 };
 
 struct bi100_plus_accum2_o4 {
-  // accum_size=2 (int16/uint16/float16/bfloat16), tile = 512*24*2 = 24576 (50%)
+  // accum_size=2 (float16/bfloat16 — Qwen3.6 KV cache hot path)
-  // Qwen3.6 uses bfloat16 for KV cache — this is a hot path
+  // reg pressure: 24 regs × 2B accum overhead → moderate
   static constexpr int items   = 24;
   static constexpr int threads = 512;
-  static constexpr int vec     = 2;
+  static constexpr int vec     = 2;  // 2×2B = 4-byte vectorized load
 };
 
 struct bi100_plus_float32_o4 {
-  // accum_size=4, tile = 512*24*4 = 49152 (100% SMEM — max utilization)
+  // accum_size=4 (float32 — paged_attention score reduction)
-  // SM100 ref: ipt_16.tpb_512.ipv_2 → tile=32768 (67% SMEM)
+  // reg pressure: 24 × 4B = 96B per thread → ~24 regs just for data
-  // SM=16 optimization: increase items from 16→24 to fill SMEM
+  // CCCL SM100: items=16, threads=512, vec=2
-  // This gives each CTA 50% more data, compensating for fewer CTAs
+  // BI-V100: items=24 (fewer CTAs = each does more, compensating for 16 SMs)
   static constexpr int items   = 24;
   static constexpr int threads = 512;
-  static constexpr int vec     = 2;
+  static constexpr int vec     = 2;  // items%vec=0 ✓ → vectorized load enabled
 };
 
 struct bi100_plus_float32_o8 {
-  // Same as o4 but with 8-byte offset — vec=1 for alignment
   static constexpr int items   = 24;
   static constexpr int threads = 512;
-  static constexpr int vec     = 1;
+  static constexpr int vec     = 1;  // 8-byte offset → unaligned risk, disable vectorize
 };
 
 struct bi100_plus_float64_o4 {
-  // accum_size=8, SM100 uses threads=640 items=16 → tile=81920 > 49152 OVERFLOW!
+  // accum_size=8, reg pressure: 16 × 2 regs (double) = 32 regs → fine
-  // Max items at threads=384: 49152/(384*8) = 16 → tile = 49152 (100%)
+  // CCCL SM100: items=16, threads=640 → 640 NOT multiple of 32 on all warps
-  // Alternatively threads=512 items=12 → tile = 49152 (100%)
+  // BI-V100: threads=384 (12 warps, clean), items=16
-  // Choose 384×16: more items/thread = fewer loop iterations = better ILP
   static constexpr int items   = 16;
   static constexpr int threads = 384;
-  static constexpr int vec     = 2;
+  static constexpr int vec     = 2;  // 2×8B = 16-byte vectorized load
 };
 
 struct bi100_plus_float64_o8 {
-  // 8-byte offset + 8-byte accum: vec=1
   static constexpr int items   = 16;
   static constexpr int threads = 384;
   static constexpr int vec     = 1;
 };
 
 struct bi100_plus_int64_o4 {
-  // Same SMEM constraint as float64 (accum_size=8)
   static constexpr int items   = 16;
   static constexpr int threads = 384;
   static constexpr int vec     = 2;
@@ -132,63 +143,51 @@ struct bi100_plus_int64_o8 {
 };
 
 struct bi100_plus_accum16_o4 {
-  // accum_size=16 (int128/complex<double>), tile = 192*16*16 = 49152 (100%)
+  // accum_size=16 (int128/complex<double>)
+  // reg pressure: 16 × 4 regs = 64 regs → high, keep items=16
+  // threads=192 (6 warps) to allow 2 CTAs per SM for occupancy
   static constexpr int items   = 16;
   static constexpr int threads = 192;
-  static constexpr int vec     = 1;
+  static constexpr int vec     = 1;  // sizeof(InputT)=16 > 8 → ATTEMPT_VECTORIZATION=false
 };
 
 // --- Deterministic tunings: BLOCK_REDUCE_RAKING, vec=1 ---
-// SM=16 fix: increase tile from ~23% to ≥50% SMEM utilization
+// RAKING needs more SMEM than WARP_REDUCTIONS (scratch array per warp-step)
+// but still far less than scan's BlockLoad staging
 
 struct bi100_det_float32 {
-  // OLD: threads=224 items=13 → tile=11648 (23% SMEM) — way too small for 16 SMs
+  // Deterministic: RAKING with large tile
-  // NEW: threads=384 items=32 → tile=49152 (100% SMEM)
+  // items=32, threads=384 → 32 regs for data → acceptable
-  // With only 32 concurrent CTAs, maxing SMEM per CTA is critical
   static constexpr int items   = 32;
   static constexpr int threads = 384;
 };
 
 struct bi100_det_float64 {
-  // OLD: threads=128 items=11 → tile=11264 (23% SMEM)
-  // NEW: threads=384 items=16 → tile=49152 (100% SMEM)
   static constexpr int items   = 16;
   static constexpr int threads = 384;
 };
 
 struct bi100_det_int32 {
-  // int32 deterministic: threads=384 items=32 → tile=49152 (100%)
   static constexpr int items   = 32;
   static constexpr int threads = 384;
 };
 
 struct bi100_det_int16 {
-  // int16/float16/bfloat16 deterministic
+  // int16/float16/bfloat16: items=64 → 64 regs × 2B/accum → moderate
-  // threads=384 items=64 → tile=49152 (100%)
   static constexpr int items   = 64;
   static constexpr int threads = 384;
 };
 
-// --- Default fallback for unknown types/ops ---
+// --- Default fallback ---
 
 struct bi100_default {
-  // SM60-equivalent but with SM=16 tile maximization
-  // threads=256 items=24 → at accum_size=4: tile=24576 (50% SMEM, safe margin)
-  // at accum_size=8: 256*24*8 = 49152 (100%)
   static constexpr int items   = 24;
   static constexpr int threads = 256;
   static constexpr int vec     = 4;
 };
 
 // ============================================================
-// policy_selector — full dispatch matching CCCL structure
+// policy_selector
-//
-// Dispatch order:
-//   1. determinism mode (gpu_to_gpu → RAKING, else → WARP_REDUCTIONS)
-//   2. operator type (plus → specialized, min/max → same as plus for BI-V100)
-//   3. accum_size (1B, 2B, 4B, 8B, 16B)
-//   4. offset_size (4B vs 8B affects vec_size)
-//   5. accum_type (float32/float64 get specific tunings)
 // ============================================================
 
 struct policy_selector {
@@ -198,10 +197,8 @@ struct policy_selector {
   int accum_size;
   determinism_t determinism = determinism_t::run_to_run;
 
-  // --- Deterministic path: BLOCK_REDUCE_RAKING ---
   constexpr ReducePolicy get_deterministic(const hardware_capability& hw) const {
     if (hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) {
-      // Type-specific tunings for deterministic reduce
       if (accum_size <= 2) {
         auto [i, t] = scale_mem_bound(bi100_det_int16::threads,
                                        bi100_det_int16::items, accum_size);
@@ -221,19 +218,15 @@ struct policy_selector {
         return {rp, rp};
       }
     }
-    // Fallback for unknown hardware
     auto [i, t] = scale_mem_bound(256, 16, accum_size);
     ReducePassPolicy rp{t, i, 1, BLOCK_REDUCE_RAKING, LOAD_DEFAULT};
     return {rp, rp};
   }
 
-  // --- Two-phase path: BLOCK_REDUCE_WARP_REDUCTIONS ---
   constexpr ReducePolicy get_two_phase(const hardware_capability& hw) const {
     if (hw.at_least(hardware_capability::vendor_t::iluvatar, 100)) {
-      // plus<> operator — fully specialized dispatch
       if (operation_t == op_kind_t::plus || operation_t == op_kind_t::min
           || operation_t == op_kind_t::max) {
-        // accum_size=1 (int8, uint8, bool)
         if (accum_size == 1) {
           auto [i, t] = scale_mem_bound(bi100_plus_accum1_o4::threads,
                                          bi100_plus_accum1_o4::items, accum_size);
@@ -241,7 +234,6 @@ struct policy_selector {
                               BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG};
           return {rp, rp};
         }
-        // accum_size=2 (int16, float16, bfloat16)
         if (accum_size == 2) {
           auto [i, t] = scale_mem_bound(bi100_plus_accum2_o4::threads,
                                          bi100_plus_accum2_o4::items, accum_size);
@@ -249,7 +241,6 @@ struct policy_selector {
                               BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG};
           return {rp, rp};
         }
-        // accum_size=4 (float32, int32)
         if (accum_size == 4) {
           int vec = (offset_size <= 4) ? bi100_plus_float32_o4::vec
                                        : bi100_plus_float32_o8::vec;
@@ -258,7 +249,6 @@ struct policy_selector {
           ReducePassPolicy rp{t, i, vec, BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG};
           return {rp, rp};
         }
-        // accum_size=8 (float64, int64)
         if (accum_size == 8) {
           if (accum_t == type_t::float64) {
             int vec = (offset_size <= 4) ? bi100_plus_float64_o4::vec
@@ -268,7 +258,6 @@ struct policy_selector {
             ReducePassPolicy rp{t, i, vec, BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG};
             return {rp, rp};
           }
-          // int64 and other 8-byte types
           int vec = (offset_size <= 4) ? bi100_plus_int64_o4::vec
                                        : bi100_plus_int64_o8::vec;
           auto [i, t] = scale_mem_bound(bi100_plus_int64_o4::threads,
@@ -276,7 +265,6 @@ struct policy_selector {
           ReducePassPolicy rp{t, i, vec, BLOCK_REDUCE_WARP_REDUCTIONS, LOAD_LDG};
           return {rp, rp};
         }
-        // accum_size=16 (int128, complex<double>)
         if (accum_size == 16) {
           auto [i, t] = scale_mem_bound(bi100_plus_accum16_o4::threads,
                                          bi100_plus_accum16_o4::items, accum_size);
@@ -286,7 +274,6 @@ struct policy_selector {
         }
       }
     }
-    // Fallback: SM60-equivalent with SM=16 tile optimization
     auto [i, t] = scale_mem_bound(bi100_default::threads,
                                    bi100_default::items, accum_size);
     ReducePassPolicy rp{t, i, bi100_default::vec,
@@ -294,7 +281,6 @@ struct policy_selector {
     return {rp, rp};
   }
 
-  // --- Main entry point ---
   constexpr ReducePolicy operator()(const hardware_capability& hw) const {
     if (determinism == determinism_t::gpu_to_gpu)
       return get_deterministic(hw);