biondizzle
9908fd64d9
Major changes from initial TileLang prototype: Kernel: - CUTLASS NVFP4 block-scaled GEMM (SM100 Blackwell, OpClassBlockScaledTensorOp) - Slot-based dispatch: L1 GEMM → SiLU+Mul per-slot → L2 GEMM → index_add scatter - 1D slot_expert_ids passed to both L1 and L2 (no 2D topk_ids rebuild) - slot_token gathered in cutlass_grouped_nvfp4_gemm when provided SF Remap (source-first): - Iterates logical (m, k_sf) source grid, uses layout_sf(make_coord(m, k_sf)) for CUTLASS dest index — no idx2crd/flatten coordinate extraction - 2D kernel launch: dim3 block(32,8), grid over (K_sf, MN) - Uses cute::cosize() for physical allocation size (not cute::size) - SFA: (MN, K_sf) row-major; SFB: (K_sf, MN) row-major (col-major) Weight transform: - UE4M3 unpack with bit reinterpret (not value cast) - Global scale folding (weight_scale_2) for gate/up split - clamp(0,448) → float8_e4m3fn, transpose (N,K)→(K,N) for CUTLASS No prepack cache: - SFB remapped per-call inside CUTLASS (~µs, not the bottleneck) - See README for why prepack cache must never return (OOM, CUDA graphs, M-dependent layout, cross-layer collisions) Stage activation: - Nearest-neighbor E2M1 quantization (no clamp, no uniform steps) - Per-tensor global scale → alpha for L2 GEMM Bug fixes: - _fold_global_scale: removed broken logical_widths branch - unpack_ue4m3_u32: int32 for CUDA bitwise, view not to, ND support - Correct expert param mapping for NVFP4 checkpoint - SiLU applied per-slot (not after summing expert paths)
38 lines
886 B
YAML
38 lines
886 B
YAML
services:
|
|
vllm:
|
|
build:
|
|
context: .
|
|
dockerfile: Dockerfile
|
|
ports:
|
|
- "8000:8000"
|
|
environment:
|
|
- OMP_NUM_THREADS=128
|
|
- CUDA_LAUNCH_BLOCKING=0
|
|
- TORCH_SHOW_CPP_STACKTRACES=0
|
|
- MEGA_MOE_DEBUG=0
|
|
- MEGA_MOE_STATIC=0
|
|
- NVFP4_DEBUG=0
|
|
- NVFP4_DEBUG_SYNC=0
|
|
- SKIP_ATTENTION=0
|
|
- MEGA_MOE_USE_CUTLASS=1
|
|
- DG_JIT_DEBUG=0
|
|
- DEEP_GEMM_JIT_DEBUG=0
|
|
command:
|
|
- /model
|
|
- --trust-remote-code
|
|
- --enable-expert-parallel
|
|
- --tensor-parallel-size=8
|
|
- --enforce-eager
|
|
- --tokenizer-mode=deepseek_v4
|
|
- --host=0.0.0.0
|
|
- --port=8000
|
|
deploy:
|
|
resources:
|
|
reservations:
|
|
devices:
|
|
- driver: nvidia
|
|
count: all
|
|
capabilities: [gpu]
|
|
volumes:
|
|
- /root/nvidia-meeting/DeepSeek-V4-Pro-NVFP4:/model:ro
|