[PixelFlat]

Description

Fix Ascend AOT compilation of apply_token_bitmask while preserving its negative-infinity masking semantics.

TileLang lowers float("-inf") to -CUDART_INF, but the Ascend compilation environment does not necessarily define that CUDA macro. The generated kernel then fails with use of undeclared identifier 'CUDART_INF'.

• Prepend a guarded CUDART_INF definition using __builtin_inff() in this kernel's AOT source generator. No CUDA headers, external compiler definitions, or compiler wrapper are required.
• Keep the kernel algorithm, dispatch ABI, and runtime allocations unchanged; do not replace -inf with a finite sentinel.
• Add a deterministic NPU regression covering FP16, BF16, and FP32: disallowed finite logits become negative infinity, allowed finite/infinite logits remain unchanged, and no NaNs are introduced.

Change Type

• Bug fix
• Test
• Build or CI

Validation

Validated on Linux ARM64 / Ascend A3 with CANN 9.0.0:

• python setup.py build --device npu --tilelang-jobs 8 passed with the original Bisheng compiler and no external infinity macro injection. All three token-bitmask dtype specializations compiled successfully.
• Negative control: removing only the compatibility header from the generated sources reproduces the original undefined-CUDART_INF error for all three dtypes with the same compiler.
• The following focused test binaries passed, with 41 tests, zero failures, and zero skips:

Test binary	Passed
apply_token_bitmask_wrapper_test	4
sampler_test	15
sampler_controls_test	19
sampler_filter_mask_test	3

• Changed-file pre-commit checks passed, including clang-format 20.1.6.
• Whole-tree Ruff 0.14.0 lint and formatting checks passed.
• git diff --check passed.

This is a single focused commit. It does not modify TileLang or operator-package dependencies.

[chatgpt-codex-connector]

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-25T13:45:16.567973Z	bb760a9	PR opened

ℹ️ 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.