| """BF16 text activation math shared by pointwise and GEMM epilogues.""" | |
| import cutlass | |
| import cutlass.cute as cute | |
| from clef_vision_norm import mul_rn | |
| 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) | |