Skip to content

[Bug] PTODSL's pto.simt_allreduce_sum/max/min rejects signless integer inputs #1580

Open

Description

@liggest

Component

PTODSL

Description

PTODSL 中,pto.simt_allreduce_sum/max/min 不接受 pto.i32,只认 f32、f16、si32、ui32,传入 pto.i32 的值直接抛出:

NotImplementedError: all_reduce: unsupported dtype i32

当前框架中 pto.i32 是合法的标量类型:pto.cast(x, pto.i32)、pto.ptr(pto.i32, "gm")
常规算术和 pto.load/store 也都接受,但编译执行到 pto.simt_allreduce_sum/max/min 就会报错

Reproduction (minimal)

from ptodsl import pto


@pto.jit(name="kernel", kernel_kind="vector", target="a5", mode="explicit")
def kernel(A: pto.ptr(pto.i32, "gm")):
    pto.init_core()
    with pto.simt(32, 1, 1):
        v = pto.cast(pto.get_tid_x(), pto.i32)
        r = pto.simt_allreduce_max(v, threads=32)
        pto.store(r, A, 0)


kernel.compile()
print("compiled")

Expected behavior

PTODSL 成功编译,功能正常
当前框架中,希望 pto.i32 行为与 pto.si32 一致

Actual behavior / error logs

Traceback (most recent call last):
  File "test_simt_allreduce_i32.py", line 14, in <module>
    kernel.compile()
  File "ptodsl/ptodsl/_jit.py", line 343, in compile
    return self._compiler.compile(**constexpr_bindings)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "ptodsl/ptodsl/_kernel_compilation.py", line 142, in compile
    compiled.build()
  File "ptodsl/ptodsl/_tracing/artifacts.py", line 36, in build
    built = self._module_factory()
            ^^^^^^^^^^^^^^^^^^^^^^
  File "ptodsl/ptodsl/_tracing/runtime.py", line 93, in build_module
    self.trace_entry(*args)
  File "ptodsl/ptodsl/_tracing/runtime.py", line 135, in trace_entry
    result = self._callback(*args, **kwargs)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "test_simt_allreduce_i32.py", line 10, in kernel
    r = pto.simt_allreduce_max(v, threads=32)
        ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "ptodsl/ptodsl/_allreduce.py", line 423, in simt_allreduce_sum
    return _simt_allreduce(value, threads=threads, scale=scale,
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "ptodsl/ptodsl/_allreduce.py", line 389, in _simt_allreduce
    raise NotImplementedError(f"all_reduce: unsupported dtype {raw_value.type}")
NotImplementedError: all_reduce: unsupported dtype i32

Git commit

0a3e017

Host platform

Linux (x86_64)

Target Ascend arch (if relevant)

a5

PTOAS build level (if relevant)

None

Activity

  1. added
    bugSomething isn't working
    tilelangIssues from tilelang
    on Sep 24, 2026
  2. changed the title [-][Bug] PTODSL's `pto.simt_allreduce_*` rejects signless `pto.i32` integer inputs[/-] [+][Bug] PTODSL's `pto.simt_allreduce_sum/max/min` rejects signless `pto.i32` integer inputs[/+] on Sep 24, 2026
  3. changed the title [-][Bug] PTODSL's `pto.simt_allreduce_sum/max/min` rejects signless `pto.i32` integer inputs[/-] [+][Bug] PTODSL's `pto.simt_allreduce_sum/max/min` rejects signless integer inputs[/+] on Sep 24, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    bugSomething isn't workingtilelangIssues from tilelang

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions