feat: add awkward_ListArray_min_range CUDA kernel

dev/generate-kernel-signatures.py

@@ -12,6 +12,7 @@
 
 
 cuda_kernels_impl = [
+    "awkward_ListArray_min_range",
     "awkward_ListArray_validity",
     "awkward_BitMaskedArray_to_ByteMaskedArray",
     "awkward_ListArray_compact_offsets",

dev/generate-tests.py

@@ -643,6 +643,7 @@ def gencpuunittests(specdict):
 
 
 cuda_kernels_tests = [
+    "awkward_ListArray_min_range",
     "awkward_ListArray_validity",
     "awkward_BitMaskedArray_to_ByteMaskedArray",
     "awkward_ListArray_compact_offsets",

kernel-test-data.json

@@ -10756,7 +10756,7 @@
         },
         {
             "name": "awkward_ListArray_min_range",
-            "status": false,
+            "status": true,
             "tests": [
                 {
                     "error": false,

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

@@ -0,0 +1,21 @@
+// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE
+
+template <typename T, typename C, typename U>
+__global__ void
+awkward_ListArray_min_range(T* tomin,
+                            const C* fromstarts,
+                            const U* fromstops,
+                            int64_t lenstarts,
+                            uint64_t invocation_index,
+                            uint64_t* err_code) {
+  if (err_code[0] == NO_ERROR) {
+    int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x;
+    int64_t shorter = fromstops[0] - fromstarts[0];
+
+    if (thread_id >=1 && thread_id < lenstarts) {
+      int64_t rangeval = fromstops[thread_id] - fromstarts[thread_id];
+      shorter = (shorter < rangeval) ? shorter : rangeval;
+      atomicMin(tomin, shorter);
+    }
+  }
+}