`hexp` is #define'd to `cutlass::fast_exp`, so the hand-written
`cutlass::fast_exp(bfloat16_t)` overload's `return ::hexp(x)` expanded into a
call to itself: bf16 `T.exp` recursed and returned garbage.
Route through float, mirroring cutlass's own `half_t fast_exp`
(cutlass/fast_math.h), so the call resolves to `cutlass::fast_exp(float)`
(= ::expf). Verified against torch.exp using the #2383 repro.
Fixes#2383.