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
|
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
|
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
|
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
|
093bfb380f
|
feat: pybind11 bindings for xllm CUDA kernels
norm.cu compiled successfully on BI-V100 (only warning: fp8 __host__ attr).
Failed at import because no PYBIND11_MODULE — now fixed.
New bindings/ directory with 4 binding files:
xllm_norm_bind.cpp → rms_norm, fused_add_rms_norm
xllm_activation_bind.cpp → silu_and_mul, gelu_and_mul, act_and_mul
xllm_rope_bind.cpp → rotary_embedding
xllm_cache_bind.cpp → reshape_paged_cache, block_copy
Build script updated: each .so = kernel .cu + binding .cpp
|
2026-08-14 10:50:13 +00:00 |
|