Commit Graph

82 Commits

Author SHA1 Message Date
Claude
16981f221e feat(muh): reduce_by_key 55→203 lines — full key_size×accum_size dispatch
From 1KB/55 lines (46× compression vs CCCL 70KB) to 203 lines:
- Add 11 type specialization structs (key=1,2,4,8 × accum=1,2,4,8)
- SM=16 tile maximization: k4_a4 hot path 100% SMEM (256*24*8=49152)
- k8_a8 also at 100% SMEM (192*16*16=49152)
- Delay halved for L2=6MB across all branches
- CCCL-matching ReduceByKeyPolicy struct with ReduceByKeyAlgorithm enum
- Dynamic SMEM fallback for unknown pair sizes
2026-08-03 07:20:23 +00:00
Claude
bdaec8da4b feat(muh): select_if SM=16 tile maximization
- Increase tiles across all elem_size branches for SM=16 (fewer CTAs need larger tiles)
- Flagged path: items increased 20-80% (e.g. elem≤2: 18→24, elem≤4: 14→18)
- Non-flagged path: items increased 30-100% (e.g. elem≤4: 18→24, elem≤8: 14→16)
- Add SMEM utilization comments for each branch (target ≥50%)
- No structural change to 3-dimension dispatch (may_alias/flagged/delay)
2026-08-03 07:19:13 +00:00
Claude
95d872e8f5 feat(muh): scan_by_key 53→284 lines — full key_size×val_size type dispatch
From 1KB/53 lines (59× compression vs CCCL 85KB) to 284 lines:
- Add 16 type specialization structs (key_size=1,2,4,8 × val_size=1,2,4,8)
- SM=16 tile maximization: k4_v4 (attention hot path) 30720→49152 (62%→100% SMEM)
- SM=16 tile increases across all small pairs (k1_v1: 3072→12288, k2_v1: 6144→15360)
- Delay halved for L2=6MB: fixed_delay values /2 (less inter-CTA contention)
- Proper CCCL-matching ScanByKeyPolicy struct with ScanByKeyAlgorithm enum
- SMEM-safe fallback with dynamic items computation from pair_size
- Macro-based dispatch (MK_POLICY) for clean type selection
2026-08-03 07:18:23 +00:00
Claude
88db0ed89c feat(muh): SM=16 tuning overhaul — reduce/scan/transform
tuning_reduce.cuh (201→311 lines):
- Add accum_size=1/2/16 branches (int8, bfloat16, int128)
- Add min/max op dispatch (same params as plus for BI-V100)
- SM=16 tile maximization: det_float32 tile 11648→49152 (23%→100% SMEM)
- SM=16 tile maximization: det_float64 tile 11264→49152 (23%→100% SMEM)
- Add float32_o8, int64_o4/o8 variants with vec_size dispatch
- Increase float32 items 16→24 (32768→49152, fill SMEM for fewer CTAs)

tuning_scan.cuh:
- Fix 1B tile from 9216→16384 (19%→33% SMEM, scan needs 2x buffer)
- Fix 2B tile from 13312→24576 (27%→100% SMEM with double buffer)
- Fix 8B_o4 tile: threads 416→384 for warp alignment, items 14→16
- Update header comments with confirmed SM=16 hardware profile
- Document lookback delay heuristic for L2=6MB

tuning_transform.cuh (128→168 lines):
- CRITICAL: bytes_in_flight 16KB→32KB (was based on 900/50=18 GB/s,
  actual is 900/16=56 GB/s — 3× error)
- Add full PrefetchPolicy struct matching CCCL upstream
- Add AsyncCopyPolicy with BI-V100 fallback (no cp.async support)
- Document CCCL cc_to_min_bytes_in_flight reference values
- Add vec_size calculation from element size (16-byte vector loads)
- Cap items_per_thread at 32 to prevent register pressure

