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_log1p(x0):
return torch.log1p(x0)
@triton.jit
def triton_log1p(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.log1p(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.log1p(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_log1p_common(param_list):
dtype, shape, ncore, xblock, xblock_sub = param_list
x0 = test_common.generate_tensor(shape, dtype).npu()
y_ref = torch_log1p(x0)
y_cal = torch.zeros(shape, dtype=eval('torch.' + dtype)).npu()
if PERF_TEST_ENABLE:
test_common.run_with_profiler(
lambda: triton_log1p[ncore, 1, 1](x0, y_cal, xblock, xblock_sub),
shape,
'log1p'
)
else:
triton_log1p[ncore, 1, 1](x0, y_cal, xblock, xblock_sub)
test_common.validate_cmp(dtype, y_cal, y_ref)
@pytest.mark.parametrize(
'param_list',
[['float32', (64,), 1, 64, 64]],
)
def test_log1p_special_values(param_list):
dtype, shape, ncore, xblock, xblock_sub = param_list
x0 = test_common.generate_tensor(shape, dtype).npu()
x0[0] = float('inf')
x0[1] = float('-inf')
x0[2] = float('nan')
x0[3] = 0.0
x0[4] = -0.0
x0[5] = 1e-38
x0[6] = -1.0
x0[7] = 1.0
y_ref = torch_log1p(x0)
y_cal = torch.zeros(shape, dtype=eval('torch.' + dtype)).npu()
triton_log1p[ncore, 1, 1](x0, y_cal, xblock, xblock_sub)
test_common.validate_cmp(dtype, y_cal, y_ref)