Skip to content

[V100/SM70] FastDeploy V100 support: custom CUDA kernels, prefill attention, and runtime bugfixes - #6306

Open
mattheliu wants to merge 54 commits into
PaddlePaddle:developfrom
mattheliu:fastdeploy_v100
Open

mattheliu wants to merge 54 commits into
PaddlePaddle:developfrom
mattheliu:fastdeploy_v100

Conversation

@mattheliu

@mattheliu mattheliu commented Feb 2, 2026 •

Copy link
Copy Markdown
Collaborator

Motivation

为 FastDeploy 添加 NVIDIA V100 GPU (SM70 架构) 支持,使其能在旧版 GPU 上进行推理部署。

V100 (Volta 架构) 不支持以下特性,需要同时适配编译系统和运行时逻辑:

特性 最低要求 V100 支持
BF16 数据类型 SM80+ (Ampere) ❌
FP8 量化 SM89+ (Ada Lovelace) ❌
cp.async 指令 SM80+ (Ampere) ❌
flash_attn_unpadded SM80+ (Ampere) ❌
tanh.approx.f32 PTX SM75+ (Turing) ❌
Marlin GEMM SM80+ (Ampere) ❌

Modifications

新增文件 (6)

文件 说明
custom_ops/gpu_ops/v100_decode_attention.cu Two-stage paged decode attention,online softmax + LSE merge,7x 加速
custom_ops/gpu_ops/v100_prefill_attention.cu Prefill attention,causal masking + paged KV cache,替代 Python fallback
custom_ops/gpu_ops/v100_rope_write_cache.cu Fused RoPE + KV cache write,支持 interleaved/neox 两种风格
fastdeploy/model_executor/layers/attention/v100_flash_attn_backend.py V100 attention backend,CUDA→Triton→Python 三级 fallback
fastdeploy/model_executor/ops/triton_ops/v100_attn_kernels.py Triton fallback kernels (write_kv_cache, decode_fused stage1/stage2)
tests/model_executor/ops/triton_ops/test_v100_attn_kernels.py V100 Triton kernel 单元测试

修改文件 (54)

编译系统 & CUDA (14)

  • setup_ops.py: SM70/80/89 编译分层,cc>=70 MoE+V100 kernels,cc>=80 BF16/Marlin/append_attn,cc>=89 FP8
  • cpp_extensions.cc: #ifdef ENABLE_BF16/FP8/APPEND_ATTENTION 条件编译守卫
  • gelu_tanh.cu: SM75 tanh fallback + block size 修复
  • moe_reduce.cu: 修复 FP16 case 错误调用 BF16 模板
  • sampling.cuh: pivot 初始化 +inf→-inf + SM70 host 编译保护
  • moe_wna16_marlin_gemm.cu / kernel.h / marlin_template.h: SM70 stub 用 PD_THROW 替代 torch::empty
  • moe_deepgemm_depermute.cu: SM70 BF16 算术通过 float 转换实现
  • stop_generation_multi_ends.cu / group_swiglu_with_masked.cu / swigluoai.cu: 跨架构兼容修复
  • auto_gen_template_instantiation.py: --skip-fp8 支持

平台 & 运行时 (4)

  • platforms/base.py: 添加 V100_FLASH_ATTN backend enum
  • platforms/cuda.py: SM 版本检测、BF16→FP16 自动降级、V100 backend 自动路由
  • config.py: BF16→FP16 dtype 降级,V100 禁用 CUDA Graph
  • gpu_model_runner.py: 抢占时 empty_cache() + host 端同步守卫

