fix: error for EmptyArray

Author: ManasviGoyal

src/awkward/_connect/cuda/cuda_kernels/awkward_ListOffsetArray_reduce_local_outoffsets_64.cu

@@ -59,17 +59,17 @@ awkward_ListOffsetArray_reduce_local_outoffsets_64_b(
     }
     __syncthreads();
 
-    for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
-      int64_t val = 0;
-      if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
-        val = temp[thread_id - stride];
+    if (thread_id < lenparents) {
+      for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
+        int64_t val = 0;
+        if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
+          val = temp[thread_id - stride];
+        }
+        __syncthreads();
+        temp[thread_id] += val;
+        __syncthreads();
       }
-      __syncthreads();
-      temp[thread_id] += val;
-      __syncthreads();
-    }
 
-    if (thread_id < lenparents) {
       int64_t parent = parents[thread_id];
       if (idx == blockDim.x - 1 || thread_id == lenparents - 1 || parents[thread_id] != parents[thread_id + 1]) {
         atomicAdd(&scan_in_array[parent], temp[thread_id]);

src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_argmax.cu

@@ -59,19 +59,19 @@ awkward_reduce_argmax_b(
     }
     __syncthreads();
 
-    for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
-      int64_t index = -1;
-      if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
-        index = temp[thread_id - stride];
-      }
-      if (index != -1 && (temp[thread_id] == -1 || fromptr[index] > fromptr[temp[thread_id]] ||
-         (fromptr[index] == fromptr[temp[thread_id]] && index < temp[thread_id]))) {
-        temp[thread_id] = index;
+    if (thread_id < lenparents) {
+      for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
+        int64_t index = -1;
+        if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
+          index = temp[thread_id - stride];
+        }
+        if (index != -1 && (temp[thread_id] == -1 || fromptr[index] > fromptr[temp[thread_id]] ||
+          (fromptr[index] == fromptr[temp[thread_id]] && index < temp[thread_id]))) {
+          temp[thread_id] = index;
+        }
+        __syncthreads();
       }
-      __syncthreads();
-    }
 
-    if (thread_id < lenparents) {
       int64_t parent = parents[thread_id];
       if (idx == blockDim.x - 1 || thread_id == lenparents - 1 || parents[thread_id] != parents[thread_id + 1]) {
         atomicExch(&atomic_toptr[parent], temp[thread_id]);

src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_argmin.cu

@@ -59,19 +59,19 @@ awkward_reduce_argmin_b(
     }
     __syncthreads();
 
-    for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
-      int64_t index = -1;
-      if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
-        index = temp[thread_id - stride];
-      }
-      if (index != -1 && (temp[thread_id] == -1 || fromptr[index] < fromptr[temp[thread_id]] ||
-         (fromptr[index] == fromptr[temp[thread_id]] && index < temp[thread_id]))) {
-        temp[thread_id] = index;
+    if (thread_id < lenparents) {
+      for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
+        int64_t index = -1;
+        if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
+          index = temp[thread_id - stride];
+        }
+        if (index != -1 && (temp[thread_id] == -1 || fromptr[index] < fromptr[temp[thread_id]] ||
+          (fromptr[index] == fromptr[temp[thread_id]] && index < temp[thread_id]))) {
+          temp[thread_id] = index;
+        }
+        __syncthreads();
       }
-      __syncthreads();
-    }
 
-    if (thread_id < lenparents) {
       int64_t parent = parents[thread_id];
       if (idx == blockDim.x - 1 || thread_id == lenparents - 1 || parents[thread_id] != parents[thread_id + 1]) {
         atomicExch(&atomic_toptr[parent], temp[thread_id]);

src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_count_64.cu

@@ -52,17 +52,17 @@ awkward_reduce_count_64_b(
     }
     __syncthreads();
 
-    for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
-      int64_t val = 0;
-      if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
-        val = temp[thread_id - stride];
+    if (thread_id < lenparents) {
+      for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
+        int64_t val = 0;
+        if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
+          val = temp[thread_id - stride];
+        }
+        __syncthreads();
+        temp[thread_id] += val;
+        __syncthreads();
       }
-      __syncthreads();
-      temp[thread_id] += val;
-      __syncthreads();
-    }
 
-    if (thread_id < lenparents) {
       int64_t parent = parents[thread_id];
       if (idx == blockDim.x - 1 || thread_id == lenparents - 1 || parents[thread_id] != parents[thread_id + 1]) {
         atomicAdd(&toptr[parent], temp[thread_id]);

src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_countnonzero.cu

@@ -54,17 +54,17 @@ awkward_reduce_countnonzero_b(
     }
     __syncthreads();
 
-    for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
-      int64_t val = 0;
-      if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
-        val = temp[thread_id - stride];
+    if (thread_id < lenparents) {
+      for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
+        int64_t val = 0;
+        if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
+          val = temp[thread_id - stride];
+        }
+        __syncthreads();
+        temp[thread_id] += val;
+        __syncthreads();
       }
-      __syncthreads();
-      temp[thread_id] += val;
-      __syncthreads();
-    }
 
