70c898ac8b
真机验证发现的问题: 1. qwen3_5.py 做 'from ex_engine.python.ix_bridge' 但包结构是 ex_engine.ix_bridge → 创建 python/ 子目录 + symlinks 2. flash_qla_sm70 只部署到 /workspace 没有到 vllm models 目录 → 显式 cp -r 到 VLLM/model_executor/models/ 3. 从 vllm v0.5.5 搬 MoE CUDA kernels (torch::Tensor API): - topk_softmax_kernels.cu (506行, CUB BlockReduce) - moe_align_block_size_kernels.cu (134行) - moe_pybind.cpp (pybind11 入口) 真机验证结果: ✓ ix_bridge import OK, available=True ✓ topk_softmax (64 experts, top8) OK — CUDA kernel 命中 ✓ ix_full_bridge silu_and_mul OK ✓ qwen3_5.py import OK ✓ ex_engine build 2/2 factors ✓ moe_topk_softmax_v3.so 编译成功 ✓ flash_qla_sm70_gdn_strided.so 编译成功 ✗ 单卡 32GB OOM (正常, 竞赛 4卡 tp=4)
50 lines
1.4 KiB
C
50 lines
1.4 KiB
C
#pragma once
|
|
|
|
#ifdef USE_ROCM
|
|
#include <hip/hip_runtime.h>
|
|
#endif
|
|
|
|
#ifndef USE_ROCM
|
|
#define WARP_SIZE 32
|
|
#else
|
|
#define WARP_SIZE warpSize
|
|
#endif
|
|
|
|
#ifndef USE_ROCM
|
|
#define VLLM_LDG(arg) __ldg(arg)
|
|
#else
|
|
#define VLLM_LDG(arg) *(arg)
|
|
#endif
|
|
|
|
#ifndef USE_ROCM
|
|
#define VLLM_SHFL_XOR_SYNC(var, lane_mask) \
|
|
__shfl_xor_sync(uint32_t(-1), var, lane_mask)
|
|
#define VLLM_SHFL_XOR_SYNC_WIDTH(var, lane_mask, width) \
|
|
__shfl_xor_sync(uint32_t(-1), var, lane_mask, width)
|
|
#else
|
|
#define VLLM_SHFL_XOR_SYNC(var, lane_mask) __shfl_xor(var, lane_mask)
|
|
#define VLLM_SHFL_XOR_SYNC_WIDTH(var, lane_mask, width) \
|
|
__shfl_xor(var, lane_mask, width)
|
|
#endif
|
|
|
|
#ifndef USE_ROCM
|
|
#define VLLM_SHFL_SYNC(var, src_lane) __shfl_sync(uint32_t(-1), var, src_lane)
|
|
#else
|
|
#define VLLM_SHFL_SYNC(var, src_lane) __shfl(var, src_lane)
|
|
#endif
|
|
|
|
#ifndef USE_ROCM
|
|
#define VLLM_SHFL_DOWN_SYNC(var, lane_delta) \
|
|
__shfl_down_sync(uint32_t(-1), var, lane_delta)
|
|
#else
|
|
#define VLLM_SHFL_DOWN_SYNC(var, lane_delta) __shfl_down(var, lane_delta)
|
|
#endif
|
|
|
|
#ifndef USE_ROCM
|
|
#define VLLM_DevFuncAttribute_SET_MaxDynamicSharedMemorySize(FUNC, VAL) \
|
|
cudaFuncSetAttribute(FUNC, cudaFuncAttributeMaxDynamicSharedMemorySize, VAL)
|
|
#else
|
|
#define VLLM_DevFuncAttribute_SET_MaxDynamicSharedMemorySize(FUNC, VAL) \
|
|
hipFuncSetAttribute(FUNC, hipFuncAttributeMaxDynamicSharedMemorySize, VAL)
|
|
#endif
|