hardware.cuh:
- Add SMEM 48KB vs 32KB disambiguation note
2026-08-03 07:16:35 +00:00
Claude
0ba4cdb025 fix: dial back max-num-seqs 256→8, revert batched-tokens and mem-util
256 concurrent seqs risks OOM: worst case with long prompts in queue
can exhaust KV cache + activation memory. 32K batched-tokens prefill
activation ≈ 20GB competes with KV cache. 0.95 mem-util leaves only
5% headroom for spikes.

Conservative start: max-num-seqs=8 (8× improvement over baseline=1).
8 seqs × 2048 avg context × 80KB/token = 1.3GB KV cache, safe.
gpu-memory-utilization and max-num-batched-tokens restored to proven
baseline values.

Optimal max-num-seqs needs real-hardware sweep: 4→8→16→32→64→128.
The value where Output TPS plateaus (KV cache saturated) is the
answer. Can't determine this without Phanthy Cloud access.
2026-08-03 06:51:33 +00:00
Claude
cdc01bbc6a fix: critical config + tuning corrections from CCCL source analysis
computility-run.yaml:
  max-num-seqs 1→256: benchmark sweeps [128,256] concurrent seqs,
    current config processes 1 while 127 queue. KV cache budget:
    256 seqs × 2048 tokens × 80KB/token = 41.9GB < 45GB available.
  max-num-batched-tokens 8192→32768: support 256 concurrent prefills.
  gpu-memory-utilization 0.9→0.95: provide KV cache headroom.

Dockerfile:
  Deploy paged_attention_v2_triton.py to vllm package path so
  try-triton-first logic in _custom_ops.py can find it. Falls back
  to PyTorch V2 automatically if Triton V2 fails (SMEM/runtime).

muh/tuning/common.cuh:
  scale_mem_bound max_smem now a parameter (default 48KB). Allows
  policy_selectors to pass hw.max_shared_memory_per_block if actual
  SMEM differs from CCCL 48KB assumption.

muh/tuning/tuning_transform.cuh:
  bytes_in_flight 16KB→32KB. Old derivation used 900/50=18 GB/s/SM
  (wrong, SM=16 confirmed). Actual per-SM BW = 56 GB/s.
  32KB is estimate pending benchmark sweep.

SM count 50→16 corrections across all affected files.
2026-08-03 06:45:54 +00:00
Claude
071fa361a3 docs: CCCL → Triton methodology transfer — parameter search translation table
Maps CCCL %RANGE% benchmark format to EngineX Triton autotune configs:
- prefix_prefill: ipt→BLOCK_M, tpb→num_warps
- triton_flash_attention: 8 configs safety filter by SMEM
- fused_moe: BLOCK_SIZE_M/N/K grid search from CCCL transform/reduce
- Execution plan: hardware confirm → grid search → filter → deploy
- Competitive advantage: systematic search vs guessing
2026-08-03 04:37:23 +00:00
Claude
e48a46a30d docs: EngineX vllm injection map — Python not C++, Triton not CUDA, 32KB SMEM claim
CRITICAL FINDINGS from enginex-vllm-bi100-qwen36-main.zip analysis:
1. No .cu files — all CUDA kernels pre-compiled in ixf_F (ixformer.functions)
2. paged_attention_v2 is NotImplementedError, use_v1=True hardcoded
3. Real tuning surface: BLOCK/NUM_WARPS in Triton, BLOCK_SIZE_M/N/K in MoE
4. _custom_ops reports SMEM=32KB (not 48KB!) — needs hardware verification
5. muh strategy shifts from C++ injection to Python parameter optimization
6. CCCL methodology still applies but targets Triton kernels not CUB dispatch
2026-08-03 03:59:49 +00:00
Claude
9bba7f4c79 docs: CCCL integration status — 462 assets inventory, 8/8 decode hot path coverage, no-clone-more verdict
- Complete CCCL asset inventory: 52 Thrust examples + 18 CUB examples + 234 Catch2 tests + 78 benchmarks + 27+27 tuning headers
- Decode hot path matrix: reduce/scan/topk/radix_sort/transform/select_if/batch_memcpy/for all covered
- SM count=16 impact analysis on tuning parameters
- muh toolchain status: 27/27 headers, gen_patch/gen_yaml/parse all functional
- Verdict: 8,900 files already sufficient, remaining 31K files are cmake/CI scaffolding
- 信创魔盒定位: muh是算法因子置换层,不是连接器适配层
2026-08-03 03:50:24 +00:00
Claude
392e644611 feat: add qwen3_5.py vllm adapter for Qwen3.6-35B-A3B
588-line vllm model implementation based on qwen3_moe.py.
Bootstrap strategy: treat ALL layers as full attention (ignoring
linear_attention optimization). Correct but suboptimal.

