Skip to content

[Docs]: Fact-check T.tile.bitwise_rshift API - #1676

Open
zhangfei-machao wants to merge 3 commits into
tile-ai:ascendc_ptofrom
zhangfei-machao:factcheck/rshift
Open

[Docs]: Fact-check T.tile.bitwise_rshift API #1676
zhangfei-machao wants to merge 3 commits into
tile-ai:ascendc_ptofrom
zhangfei-machao:factcheck/rshift

Conversation

@zhangfei-machao

Copy link
Copy Markdown
  1. 文件改动
    文件 类型 改动说明
    docs/language/tile_bitwise_rshift.md 新增 校验后的 API 文档,包含函数原型、参数说明、dtype 支持矩阵、约束条件、已知限制
    tilelang/language/ascend_tile.py 修改 bitwise_rshift docstring 扩充:补充移位语义(有符号算术/无符号逻辑)、支持 dtype、UB scope、scalarValue 自动类型转换
    testing/python/language/test_tilelang_ascend_language_elementwise.py 修改 新增 112 个参数化测试用例,覆盖签名语义、dtype、shape、边界值、非对齐、不支持类型
  2. 测试用例
    2.1 约束测试(动态验证)
    验证项 测什么
    有符号算术右移 int16/int32 负值输入,shift=1/2/4,AscendC+PTO
    无符号逻辑右移 uint16/uint32 高位为 1 的值,shift=1/4,AscendC+PTO
    1D shape 支持 int16/uint16/int32/uint32,1D buffer,shift=1/4/8
    边界值 shift=0 int16/uint16/int32/uint32,AscendC+PTO
    边界值 shift=max-1 int16 shift=15,int32 shift=31
    边界值 shift=max int16 shift=16,int32 shift=32
    scalarValue 类型不匹配 Python int(int32 IR)传入 int16/uint16 操作
    非对齐 shape (PTO) shape=(1024,100),4 dtype,shift=1/4
    非对齐 shape (AscendC) 同上,AscendC 后端
    int8/uint8 不支持 编译 int8/uint8 kernel
    int64/uint64 不支持 编译 int64/uint64 kernel
    float16/float32 不支持 编译 float kernel
    2.2 正式测试文件
    测试函数 用例数 测什么
    test_bitwise_rshift_signed_semantics 12 有符号类型算术右移语义(int16/int32,shift=1/2/4,AscendC+PTO)
    test_bitwise_rshift_unsigned_high_bits 8 无符号类型逻辑右移语义(uint16/uint32,高位为 1,shift=1/4,AscendC+PTO)
    test_bitwise_rshift_1d 24 1D buffer 支持(4 dtype × shift=1/4/8 × AscendC+PTO)
    test_bitwise_rshift_boundary_int16 16 16-bit 边界值(int16/uint16,shift=0/1/15/16,AscendC+PTO)
    test_bitwise_rshift_boundary_int32 16 32-bit 边界值(int32/uint32,shift=0/1/31/32,AscendC+PTO)
    test_bitwise_rshift_non_aligned 8 非对齐 shape PTO 精度(4 dtype,shift=1/4)
    test_bitwise_rshift_non_aligned_ascendc_xfail 8 非对齐 shape AscendC 精度错误(预期失败)
    test_bitwise_rshift_int8_xfail 4 int8/uint8 编译失败(预期失败)
    test_bitwise_rshift_int64_skip 4 int64/uint64 编译器段错误(跳过)
    test_bitwise_rshift_float_xfail 4 float16/float32 编译失败(预期失败)
    test_bitwise_rshift_scalar_dtype_mismatch 8 scalarValue dtype 与 dst 不一致时 codegen 自动转换
  3. 测试结果
  • 测试环境:Ascend 910B3、CANN 8.5.2、torch 2.7.1+cpu、torch_npu 2.7.1.post4、tilelang 0.1.4、Python 3.11.4
  • 约束测试:92 PASSED + 16 XFAILED + 4 SKIPPED(全部预期行为)
  • 正式测试:108 PASSED(含 16 既有 + 92 新增运行时用例),16 XFAILED,4 SKIPPED,0 FAILED
  • 必要回归测试:test_bitwise_xor 4 PASSED
    3.1 Pytest 标签
    本次新增测试未使用额外 pytest 标签。
  1. 校验过程中发现的问题
  2. scalarValue 类型约束错误:原文档声称"scalarValue 的数据类型需与 dst 元素类型一致(Ascend C 约束)"。核查 codegen(TshCodegen / ShiftOpCodegen)发现当 scalarValue 的 dtype 与 src0 不一致时,codegen 自动插入类型转换 int16(scalarValue)。真机测试(8 用例,AscendC+PTO,4 dtype)全部通过。文档修正为"scalarValue 的 dtype 不要求与 dst 一致(codegen 自动转换)"。
  3. int8/uint8 不支持:原文档在 A5 dtype 表中列出 int8/uint8,但未说明 A2/A3 不支持。真机验证(910B3)表明 AscendC ShiftRight 无 int8/uint8 特化,PTO TShiftCheck static_assert 仅允许 int16/uint16/int32/uint32,两者均编译失败。文档新增 A2/A3 完整 dtype 支持矩阵并标注 int8/uint8 不支持。
  4. int64/uint64 编译器段错误:原文档声称 A5 支持 int64/uint64。真机验证(910B3/A2A3)发现 int64/uint64 导致 TVM 编译器 segfault 而非优雅编译错误,已用 pytest.mark.skip 标记规避。文档标注为不支持。
  5. A5 支持范围未经真机验证:原文档声称 A5 支持 int8/uint8/int16/uint16/int32/uint32/int64/uint64。经核查 PTO a5 TSHRS_IMPL static_assert 允许 int8-uint32,但不包含 int64/uint64。当前环境(A2/A3)无法验证 A5 行为,文档标注为"未经真机验证"。
  6. 非对齐 shape 精度问题:原文档声称"操作数地址需 32 字节对齐"。真机验证发现 AscendC 后端在非对齐 shape(如 1024×100)下产生精度错误,PTO 后端不受影响。与 bitwise_lshift 表现一致。文档新增"已知限制"章节说明。
  7. 移位语义验证:原文档声称"无符号逻辑右移、有符号算术右移"。真机验证(910B3,20 用例)确认该描述正确,与 bitwise_lshift(左移为逻辑移位)不同,右移根据 C++ 模板参数 T 的符号性自动区分逻辑/算术右移。
  8. 未解决问题与风险
  • int64/uint64 在 A2/A3 导致 TVM 编译器段错误,根因未排查,已用 skip 标记规避
  • A5 平台所有 dtype 支持声明未经真机验证(无 A5 硬件)
  • AscendC 非对齐 shape 精度错误根因未深入排查,与 lshift 一致
  1. 验证
  • python -m pytest ... -k "rshift" -vv:通过,108 passed, 4 skipped, 16 xfailed, 0 failed(873.79s)
  • python -m pytest ... -k "bitwise_xor and not slice" -vv:通过,4 passed(42.09s)
  • python -c "import ast; ast.parse(open('...elementwise.py').read())":通过
  • python -c "import ast; ast.parse(open('ascend_tile.py').read())":通过
  • git diff --check:通过,无空白错误
    汇总:
  • 编译验证:通过(所有有效 dtype kernel 编译成功)
  • NPU 运行:通过(910B3 真机运行成功)
  • 精度验证:通过(rtol=0, atol=0,与 PyTorch 参考实现完全一致)
  • 目标测试:128 collected, 108 passed, 4 skipped, 16 xfailed
  • 回归测试:4 passed(test_bitwise_xor)
  • 语法检查:通过
  • Ruff/格式检查:未执行(环境不可用)