-    if (thread_id < lenparents) {
       int64_t parent = parents[thread_id];
       if (idx == blockDim.x - 1 || thread_id == lenparents - 1 || parents[thread_id] != parents[thread_id + 1]) {
         atomicAdd(&toptr[parent], temp[thread_id]);

src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_max.cu

@@ -55,18 +55,18 @@ awkward_reduce_max_b(
     }
     __syncthreads();
 
-    for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
-      T val = identity;
+    if (thread_id < lenparents) {
+      for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
+        T val = identity;
 
-      if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
-        val = temp[idx - stride];
+        if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
+          val = temp[idx - stride];
+        }
+        __syncthreads();
+        temp[thread_id] = val > temp[thread_id] ? val : temp[thread_id];
+        __syncthreads();
       }
-      __syncthreads();
-      temp[thread_id] = val > temp[thread_id] ? val : temp[thread_id];
-      __syncthreads();
-    }
 
-    if (thread_id < lenparents) {
       int64_t parent = parents[thread_id];
       if (idx == blockDim.x - 1 || thread_id == lenparents - 1 || parents[thread_id] != parents[thread_id + 1]) {
         atomicMax(&toptr[parent], temp[thread_id]);

src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_min.cu

@@ -56,17 +56,17 @@ awkward_reduce_min_b(
     }
     __syncthreads();
 
-    for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
-      T val = identity;
-      if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
-        val = temp[thread_id - stride];
+    if (thread_id < lenparents) {
+      for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
+        T val = identity;
+        if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
+          val = temp[thread_id - stride];
+        }
+        __syncthreads();
+        temp[thread_id] = val < temp[thread_id] ? val : temp[thread_id];
+        __syncthreads();
       }
-      __syncthreads();
-      temp[thread_id] = val < temp[thread_id] ? val : temp[thread_id];
-      __syncthreads();
-    }
 
-    if (thread_id < lenparents) {
       int64_t parent = parents[thread_id];
       if (idx == blockDim.x - 1 || thread_id == lenparents - 1 || parents[thread_id] != parents[thread_id + 1]) {
         atomicMin(&toptr[parent], temp[thread_id]);

src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_prod.cu

@@ -59,17 +59,17 @@ awkward_reduce_prod_b(
     }
     __syncthreads();
 
-    for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
-      T val = 1;
-      if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
-        val = temp[thread_id - stride];
+    if (thread_id < lenparents) {
+      for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
+        T val = 1;
+        if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
+          val = temp[thread_id - stride];
+        }
+        __syncthreads();
+        temp[thread_id] *= val;
+        __syncthreads();
       }
-      __syncthreads();
-      temp[thread_id] *= val;
-      __syncthreads();
-    }
 
-    if (thread_id < lenparents) {
       int64_t parent = parents[thread_id];
       if (idx == blockDim.x - 1 || thread_id == lenparents - 1 || parents[thread_id] != parents[thread_id + 1]) {
         atomicMul(&atomic_toptr[parent], temp[thread_id]);

src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_prod_bool.cu

@@ -59,17 +59,17 @@ awkward_reduce_prod_bool_b(
     }
     __syncthreads();
 
-    for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
-      T val = 1;
-      if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
-        val = temp[thread_id - stride];
+    if (thread_id < lenparents) {
+      for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
+        T val = 1;
+        if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
+          val = temp[thread_id - stride];
+        }
+        __syncthreads();
+        temp[thread_id] &= (val != 0);
+        __syncthreads();
       }
-      __syncthreads();
-      temp[thread_id] &= (val != 0);
-      __syncthreads();
-    }
 
-    if (thread_id < lenparents) {
       int64_t parent = parents[thread_id];
       if (idx == blockDim.x - 1 || thread_id == lenparents - 1 || parents[thread_id] != parents[thread_id + 1]) {
         atomicAnd(&atomic_toptr[parent], temp[thread_id]);

src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_sum.cu

@@ -54,17 +54,17 @@ awkward_reduce_sum_b(
     }
     __syncthreads();
 
-    for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
-      T val = 0;
-      if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
-        val = temp[thread_id - stride];
+    if (thread_id < lenparents) {
+      for (int64_t stride = 1; stride < blockDim.x; stride *= 2) {
+        T val = 0;
+        if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) {
+          val = temp[thread_id - stride];
+        }
+        __syncthreads();
+        temp[thread_id] += val;
+        __syncthreads();
       }
-      __syncthreads();
-      temp[thread_id] += val;
-      __syncthreads();
-    }
 
-    if (thread_id < lenparents) {
       int64_t parent = parents[thread_id];
       if (idx == blockDim.x - 1 || thread_id == lenparents - 1 || parents[thread_id] != parents[thread_id + 1]) {
         atomicAdd(&toptr[parent], temp[thread_id]);