Key adaptations:
- _get_text_config(): unwrap composite config -> text_config
- Shared expert support (shared_expert_intermediate_size)
- Skip linear attention weights (conv1d, delta_net, gated_delta)
- Skip vision encoder and MTP weights
- QK norm (Qwen3 style)
- Partial rotary embedding (rope_pct=0.25)

Includes deploy.sh and run_baseline.sh for server deployment.
2026-08-01 13:54:33 +00:00
root
6beb497447 fix(hardware): SM count 50→16 confirmed on Phanthy Cloud BI-V100
ixsmi + torch.cuda.get_device_properties confirmed:
- multi_processor_count: 16 (not 50 as in spec sheet)
- compute_capability: 7.0 (Volta-compatible)
- max_threads_per_SM: 8192
- total_memory: 32GB per GPU
- SM clock: 1500MHz (max 2500MHz)

Impact: bandwidth_per_sm = 900/16 = 56.25 GB/s (was 18 GB/s at 50 SM)
All occupancy and tile-size calculations need revision.
2026-08-01 13:39:28 +00:00
Claude
fee8f1b9e4 docs: document Qwen3.6-35B-A3B bootstrap failure and architecture analysis
vllm 0.6.3 KeyError on qwen3_5_moe model type.
Model is hybrid linear+full attention MoE with 256 experts (top-8).
enginex-vllm-bi100-qwen36-main.zip in repo likely contains the fix.
2026-08-01 13:16:01 +00:00
Claude
9b0d1c283c docs: add competition server profile (4×BI-V100, Qwen3.6-35B-A3B)
Hardware: 4× Iluvatar BI-V100 32GB, Xeon Gold 6530, 503GB RAM
Software: vllm 0.6.3+corex.3.2.3, torch 2.1.0+corex.3.2.3
Model: Qwen3.6-35B-A3B at /root/public-storage/models/Qwen/
Benchmark: benchmark_server_v0.5.0.py with automated sweeps
2026-08-01 13:12:59 +00:00
dylanyunlon
79730ea907 test: add SMEM safety validator for all 26 tuning algorithms
191 combinations tested: algorithm × type_size × (key,value) pairs.
Verifies every policy_selector output satisfies tile ≤ 49152 bytes.
Exit code 0 = all safe, 1 = overflow detected.

Usage: python3 muh/tests/test_smem_safety.py [--verbose]
2026-08-01 12:37:14 +08:00
dylanyunlon
0154a3b297 fix(tuning_batched_topk): force bits=8, fix SMEM overflow
Previous version used base topk policy's bits (11 for key>=2B),
causing SMEM overflow: 512*4*key_size + 2048*4*batches > 49152.

Fix: force bits=8 (same as radix_sort decision for BI-V100).
SMEM: 512*4*key_size + 256*4*batches = manageable.
Also adds while-loop SMEM check on max_batches.

Detected by test_smem_safety.py: 3 overflows at key_size=2,4,8.
2026-08-01 12:36:57 +08:00
dylanyunlon
2c5e77f370 feat(gen_patch): add TUNING_REGISTRY for all 26 algorithms
Registers all 26 CUB algorithms with metadata:
- 6 'injection' mode: have VLLM_INJECTION_POINTS (reduce/scan/topk/transform/batch_memcpy/for)
- 20 'library' mode: used via CCCL device API, no direct #define injection
- struct_mode: 'named' (bi100_* structs) vs 'inline' (policy_selector returns)

