Skip to content

[NPU][Bugfix] use CANN libdevice for DSpark Markov sampling - #43849

Merged
sglang-npu-bot merged 1 commit into
sgl-project:mainfrom
AndyLi429:fix/npu-dspark-markov-sample
Oct 14, 2026
Merged

sglang-npu-bot merged 1 commit into
sgl-project:mainfrom
AndyLi429:fix/npu-dspark-markov-sample

Conversation

@AndyLi429

@AndyLi429 AndyLi429 commented Oct 12, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

DeepSeek V4 Flash with DSpark fails during NPU draft graph capture in [this CI job](https://github.com/sgl-project/sglang/actions/runs/38233348303/job/114772999125).

MarkovSampleStep calls -libdevice.log1p(-uniform), but the Ascend backend does not map the generic Triton libdevice stub to its CANN implementation. The call returns None, causing compilation to fail with:

AST unary operator '__neg__' is not (currently) implemented on type NoneType

The sampling kernel was introduced in #43621 and integrated into draft graph capture in #43668.

Modifications

  • Import triton.language.extra.cann.libdevice on NPU in markov_sample.py.
  • Retain the existing generic libdevice import on other platforms.

Accuracy Tests

  • Pre-commit checks passed.
  • Host validation with mocked platform/Triton imports and real CPU PyTorch passed: library selection, greedy special values, deterministic sampling, and noise-offset advancement.
  • Pending follow-up: validate NPU compilation, graph capture/replay, and rerun the failing accuracy job.

Speed Tests and Profiling

Pending follow-up: verify NPU inference performance after restoring successful compilation. No performance improvement is claimed.

Checklist

Review and Merge Process

  1. Ping Merge Oncalls to start the process. See the PR Merge Process.
  2. Get approvals from CODEOWNERS and other reviewers.
  3. Trigger CI tests with comments or contact authorized users to do so.
    • Common commands include /tag-and-rerun-ci, /tag-run-ci-label, /rerun-failed-ci
  4. After green CI and required approvals, ask Merge Oncalls or people with Write permission to merge the PR.

CI States

Latest PR Test (Base): ✅ Run #38245525089
Latest PR Test (Extra): ❌ Run #38245524642
Latest PR Test (AMD ROCm 10): ❌ Run #38245525423

Select the Ascend libdevice implementation on NPU so log1p lowers to the CANN intrinsic instead of executing the generic Triton stub. This fixes draft graph capture failing with a NoneType unary negation error. Keep the existing libdevice selection on other platforms.
@AndyLi429
AndyLi429 marked this pull request as ready for review October 12, 2026 12:18
@chatgpt-codex-connector

chatgpt-codex-connector Bot commented Oct 12, 2026 •

Copy link
Copy Markdown

Codex Review Summary

This comment shows the latest Codex review activity on this pull request.

Review Status Commit Review trigger
📝 Code Review ✅ Completed 2026-10-12T12:26:51.591965Z 8e4fd24 Draft marked ready
ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

  • Open a pull request for review
  • Mark a draft as ready
  • Comment "@codex review" or "@codex security review".

Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings.

@sglang-npu-bot

Copy link
Copy Markdown
Collaborator

/tag-and-rerun-ci

@github-actions github-actions Bot added the run-ci CI: run the baseline test suite on this PR label Oct 12, 2026
@AndyLi429

Copy link
Copy Markdown
Contributor Author

fix base-c-test-acc-16-npu-a3 test_npu_deepseek_v4_flash_w8a8_8p_gpqa.py

@AndyLi429

Copy link
Copy Markdown
Contributor Author

/rerun-failed-ci

1 similar comment
@AndyLi429

Copy link
Copy Markdown
Contributor Author

/rerun-failed-ci

@sglang-npu-bot
sglang-npu-bot merged commit 05bb5af into sgl-project:main Oct 14, 2026
483 of 552 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

jit-kernel run-ci CI: run the baseline test suite on this PR

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants