div

1. OP 概述

简介：除法，四则运算 ‘/’，无tl.div方法

底层实现与fdiv算子相同，只不过fdiv明确限制入参必须是float类型，‘/’无此限制，它会把非浮点型转换为浮点型再做计算

2. OP 规格

2.1 参数说明

参数名	类型	说明
self	tensor or Number	第一个入参 ，被除数
other	tensor or Number	第二个入参 ，除数

返回值： tl.tensor：除法结果 返回结果类型：总是返回浮点型

输入类型	处理方式	结果类型
int / int	两个都转成 float32	float32
int / float	int 转 float	float 类型
float / float	统一到更高精度 float	高精度 float
float / int	int 转 float	float 类型

2.2 支持规格

2.2.1 DataType 支持

	uint8	int8	uint16	int16	uint32	int32	uint64	int64	fp16	fp32	bf16	bool/int1
GPU	√	√	√	√	√	√	√	√	√	√	√	√
Ascend A2/A3	×	√	×	√	×	√	×	√	√	√	√	√

2.2.2 Shape 支持

	支持维度范围
GPU	无限制
Ascend A2/A3	无限制

结论：在 Shape 方面，GPU 与 Ascend 平台无差异。

2.3 特殊限制说明

Ascend A3 对比 GPU 缺失uint8、uint16、uint32、uint64、fp64的支持

2.4 使用方法

以下示例实现了对输入张量 in_ptr0, in_ptr1 做除法计算：

@triton.jit
def triton_div(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)