Also adds coverage reporting to generate_patches().
2026-08-01 12:33:14 +08:00
dylanyunlon
84c18150e6 fix(tuning_select_if): restore 3 collapsed dispatch dimensions
Previous version collapsed 77 CCCL specializations into 4 if/else
branches by elem_size only, losing:

1. may_alias dimension: now dispatches LOAD_CA (alias-safe) vs
   LOAD_DIRECT+LOAD_LDG (no-alias, ~5-10% faster for common case).
   CCCL SM100 no-alias small-type uses BLOCK_LOAD_DIRECT.

2. has_flags dimension: flagged path now gets 2-4 fewer items_per_thread
   because flag array takes additional SMEM. SMEM check includes flag_tile.

3. delay dimension: type-size-dependent delays instead of fixed(350,450).
   Scaled from CCCL SM100 benchmarks: ns*0.5, l2w*0.6 for BI-V100 L2.

SMEM check: input_tile + output_scatter + flag_tile ≤ 48KB.
2026-08-01 02:26:48 +08:00
dylanyunlon
9287700964 fix(tuning_radix_sort): remove invented portioned_smem_per_warp field
The previous version had a `portioned_smem_per_warp` field that doesn't
exist in CCCL. The actual CCCL RadixSortOnesweepPolicy has:
  threads, items, store_algorithm, rank_algorithm, scan_algorithm,
  rank_private_partitions, radix_bits

Also adds proper SMEM calculation:
  total = max(keys_tile, values_tile, rank_smem) + offsets
  with 2KB headroom for kernel stack/locals.

