triton.language.mod

triton.language.mod(x, y)

Element-wise remainder of x divided by y, the % operator on tensors.

• float % float — floating-point remainder (frem)

• int % int — integer remainder; signed integers use srem, unsigned integers use urem. Operands with mismatched signedness raise TypeError.

参数:

• x (Block or scalar number) -- the dividend

• y (Block or scalar number) -- the divisor

示例

import triton
import triton.language as tl
import torch


def torch_mod(x0, x1):
    res = x0 % x1
    return res


@triton.jit
def triton_mod(in_ptr0, in_ptr1, out_ptr0, XBLOCK: tl.constexpr, XBLOCK_SUB: tl.constexpr):
    offset = tl.program_id(0) * XBLOCK
    base1 = tl.arange(0, XBLOCK_SUB)
    loops1: tl.constexpr = (XBLOCK + XBLOCK_SUB - 1) // XBLOCK_SUB
    for loop1 in range(loops1):
        x0 = offset + (loop1 * XBLOCK_SUB) + base1
        tmp0 = tl.load(in_ptr0 + (x0), None)
        tmp1 = tl.load(in_ptr1 + (x0), None)
        tmp2 = tmp0 % tmp1
        tl.store(out_ptr0 + (x0), tmp2, None)


def test_mod():
    param_list = ['int32', (2, 4096, 8), 2, 32768, 1024]
    dtype, shape, ncore, xblock, xblock_sub = param_list
    x0 = torch.randint(1, 10, size=shape, dtype=eval('torch.' + dtype)).npu()
    x1 = torch.randint(1, 10, size=shape, dtype=eval('torch.' + dtype)).npu()

    torch_res = torch_mod(x0, x1)
    triton_res = torch.empty_like(x0)
    triton_mod[ncore, 1, 1](x0, x1, triton_res, xblock, xblock_sub)

    torch.testing.assert_close(torch_res, triton_res, rtol=0, atol=0, equal_nan=True)


if __name__ == '__main__':
    test_mod()

数据类型支持

平台	uint8	int8	uint16	int16	uint32	int32	uint64	int64	fp16	fp32	fp64	bf16	fp8e(e4m3)	fp8e5(e5m2)	bool
Ascend A2/A3	√	√	×	√	×	√	×	√	√	√	×	√	×	×	√
Ascend 950	√	√	√	√	√	√	√	√	√	√	×	√	√	√	√