import os
import pytest
import triton
import torch
import torch_npu
import triton.language as tl
import triton.language.extra.cann.libdevice as libdevice
import test_common
PERF_TEST_ENABLE = os.getenv('PERF_TEST_ENABLE', 'False').lower() == 'true'
default_param_list = test_common.make_default_param_list(['float32'])
full_param_list = test_common.make_full_param_list(['float32'])
@triton.jit
def triton_asin(in_ptr0, out_ptr0, XBLOCK: tl.constexpr, XBLOCK_SUB: tl.constexpr):
xoffset = tl.program_id(0) * XBLOCK
base = tl.arange(0, XBLOCK_SUB)
loops: tl.constexpr = XBLOCK // XBLOCK_SUB
for loop in range(loops):
xindex = xoffset + loop * XBLOCK_SUB + base
x0 = tl.load(in_ptr0 + xindex)
y = libdevice.asin(x0)
tl.store(out_ptr0 + xindex, y)
remaining: tl.constexpr = XBLOCK % XBLOCK_SUB
if remaining > 0:
rem_xindex = xoffset + loops * XBLOCK_SUB + base
mask = base < remaining
x0 = tl.load(in_ptr0 + rem_xindex, mask=mask)
y = libdevice.asin(x0)
tl.store(out_ptr0 + rem_xindex, y, mask=mask)
@pytest.mark.parametrize(
'param_list',
default_param_list if not PERF_TEST_ENABLE else full_param_list,
)
def test_asin(param_list):
dtype, shape, ncore, xblock, xblock_sub = param_list
x = test_common.generate_tensor(shape, dtype).npu()
x[0, 0] = 0.0
x[0, 1] = 0.5
x[0, 2] = -0.5
x[0, 3] = 1.0
x[0, 4] = -1.0
x[0, 5] = 0.707
x[0, 6] = 0.866
x[0, 7] = 1.1
x[0, 8] = -1.1
y_ref = torch.asin(x)
y_cal = torch.zeros_like(y_ref)
if PERF_TEST_ENABLE:
test_common.run_with_profiler(
lambda: triton_asin[ncore, 1, 1](x, y_cal, xblock, xblock_sub),
shape,
'asin'
)
else:
triton_asin[ncore, 1, 1](x, y_cal, xblock, xblock_sub)
valid_mask = (x >= -1) & (x <= 1)
if torch.any(valid_mask):
valid_y = y_cal[valid_mask]
valid_expected = y_ref[valid_mask]
torch.testing.assert_close(valid_y, valid_expected, rtol=1e-3, atol=1e-3)
invalid_mask = (x < -1) | (x > 1)
if torch.any(invalid_mask):
invalid_y = y_cal[invalid_mask]
assert torch.all(torch.isnan(invalid_y)), "Invalid inputs should return NaN"