rank_private_partitions set to 1 to minimize SMEM pressure.
2026-08-01 02:26:46 +08:00
dylanyunlon
2bc3263793 [muh] add tuning_radix_sort.cuh: BI-V100 tuning for radix_sort
Translated from CCCL with SMEM overflow protection.
All SM100 values checked against 48KB limit.
2026-08-01 02:11:05 +08:00
dylanyunlon
437fc3ea20 [muh] add tuning_unique_by_key.cuh: BI-V100 tuning for unique_by_key
Translated from CCCL with SMEM overflow protection.
All SM100 values checked against 48KB limit.
2026-08-01 02:11:03 +08:00
dylanyunlon
58de86d817 [muh] add tuning_select_if.cuh: BI-V100 tuning for select_if
Translated from CCCL with SMEM overflow protection.
All SM100 values checked against 48KB limit.
2026-08-01 02:11:01 +08:00
dylanyunlon
91f9a3a0e5 [muh] add tuning_scan_by_key.cuh: BI-V100 tuning for scan_by_key
Translated from CCCL with SMEM overflow protection.
All SM100 values checked against 48KB limit.
2026-08-01 02:10:59 +08:00
dylanyunlon
0ec355cf74 [muh] add tuning_reduce_by_key.cuh: BI-V100 tuning for reduce_by_key
Translated from CCCL with SMEM overflow protection.
All SM100 values checked against 48KB limit.
2026-08-01 02:10:57 +08:00
dylanyunlon
915c4aff56 [muh] add tuning_segmented_sort.cuh: BI-V100 tuning for segmented_sort
Translated from CCCL with SMEM overflow protection.
All SM100 values checked against 48KB limit.
2026-08-01 02:10:55 +08:00
dylanyunlon
561a82c849 [muh] add tuning_three_way_partition.cuh: BI-V100 tuning for three_way_partition
Translated from CCCL with SMEM overflow protection.
All SM100 values checked against 48KB limit.
2026-08-01 02:10:54 +08:00
dylanyunlon
aeb270578f [muh] add tuning_rle_non_trivial_runs.cuh: BI-V100 tuning for rle_non_trivial_runs
Translated from CCCL with SMEM overflow protection.
All SM100 values checked against 48KB limit.
2026-08-01 02:10:52 +08:00
dylanyunlon
f3ae28bbb5 [muh] add tuning_rle_encode.cuh: BI-V100 tuning for rle_encode
Translated from CCCL with SMEM overflow protection.
All SM100 values checked against 48KB limit.
2026-08-01 02:10:50 +08:00
dylanyunlon
f838c22bad [muh] add tuning_histogram.cuh: BI-V100 tuning for histogram
Translated from CCCL with SMEM overflow protection.
All SM100 values checked against 48KB limit.
2026-08-01 02:10:48 +08:00
dylanyunlon
3a8030224b [muh] add tuning_segmented_radix_sort.cuh: BI-V100 tuning header for segmented_radix_sort
Translated from CCCL cub/device/dispatch/tuning/tuning_segmented_radix_sort.cuh.
Uses hardware_capability dispatch instead of compute_capability.
2026-08-01 02:08:44 +08:00
dylanyunlon
c3e9e5b6f5 [muh] add tuning_batched_topk.cuh: BI-V100 tuning header for batched_topk
Translated from CCCL cub/device/dispatch/tuning/tuning_batched_topk.cuh.
Uses hardware_capability dispatch instead of compute_capability.
2026-08-01 02:08:42 +08:00
dylanyunlon
0344a3fbb8 [muh] add tuning_transform_tile.cuh: BI-V100 tuning header for transform_tile
Translated from CCCL cub/device/dispatch/tuning/tuning_transform_tile.cuh.
Uses hardware_capability dispatch instead of compute_capability.
2026-08-01 02:08:40 +08:00
dylanyunlon
3fee54f4d7 [muh] add tuning_merge_sort.cuh: BI-V100 tuning header for merge_sort
Translated from CCCL cub/device/dispatch/tuning/tuning_merge_sort.cuh.
Uses hardware_capability dispatch instead of compute_capability.
2026-08-01 02:08:38 +08:00
dylanyunlon
af8fc0caeb [muh] add tuning_merge.cuh: BI-V100 tuning header for merge
Translated from CCCL cub/device/dispatch/tuning/tuning_merge.cuh.
Uses hardware_capability dispatch instead of compute_capability.
2026-08-01 02:08:37 +08:00
dylanyunlon
c0bfc8c93d [muh] add tuning_segmented_scan.cuh: BI-V100 tuning header for segmented_scan
Translated from CCCL cub/device/dispatch/tuning/tuning_segmented_scan.cuh.
Uses hardware_capability dispatch instead of compute_capability.
2026-08-01 02:08:35 +08:00
dylanyunlon
105dd96b52 [muh] add tuning_segmented_reduce.cuh: BI-V100 tuning header for segmented_reduce
Translated from CCCL cub/device/dispatch/tuning/tuning_segmented_reduce.cuh.
Uses hardware_capability dispatch instead of compute_capability.
2026-08-01 02:08:33 +08:00
dylanyunlon
81e4a907c0 [muh] add tuning_find_bound_sorted_values.cuh: BI-V100 tuning header for find_bound_sorted_values
Translated from CCCL cub/device/dispatch/tuning/tuning_find_bound_sorted_values.cuh.
Uses hardware_capability dispatch instead of compute_capability.
2026-08-01 02:08:31 +08:00
dylanyunlon
25f7a636a9 [muh] add tuning_find.cuh: BI-V100 tuning header for find
Translated from CCCL cub/device/dispatch/tuning/tuning_find.cuh.
Uses hardware_capability dispatch instead of compute_capability.
2026-08-01 02:08:29 +08:00
dylanyunlon
eaca73a390 [muh] add tuning_adjacent_difference.cuh: BI-V100 tuning header for adjacent_difference
Translated from CCCL cub/device/dispatch/tuning/tuning_adjacent_difference.cuh.
Uses hardware_capability dispatch instead of compute_capability.
2026-08-01 02:08:28 +08:00
dylanyunlon
52c5ca7ce5 refactor(muh_dispatch): read-once from C++ headers, not write-twice
Replaces hand-written reduce_threads=512, reduce_items=16 with
_read_reduce_config(accum_size) that reads from tuning_reduce.cuh
via gen_patch.extract_bi100_structs().

Architecture change:
  OLD: hand-write values in Python + verify_against_headers() asserts equal
  NEW: _read_reduce_config() reads from C++ header (single source of truth)
       Falls back to compiled-in defaults only when headers not on disk
       (deployed container), with RuntimeWarning.

