import os
import triton
import triton.language as tl
from triton.language.extra.cann.libdevice import pow
import torch
import torch_npu
import pytest
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_pow(in_ptr0, in_ptr1, 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)
x1 = tl.load(in_ptr1 + xindex)
y = pow(x0, x1)
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)
x1 = tl.load(in_ptr1 + rem_xindex, mask=mask)
y = pow(x0, x1)
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_pow(param_list):
dtype, shape, ncore, xblock, xblock_sub = param_list
x0 = test_common.generate_tensor(shape, dtype).npu()
x1 = test_common.generate_tensor(shape, dtype).npu()
y_ref = torch.pow(x0, x1)
y_cal = torch.zeros_like(y_ref)
if PERF_TEST_ENABLE:
test_common.run_with_profiler(
lambda: triton_pow[ncore, 1, 1](x0, x1, y_cal, xblock, xblock_sub),
shape,
'pow'
)
else:
triton_pow[ncore, 1, 1](x0, x1, y_cal, xblock, xblock_sub)
test_common.validate_cmp(dtype, y_cal, y_ref)