@github-actions

Copy link
Copy Markdown

👋 Hi! Thank you for contributing to the TileLang project.

Please remember to run bash format.sh in the root directory of the project to ensure your changes are properly linted and formatted. This will help ensure your contribution passes the format check.

We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work!

🚀

Mark 9 test functions as low_priority to reduce PR CI runtime:
- test_bitwise_rshift_1d (24 cases, shape variant)
- test_bitwise_rshift_boundary_int16/int32 (32 cases, boundary edge)
- test_bitwise_rshift_non_aligned/_ascendc_xfail (16 cases, non-aligned)
- test_bitwise_rshift_int8_xfail (4 cases, unsupported dtype)
- test_bitwise_rshift_int64_skip (4 cases, segfault skip)
- test_bitwise_rshift_float_xfail (4 cases, unsupported dtype)
- test_bitwise_rshift_scalar_dtype_mismatch (8 cases, dtype edge case)

Keep 2 core semantic tests in PR CI:
- test_bitwise_rshift_signed_semantics (12 cases, arithmetic shift)
- test_bitwise_rshift_unsigned_high_bits (8 cases, logical shift)

PR CI: 26 rshift tests (down from 128), ~3.7min runtime.
@zhangfei-machao

Copy link
Copy Markdown
Author

/re-test

@github-actions

Copy link
Copy Markdown

🔄 Re-running failed jobs

Original workflow run: View details

Only the failed jobs will be re-executed.

@zhangfei-machao

Copy link
Copy Markdown
Author

/re-test

@github-actions

Copy link
Copy Markdown

🔄 Re-running failed jobs

Original workflow run: View details

Only the failed jobs will be re-executed.

@zhangfei-machao

Copy link
Copy Markdown
Author

/re-test

@github-actions

Copy link
Copy Markdown

🔄 Re-running failed jobs

Original workflow run: View details

Only the failed jobs will be re-executed.

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.

1 participant