Skip to content

[Docs](fix) Repair A5 libdevice API example tests from #1782 #1886

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
hockyy wants to merge 1 commit into
base: main
Choose a base branch
from
Open

[Docs](fix) Repair A5 libdevice API example tests from #1782 #1886

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.associative_scan.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,7 @@ def test_associative_scan_2d():

expected = torch.cumsum(x.cpu(), dim=dim)

assert torch.allclose(output.cpu(), expected.cpu())
assert torch.allclose(output.cpu(), expected.cpu(), rtol=1e-3, atol=1e-3)


if __name__ == "__main__":
Expand Down
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.brev.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,7 @@ def _brev(x):
@pytest.mark.skipif(not triton_enable_libdevice_simt(), reason=_SIMT_SKIP_MSG)
def test_brev():
x0 = (torch.randint(1, 16, (8, ))).to(torch.int32).npu()
expected = torch.tensor([_brev(*v) for v in zip(x0.tolist())], dtype=torch.int32).npu()
expected = torch.tensor([_brev(v) & 0xFFFFFFFF for v in x0.tolist()], dtype=torch.int64).to(torch.int32).npu()
output = torch.empty(8, dtype=torch.int32, device='npu')
triton_kernel[(1, )](x0, output, 8, XBLOCK=8, XBLOCK_SUB=8, force_simt_only=True)
torch.testing.assert_close(output, expected, rtol=0, atol=0)
Expand Down
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.cbrt.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,7 +32,7 @@ def triton_kernel(input0, output, n_elements, XBLOCK: tl.constexpr, XBLOCK_SUB:
@pytest.mark.skipif(not triton_enable_libdevice_simt(), reason=_SIMT_SKIP_MSG)
def test_cbrt():
x0 = (torch.rand((8, )) * 3.0 - 1.5).to(torch.float32).npu()
expected = (torch.cbrt(x0)).to(torch.float32).npu()
expected = (torch.sign(x0) * torch.pow(torch.abs(x0), 1.0 / 3.0)).to(torch.float32).npu()
output = torch.empty(8, dtype=torch.float32, device='npu')
triton_kernel[(1, )](x0, output, 8, XBLOCK=8, XBLOCK_SUB=8, force_simt_only=True)
torch.testing.assert_close(output, expected, rtol=1e-03, atol=1e-03, equal_nan=True)
Expand Down
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.erfcinv.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,7 +32,7 @@ def triton_kernel(input0, output, n_elements, XBLOCK: tl.constexpr, XBLOCK_SUB:
@pytest.mark.skipif(not triton_enable_libdevice_simt(), reason=_SIMT_SKIP_MSG)
def test_erfcinv():
x0 = (torch.rand((8, )) * 0.9 + 0.05).to(torch.float32).npu()
expected = (torch.erfc(torch.erfcinv(x0))).to(torch.float32).npu()
expected = torch.erfinv(1.0 - x0).to(torch.float32).npu()
output = torch.empty(8, dtype=torch.float32, device='npu')
triton_kernel[(1, )](x0, output, 8, XBLOCK=8, XBLOCK_SUB=8, force_simt_only=True)
torch.testing.assert_close(output, expected, rtol=1e-03, atol=1e-03, equal_nan=True)
Expand Down
6 changes: 3 additions & 3 deletions docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.ffs.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,11 +30,11 @@ def triton_kernel(input0, output, n_elements, XBLOCK: tl.constexpr, XBLOCK_SUB:


def _ffs(x):
"""Index of the least significant set bit (0-based), 32 if none."""
"""1-based index of the least significant set bit; 0 if none (CUDA ffs)."""
v = int(x) & 0xFFFFFFFF
if v == 0:
return 32
return (v & -v).bit_length() - 1
return 0
return (v & -v).bit_length()


@pytest.mark.skipif(not triton_enable_libdevice_simt(), reason=_SIMT_SKIP_MSG)
Expand Down
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.rcbrt.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,7 +32,7 @@ def triton_kernel(input0, output, n_elements, XBLOCK: tl.constexpr, XBLOCK_SUB:
@pytest.mark.skipif(not triton_enable_libdevice_simt(), reason=_SIMT_SKIP_MSG)
def test_rcbrt():
x0 = (torch.rand((8, )) + 0.1).to(torch.float32).npu()
expected = (1.0 / torch.cbrt(x0)).to(torch.float32).npu()
expected = (torch.sign(x0) * torch.pow(torch.abs(x0), -1.0 / 3.0)).to(torch.float32).npu()
output = torch.empty(8, dtype=torch.float32, device='npu')
triton_kernel[(1, )](x0, output, 8, XBLOCK=8, XBLOCK_SUB=8, force_simt_only=True)
torch.testing.assert_close(output, expected, rtol=1e-03, atol=1e-03, equal_nan=True)
Expand Down
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.rhadd.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,7 +34,7 @@ def triton_kernel(input0, input1, output, n_elements, XBLOCK: tl.constexpr, XBLO
def test_rhadd():
x0 = (torch.randint(1, 16, (8, ))).to(torch.int32).npu()
x1 = (torch.randint(1, 16, (8, ))).to(torch.int32).npu()
expected = ((x0 + x1) // 2).to(torch.int32).npu()
expected = ((x0.to(torch.int64) + x1.to(torch.int64) + 1) // 2).to(torch.int32).npu()
output = torch.empty(8, dtype=torch.int32, device='npu')
triton_kernel[(1, )](x0, x1, output, 8, XBLOCK=8, XBLOCK_SUB=8, force_simt_only=True)
torch.testing.assert_close(output, expected, rtol=0, atol=0)
Expand Down
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.uint2float_rd.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ def triton_kernel(input0, output, n_elements, XBLOCK: tl.constexpr, XBLOCK_SUB:
x0 = offset + (loop * XBLOCK_SUB) + base
mask = x0 < n_elements
tmp0 = tl.load(input0 + (x0), mask=mask)
tmp1 = libdevice.uint2float_rd(tmp0)
tmp1 = libdevice.uint2float_rd(tmp0.to(tl.uint32))
tl.store(output + (x0), tmp1, mask=mask)


Expand Down
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.uint2float_rn.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ def triton_kernel(input0, output, n_elements, XBLOCK: tl.constexpr, XBLOCK_SUB:
x0 = offset + (loop * XBLOCK_SUB) + base
mask = x0 < n_elements
tmp0 = tl.load(input0 + (x0), mask=mask)
tmp1 = libdevice.uint2float_rn(tmp0)
tmp1 = libdevice.uint2float_rn(tmp0.to(tl.uint32))
tl.store(output + (x0), tmp1, mask=mask)


Expand Down
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.uint2float_ru.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ def triton_kernel(input0, output, n_elements, XBLOCK: tl.constexpr, XBLOCK_SUB:
x0 = offset + (loop * XBLOCK_SUB) + base
mask = x0 < n_elements
tmp0 = tl.load(input0 + (x0), mask=mask)
tmp1 = libdevice.uint2float_ru(tmp0)
tmp1 = libdevice.uint2float_ru(tmp0.to(tl.uint32))
tl.store(output + (x0), tmp1, mask=mask)


Expand Down
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.uint2float_rz.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ def triton_kernel(input0, output, n_elements, XBLOCK: tl.constexpr, XBLOCK_SUB:
x0 = offset + (loop * XBLOCK_SUB) + base
mask = x0 < n_elements
tmp0 = tl.load(input0 + (x0), mask=mask)
tmp1 = libdevice.uint2float_rz(tmp0)
tmp1 = libdevice.uint2float_rz(tmp0.to(tl.uint32))
tl.store(output + (x0), tmp1, mask=mask)


Expand Down
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.ull2float_rd.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ def triton_kernel(input0, output, n_elements, XBLOCK: tl.constexpr, XBLOCK_SUB:
x0 = offset + (loop * XBLOCK_SUB) + base
mask = x0 < n_elements
tmp0 = tl.load(input0 + (x0), mask=mask)
tmp1 = libdevice.ull2float_rd(tmp0)
tmp1 = libdevice.ull2float_rd(tmp0.to(tl.uint64))
tl.store(output + (x0), tmp1, mask=mask)


Expand Down
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.ull2float_rn.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ def triton_kernel(input0, output, n_elements, XBLOCK: tl.constexpr, XBLOCK_SUB:
x0 = offset + (loop * XBLOCK_SUB) + base
mask = x0 < n_elements
tmp0 = tl.load(input0 + (x0), mask=mask)
tmp1 = libdevice.ull2float_rn(tmp0)
tmp1 = libdevice.ull2float_rn(tmp0.to(tl.uint64))
tl.store(output + (x0), tmp1, mask=mask)


Expand Down
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.ull2float_ru.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ def triton_kernel(input0, output, n_elements, XBLOCK: tl.constexpr, XBLOCK_SUB:
x0 = offset + (loop * XBLOCK_SUB) + base
mask = x0 < n_elements
tmp0 = tl.load(input0 + (x0), mask=mask)
tmp1 = libdevice.ull2float_ru(tmp0)
tmp1 = libdevice.ull2float_ru(tmp0.to(tl.uint64))
tl.store(output + (x0), tmp1, mask=mask)


Expand Down
2 changes: 1 addition & 1 deletion docs/zh/python-api/_examples/triton.language.extra.cann.libdevice.ull2float_rz.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ def triton_kernel(input0, output, n_elements, XBLOCK: tl.constexpr, XBLOCK_SUB:
x0 = offset + (loop * XBLOCK_SUB) + base
mask = x0 < n_elements
tmp0 = tl.load(input0 + (x0), mask=mask)
tmp1 = libdevice.ull2float_rz(tmp0)
tmp1 = libdevice.ull2float_rz(tmp0.to(tl.uint64))
tl.store(output + (x0), tmp1, mask=mask)


Expand Down
Loading