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_log10(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.log10(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.log10(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_log10(param_list):
dtype, shape, ncore, xblock, xblock_sub = param_list
x = test_common.generate_tensor(shape, dtype).npu()
x[0, 0] = 1.0
x[0, 1] = 10.0
x[0, 2] = 100.0
x[0, 3] = 0.1
x[0, 4] = 0.0
x[0, 5] = -1.0
x[0, 6] = 2.0
y_ref = torch.log10(x)
y_cal = torch.zeros_like(y_ref)
if PERF_TEST_ENABLE:
test_common.run_with_profiler(
lambda: triton_log10[ncore, 1, 1](x, y_cal, xblock, xblock_sub),
shape,
'log10'
)
else:
triton_log10[ncore, 1, 1](x, y_cal, xblock, xblock_sub)
valid_mask = (x > 0)
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)
negative_mask = (x < 0)
if torch.any(negative_mask):
negative_y = y_cal[negative_mask]
assert torch.all(torch.isnan(negative_y)), "Negative inputs should return NaN"
zero_mask = (x == 0)
if torch.any(zero_mask):
zero_y = y_cal[zero_mask]
assert torch.all(torch.isinf(zero_y) & (zero_y < 0)), "Zero inputs should return -inf"