No hand-written tuning values remain in the normal code path.
verify_against_headers() removed — there is nothing to verify
when there is only one copy of the truth.
2026-08-01 01:30:42 +08:00
dylanyunlon
482aabdea3 fix(muh_kernel_map): add threads >= 32 floor in Python scale_mem_bound
Mirrors the C++ fix in common.cuh.
2026-08-01 01:29:57 +08:00
dylanyunlon
142568072a fix(common.cuh): add threads >= 32 floor in scale_mem_bound
Defensive guard: if SMEM cap computes max_threads_by_smem < 32
(or rounds to 0), floor at 32 (one warp). Prevents launching
0 threads which is undefined behavior.
2026-08-01 01:29:47 +08:00
dylanyunlon
03f6a59ebf fix(muh_dispatch): add verify_against_headers() to close the loop
Adds verification that hand-written values in muh_dispatch.py
(reduce_threads=512, reduce_items=16, etc.) match the C++ headers
(bi100_float32_plus_o4 in tuning_reduce.cuh).

Previously: muh_dispatch.py had hand-coded values with no link to
the C++ source of truth. gen_patch.py reads from C++ headers,
but muh_dispatch.py was a separate copy that could diverge.

Now: verify_against_headers() calls gen_patch.extract_bi100_structs()
and compares. Self-test prints mismatches if any exist.
2026-08-01 00:32:16 +08:00
dylanyunlon
3a2b67c166 fix(tuning_reduce): auto [t,i] → auto [i,t] matching CCCL scaling_result
scale_mem_bound now returns {items, threads} (items-first) to match
CCCL's scaling_result struct. All 7 call sites in this file updated.

Previously: auto [t, i] bound threads→t, items→i
Now:        auto [i, t] binds items→i, threads→t

The ReducePassPolicy{t, i, ...} constructors remain correct because
they take (threads, items, ...) — t is threads, i is items in both cases.
The old code worked by accident (two reversals canceling out).
2026-08-01 00:31:41 +08:00
dylanyunlon
ec1c85cd9a fix(common.cuh): scale_mem_bound — 3 bugs vs CCCL original
1. Return order: {threads, items} → {items, threads} matching CCCL scaling_result
2. Upper clamp: nominal*1 → nominal*2 (CCCL allows small types to double items)
3. Add threads SMEM cap: min(nominal, round_up(48KB/(ts*items), 32))

Verified against all 8 test vectors from CCCL catch2_test_util_arch.cu.
The old code was only safe because current bi100_* structs don't hit the
edge cases — but any future CCCL code copy would silently produce wrong
values.
2026-08-01 00:31:22 +08:00
dylanyunlon
3ebc37d80d [muh] fix scale_mem_bound: 3 bugs vs CCCL util_arch.cuh
1. Return order: (items, threads) not (threads, items) — matches CCCL scaling_result
2. Items clamp upper bound: nominal*2, not nominal*1 — allows small types to double
3. Threads SMEM cap: min(nominal, round_up(max_smem/(type*items), 32)) — prevents SMEM overflow

Verified against all 18 CCCL test cases in catch2_test_util_arch.cu (was 4/14, now 18/18).

Note: C++ tuning headers (tuning_reduce.cuh etc.) have corresponding auto [t, i] destructuring
that also needs to flip to auto [i, t]. The bi100_* struct values themselves are correct
(hand-derived from SMEM constraints), but the policy_selector callers of scale_mem_bound
will produce wrong destructuring. Tracked in project/6 as separate fix item.
2026-08-01 00:00:02 +08:00
Claude
173c6afe09 [muh] kernel_map: full vllm→CCCL mapping with SMEM overflow detection
muh_kernel_map.py maps every vllm kernel to its CCCL algorithm(s):
  paged_attention_v1 → reduce (compound: summary_statistics pattern)
  paged_attention_v2 → reduce + scan (two-pass partitioned)
  sampling_topk → topk + radix_sort
  activation_kernels → transform (SiLU/GELU)
  layernorm_kernels → reduce + transform (variance + normalize)
  rotary_embedding → for_each + transform (RoPE)
  cache_kernels → batch_memcpy (KV block copy)