Attention & 模型层 (22)

  • attention/__init__.py: 注册 V100FlashAttentionBackend
  • attention/mla_attention_backend.py / native_paddle_backend.py: SM 兼容性修复
  • attention/ops/*.py (5 个): SM80+ 专属 ops 添加 try-except 保护
  • moe/moe.py: Marlin/Triton → CUTLASS fallback (SM<80)
  • moe/fused_moe_cutlass_backend.py / fused_moe_deepgemm_backend.py: API 适配
  • quantization/__init__.py: FP8 → INT8/INT4 自动 fallback
  • quantization/block_wise_fp8.py / mix_quant.py / weight_only.py: SM89+ 守卫
  • embeddings.py / linear.py / lm_head.py / normalization.py / load_weight_utils.py: BF16 weight → FP16 cast
  • ernie4_5_moe.py / pre_and_post_process.py / utils.py / token_processor.py: 运行时兼容修复
  • triton_ops/__init__.py: V100 kernel 导出
  • engine/ (4 个) / entrypoints/ (2 个) / worker/ (2 个): 运行时兼容修复

测试 (6)

  • test_attention_layer.py / test_fusedmoe.py / test_w4afp8.py: FP8 SM89+ skip 装饰器
  • test_ffn.py: SM 版本自动选择 dtype
  • test_platforms.py: 平台检测测试

SM70/SM75 Fallback 策略

功能 原始 Fallback 原因
数据类型 BF16 FP16 SM80+
Attention APPEND/MLA/FLASH_ATTN V100_FLASH_ATTN cp.async SM80+
CUDA Graph 启用 禁用 Python 后端不兼容
MoE Backend Marlin/Triton CUTLASS SM80+
量化 block_wise_fp8/w4afp8 wint8/wint4 FP8 SM89+

V100 Attention 架构

Decode:  CUDA C++ kernel → Triton v100_decode_fused → Python SDPA
Prefill: CUDA prefill kernel → Python SDPA

Usage or Command

# 编译
cd custom_ops && python setup_ops.py install

# 推理
python -c "
from fastdeploy import LLM, SamplingParams
llm = LLM(model='PaddlePaddle/ERNIE-4.5-0.3B-Paddle', max_model_len=2048, max_num_seqs=4)
outputs = llm.chat([[{'role': 'user', 'content': 'Hello'}]], SamplingParams(max_tokens=50))
print(outputs[0].outputs.text)
"

Accuracy Tests

已验证模型(V100, SM70):

  • ✅ ERNIE-4.5-0.3B (interleaved RoPE):32 并发,~550 tok/s,零 CUDA error
  • ✅ Qwen3-0.6B (neox RoPE):正常推理
  • ✅ C-Eval 1346 题,准确率 30.31%(与 0.3B 量级预期一致)

本 PR 为硬件兼容性支持,不影响现有 SM80+/SM90 平台计算逻辑。

Checklist

  • Add at least a tag in the PR title.
  • Format your code, run pre-commit before commit.
  • Add unit tests.
    • test_v100_attn_kernels.py 覆盖 V100 Triton kernel
    • SM 版本 skip 装饰器确保测试在不支持硬件上正确跳过
  • Provide accuracy results.
  • If the current PR is submitting to the release branch, make sure the PR has been submitted to the develop branch, then cherry-pick it to the release branch with the [Cherry-Pick] PR tag.

@paddle-bot

paddle-bot Bot commented Feb 2, 2026

Copy link
Copy Markdown

Thanks for your contribution!

@mattheliu

Copy link
Copy Markdown
Collaborator Author

/Re-run failed jobs

@mattheliu
mattheliu marked this pull request as ready for review February 4, 2026 05:47
@codecov-commenter

codecov-commenter commented Feb 4, 2026 •

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 25.64470% with 519 lines in your changes missing coverage. Please review.
⚠️ Please upload report for BASE (develop@81acdb6). Learn more about missing BASE report.

Files with missing lines Patch % Lines
...ecutor/layers/attention/v100_flash_attn_backend.py 12.54% 271 Missing and 1 partial ⚠️
...model_executor/ops/triton_ops/v100_attn_kernels.py 30.00% 97 Missing and 1 partial ⚠️
...loy/model_executor/layers/quantization/__init__.py 27.02% 25 Missing and 2 partials ⚠️
...oy/model_executor/layers/quantization/mix_quant.py 25.00% 20 Missing and 1 partial ⚠️
..._executor/layers/moe/fused_moe_deepgemm_backend.py 6.25% 15 Missing ⚠️
fastdeploy/config.py 36.36% 7 Missing and 7 partials ⚠️
...executor/layers/attention/native_paddle_backend.py 31.57% 12 Missing and 1 partial ⚠️
fastdeploy/platforms/cuda.py 70.45% 10 Missing and 3 partials ⚠️
.../model_executor/layers/quantization/weight_only.py 21.42% 7 Missing and 4 partials ⚠️
fastdeploy/model_executor/layers/moe/moe.py 22.22% 4 Missing and 3 partials ⚠️
... and 10 more
Additional details and impacted files
@@            Coverage Diff             @@
##             develop    #6306   +/-   ##
==========================================
  Coverage           ?   70.76%           
==========================================
  Files              ?      394           
  Lines              ?    54353           
  Branches           ?     8518           
==========================================
  Hits               ?    38461           
  Misses             ?    13109           
  Partials           ?     2783           
Flag Coverage Δ
GPU 70.76% <25.64%> (?)

Flags with carried forward coverage won't be shown. Click here to find out more.

☔ View full report in Codecov by Sentry.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

mattheliu and others added 28 commits April 1, 2026 12:50
Stage1 was storing unnormalized acc, but stage2 merge formula
expects normalized partials. Divide acc by l_i before storing.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
RoPE is a lightweight memory-bound op where Triton kernel launch
overhead dominates at small token counts. Use Paddle vectorized ops
instead, which benchmarks show are 10x faster for typical batch sizes.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
…partial_out

When num_kv_splits > actual KV blocks, empty splits in stage1 leave
partial_out uninitialized. In stage2, exp(-inf - (-inf)) produces 0,
but 0 * NaN (IEEE 754) = NaN, corrupting the entire output. This caused
all decode tokens to be <unk> (token_id=0) on V100.

Fix:
- Initialize partial_out with zeros instead of empty
- Add is_valid guard in stage2 to skip loading from empty splits
- Add decode tests with small kv_len (7, 3) to cover this edge case

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
1. Add FD_V100_USE_PYTHON_ATTN=1 env var to force Python/Paddle fallback
   when Triton kernels produce wrong results on SM70. This helps isolate
   whether decode issues are caused by Triton codegen problems.

2. Replace _simple_attention_forward with zeros return for dummy/profile
   runs. The naive attention computes O(n^2) score matrix (e.g. 16 heads *
   8192 * 8192 = 4GB), causing OOM on V100 32GB for larger models like
   Qwen3-0.6B. Since V100 Triton attention uses tiled flash-decoding with
   O(1) extra memory, returning zeros gives a more accurate memory estimate.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Unit tests pass but model produces wrong output. The KV cache is
pre-populated in tests, but in the model pipeline v100_write_kv_cache
writes KV then v100_paged_attention immediately reads it. If Triton
kernels run on a different CUDA stream than Paddle ops, the attention
kernel may read stale/zero cache data.

Add paddle.device.cuda.synchronize() between steps 4 and 5 to test
this hypothesis. If this fixes the model output, we need proper stream
synchronization instead of a full device sync.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
…on Qwen3/ERNIE

- Fix extend_attention kernel: cast tl.dot inputs to fp16 with fp32 accumulator
  for SM70 Tensor Core compatibility (fp32 tl.dot causes GPU hang on V100)
- Refactor _triton_forward to use proven-correct Python data prep with Triton
  paged attention only, ensuring correctness on both Qwen3-0.6B and ERNIE-4.5-0.3B
- Add vectorized Paddle helpers (_paddle_compute_positions, _paddle_compute_total_seq_lens,
  _paddle_write_kv_to_block_cache) for future performance optimization
- Add unit tests for ERNIE-4.5 config (num_heads=8, kv_num_heads=2) and multi-split decode

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
…riton OOM

Replace Triton paged attention with Python/Paddle cuBLAS SDPA for both decode
and prefill phases. The Triton JIT compilation memory overhead was causing OOM
on V100 32GB during prefill (KV cache ~25GB leaves insufficient headroom).

This establishes a correct, working baseline where both _triton_forward and
_python_forward use identical logic. Triton kernels are imported for upcoming
decode-only Triton optimization.

Benchmarks (Qwen3-0.6B, batch=1, 64 tokens):
- Before: Triton path ~28s (due to .item() overhead), Python path ~6s
- After: Both paths ~6s, correct output, no OOM

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
…riton

Add a CUDA C++ custom op (v100_decode_attention) that replaces Triton
flash-decoding kernels for V100 decode attention. This eliminates the
torch_proxy kernel launch overhead (~1.5ms per Triton launch × 28 layers),
reducing 512-token generation from 153s to 21.5s (7x speedup, 9.5 tok/s).

Three CUDA kernels in v100_decode_attention.cu:
- v100_write_kv_cache_kernel: vectorized KV cache write
- v100_decode_attn_stage1_kernel: flash-decoding with online softmax
- v100_decode_attn_stage2_kernel: LSE-based partial output merge

Backend changes (v100_flash_attn_backend.py):
- CUDA C++ kernel as primary decode path (~0.01ms launch overhead)
- Triton v100_decode_fused as fallback (when CUDA op unavailable)
- Cross-layer caching: positions, seq_lens, q_start_locs computed once
  at layer 0 and reused across all 28 layers
- Small KV (≤2 blocks) still uses Python SDPA path

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
…backend

After adding the CUDA C++ decode attention kernel, several Triton functions
became dead code. This commit removes them and updates tests accordingly:

- v100_attn_kernels.py: delete 6 unused functions (compute_positions,
  fused_rope, decode_attn_stage1, decode_attention, extend_attention,
  paged_attention), keeping only write_kv_cache, decode_fused, stage2
- v100_flash_attn_backend.py: remove unused Triton imports, keep only
  v100_decode_fused
- test_v100_attn_kernels.py: delete tests for removed functions, rewrite
  TestDecodeAttention as TestDecodeFusedAttention testing v100_decode_fused

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
The __init__.py imported v100_compute_positions, v100_decode_attention,
v100_extend_attention, v100_fused_rope, v100_paged_attention which were
deleted in the dead code cleanup. This caused ImportError that silently
disabled ALL Triton ops (including moe_wint2_ffn_kernel) on every GPU.

Fix: separate V100 Triton imports into their own try/except block so
failures don't affect other Triton ops. Import only the 2 surviving
functions: v100_decode_fused and v100_write_kv_cache.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Remove unused methods (_paddle_compute_positions, _paddle_compute_total_seq_lens,
_paddle_write_kv_to_block_cache, _simple_attention_forward), unused imports
(scaled_dot_product_attention), unused instance vars (use_speculate, rope_3d,
_use_fp16), and unused metadata fields (cu_seqlens_k, max_len_tensor_cpu_decoder).

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
- Add PD_CHECK(head_dim <= THREADS*4) in v100_decode_attention.cu to
  prevent silent buffer overflow for large head_dim models
- Unify block_wise_fp8 fallback to wint8 in quantization/__init__.py,
  consistent with mix_quant.py (was incorrectly disabling quantization)
- Remove 4 debug print statements from test_ffn.py
- Extract duplicated _check_fp8_support() from 3 test files into
  shared check_fp8_support() in tests/utils.py

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
…it__.py)

The `from tests.utils import check_fp8_support` fails on machines where
tests/ is not a Python package. Revert to inline _check_fp8_support()
in each test file.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
Prevents import failure on non-V100 environments where V100-specific
dependencies might not be available.

Generated with [Claude Code](https://claude.ai/code)
via [Happy](https://happy.engineering)

Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Happy <yesreply@happy.engineering>
SM70 (V100) and SM75 (Turing) don't support BF16 MMA instructions.
When compiling purely for these architectures, the generated
kernel_bf16_*.cu files cause undefined symbol errors at .so load time
because BF16 Marlin template instantiations reference SM80+ intrinsics.

Filter out kernel_bf16 files in setup_ops.py when cc < 80.

Co-Authored-By: Claude <noreply@anthropic.com>
When compiling purely for SM70, the host-side BF16 dispatch in
moe_wna16_marlin_gemm.cu references Marlin<nv_bfloat16> kernel
symbols that don't exist (SM70 has no BF16 MMA support).

Two-part fix:
1. setup_ops.py: Add -DMARLIN_DISABLE_BF16 flag when cc < 80
2. moe_wna16_marlin_gemm.cu: Guard BF16 dispatch with #ifdef

SM80+ builds are unaffected (MARLIN_DISABLE_BF16 not defined).

Co-Authored-By: Claude <noreply@anthropic.com>
…tly used BFLOAT16 template

The FLOAT16 switch case in MoeExpertReduceKernel was calling
MoeReduceKernel<paddle::DataType::BFLOAT16> instead of FLOAT16,
causing dtype mismatch errors on V100 (SM70) and any FP16 inference.

Co-Authored-By: Claude <noreply@anthropic.com>
1. CacheConfig BF16 leak: cache_dtype was not adjusted for V100 hardware,
   causing BF16 KV cache allocation on SM70 which doesn't support BF16.
   Added hardware check in CacheConfig.__init__() and after read_from_config().

2. V100 attention backend BF16 passthrough: init_attention_metadata() only
   warned about BF16 but still set metadata._dtype=bfloat16. Now forces
   FP16 for correctness on V100.

3. MoE SwiGLU hardcoded BF16: swigluoai.cu and group_swiglu_with_masked.cu
   had PD_CHECK(dtype==BFLOAT16) which crashes when model dtype is auto-
   downgraded to FP16 on V100. Added FP16 dispatch path.

Generated with [Claude Code](https://claude.ai/code)

Co-Authored-By: Claude <noreply@anthropic.com>
Add v100_rope_write_cache.cu - a fused CUDA kernel that combines:
1. RoPE application to Q and K
2. KV cache write to paged block cache

This replaces the Python implementations (_python_apply_rope_to_qk and
_python_write_kv_to_block_cache) which were the main performance bottleneck
on V100. Expected speedup: 10-50x for data preparation phase.

Key changes:
- custom_ops/gpu_ops/v100_rope_write_cache.cu: New fused kernel
- custom_ops/gpu_ops/v100_decode_attention.cu: Add skip_kv_write param
- v100_flash_attn_backend.py: Integrate CUDA kernel, add batched SDPA
- v100_attn_kernels.py: Add skip_kv_write to Triton path
- setup_ops.py: Add new kernel to build

V100 architecture now matches A100/H100 pattern:
- v100_rope_write_cache (RoPE + KV write) - NEW
- v100_decode_attention (decode attention) - existing
- Paddle SDPA (prefill attention) - existing

Co-Authored-By: Claude <noreply@anthropic.com>
…ion bugs

1. Rewrite v100_rope_write_cache.cu: single kernel, PD_BUILD_STATIC_OP
   with SetInplaceMap, correct cos/sin stride (pos*rotary_dim+d),
   NeoX-only RoPE, float4 vectorized V copy, zero-waste 2D grid
2. Update Python _cuda_rope_write_cache() to match new inplace interface:
   pre-allocate q_out/k_out, pass as first two args, non-NeoX fallback
3. Fix batched SDPA query indexing: use batch_id->token_idx mapping
   instead of list index (wrong when batch_ids have gaps)
4. Add .contiguous() on cos/sin slices for CUDA memory safety
5. Add default value for skip_kv_write in v100_decode_attention.cu

Co-Authored-By: Claude <noreply@anthropic.com>
…utput

Root cause: data processors (ernie4_5_processor.py, text_processor.py)
converted temperature < _SAMPLING_EPS (1e-5) to temperature=1, but
forgot to set top_p to a small value. With temperature=1 and top_p=0.8
from generation_config.json, the sampler used random top-p sampling.

Fix: when temperature is near-zero (greedy), set top_p=_SAMPLING_EPS
in addition to temperature=1. This forces paddle.tensor.top_p_sampling
to select only the top token (argmax behavior).

Impact: All benchmarks using temperature=0 (IFEval, BBH, ZebraLogic,
LiveCodeBench) were producing wrong random answers instead of greedy
predictions.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
When max_dec_len is reached, stop_flags was set True BEFORE calling
set_stop_value_multi_ends. The CUDA kernel sees stop_flags=True and
replaces sampled_token_ids with EOS, discarding the actual generated
token. This caused max_tokens=1 requests to always return EOS.

Fix: move the stop_flags |= length_cond assignment to AFTER the kernel
call, so the sampled token is preserved for length-triggered stops.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
- setup_ops.py: keep cc>=70 MoE build block + upstream cc>=75 SM75 ext ops
- gpu_model_runner.py: keep V100 BFC flush (empty_cache) + adapt to upstream interface
- ernie4_5_processor.py / text_processor.py: upstream refactored to super().__init__(),
  greedy sampling fix (top_p=_SAMPLING_EPS when temperature=0) preserved via base class

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
- black/isort: auto-format Python files
- clang-format: auto-format CUDA C++ files
- Fix undefined names: logger (weight_only.py), layers_are_grouped
  (load_weight_utils.py), x_fp8->x (fused_moe_deepgemm_backend.py)
- Fix unused variables: q_len/kv_len (v100_flash_attn_backend.py),
  _e (ernie4_5_moe.py), _exec_cost/start_execute_time (worker_process.py)
- Restore append_attention.py from upstream (out_swa undefined)
- Restore _maybe_view_bf16_as_fp16 in load_weight_utils.py for V100
- Add noqa comments for intentional patterns

Co-Authored-By: mattheliu <leonliuzx@outlook.com>
append_attention CUDA op is not compiled on SM70 (V100). Without the
try-except guard, importing this module crashes on V100 at startup.

Co-Authored-By: mattheliu <leonliuzx@outlook.com>

@PaddlePaddle-bot PaddlePaddle-bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🤖 AI Code Review | {{TIMESTAMP}}

📋 Review 摘要

PR 概述:为 FastDeploy 添加 NVIDIA V100 GPU (SM70 架构) 支持,包括自定义 CUDA kernel、BF16→FP16 降级、Attention backend fallback 等。

变更范围:custom_ops/gpu_ops/, fastdeploy/model_executor/layers/, fastdeploy/platforms/

影响面 Tag:OP Engine Models


📝 PR 规范检查

标题问题:使用了非官方 Tag [V100/SM70],应使用官方 Tag 列表中的标签。

标题建议(可直接复制):

  • [Feature][OP] Add NVIDIA V100 (SM70) GPU support with custom CUDA kernels
  • 或 [Feature] FastDeploy V100 support: custom CUDA kernels and runtime fallbacks

问题

级别 文件 概述
🔴 Bug custom_ops/setup_ops.py:426 ENABLE_APPEND_ATTENTION 宏未定义,导致 AppendAttention 等函数无法编译

总体评价

PR 整体设计良好,V100 fallback 策略清晰,新增的 CUDA kernel 实现正确。但存在一个关键的编译配置问题:cpp_extensions.cc 中使用 ENABLE_APPEND_ATTENTION 宏保护 AppendAttention 相关代码,但 setup_ops.py 中未定义该宏,会导致 SM80+ 设备上这些函数无法被正确编译和调用。修复此问题后可以合入。

Comment thread custom_ops/setup_ops.py
if cc >= 70:
nvcc_compile_args += [
"-Igpu_ops/moe",
"-DENABLE_BF16",

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔴 Bug 条件编译宏 ENABLE_APPEND_ATTENTION 未定义

cpp_extensions.cc 中使用 #ifdef ENABLE_APPEND_ATTENTION 保护了以下函数的声明和 pybind 注册:

  • AppendAttention
  • PreCacheLenConcat
  • GetBlockShapeAndSplitKVBlock

但 setup_ops.py 中从未定义 ENABLE_APPEND_ATTENTION 宏(第 486-487 行仅定义了 ENABLE_SM80_EXT_OPS)。

影响:这会导致 SM80+ 设备上 AppendAttention 等函数的 Python 绑定不会被编译,运行时调用会抛出 AttributeError。

建议修复:在 cc >= 80 分支中添加宏定义:

if cc >= 80:
    cc_compile_args += ["-DENABLE_SM80_EXT_OPS", "-DENABLE_APPEND_ATTENTION"]
    nvcc_compile_args += ["-DENABLE_SM80_EXT_OPS", "-DENABLE_APPEND_ATTENTION"]

This branch had an error being deployed

1 failed deployment
Metax_ci — 03b4fb90 Deployed Apr 1, 2026 by mattheliu via Trigger Jenkins for PR #6377
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants