dylanyunlong/project_6
1
0
Fork 0
Code Issues Pull Requests Actions Projects Releases Wiki Activity
656 Commits 1 Branch 0 Tags
d8d241bf9f867ff270e7f26215fee7ee29023e9c
Commit Graph

656 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
Claude
d8d241bf9f fix: corex_batched_gemm_kernel — use OpClassTensorOp + arch::Cu10 + FP32 accumulator 
Root cause of 25ms (vs expected 2.5ms):
1. ElementAccumulator was half_t → now float (FP32 accumulation)
2. Missing OpClassTensorOp → was defaulting to OpClassSimt (CUDA cores only)
3. Missing arch::Cu10 → was defaulting to arch::Sm61

With these fixes it should use __ivcorex_matrix_mad_f32x4_f16x4 (TCU)
same as moe_cutlass_batched.cu which benchmarked at 2.462ms.
 2026-08-15 12:34:01 +00:00
dylan
3481f2903f fix(build): cutlass.h lives under tensorflow/include on this image 2026-08-15 12:03:12 +00:00
dylan
f41900c06b fix(build): auto-find cutlass/cutlass.h under COREX_ROOT 2026-08-15 12:02:23 +00:00
dylan
bfa18cd5b4 fix(build): use CoreX clang++ instead of nvcc — match working build scripts 2026-08-15 12:01:17 +00:00
dylan
04cc9b88af fix(build): add CUDA include path to g++ step in build_corex_batched_gemm.sh 2026-08-15 11:58:53 +00:00
dylan
ddcfbad431 feat: pybind wrapper for CUTLASS batched GEMM → MoE decode path 
Based on verified result (issue #68):
  CUTLASS Cu10 TensorOp batched: 2.462ms (8 experts, 1 launch)
  vs 8× torch.matmul: 4.6ms (8 launches)
  vs Python F.linear loop: 10.36ms

New files:
  ex_engine/xllm_kernels/cuda/bindings/corex_batched_gemm_bind.cpp
    pybind11 wrapper: batched_gemm_fp16() + moe_decode_fused()
  ex_engine/xllm_kernels/cuda/corex_batched_gemm_kernel.cu
    CUTLASS GemmBatched<half> kernel (from cat_files/batched_gemm.cu)
  qwen3_6_scripts/build_corex_batched_gemm.sh
    Build script for BI-V100 (ivcore10)

Modified:
  qwen3_6_scripts/qwen3_5.py
    import corex_batched_gemm + _USE_COREX_BATCHED_GEMM flag
    Tier 1.5 in MoE decode: after corex_direct_routed, before corex_gather

Build on device: bash qwen3_6_scripts/build_corex_batched_gemm.sh
Output: prebuilt/corex-3.2.3-ivcore10/corex_batched_gemm.so
 2026-08-15 11:54:26 +00:00
dylan
a875fa5d4c Revert "data: cat SGEMM files from 3 repos into cat_files/" 
This reverts commit 7cfa87b5ac.
 2026-08-15 11:48:21 +00:00
Claude
36676f2d1b data: complete SGEMM upstream from 3 repos (siboehm+wangzyon+edtallison) + xllm fused_qknorm_rope + xattention kernels 
SGEMM repos (upstream_ref/sgemm_cuda/, 41 files):
  siboehm/SGEMM_CUDA: kernel 1-12, runner, CMake, cuBLAS benchmark
  wangzyon/NVIDIA_SGEMM_PRACTICE: kernel 1-7 (Chinese comments), utils
  edtallison/sgemm-cuda: kernel 01-09 (learning notes), Makefile

xllm kernels (ex_engine/xllm_kernels/cuda/):
  fused_qknorm_rope.cu + bind — saves 128 kernel launches/fwd
  xattention/ — 6 files from upstream xllm
  headers: corex_compat_utils.h, topk_last_dim.cuh
  ilu/CMakeLists.txt

SO_BUILD_MANIFEST.md — complete .so inventory and call chain analysis
 2026-08-15 07:00:09 +00:00
dylan
7cfa87b5ac data: cat SGEMM files from 3 repos into cat_files/ 
siboehm/SGEMM_CUDA (19 files):
  siboehm_sgemm.cu, siboehm_runner.cu, siboehm_runner.cuh, siboehm_kernels.cuh
  siboehm_cuBLAS_sgemm.cu, siboehm_simplest_kernel.cu, siboehm_CMakeLists.txt
  siboehm_{1_naive..12_kernel_double_buffering}.cuh

wangzyon/NVIDIA_SGEMM_PRACTICE (12 files):
  wangzyon_sgemm.cu, wangzyon_utils.cu, wangzyon_utils.cuh, wangzyon_kernel.cuh
  wangzyon_CMakeLists.txt, wangzyon_kernel_{1..7}.cuh

edtallison/sgemm-cuda (19 files):
  edtallison_sgemm.cu, edtallison_runner.cu, edtallison_runner.cuh
  edtallison_kernels.cuh, edtallison_cuBLAS_sgemm.cu, edtallison_simplest_kernel.cu
  edtallison_CMakeLists.txt, edtallison_{01_naive..12_kernel_double_buffering}.cuh

cat_files/ total: 25 → 75 files
 2026-08-15 06:59:18 +00:00
dylan
284804ac53 data: cat 3 SGEMM repos — siboehm, wangzyon, edtallison (full clone, no --depth) 
Sources:
  siboehm/SGEMM_CUDA        → upstream_ref/sgemm_siboehm/     (25 files)
  wangzyon/NVIDIA_SGEMM_PRACTICE → upstream_ref/nvidia_sgemm_practice/ (23 files, filled gaps)
  edtallison/sgemm-cuda      → upstream_ref/sgemm_edtallison/  (41 files)

All files cat'd one by one from git clone (no --depth).
These are the 3 public SGEMM repos that can compile on CUDA 10.2 + CoreX ivcore10.

Key files for BI-V100 porting:
  kernel 10 (warp tiling) — already proven on device with WARPSIZE=64
  kernel 11/12 (double buffering) — next optimization target
  sgemm.cu + runner.cu — complete build+benchmark harness
  CMakeLists.txt — build system reference
 2026-08-15 06:58:07 +00:00
dylan
854fb93a8e test: add test_ex_engine_cuda.py — test all 22 prebuilt .so on BI-V100 
Tests xllm_norm, xllm_activation, xllm_rope, xllm_moe, ix_full_bridge,
and 11 corex_*.so modules with correctness checks against PyTorch reference.

Run: python3 test_ex_engine_cuda.py
 2026-08-15 06:28:34 +00:00
dylan
e8f0948fe1 feat: ix_ops integration layer — wire ix_full_bridge.so into vllm hot path 
Architecture (CCCL dispatch pattern):
  base_image ixformer::infer → ix_full_bridge.so → ix_ops.py → vllm patches

New files:
  ex_engine/python/ix_ops.py          — Python API for all 14 ixformer::infer ops
  ex_engine/python/patch_vllm_ops.py  — monkey-patch vllm GemmaRMSNorm, SiluAndMul
  ex_engine/deploy_ix_bridge.sh       — build-time deployment script

Modified:
  qwen3_6_scripts/patch_ops.sh        — integrated ix_bridge deployment + startup hook

Call chain: DecoderLayer.forward → GemmaRMSNorm → ix_ops.fused_add_rms_norm
            → ixformer::infer::residual_rms_norm (fused C++ kernel)
 2026-08-15 06:15:17 +00:00
root
109d29fa60 Merge remote-tracking branch 'modelhub/main' 2026-08-15 05:58:37 +00:00
Claude
045ea5df79 feat: Cu10 TensorOp batched HGEMM via Iluvatar CUTLASS framework 
From corex-samples batched_gemm.cu, changed:
  float → half_t, OpClassSimt → OpClassTensorOp, Sm61 → Cu10
Uses __ivcorex_matrix_mad_f32x4_f16x4 via mma_cu10.h
Default config: TB<128,128,32> Warp<32,32,32> Inst<16,16,16>

Standalone test: correctness + perf for MoE decode (8 × 1x4096@4096x11008)
 2026-08-15 05:44:15 +00:00
root
30f98c0674 data: Cu10 CUTLASS part 2 — tensorop example, arch.h, cutlass.h 2026-08-15 05:41:30 +00:00
Claude
f006ab1a01 test: cat tensorop GEMM example + arch.h + cutlass.h from corex-samples 2026-08-15 05:41:03 +00:00
root
b922d694dc data: Cu10 CUTLASS headers from corex-samples 2026-08-15 05:32:10 +00:00
Claude
1d36754efc fix: cat_cutlass_cu10.sh writes to cat_files/ directory instead of stdout 2026-08-15 05:31:33 +00:00
Claude
4abb4df215 test: cat Cu10 CUTLASS files — mma_cu10.h, iluvatar_mma.hpp, batched_gemm.cu, default_mma_core_cu10.h 2026-08-15 05:30:15 +00:00
Claude
6b9086c3a9 test: probe Cu10 CUTLASS fork — find mma_cu10.h, tensor op files, batched_gemm example 2026-08-15 05:27:43 +00:00
Claude
a465dd1d75 fix: use F.silu in test script for old corex torch 2026-08-15 05:22:31 +00:00
Claude
a12d070d82 fix: replace torch::silu with x*sigmoid(x) for old corex torch 2026-08-15 05:22:25 +00:00
Claude
c840c9159f feat: moe_tcu_dispatch.cpp — C++ MoE expert loop via torch::mm (TCU kernel) 
torch profiler confirmed: torch.mm launches Gemm_tcu_bi_kernel::gemm_h_h_tcu_25
which is BI-V100 TCU (Tensor Compute Unit) hardware-accelerated GEMM.
0.58ms per call vs our custom kernel 7.7ms — TCU is 13x faster.

Python for-loop overhead measured: 0.892 ms/expert = 7.1 ms for 8 experts.
This C++ dispatch eliminates that overhead while using the same TCU kernel.

Three entry points:
- moe_decode: full MoE forward (FC1 + SiLU*mul + FC2) for decode
- moe_prefill: group-by-expert MoE forward for prefill
- moe_expert_gemm_tcu: raw GEMM loop for benchmarking
 2026-08-15 05:20:05 +00:00
Claude
9514092980 test: probe torch.matmul backend + ixformer.matmul/linear + Python loop overhead 2026-08-15 05:17:13 +00:00
Claude
395b3e4042 test: clean rebuild + debug output for kernel 10 correctness 2026-08-15 05:14:18 +00:00
Claude
a8ca42b59c perf: hgemm_warptiling Config B — beats cublas on MoE-sized GEMM (0.7x) 
probe_k10_configs.sh results on BI-V100:
  256x4096 @ 4096x11008:
    cublas:   10.554 ms
    Config B:  7.649 ms (0.7x cublas — FASTER)
    Config A:  2308 ms  (old broken config)

Config B: BM128 BN128 BK16 WM64 WN64 WNITER2 TM8 TN4 NT128
Root cause of Config A slowness: WN=128 WNITER=4 caused
excessive register pressure and smem bank conflicts.
 2026-08-15 05:11:32 +00:00
Claude
21417319bc test: sweep 6 kernel 10 configs + cublas baseline — find best params for warp64 2026-08-15 05:08:11 +00:00
Claude
27bb8d28df test: probe kernel 10 perf with CUDA events — isolate bottleneck 2026-08-14 17:23:41 +00:00
Claude
2b12fe687e feat: hgemm_warptiling.cu — siboehm kernel 10 ported to WARPSIZE=64 FP16 
1:1 from upstream_ref/sgemm_cuda/10_kernel_warptiling.cuh.
3 changes: WARPSIZE 32→64, float→__half, FP32 accumulator.

Launch config (confirmed by probe_warp64.sh):
  NUM_THREADS=128, 2 warps of 64
  BM=128 BN=128 BK=16 WM=64 WN=128 WNITER=4 TM=4 TN=4
  WMITER=2, WSUBM=32, WSUBN=32, threads_per_warp=64 ✓
 2026-08-14 17:05:59 +00:00
Claude
11b8a98eea test: probe warp_size=64 behavior + kernel 10 warp tiling with WARPSIZE=64 on BI-V100 2026-08-14 17:02:55 +00:00
Claude
1af7e7cf48 fix: use c10::cuda::getCurrentCUDAStream().stream() for corex torch 2026-08-14 16:49:47 +00:00
Claude
3bee73207e fix: add cuda_runtime.h to hgemm_bind.cpp for cudaStream_t 2026-08-14 16:33:29 +00:00
Claude
09e5261ba6 refactor: hgemm_blocktiling.cu — strict 1:1 from siboehm kernel 6 
Only 3 changes from upstream_ref/sgemm_cuda/6_kernel_vectorize.cuh:
1. float → __half for A/B/C data and shared memory
2. float4 vectorized load → 4 scalar half loads (float4 needs 16-byte align)
3. threadResults accumulator stays float (FP32 accumulation)

Everything else identical: same shared mem layout, same indexing,
same A-transpose-while-loading, same thread tile computation.
No WARPSIZE. No cooperative_groups. No cuda::barrier.
 2026-08-14 16:24:22 +00:00
Claude
ab42fc1fd7 feat: hgemm_blocktiling.cu — FP16 GEMM kernel for MoE expert dispatch on BI-V100 
Adapted from siboehm/SGEMM_CUDA kernel 6 (vectorize + A transpose)
and wangzyon/NVIDIA_SGEMM_PRACTICE kernel 6 (mysgemm_v6).

Key design decisions:
- FP16 data with FP32 accumulation (avoid precision loss)
- No WARPSIZE dependency (safe for BI-V100 warp_size=64)
- Boundary checks for non-aligned M/N/K (MoE expert token counts vary)
- BM=128 BN=128 BK=8 TM=8 TN=8 (256 threads, fits BI-V100 128KB smem)
- A transpose in shared memory for coalesced reads

Two entry points:
1. hgemm(A, B) — standalone FP16 GEMM
2. moe_expert_gemm(input, weights, expert_counts) — MoE prefill path
   loops over experts with variable token counts

For decode (M=1), use cublasHgemmStridedBatched (confirmed working).

Upstream refs: upstream_ref/sgemm_cuda/6_kernel_vectorize.cuh
              upstream_ref/nvidia_sgemm_practice/kernel_6.cuh
 2026-08-14 16:22:00 +00:00
Claude
9ca33cf4d5 upstream: add GEMM kernel references from 4 repos for BI-V100 porting 
Sources (all CUDA 10.2 compatible, no CUTLASS/Triton dependency):
- leimao/CUDA-GEMM-Optimization: v00-v07, fp16 WMMA variant, double buffered
- siboehm/SGEMM_CUDA: kernel 1-12, warp tiling + double buffering
- wangzyon/NVIDIA_SGEMM_PRACTICE: kernel 1-7
- edtallison/sgemm-cuda: kernel 1-12 (reimplementation with notes)

Key porting issue: ALL kernels hardcode WARPSIZE=32.
BI-V100 has warp_size=64. Need to:
1. Replace all 32U / WARPSIZE constants with 64
2. Adjust warp subtile decomposition (WMITER, WNITER, WSUBM, WSUBN)
3. Adjust shared memory bank conflict avoidance (may have different bank count)
4. Test __shfl_down_sync with mask=0xFFFFFFFFFFFFFFFF (64-bit)
 2026-08-14 15:11:57 +00:00
Claude
29ecc2e602 feat: moe_expert_gemm.cpp — C++ loop over experts via ixformer_linear (replaces Python for-loop) 
Key difference from the reverted batched approach:
- Does NOT use torch::mm in a C++ loop (that was the reverted commit)
- Uses ixformer_torch_ext::ixformer_linear — the base image's optimized GEMM
- Same kernel the competitor (sub 168) uses via corex_moe.py
- Eliminates Python interpreter + dispatcher overhead per expert
- Links against _ixformer_torch.cpython-310.so (already in base image)

Decode: 1 Python call → 8 C++ ixformer_linear (vs 8 Python F.linear)
Prefill: 1 Python call → 64 C++ ixformer_linear (vs 64 Python F.linear)
 2026-08-14 12:09:50 +00:00
Claude
0ace44e293 test: cat ixinfer.h + functional batched GEMM test on BI-V100 2026-08-14 12:07:00 +00:00
Claude
bfc4de2cf3 test: probe cuinfer + cublasLt + cublas batched GEMM APIs for MoE expert dispatch 2026-08-14 12:03:52 +00:00
Claude
d6958070cb test: probe CUTLASS + __CUDA_ARCH__ on BI-V100 for grouped GEMM feasibility 2026-08-14 11:59:51 +00:00
claude
50a249e0a3 Revert "feat: batched MoE expert GEMM — replaces Python for-loop" 
This reverts commit 06d7713db6.
 2026-08-14 11:47:37 +00:00
claude
06d7713db6 feat: batched MoE expert GEMM — replaces Python for-loop 
ixformer probe results:
  ✗ moe_w16a16_group_gemm NOT in ixformer .so
  ✗ CUTLASS grouped GEMM needs cuda/std (variadic function error on corex)
  ✓ ixformer_linear EXISTS (fused matmul)
  ✓ torch.mm works (uses corex cublas)

Solution: moe_batched_gemm.cu
  - C++ loop over experts (eliminates Python overhead)
  - torch::mm for GEMM (corex cublas, not F.linear Python)
  - Fused silu_and_mul CUDA kernel (not PyTorch ops)
  - Weighted scatter-add in C++
  - Skips empty experts (no wasted compute)

Integration in qwen3_5.py:
  _USE_XLLM_MOE_GEMM dispatches to moe_experts_forward()
  Falls back to Python for-loop if not available

Build: bash qwen3_6_scripts/build_xllm_kernels.sh
 2026-08-14 11:43:46 +00:00
claude
93353a1414 test: probe ixformer .so symbols for MoE grouped GEMM 
CUTLASS grouped GEMM (example 24) requires SM80 Tensor Core + cuda/std headers.
Cannot compile on corex (same issue as CCCL 3.6 variadic functions).

Alternative path: ix_moe_bridge.so calls ixformer::infer::moe_w16a16_group_gemm
which is BI-V100 optimized grouped GEMM already in the base image.

This probe script checks if the MoE functions exist in ixformer .so
before attempting to build ix_moe_bridge.so.

Run: bash qwen3_6_scripts/probe_ixformer_symbols.sh
 2026-08-14 11:40:13 +00:00
claude
865c18f852 feat: integrate xllm_moe into qwen3_5.py MoE hot path 
xllm_moe.so provides 3 fused CUDA kernels compiled for ivcore10:
  - moe_fused_topk: CUB topk + softmax (replaces corex_moe_topk_softmax)
  - moe_compute_index: histogram + prefix_sum + place (replaces corex_moe_index_combine)
  - moe_combine_result: reorder + weighted sum (available but not yet wired to output)

Dispatch priority in _pure_pytorch_experts():
  Tier 0: xllm_moe (if available)
  Tier 1: corex_moe_* individual .so
  Tier 2: PyTorch fallback

Integration points:
  1. Topk routing: xllm_moe.moe_fused_topk → corex_moe_topk_softmax → torch.topk
  2. Index computation: xllm_moe.moe_compute_index → corex_moe_index_combine → torch.argsort
  3. Expert loop: still Python F.linear (next target: batch GEMM)

patch_ops.sh already deploys all prebuilt/*.so including xllm_moe.so
 2026-08-14 11:37:21 +00:00
root
e147c283e3 prebuilt: 5 xllm CUDA kernel .so for BI-V100 (ivcore10) 
All 12/12 AST call chain tests passed:
  xllm_norm.so       (1.4MB) rms_norm max_err=0.000000, fused_add_rms_norm err=0.001953
  xllm_activation.so (1.3MB) silu_and_mul err=0.001953, gelu_and_mul err=0.001953
  xllm_rope.so       (1.2MB) rotary_embedding q_diff=1360 k_diff=1404
  xllm_cache.so      (1.3MB) reshape_paged_cache err=0.000000
  xllm_moe.so        (?.?MB) fused_topk weight_sum_err<0.01, compute_index total=64, combine err<0.1

vs ixformer: rms_norm err=0.000000, silu_and_mul err=0.000000

Compiled with corex clang/16 --cuda-gpu-arch=ivcore10
Using corex CUB (/usr/local/corex/include/cub/), NOT CCCL 3.6
 2026-08-14 11:34:23 +00:00
claude
31d3ee99bb fix: MoE kernel include paths — device_utils.cuh + arch_condition.h 
Fixed xllm internal paths to our headers/ directory:
  kernels/cuda/device_utils.cuh → device_utils.cuh
  core/kernels/cuda/device_utils.cuh → device_utils.cuh
  core/kernels/cuda/arch_condition.h → arch_condition.h (copied)
 2026-08-14 11:31:54 +00:00
claude
df6a0f5d47 fix: remove cuda/functional from MoE topk kernels (not available on corex) 2026-08-14 11:29:19 +00:00
claude
a50adefdfc feat: xllm MoE CUDA kernels — fused_topk + compute_index + combine 
3 MoE kernel files adapted for corex:
  moe_fused_topk.cu: LOG(FATAL)→TORCH_CHECK, +torch/extension.h
  moe_compute_index.cu: CHECK_LE→TORCH_CHECK, uses cub::BlockScan (corex CUB)
  moe_combine.cu: fixed duplicate include, +torch/extension.h

New pybind binding: xllm_moe_bind.cpp
  → moe_fused_topk(gating, topk, renormalize, bias, scoring_func)
  → moe_compute_index(expert_id, num_experts)
  → moe_combine_result(gemm2, weights, N, topk)

AST verification added for all 3 functions
 2026-08-14 11:23:49 +00:00
claude
49cd7def89 fix: slot_ids dtype long→int32 in cache test 2026-08-14 11:20:07 +00:00
claude
1f51feee05 fix: cos_sin_cache dtype float→half in rope test 2026-08-14 11:18:02 +00:00
claude
7fc0c1defa test: AST call chain verification for all 4 xllm kernel .so 
verify_ast_chain.py tests every exported function:
  xllm_norm.so:       rms_norm (vs PyTorch), fused_add_rms_norm (vs PyTorch)
  xllm_activation.so: silu_and_mul (vs F.silu), gelu_and_mul (vs F.gelu)
  xllm_rope.so:       rotary_embedding (verify rotation applied)
  xllm_cache.so:      reshape_paged_cache (verify slot write)
  vs_ixformer:        compare our .so output vs base image ixformer

All 4 .so compiled successfully:
  xllm_norm.so       1.4MB  [rms_norm, fused_add_rms_norm]
  xllm_activation.so 1.3MB  [silu_and_mul, gelu_and_mul, act_and_mul]
  xllm_rope.so       1.2MB  [rotary_embedding]
  xllm_cache.so      1.3MB  [block_copy, reshape_paged_cache]
 2026-08-14 11:14:58 +00:00