tilelang.contrib.ptodsl.simd_inst

Functions

vdiv_precise_f32(src0, src1, mask)	Emit precise f32 vector division matching Ascend vdiv_precise.
vexp_1ulp_ftz_false(src, mask)	Emit subnormal-preserving f32 vexp matching Ascend vexp_1ulp_ftz_false.
vln_1ulp_ftz_false(src, mask)	Emit subnormal-preserving f32 vln matching Ascend vln_1ulp_ftz_false.
vsqrt_0ulp_ftz_false(src, mask)	Emit subnormal-preserving f32 vsqrt matching Ascend vsqrt_0ulp_ftz_false.

Module Contents

tilelang.contrib.ptodsl.simd_inst.vdiv_precise_f32(src0, src1, mask)

Emit precise f32 vector division matching Ascend vdiv_precise.

tilelang.contrib.ptodsl.simd_inst.vexp_1ulp_ftz_false(src, mask)

Emit subnormal-preserving f32 vexp matching Ascend vexp_1ulp_ftz_false.

Normal outputs pass through; outputs that would land in the subnormal range are computed as (e^(x/2))^2 so the SFU input stays normal. Self-consistent at x < -174.7: e^(x/2) itself is flushed to 0, squaring yields 0, and the true e^x < 2^-252 correctly rounds to 0.

tilelang.contrib.ptodsl.simd_inst.vln_1ulp_ftz_false(src, mask)

Emit subnormal-preserving f32 vln matching Ascend vln_1ulp_ftz_false.

Positive subnormal inputs are scaled by 2^23 before VLN and compensated by -ln(2^23) (the scaled value is exact: the subnormal mantissa shifts into the normal range). Other inputs keep hardware semantics (0 -> -inf, negative -> NaN).

tilelang.contrib.ptodsl.simd_inst.vsqrt_0ulp_ftz_false(src, mask)

Emit subnormal-preserving f32 vsqrt matching Ascend vsqrt_0ulp_ftz_false.

Replica of CANN SqrtFastInverseImpl (PRECISION_0ULP_FTZ_FALSE / FAST_INVERSE), chosen over the 1ULP scale/unscale variant because the latter mis-rounds 0x007fffff to +0. Inputs < 1 are scaled by 2^24 so the chain stays in the normal range, then unscaled by 2^-12; a 1/sqrt initial value plus a Newton step and a second residual correction gives correct rounding; +-0 and +inf pass through.