Found 5 lookahead SMEM overflows — documented in SPECIALIZATION_ANALYSIS.md.
These are non-functional (BI-V100 lacks warpspeed pipeline) but the
dispatch correctly falls back to lookback.

The competitive moat:
  Others: tune 5 vllm launch params → hours
  Us: tune 7 CUB primitive dimensions per algorithm × 6 algorithms,
      constrained by SMEM/occupancy/L2, with CCCL benchmark protocol
2026-07-31 11:13:33 +00:00
dylanyunlon
e69c46d0b7 [muh] add muh_dispatch.py — CCCL-style type-dispatched kernel config for BI-V100
This is the key differentiator vs parameter brute-force.

Everyone else hardcodes BLOCK_SIZE=64, NUM_WARPS=4, PARTITION_SIZE=512.
muh_dispatch replaces these with type-dispatched values derived from
CCCL's policy_selector architecture.

Dispatch axes (matching CCCL type_t × op_kind_t × offset_size):
  - dtype → determines accum_size, SMEM per element
  - head_dim → determines tile width, SMEM constraint
  - max_seq_len → determines V1/V2 threshold (single_tile vs multi_tile)
  - num_kv_heads → determines GQA ratio (memory access pattern)

Output: AttentionConfig struct with all kernel parameters.
CCCL reference: ReducePolicy{multi_tile, single_tile} pattern.

Example type dispatches for Qwen3.6 on BI-V100:
  bf16 h128 100K → partition=512, v1_thresh=8192, reduce(512,16,vec=4)
  bf16 h256 100K → partition=256, v1_thresh=8192, reduce(512,16,vec=4)
  fp32 h128 32K  → partition=256, v1_thresh=8192, reduce(512,16,vec=4)
  bf16 h128 2K   → v1_thresh=2049 (always V1, skip V2 overhead)
2026-07-31 19:11:46 +08:00
dylanyunlon
d14b0c19e4 [docs] add CCCL SM100 vs muh BI-V100 specialization parity analysis 2026-07-31 18:35:50 +08:00
dylanyunlon
35ef79c5f8 [muh] scan: add bi100_lookback_1B_o8 — close SM100 parity gap (7/7 lookback branches)
CCCL SM100 scan lookback has 7 type-specialized branches:
  offset_size=4: 1B, 2B, 4B, 8B
  offset_size=8: 1B, 4B, 8B

muh BI-V100 previously had 6 (missing o8_1B).
This commit adds the o8_1B branch derived from SM100 ref:
  ipt_14.tpb_384.ns_228.dcid_7.l2w_775 → 1.107x
  BI-V100 delay: halved ns (L2 6MB vs 50MB): backon(114, 465)
  nominal_tile = 384*14*4 = 21504 ≤ 49152 ✓

Now: 7/7 lookback + 6/6 lookahead = 13/13 SM100 parity.
2026-07-31 18:35:01 +08:00
Claude
c5a0d61851 sync: align with enginex-vllm-bi100-qwen36 baseline (1902c81f)
Synced files from EngineX baseline zip (2026-06-30):
- ADD paged_attn.py (root): production paged attention with PyTorch fallback
- ADD launch_service: BI-V100 server startup script with env configuration
- SYNC computility-run.yaml: gpu_memory=0.9, batched_tokens=8192, seq_capture=32768
- SYNC qwen3_6_scripts/paged_attn.py: +311 lines, Triton bypass docs, _forward_decode_pytorch shape docs
- SYNC qwen3_6_scripts/qwen3_5.py: -72 lines, revert optimized MoE prefill to baseline (untested on BI-V100)
- KEEP Dockerfile: repo version has V2/Triton/head256 optimization patches not in baseline

Baseline commit: 1902c81fdd373943f17f5983eb8750758c7f4a69
Source: enginex-vllm-bi100-qwen36-main.zip (dev.modelhub.org.cn)
2026-07-31 09:43:58 +00:00