From 1873096a645b10de4d5ccecf36f0902b79b70550 Mon Sep 17 00:00:00 2001 From: Manasvi Goyal Date: Thu, 7 Dec 2023 12:05:13 +0100 Subject: [PATCH 1/4] refactor: remove obsolete cuda_kernels --- dev/generate-kernel-signatures.py | 2 -- dev/generate-tests.py | 2 -- .../awkward_ByteMaskedArray_getitem_carry.cu | 25 ------------------- .../awkward_ByteMaskedArray_mask.cu | 17 ------------- 4 files changed, 46 deletions(-) delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_ByteMaskedArray_getitem_carry.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_ByteMaskedArray_mask.cu diff --git a/dev/generate-kernel-signatures.py b/dev/generate-kernel-signatures.py index b043911fc3..838269c827 100644 --- a/dev/generate-kernel-signatures.py +++ b/dev/generate-kernel-signatures.py @@ -20,7 +20,6 @@ "awkward_ListOffsetArray_flatten_offsets", "awkward_IndexedArray_overlay_mask", "awkward_IndexedArray_mask", - "awkward_ByteMaskedArray_mask", "awkward_zero_mask", "awkward_IndexedArray_fill_count", "awkward_UnionArray_fillna", @@ -40,7 +39,6 @@ "awkward_IndexedArray_simplify", "awkward_UnionArray_validity", "awkward_index_carry", - "awkward_ByteMaskedArray_getitem_carry", "awkward_IndexedArray_validity", "awkward_ByteMaskedArray_overlay_mask", "awkward_NumpyArray_reduce_mask_ByteMaskedArray_64", diff --git a/dev/generate-tests.py b/dev/generate-tests.py index 432d8c5c59..2d031444ff 100644 --- a/dev/generate-tests.py +++ b/dev/generate-tests.py @@ -534,7 +534,6 @@ def gencpukerneltests(specdict): "awkward_ListOffsetArray_flatten_offsets", "awkward_IndexedArray_overlay_mask", "awkward_IndexedArray_mask", - "awkward_ByteMaskedArray_mask", "awkward_zero_mask", "awkward_IndexedArray_fill_count", "awkward_UnionArray_fillna", @@ -554,7 +553,6 @@ def gencpukerneltests(specdict): "awkward_IndexedArray_simplify", "awkward_UnionArray_validity", "awkward_index_carry", - "awkward_ByteMaskedArray_getitem_carry", "awkward_IndexedArray_validity", "awkward_ByteMaskedArray_overlay_mask", "awkward_NumpyArray_reduce_mask_ByteMaskedArray_64", diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_ByteMaskedArray_getitem_carry.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_ByteMaskedArray_getitem_carry.cu deleted file mode 100644 index 0300265f4e..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_ByteMaskedArray_getitem_carry.cu +++ /dev/null @@ -1,25 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -enum class BYTEMASKEDARRAY_GETITEM_CARRY_ERRORS { - IND_OUT_OF_RANGE // message: "index out of range" -}; - -template -__global__ void -awkward_ByteMaskedArray_getitem_carry(T* tomask, - const C* frommask, - int64_t lenmask, - const U* fromcarry, - int64_t lencarry, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < lencarry) { - if (fromcarry[thread_id] >= lenmask) { - RAISE_ERROR(BYTEMASKEDARRAY_GETITEM_CARRY_ERRORS::IND_OUT_OF_RANGE) - } - tomask[thread_id] = frommask[fromcarry[thread_id]]; - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_ByteMaskedArray_mask.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_ByteMaskedArray_mask.cu deleted file mode 100644 index 88f9f6ac3b..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_ByteMaskedArray_mask.cu +++ /dev/null @@ -1,17 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_ByteMaskedArray_mask(T* tomask, - const C* frommask, - int64_t length, - bool validwhen, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < length) { - tomask[thread_id] = (frommask[thread_id] != 0) != validwhen; - } - } -} From 5d62b767a7cf68985230f919ea774f05c1477f8d Mon Sep 17 00:00:00 2001 From: Manasvi Goyal Date: Thu, 7 Dec 2023 13:10:11 +0100 Subject: [PATCH 2/4] refactor: remove rest of the obsolete cuda_kernels --- dev/generate-kernel-signatures.py | 20 ------ dev/generate-tests.py | 20 ------ src/awkward/_connect/cuda/__init__.py | 2 - .../cuda_kernels/awkward_Index_to_Index64.cu | 16 ----- ...edArray_getitem_nextcarry_outindex_mask.cu | 70 ------------------- .../cuda_kernels/awkward_IndexedArray_mask.cu | 16 ----- ...tOffsetArray_reduce_global_startstop_64.cu | 16 ----- .../awkward_NumpyArray_contiguous_init.cu | 16 ----- .../awkward_NumpyArray_contiguous_next.cu | 21 ------ .../awkward_NumpyArray_fill_tobool.cu | 17 ----- ...ward_NumpyArray_getitem_boolean_nonzero.cu | 56 --------------- .../awkward_NumpyArray_getitem_next_array.cu | 24 ------- ..._NumpyArray_getitem_next_array_advanced.cu | 20 ------ .../awkward_NumpyArray_getitem_next_at.cu | 18 ----- .../awkward_NumpyArray_getitem_next_range.cu | 23 ------ ..._NumpyArray_getitem_next_range_advanced.cu | 26 ------- .../cuda/cuda_kernels/awkward_carry_arange.cu | 15 ---- .../cuda/cuda_kernels/awkward_combinations.cu | 18 ----- .../awkward_content_reduce_zeroparents_64.cu | 15 ---- .../cuda/cuda_kernels/awkward_index_carry.cu | 26 ------- .../awkward_index_carry_nocheck.cu | 17 ----- .../awkward_regularize_arrayslice.cu | 26 ------- .../cuda/cuda_kernels/awkward_zero_mask.cu | 15 ---- 23 files changed, 513 deletions(-) delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_Index_to_Index64.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_IndexedArray_getitem_nextcarry_outindex_mask.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_IndexedArray_mask.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_ListOffsetArray_reduce_global_startstop_64.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_contiguous_init.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_contiguous_next.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_fill_tobool.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_boolean_nonzero.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_array.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_array_advanced.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_at.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_range.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_range_advanced.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_carry_arange.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_combinations.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_content_reduce_zeroparents_64.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_index_carry.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_index_carry_nocheck.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_regularize_arrayslice.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_zero_mask.cu diff --git a/dev/generate-kernel-signatures.py b/dev/generate-kernel-signatures.py index 838269c827..fec3960a8c 100644 --- a/dev/generate-kernel-signatures.py +++ b/dev/generate-kernel-signatures.py @@ -19,35 +19,19 @@ "awkward_Identities32_to_Identities64", "awkward_ListOffsetArray_flatten_offsets", "awkward_IndexedArray_overlay_mask", - "awkward_IndexedArray_mask", - "awkward_zero_mask", "awkward_IndexedArray_fill_count", "awkward_UnionArray_fillna", "awkward_localindex", - "awkward_content_reduce_zeroparents_64", - "awkward_ListOffsetArray_reduce_global_startstop_64", "awkward_IndexedArray_reduce_next_fix_offsets_64", - "awkward_Index_to_Index64", - "awkward_carry_arange", - "awkward_index_carry_nocheck", - "awkward_NumpyArray_contiguous_init", - "awkward_NumpyArray_getitem_next_array_advanced", - "awkward_NumpyArray_getitem_next_at", "awkward_RegularArray_getitem_next_array_advanced", "awkward_ByteMaskedArray_toIndexedOptionArray", - "awkward_combinations", # ? "awkward_IndexedArray_simplify", "awkward_UnionArray_validity", - "awkward_index_carry", "awkward_IndexedArray_validity", "awkward_ByteMaskedArray_overlay_mask", "awkward_NumpyArray_reduce_mask_ByteMaskedArray_64", "awkward_RegularArray_getitem_carry", - "awkward_NumpyArray_getitem_next_array", "awkward_RegularArray_localindex", - "awkward_NumpyArray_contiguous_next", - "awkward_NumpyArray_getitem_next_range", - "awkward_NumpyArray_getitem_next_range_advanced", "awkward_RegularArray_getitem_next_range", "awkward_RegularArray_getitem_next_range_spreadadvanced", "awkward_RegularArray_getitem_next_array", @@ -56,10 +40,8 @@ "awkward_RegularArray_getitem_jagged_expand", "awkward_ListArray_getitem_jagged_expand", "awkward_ListArray_getitem_next_array", - "awkward_NumpyArray_fill_tobool", "awkward_NumpyArray_reduce_adjust_starts_64", "awkward_NumpyArray_reduce_adjust_starts_shifts_64", - "awkward_regularize_arrayslice", "awkward_RegularArray_getitem_next_at", "awkward_BitMaskedArray_to_IndexedOptionArray", "awkward_ByteMaskedArray_getitem_nextcarry", @@ -71,7 +53,6 @@ "awkward_IndexedArray_flatten_nextcarry", "awkward_IndexedArray_getitem_nextcarry", "awkward_IndexedArray_getitem_nextcarry_outindex", - "awkward_IndexedArray_getitem_nextcarry_outindex_mask", "awkward_IndexedArray_reduce_next_64", "awkward_IndexedArray_reduce_next_nonlocal_nextshifts_64", "awkward_IndexedArray_reduce_next_nonlocal_nextshifts_fromshifts_64", @@ -79,7 +60,6 @@ "awkward_ListOffsetArray_rpad_and_clip_axis1", # "awkward_ListOffsetArray_rpad_axis1", "awkward_MaskedArray_getitem_next_jagged_project", - "awkward_NumpyArray_getitem_boolean_nonzero", "awkward_UnionArray_project", "awkward_reduce_argmax", "awkward_reduce_argmax_bool_64", diff --git a/dev/generate-tests.py b/dev/generate-tests.py index 2d031444ff..1438849376 100644 --- a/dev/generate-tests.py +++ b/dev/generate-tests.py @@ -533,35 +533,19 @@ def gencpukerneltests(specdict): "awkward_Identities32_to_Identities64", "awkward_ListOffsetArray_flatten_offsets", "awkward_IndexedArray_overlay_mask", - "awkward_IndexedArray_mask", - "awkward_zero_mask", "awkward_IndexedArray_fill_count", "awkward_UnionArray_fillna", "awkward_localindex", - "awkward_content_reduce_zeroparents_64", - "awkward_ListOffsetArray_reduce_global_startstop_64", "awkward_IndexedArray_reduce_next_fix_offsets_64", - "awkward_Index_to_Index64", - "awkward_carry_arange", - "awkward_index_carry_nocheck", - "awkward_NumpyArray_contiguous_init", - "awkward_NumpyArray_getitem_next_array_advanced", - "awkward_NumpyArray_getitem_next_at", "awkward_RegularArray_getitem_next_array_advanced", "awkward_ByteMaskedArray_toIndexedOptionArray", - "awkward_combinations", # ? "awkward_IndexedArray_simplify", "awkward_UnionArray_validity", - "awkward_index_carry", "awkward_IndexedArray_validity", "awkward_ByteMaskedArray_overlay_mask", "awkward_NumpyArray_reduce_mask_ByteMaskedArray_64", "awkward_RegularArray_getitem_carry", - "awkward_NumpyArray_getitem_next_array", "awkward_RegularArray_localindex", - "awkward_NumpyArray_contiguous_next", - "awkward_NumpyArray_getitem_next_range", - "awkward_NumpyArray_getitem_next_range_advanced", "awkward_RegularArray_getitem_next_range", "awkward_RegularArray_getitem_next_range_spreadadvanced", "awkward_RegularArray_getitem_next_array", @@ -570,10 +554,8 @@ def gencpukerneltests(specdict): "awkward_RegularArray_getitem_jagged_expand", "awkward_ListArray_getitem_jagged_expand", "awkward_ListArray_getitem_next_array", - "awkward_NumpyArray_fill_tobool", "awkward_NumpyArray_reduce_adjust_starts_64", "awkward_NumpyArray_reduce_adjust_starts_shifts_64", - "awkward_regularize_arrayslice", "awkward_RegularArray_getitem_next_at", "awkward_BitMaskedArray_to_IndexedOptionArray", "awkward_ByteMaskedArray_getitem_nextcarry", @@ -585,7 +567,6 @@ def gencpukerneltests(specdict): "awkward_IndexedArray_flatten_nextcarry", "awkward_IndexedArray_getitem_nextcarry", "awkward_IndexedArray_getitem_nextcarry_outindex", - "awkward_IndexedArray_getitem_nextcarry_outindex_mask", "awkward_IndexedArray_reduce_next_64", "awkward_IndexedArray_reduce_next_nonlocal_nextshifts_64", "awkward_IndexedArray_reduce_next_nonlocal_nextshifts_fromshifts_64", @@ -593,7 +574,6 @@ def gencpukerneltests(specdict): "awkward_ListOffsetArray_rpad_and_clip_axis1", # "awkward_ListOffsetArray_rpad_axis1", "awkward_MaskedArray_getitem_next_jagged_project", - "awkward_NumpyArray_getitem_boolean_nonzero", "awkward_UnionArray_project", "awkward_reduce_argmax", "awkward_reduce_argmax_bool_64", diff --git a/src/awkward/_connect/cuda/__init__.py b/src/awkward/_connect/cuda/__init__.py index 4806be58e5..142eb7716b 100644 --- a/src/awkward/_connect/cuda/__init__.py +++ b/src/awkward/_connect/cuda/__init__.py @@ -78,14 +78,12 @@ def fetch_template_specializations(kernel_dict): "awkward_IndexedArray_flatten_nextcarry", "awkward_IndexedArray_getitem_nextcarry", "awkward_IndexedArray_getitem_nextcarry_outindex", - "awkward_IndexedArray_getitem_nextcarry_outindex_mask", "awkward_ListArray_compact_offsets", "awkward_IndexedArray_reduce_next_64", "awkward_IndexedArray_reduce_next_nonlocal_nextshifts_64", "awkward_IndexedArray_reduce_next_nonlocal_nextshifts_fromshifts_64", "awkward_IndexedOptionArray_rpad_and_clip_mask_axis1", "awkward_MaskedArray_getitem_next_jagged_project", - "awkward_NumpyArray_getitem_boolean_nonzero", "awkward_UnionArray_project", "awkward_reduce_count_64", "awkward_reduce_sum", diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_Index_to_Index64.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_Index_to_Index64.cu deleted file mode 100644 index 7bb1382e72..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_Index_to_Index64.cu +++ /dev/null @@ -1,16 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_Index_to_Index64(T* toptr, - const C* fromptr, - int64_t length, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < length) { - toptr[thread_id] = (int64_t)(fromptr[thread_id]); - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_IndexedArray_getitem_nextcarry_outindex_mask.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_IndexedArray_getitem_nextcarry_outindex_mask.cu deleted file mode 100644 index 4b7bb5a562..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_IndexedArray_getitem_nextcarry_outindex_mask.cu +++ /dev/null @@ -1,70 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -enum class INDEXEDARRAY_GETITEM_NEXTCARRY_OUTINDEX_MASK_ERRORS { - IND_OUT_OF_RANGE, // message: "index out of range" -}; - -// BEGIN PYTHON -// def f(grid, block, args): -// (tocarry, toindex, fromindex, lenindex, lencontent, invocation_index, err_code) = args -// scan_in_array = cupy.empty(lenindex, dtype=cupy.int64) -// cuda_kernel_templates.get_function(fetch_specialization(["awkward_IndexedArray_getitem_nextcarry_outindex_mask_a", tocarry.dtype, toindex.dtype, fromindex.dtype]))(grid, block, (tocarry, toindex, fromindex, lenindex, lencontent, scan_in_array, invocation_index, err_code)) -// scan_in_array = inclusive_scan(grid, block, (scan_in_array, invocation_index, err_code)) -// cuda_kernel_templates.get_function(fetch_specialization(["awkward_IndexedArray_getitem_nextcarry_outindex_mask_b", tocarry.dtype, toindex.dtype, fromindex.dtype]))(grid, block, (tocarry, toindex, fromindex, lenindex, lencontent, scan_in_array, invocation_index, err_code)) -// out["awkward_IndexedArray_getitem_nextcarry_outindex_mask_a", {dtype_specializations}] = None -// out["awkward_IndexedArray_getitem_nextcarry_outindex_mask_b", {dtype_specializations}] = None -// END PYTHON - -template -__global__ void -awkward_IndexedArray_getitem_nextcarry_outindex_mask_a(T* tocarry, - C* toindex, - const U* fromindex, - int64_t lenindex, - int64_t lencontent, - int64_t* scan_in_array, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < lenindex) { - C j = fromindex[thread_id]; - if (j >= lencontent) { - RAISE_ERROR( - INDEXEDARRAY_GETITEM_NEXTCARRY_OUTINDEX_MASK_ERRORS::IND_OUT_OF_RANGE) - } else if (j < 0) { - scan_in_array[thread_id] = 0; - } else { - scan_in_array[thread_id] = 1; - } - } - } -} - -template -__global__ void -awkward_IndexedArray_getitem_nextcarry_outindex_mask_b(T* tocarry, - C* toindex, - const U* fromindex, - int64_t lenindex, - int64_t lencontent, - int64_t* scan_in_array, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - - if (thread_id < lenindex) { - C j = fromindex[thread_id]; - if (j >= lencontent) { - RAISE_ERROR( - INDEXEDARRAY_GETITEM_NEXTCARRY_OUTINDEX_MASK_ERRORS::IND_OUT_OF_RANGE) - } else if (j < 0) { - toindex[thread_id] = -1; - } else { - tocarry[scan_in_array[thread_id] - 1] = j; - toindex[thread_id] = (C)(scan_in_array[thread_id] - 1); - } - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_IndexedArray_mask.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_IndexedArray_mask.cu deleted file mode 100644 index 68d5b1578f..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_IndexedArray_mask.cu +++ /dev/null @@ -1,16 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_IndexedArray_mask(T* tomask, - const C* fromindex, - int64_t length, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < length) { - tomask[thread_id] = fromindex[thread_id] < 0; - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_ListOffsetArray_reduce_global_startstop_64.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_ListOffsetArray_reduce_global_startstop_64.cu deleted file mode 100644 index 610f9222fd..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_ListOffsetArray_reduce_global_startstop_64.cu +++ /dev/null @@ -1,16 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_ListOffsetArray_reduce_global_startstop_64(T* globalstart, - C* globalstop, - const U* offsets, - int64_t length, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - globalstart[0] = offsets[0]; - globalstop[0] = offsets[length]; - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_contiguous_init.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_contiguous_init.cu deleted file mode 100644 index bb27505abf..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_contiguous_init.cu +++ /dev/null @@ -1,16 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_NumpyArray_contiguous_init(T* toptr, - int64_t skip, - int64_t stride, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < skip) { - toptr[thread_id] = (thread_id * stride); - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_contiguous_next.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_contiguous_next.cu deleted file mode 100644 index fc9ee70ca4..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_contiguous_next.cu +++ /dev/null @@ -1,21 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_NumpyArray_contiguous_next(T* topos, - const C* frompos, - int64_t length, - int64_t skip, - int64_t stride, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = (blockIdx.x * blockDim.x + threadIdx.x) / skip; - int64_t thready_id = (blockIdx.x * blockDim.x + threadIdx.x) % skip; - - if (thread_id < length) { - topos[(thread_id * skip) + thready_id] = - frompos[thread_id] + (thready_id * stride); - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_fill_tobool.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_fill_tobool.cu deleted file mode 100644 index 94b63bc24e..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_fill_tobool.cu +++ /dev/null @@ -1,17 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_NumpyArray_fill_tobool(T* toptr, - int64_t tooffset, - const C* fromptr, - int64_t length, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < length) { - toptr[tooffset + thread_id] = fromptr[thread_id] > 0 ? true : false; - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_boolean_nonzero.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_boolean_nonzero.cu deleted file mode 100644 index 14c70768bb..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_boolean_nonzero.cu +++ /dev/null @@ -1,56 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -// BEGIN PYTHON -// def f(grid, block, args): -// (toptr, fromptr, length, stride, invocation_index, err_code) = args -// scan_in_array = cupy.empty(length, dtype=cupy.int64) -// cuda_kernel_templates.get_function(fetch_specialization(["awkward_NumpyArray_getitem_boolean_nonzero_a", toptr.dtype, fromptr.dtype]))(grid, block, (toptr, fromptr, length, stride, scan_in_array, invocation_index, err_code)) -// scan_in_array = inclusive_scan(grid, block, (scan_in_array, invocation_index, err_code)) -// cuda_kernel_templates.get_function(fetch_specialization(["awkward_NumpyArray_getitem_boolean_nonzero_b", toptr.dtype, fromptr.dtype]))(grid, block, (toptr, fromptr, length, stride, scan_in_array, invocation_index, err_code)) -// out["awkward_NumpyArray_getitem_boolean_nonzero_a", {dtype_specializations}] = None -// out["awkward_NumpyArray_getitem_boolean_nonzero_b", {dtype_specializations}] = None -// END PYTHON - -template -__global__ void -awkward_NumpyArray_getitem_boolean_nonzero_a(T* toptr, - const C* fromptr, - int64_t length, - int64_t stride, - int64_t* scan_in_array, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - - if (thread_id < length) { - if (thread_id % stride == 0) { - if (fromptr[thread_id] != 0) { - scan_in_array[thread_id] = 1; - } - } else { - scan_in_array[thread_id] = 0; - } - } - } -} - -template -__global__ void -awkward_NumpyArray_getitem_boolean_nonzero_b(T* toptr, - const C* fromptr, - int64_t length, - int64_t stride, - int64_t* scan_in_array, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - - if (thread_id < length && thread_id % stride == 0) { - if (fromptr[thread_id] != 0) { - toptr[scan_in_array[thread_id] - 1] = thread_id; - } - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_array.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_array.cu deleted file mode 100644 index f07e836d2e..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_array.cu +++ /dev/null @@ -1,24 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_NumpyArray_getitem_next_array(T* nextcarryptr, - C* nextadvancedptr, - const U* carryptr, - const V* flatheadptr, - int64_t lencarry, - int64_t lenflathead, - int64_t skip, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = (blockIdx.x * blockDim.x + threadIdx.x) / lenflathead; - int64_t thready_id = (blockIdx.x * blockDim.x + threadIdx.x) % lenflathead; - - if (thread_id < lencarry) { - nextcarryptr[(thread_id * lenflathead) + thready_id] = - (skip * carryptr[thread_id]) + flatheadptr[thready_id]; - nextadvancedptr[(thread_id * lenflathead) + thready_id] = thready_id; - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_array_advanced.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_array_advanced.cu deleted file mode 100644 index 9c60e0550e..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_array_advanced.cu +++ /dev/null @@ -1,20 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_NumpyArray_getitem_next_array_advanced(T* nextcarryptr, - const C* carryptr, - const U* advancedptr, - const W* flatheadptr, - int64_t lencarry, - int64_t skip, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < lencarry) { - nextcarryptr[thread_id] = - (skip * carryptr[thread_id]) + flatheadptr[advancedptr[thread_id]]; - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_at.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_at.cu deleted file mode 100644 index 6f8cae7c0b..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_at.cu +++ /dev/null @@ -1,18 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_NumpyArray_getitem_next_at(T* nextcarryptr, - const C* carryptr, - int64_t lencarry, - int64_t skip, - int64_t at, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < lencarry) { - nextcarryptr[thread_id] = (skip * carryptr[thread_id]) + at; - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_range.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_range.cu deleted file mode 100644 index be31e2a504..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_range.cu +++ /dev/null @@ -1,23 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_NumpyArray_getitem_next_range(T* nextcarryptr, - const C* carryptr, - int64_t lencarry, - int64_t lenhead, - int64_t skip, - int64_t start, - int64_t step, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = (blockIdx.x * blockDim.x + threadIdx.x) / lenhead; - int64_t thready_id = (blockIdx.x * blockDim.x + threadIdx.x) % lenhead; - - if (thread_id < lencarry) { - nextcarryptr[(thread_id * lenhead) + thready_id] = - ((skip * carryptr[thread_id]) + start) + (thready_id * step); - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_range_advanced.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_range_advanced.cu deleted file mode 100644 index 1030865667..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_getitem_next_range_advanced.cu +++ /dev/null @@ -1,26 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_NumpyArray_getitem_next_range_advanced(T* nextcarryptr, - C* nextadvancedptr, - const U* carryptr, - const V* advancedptr, - int64_t lencarry, - int64_t lenhead, - int64_t skip, - int64_t start, - int64_t step, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x / lenhead; - int64_t thready_id = blockIdx.x * blockDim.x + threadIdx.x % lenhead; - if (thread_id < lencarry) { - nextcarryptr[((thread_id * lenhead) + thready_id)] = - (((skip * carryptr[thread_id]) + start) + (thready_id * step)); - nextadvancedptr[((thread_id * lenhead) + thready_id)] = - advancedptr[thread_id]; - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_carry_arange.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_carry_arange.cu deleted file mode 100644 index b0ad81c0d4..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_carry_arange.cu +++ /dev/null @@ -1,15 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_carry_arange(T* toptr, - int64_t length, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < length) { - toptr[thread_id] = thread_id; - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_combinations.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_combinations.cu deleted file mode 100644 index da9faee2ae..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_combinations.cu +++ /dev/null @@ -1,18 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -enum class COMBINATIONS_ERRORS { - FIXME // message: "FIXME: awkward_combinations" -}; - -template -__global__ void -awkward_combinations(T* toindex, - int64_t n, - bool replacement, - int64_t singlelen, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - RAISE_ERROR(COMBINATIONS_ERRORS::FIXME) - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_content_reduce_zeroparents_64.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_content_reduce_zeroparents_64.cu deleted file mode 100644 index 4facbfea2d..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_content_reduce_zeroparents_64.cu +++ /dev/null @@ -1,15 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_content_reduce_zeroparents_64(T* toparents, - int64_t length, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < length) { - toparents[thread_id] = 0; - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_index_carry.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_index_carry.cu deleted file mode 100644 index 747e92d9ea..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_index_carry.cu +++ /dev/null @@ -1,26 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -enum class INDEX_CARRY_ERRORS { - IND_OUT_OF_RANGE // message: "index out of range" -}; - -template -__global__ void -awkward_index_carry(T* toindex, - const C* fromindex, - const U* carry, - int64_t lenfromindex, - int64_t length, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < length) { - U j = carry[thread_id]; - if (j > lenfromindex) { - RAISE_ERROR(INDEX_CARRY_ERRORS::IND_OUT_OF_RANGE) - } - toindex[thread_id] = fromindex[(int64_t)j]; - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_index_carry_nocheck.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_index_carry_nocheck.cu deleted file mode 100644 index 5421394d0c..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_index_carry_nocheck.cu +++ /dev/null @@ -1,17 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_index_carry_nocheck(T* toindex, - const C* fromindex, - const U* carry, - int64_t length, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < length) { - toindex[thread_id] = fromindex[(int64_t)carry[thread_id]]; - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_regularize_arrayslice.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_regularize_arrayslice.cu deleted file mode 100644 index 551b21300b..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_regularize_arrayslice.cu +++ /dev/null @@ -1,26 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -enum class REGULARIZE_ARRYSLICE { - IND_OUT_OF_RANGE, // message: "index out of range" -}; - -template -__global__ void -awkward_regularize_arrayslice(T* flatheadptr, - int64_t lenflathead, - int64_t length, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < lenflathead) { - if (flatheadptr[thread_id] < -length || - flatheadptr[thread_id] >= length) { - RAISE_ERROR(REGULARIZE_ARRYSLICE::IND_OUT_OF_RANGE) - } - if (flatheadptr[thread_id] < 0) { - flatheadptr[thread_id] += length; - } - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_zero_mask.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_zero_mask.cu deleted file mode 100644 index 566be0967e..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_zero_mask.cu +++ /dev/null @@ -1,15 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -template -__global__ void -awkward_zero_mask(T* tomask, - int64_t length, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < length) { - tomask[thread_id] = 0; - } - } -} From 659ef9e30ea7720f17ad37ad4c0670be7704c8d5 Mon Sep 17 00:00:00 2001 From: Manasvi Goyal Date: Thu, 7 Dec 2023 13:13:00 +0100 Subject: [PATCH 3/4] fix: kernel name in awkward_ListArray_getitem_next_at.cpp --- .../src/cpu-kernels/awkward_ListArray_getitem_next_at.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/awkward-cpp/src/cpu-kernels/awkward_ListArray_getitem_next_at.cpp b/awkward-cpp/src/cpu-kernels/awkward_ListArray_getitem_next_at.cpp index 47dfe12d4e..bc57e5de8d 100644 --- a/awkward-cpp/src/cpu-kernels/awkward_ListArray_getitem_next_at.cpp +++ b/awkward-cpp/src/cpu-kernels/awkward_ListArray_getitem_next_at.cpp @@ -1,6 +1,6 @@ // BSD 3-Clause License; see https://github.com/scikit-hep/awkward/blob/main/LICENSE -#define FILENAME(line) FILENAME_FOR_EXCEPTIONS_C("src/cpu-kernels/awkward_NumpyArray_getitem_next_at.cpp", line) +#define FILENAME(line) FILENAME_FOR_EXCEPTIONS_C("src/cpu-kernels/awkward_ListArray_getitem_next_at.cpp", line) #include "awkward/kernels.h" From 46f7f5454ec8775309ba9ba34062f72b5382b8b0 Mon Sep 17 00:00:00 2001 From: Manasvi Goyal Date: Fri, 8 Dec 2023 09:23:32 +0100 Subject: [PATCH 4/4] remove awkward_reduce_argmin_bool_64.cu and awkward_reduce_argmax_bool_64.cu --- dev/generate-kernel-signatures.py | 2 - dev/generate-tests.py | 2 - src/awkward/_connect/cuda/__init__.py | 2 - .../awkward_reduce_argmax_bool_64.cu | 49 ------------------- .../awkward_reduce_argmin_bool_64.cu | 49 ------------------- 5 files changed, 104 deletions(-) delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_argmax_bool_64.cu delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_argmin_bool_64.cu diff --git a/dev/generate-kernel-signatures.py b/dev/generate-kernel-signatures.py index fec3960a8c..10d7b7a6ed 100644 --- a/dev/generate-kernel-signatures.py +++ b/dev/generate-kernel-signatures.py @@ -62,9 +62,7 @@ "awkward_MaskedArray_getitem_next_jagged_project", "awkward_UnionArray_project", "awkward_reduce_argmax", - "awkward_reduce_argmax_bool_64", "awkward_reduce_argmin", - "awkward_reduce_argmin_bool_64", "awkward_reduce_count_64", "awkward_reduce_max", "awkward_reduce_min", diff --git a/dev/generate-tests.py b/dev/generate-tests.py index 1438849376..4ae9a8cf9c 100644 --- a/dev/generate-tests.py +++ b/dev/generate-tests.py @@ -576,9 +576,7 @@ def gencpukerneltests(specdict): "awkward_MaskedArray_getitem_next_jagged_project", "awkward_UnionArray_project", "awkward_reduce_argmax", - "awkward_reduce_argmax_bool_64", "awkward_reduce_argmin", - "awkward_reduce_argmin_bool_64", "awkward_reduce_count_64", "awkward_reduce_max", "awkward_reduce_min", diff --git a/src/awkward/_connect/cuda/__init__.py b/src/awkward/_connect/cuda/__init__.py index 142eb7716b..5bc7a8ba95 100644 --- a/src/awkward/_connect/cuda/__init__.py +++ b/src/awkward/_connect/cuda/__init__.py @@ -92,9 +92,7 @@ def fetch_template_specializations(kernel_dict): "awkward_reduce_sum_bool", "awkward_reduce_prod_bool", "awkward_reduce_argmax", - "awkward_reduce_argmax_bool_64", "awkward_reduce_argmin", - "awkward_reduce_argmin_bool_64", "awkward_reduce_countnonzero", "awkward_reduce_max", "awkward_reduce_min", diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_argmax_bool_64.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_argmax_bool_64.cu deleted file mode 100644 index d9cea5a3b3..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_argmax_bool_64.cu +++ /dev/null @@ -1,49 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -// BEGIN PYTHON -// def f(grid, block, args): -// (toptr, fromptr, parents, lenparents, outlength, invocation_index,err_code) = args -// cuda_kernel_templates.get_function(fetch_specialization(["awkward_reduce_argmax_bool_64_a", toptr.dtype, fromptr.dtype, parents.dtype]))(grid, block, (toptr, fromptr, parents, lenparents,outlength, invocation_index,err_code)) -// cuda_kernel_templates.get_function(fetch_specialization(["awkward_reduce_argmax_bool_64_b", toptr.dtype, fromptr.dtype, parents.dtype]))(grid, block, (toptr, fromptr, parents, lenparents,outlength, invocation_index,err_code)) -// out["awkward_reduce_argmax_bool_64_a", {dtype_specializations}] = None -// out["awkward_reduce_argmax_bool_64_b", {dtype_specializations}] = None -// END PYTHON - -template -__global__ void -awkward_reduce_argmax_bool_64_a(T* toptr, - const C* fromptr, - const U* parents, - int64_t lenparents, - int64_t outlength, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < outlength) { - toptr[thread_id] = -1; - } - } -} - -template -__global__ void -awkward_reduce_argmax_bool_64_b(T* toptr, - const C* fromptr, - const U* parents, - int64_t lenparents, - int64_t outlength, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - - if (thread_id < lenparents) { - int64_t parent = parents[thread_id]; - if (toptr[parent] == -1 || - (fromptr[thread_id] != 0) > (fromptr[toptr[parent]] != 0)) { - toptr[parent] = thread_id; - } - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_argmin_bool_64.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_argmin_bool_64.cu deleted file mode 100644 index c3d1be7fb2..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_argmin_bool_64.cu +++ /dev/null @@ -1,49 +0,0 @@ -// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE - -// BEGIN PYTHON -// def f(grid, block, args): -// (toptr, fromptr, parents, lenparents, outlength, invocation_index,err_code) = args -// cuda_kernel_templates.get_function(fetch_specialization(["awkward_reduce_argmin_bool_64_a", toptr.dtype, fromptr.dtype, parents.dtype]))(grid, block, (toptr, fromptr, parents, lenparents,outlength, invocation_index,err_code)) -// cuda_kernel_templates.get_function(fetch_specialization(["awkward_reduce_argmin_bool_64_b", toptr.dtype, fromptr.dtype, parents.dtype]))(grid, block, (toptr, fromptr, parents, lenparents,outlength, invocation_index,err_code)) -// out["awkward_reduce_argmin_bool_64_a", {dtype_specializations}] = None -// out["awkward_reduce_argmin_bool_64_b", {dtype_specializations}] = None -// END PYTHON - -template -__global__ void -awkward_reduce_argmin_bool_64_a(T* toptr, - const C* fromptr, - const U* parents, - int64_t lenparents, - int64_t outlength, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id < outlength) { - toptr[thread_id] = -1; - } - } -} - -template -__global__ void -awkward_reduce_argmin_bool_64_b(T* toptr, - const C* fromptr, - const U* parents, - int64_t lenparents, - int64_t outlength, - uint64_t invocation_index, - uint64_t* err_code) { - if (err_code[0] == NO_ERROR) { - int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - - if (thread_id < lenparents) { - int64_t parent = parents[thread_id]; - if (toptr[parent] == -1 || - (fromptr[thread_id] != 0) > (fromptr[toptr[parent]] != 0)) { - toptr[parent] = thread_id; - } - } - } -}