[Hopper CuTeDSL] Fix max reduction in fmha kernel

[Aladoro]

The old reduction logic seemed to have mistakenly called cutlass.max rather than cute.arch.fmax during one of the in-place reductions:

...
 # Linear reduction is faster here, as well
  for j in cutlass.range_constexpr(cute.size(acc_qk_mn, mode=[1])):
      s_max[i] = cute.arch.fmax(s_max[i], acc_qk_mn[i, j])
      # old code: s_max[i] = cutlass.max(s_max[i], acc_qk_mn[i, j])
...

This seems to have been accidental (the kernel previously already uses the correct cute.arch.fmax for the first iteration), and seems very costly, emitting several operations rather than just nvmm fmax:

setp.le.f32
setp.nan.f32
selp.f32
selp.f32

Validation

Environment: NVIDIA H100 SxM, CUDA 12.9

Before the fix (default shapes/dtypes): 3884.84 us, 566.05 TFLOP/s
After: 3440.17 us, 639.22 TFLOP/s