import pytest
import triton
import triton.language as tl
import triton.language.extra.cann.libdevice as libdevice
import test_common
import torch
import os
PERF_TEST_ENABLE = os.getenv('PERF_TEST_ENABLE', 'False').lower() == 'true'
def torch_atan(x0):
return torch.atan(x0)
@triton.jit
def triton_atan(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.atan(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.atan(x0)
tl.store(out_ptr0 + rem_xindex, y, mask=mask)
default_param_list = test_common.make_default_param_list(['float32'])
full_param_list = test_common.make_full_param_list(['float32'])
@pytest.mark.parametrize(
'param_list',
default_param_list if not PERF_TEST_ENABLE else full_param_list,
)
def test_atan(param_list):
dtype, shape, ncore, xblock, xblock_sub = param_list
x0 = test_common.generate_tensor(shape, dtype).npu()
y_ref = torch_atan(x0)
y_cal = torch.zeros_like(y_ref)
if PERF_TEST_ENABLE:
test_common.run_with_profiler(
lambda: triton_atan[ncore, 1, 1](x0, y_cal, xblock, xblock_sub),
shape,
'atan'
)
else:
triton_atan[ncore, 1, 1](x0, y_cal, xblock, xblock_sub)
test_common.validate_cmp(dtype, y_cal, y_ref)