"""BF16 text activation math shared by pointwise and GEMM epilogues.""" import cutlass import cutlass.cute as cute from clef_vision_norm import mul_rn @cute.jit def silu_bf16(value): """Input is BF16-rounded; validate the complete finite BF16 domain on SM90.""" denominator = 1.0 + cute.math.exp(-value, fastmath=True) activated = cutlass.Float32(0) if value > -87.0: inverse = cute.math.rcp(denominator, fastmath=True) activated = mul_rn(value, inverse) else: # Preserve tiny negative activations when the reciprocal would flush to zero. activated = value / denominator return activated.to(cutlass.BFloat16)