Suboptimal warp reductions

It is preferred to implement reductions with __shfl_xor_sync as oppose to __shfl_down_sync. This way all lanes will have the final value and you no longer need to broadcast final value to other lanes.

```
// warp-level reduction for finding the maximum value
__device__ float warpReduceMax(float val) {
for (int offset = 16; offset > 0; offset /= 2) {
val = fmaxf(val, __shfl_xor_sync(0xFFFFFFFF, val, offset));
}
return val;
}

// warp-level reduction for summing values
__device__ float warpReduceSum(float val) {
for (int offset = 16; offset > 0; offset /= 2) {
val += __shfl_xor_sync(0xFFFFFFFF, val, offset);
}
return val;
}
```

Start by locating the warp-level reduction implementations in the repository and reviewing their use of __shfl_down_sync. Compare them with the __shfl_xor_sync examples in this issue, then verify that the CUDA training code still builds and the reduction results remain correct.

Written by the indexing model from the issue text.