Fuse online softmax: compute running max and sum in one tree reduction

[pranathi000]

This change fuses the two softmax reductions into one.

Currently both softmaxKernel and softmaxKernelDecode do two full tree reductions per row: the first finds the maximum, and the second sums the exponentials, each with its own __syncthreads() loop. This PR combines them into a single reduction using the online softmax formulation already used per-token in pagedAttentionKernel.

Each thread now carries a (max, sum) pair up the tree and merges them together at every level:

m = max(m_a, m_b)
d = d_a * exp(m_a - m) + d_b * exp(m_b - m)

The exp(m_a - m) term rescales a partial sum whenever a larger max appears, which keeps the computation numerically stable. When the reduction reaches the root, both the row max and the normalized denominator are already available, so the separate max broadcast and the intermediate exponentiation pass are no longer needed.

This roughly halves the number of __syncthreads() barriers per row. It adds a couple of expf calls per merge, but on GPU the reduction in barriers is the better trade-off.
Kernel signatures, launch configurations, and the existing 1024-thread guard are unchanged. Shared memory usage goes from a single float[1024] (plus one scalar) to two float[1024] arrays.

[jmaczan]

Hi there, thanks for the contribution! Looks good to me, and I love that you actually implemented the proper online softmax. Please include add some tests or at least show some proofs that it works to make sure there is no regression, and then we can merge. Thank you!