Revert "[ATen] Fix CUDA reduction warp shuffle order (pytorch#164790)"

This reverts commit 8e1f409.

Reverted pytorch#164790 on behalf of https://github.com/jeffdaily due to broke cuda and rocm ci ([comment](pytorch#164790 (comment)))

Author: pytorchmergebot

aten/src/ATen/native/cuda/Reduce.cuh

@@ -655,17 +655,12 @@ struct ReduceOp {
     }
 
     __syncthreads();
-    // Warp-level reduction for remaining threads
-    // For non-power-of-2 sizes, we start from the next power-of-2 divided by 2
-    // and use a boundary check to avoid out-of-bounds access
-    for (size_t offset = warpSize / 2; offset > 0; offset >>= 1) {
+
+    for (int offset = 1; offset < dim_x; offset <<= 1) {
       #pragma unroll
       for (int i = 0; i < output_vec_size; i++) {
         arg_t other = ops.warp_shfl_down(value[i], offset);
-        // Only combine if the source thread (threadIdx.x + offset) is within bounds
-        if (threadIdx.x + offset < dim_x) {
-          value[i] = ops.combine(value[i], other);
-        }
+        value[i] = ops.combine(value[i], other);
       }
     }
     return value;

aten/src/ATen/native/cuda/reduction_template.cuh

@@ -466,13 +466,11 @@ struct ReduceJitOp {
 
     __syncthreads();
 
-    for (size_t offset = warpSize / 2; offset > 0; offset >>= 1) {
+    for (int offset = 1; offset < dim_x; offset <<= 1) {
       #pragma unroll
       for (int i = 0; i < output_vec_size; i++) {
         arg_t other = reducer::warp_shfl_down(value[i], offset);
-        if (threadIdx.x + offset < dim_x) {
-          value[i] = reducer::combine(value[i], other);
-        }
+        value[i] = reducer::combine(value[i], other);
       }
     }
     return value;