From 8a701a580df2720f9a8c45eaa458b59bfa906358 Mon Sep 17 00:00:00 2001 From: Manasvi Goyal <55101825+ManasviGoyal@users.noreply.github.com> Date: Sat, 20 Apr 2024 21:38:58 +0200 Subject: [PATCH] feat: add boolean cases for CPU kernels (#3059) * feat: add awkward_ListArray_getitem_jagged_apply.cu and awkward_ListArray_getitem_jagged_shrink * fix: use cupy.ones instead of multiple kernels * feat: add awkward_NumpyArray_rearrange_shifted CUDA kernel * feat: add new CPU kernel for awkward_unique_bool_64 * fix: test outputs * feat: add bool case for awkward_unique * fix: tolength assignment * fix: test cases * feat: add bool case for some kernels * fix: test case errors * fix: remove templates for bool * fix: segfault in ctypes double pointer * fix: assume sorted array for awkward_unique tests * fix: clean up PR * fix: remove awkward_NumpyArray_fill kernel * fix: remove awkward_unique and awkward_unique_copy kernels * fix: lint error * fix: unexpected keyword argument equal_nan error --- .../cpu-kernels/awkward_NumpyArray_fill.cpp | 826 ------------------ .../awkward_NumpyArray_rearrange_shifted.cpp | 18 +- .../awkward_NumpyArray_subrange_equal.cpp | 14 - ...awkward_NumpyArray_subrange_equal_bool.cpp | 37 + .../src/cpu-kernels/awkward_unique.cpp | 121 --- .../src/cpu-kernels/awkward_unique_copy.cpp | 145 --- .../src/cpu-kernels/awkward_unique_ranges.cpp | 14 - .../awkward_unique_ranges_bool.cpp | 26 + dev/generate-kernel-signatures.py | 2 +- dev/generate-tests.py | 26 +- kernel-specification.yml | 812 +---------------- kernel-test-data.json | 425 +++++++-- src/awkward/_connect/cuda/__init__.py | 3 +- .../cuda_kernels/awkward_NumpyArray_fill.cu | 18 - .../awkward_NumpyArray_rearrange_shifted.cu | 62 ++ ...d_RegularArray_reduce_local_nextparents.cu | 28 +- ...egularArray_reduce_nonlocal_preparenext.cu | 30 +- .../cuda_kernels/awkward_sorting_ranges.cu | 7 +- src/awkward/_nplikes/array_module.py | 25 +- src/awkward/contents/indexedarray.py | 40 +- src/awkward/contents/numpyarray.py | 136 ++- .../kernels_to_tests.json | 29 - .../tests_to_kernels.json | 29 +- tests/test_3059_boolean_kernels.py | 117 +++ 24 files changed, 793 insertions(+), 2197 deletions(-) delete mode 100644 awkward-cpp/src/cpu-kernels/awkward_NumpyArray_fill.cpp create mode 100644 awkward-cpp/src/cpu-kernels/awkward_NumpyArray_subrange_equal_bool.cpp delete mode 100644 awkward-cpp/src/cpu-kernels/awkward_unique.cpp delete mode 100644 awkward-cpp/src/cpu-kernels/awkward_unique_copy.cpp create mode 100644 awkward-cpp/src/cpu-kernels/awkward_unique_ranges_bool.cpp delete mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_fill.cu create mode 100644 src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_rearrange_shifted.cu create mode 100644 tests/test_3059_boolean_kernels.py diff --git a/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_fill.cpp b/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_fill.cpp deleted file mode 100644 index 02f4e7be2f..0000000000 --- a/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_fill.cpp +++ /dev/null @@ -1,826 +0,0 @@ -// 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_fill.cpp", line) - -#include "awkward/kernels.h" - -template -ERROR -awkward_NumpyArray_fill(TO* toptr, - int64_t tooffset, - const FROM* fromptr, - int64_t length) { - for (int64_t i = 0; i < length; i++) { - toptr[tooffset + i] = (TO)fromptr[i]; - } - return success(); -} -ERROR -awkward_NumpyArray_fill_toint8_fromint8(int8_t* toptr, - int64_t tooffset, - const int8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint16_fromint8(int16_t* toptr, - int64_t tooffset, - const int8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint32_fromint8(int32_t* toptr, - int64_t tooffset, - const int8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint64_fromint8(int64_t* toptr, - int64_t tooffset, - const int8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint8_fromint8(uint8_t* toptr, - int64_t tooffset, - const int8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint16_fromint8(uint16_t* toptr, - int64_t tooffset, - const int8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint32_fromint8(uint32_t* toptr, - int64_t tooffset, - const int8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint64_fromint8(uint64_t* toptr, - int64_t tooffset, - const int8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat32_fromint8(float* toptr, - int64_t tooffset, - const int8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat64_fromint8(double* toptr, - int64_t tooffset, - const int8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} - -ERROR -awkward_NumpyArray_fill_toint8_fromint16(int8_t* toptr, - int64_t tooffset, - const int16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint16_fromint16(int16_t* toptr, - int64_t tooffset, - const int16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint32_fromint16(int32_t* toptr, - int64_t tooffset, - const int16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint64_fromint16(int64_t* toptr, - int64_t tooffset, - const int16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint8_fromint16(uint8_t* toptr, - int64_t tooffset, - const int16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint16_fromint16(uint16_t* toptr, - int64_t tooffset, - const int16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint32_fromint16(uint32_t* toptr, - int64_t tooffset, - const int16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint64_fromint16(uint64_t* toptr, - int64_t tooffset, - const int16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat32_fromint16(float* toptr, - int64_t tooffset, - const int16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat64_fromint16(double* toptr, - int64_t tooffset, - const int16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} - -ERROR -awkward_NumpyArray_fill_toint8_fromint32(int8_t* toptr, - int64_t tooffset, - const int32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint16_fromint32(int16_t* toptr, - int64_t tooffset, - const int32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint32_fromint32(int32_t* toptr, - int64_t tooffset, - const int32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint64_fromint32(int64_t* toptr, - int64_t tooffset, - const int32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint8_fromint32(uint8_t* toptr, - int64_t tooffset, - const int32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint16_fromint32(uint16_t* toptr, - int64_t tooffset, - const int32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint32_fromint32(uint32_t* toptr, - int64_t tooffset, - const int32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint64_fromint32(uint64_t* toptr, - int64_t tooffset, - const int32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat32_fromint32(float* toptr, - int64_t tooffset, - const int32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat64_fromint32(double* toptr, - int64_t tooffset, - const int32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} - -ERROR -awkward_NumpyArray_fill_toint8_fromint64(int8_t* toptr, - int64_t tooffset, - const int64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint16_fromint64(int16_t* toptr, - int64_t tooffset, - const int64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint32_fromint64(int32_t* toptr, - int64_t tooffset, - const int64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint64_fromint64(int64_t* toptr, - int64_t tooffset, - const int64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint8_fromint64(uint8_t* toptr, - int64_t tooffset, - const int64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint16_fromint64(uint16_t* toptr, - int64_t tooffset, - const int64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint32_fromint64(uint32_t* toptr, - int64_t tooffset, - const int64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint64_fromint64(uint64_t* toptr, - int64_t tooffset, - const int64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat32_fromint64(float* toptr, - int64_t tooffset, - const int64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat64_fromint64(double* toptr, - int64_t tooffset, - const int64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} - -ERROR -awkward_NumpyArray_fill_toint8_fromuint8(int8_t* toptr, - int64_t tooffset, - const uint8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint16_fromuint8(int16_t* toptr, - int64_t tooffset, - const uint8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint32_fromuint8(int32_t* toptr, - int64_t tooffset, - const uint8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint64_fromuint8(int64_t* toptr, - int64_t tooffset, - const uint8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint8_fromuint8(uint8_t* toptr, - int64_t tooffset, - const uint8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint16_fromuint8(uint16_t* toptr, - int64_t tooffset, - const uint8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint32_fromuint8(uint32_t* toptr, - int64_t tooffset, - const uint8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint64_fromuint8(uint64_t* toptr, - int64_t tooffset, - const uint8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat32_fromuint8(float* toptr, - int64_t tooffset, - const uint8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat64_fromuint8(double* toptr, - int64_t tooffset, - const uint8_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} - -ERROR -awkward_NumpyArray_fill_toint8_fromuint16(int8_t* toptr, - int64_t tooffset, - const uint16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint16_fromuint16(int16_t* toptr, - int64_t tooffset, - const uint16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint32_fromuint16(int32_t* toptr, - int64_t tooffset, - const uint16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint64_fromuint16(int64_t* toptr, - int64_t tooffset, - const uint16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint8_fromuint16(uint8_t* toptr, - int64_t tooffset, - const uint16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint16_fromuint16(uint16_t* toptr, - int64_t tooffset, - const uint16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint32_fromuint16(uint32_t* toptr, - int64_t tooffset, - const uint16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint64_fromuint16(uint64_t* toptr, - int64_t tooffset, - const uint16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat32_fromuint16(float* toptr, - int64_t tooffset, - const uint16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat64_fromuint16(double* toptr, - int64_t tooffset, - const uint16_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} - -ERROR -awkward_NumpyArray_fill_toint8_fromuint32(int8_t* toptr, - int64_t tooffset, - const uint32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint16_fromuint32(int16_t* toptr, - int64_t tooffset, - const uint32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint32_fromuint32(int32_t* toptr, - int64_t tooffset, - const uint32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint64_fromuint32(int64_t* toptr, - int64_t tooffset, - const uint32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint8_fromuint32(uint8_t* toptr, - int64_t tooffset, - const uint32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint16_fromuint32(uint16_t* toptr, - int64_t tooffset, - const uint32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint32_fromuint32(uint32_t* toptr, - int64_t tooffset, - const uint32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint64_fromuint32(uint64_t* toptr, - int64_t tooffset, - const uint32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat32_fromuint32(float* toptr, - int64_t tooffset, - const uint32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat64_fromuint32(double* toptr, - int64_t tooffset, - const uint32_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} - -ERROR -awkward_NumpyArray_fill_toint8_fromuint64(int8_t* toptr, - int64_t tooffset, - const uint64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint16_fromuint64(int16_t* toptr, - int64_t tooffset, - const uint64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint32_fromuint64(int32_t* toptr, - int64_t tooffset, - const uint64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint64_fromuint64(int64_t* toptr, - int64_t tooffset, - const uint64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint8_fromuint64(uint8_t* toptr, - int64_t tooffset, - const uint64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint16_fromuint64(uint16_t* toptr, - int64_t tooffset, - const uint64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint32_fromuint64(uint32_t* toptr, - int64_t tooffset, - const uint64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint64_fromuint64(uint64_t* toptr, - int64_t tooffset, - const uint64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat32_fromuint64(float* toptr, - int64_t tooffset, - const uint64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat64_fromuint64(double* toptr, - int64_t tooffset, - const uint64_t* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} - -ERROR -awkward_NumpyArray_fill_toint8_fromfloat32(int8_t* toptr, - int64_t tooffset, - const float* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint16_fromfloat32(int16_t* toptr, - int64_t tooffset, - const float* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint32_fromfloat32(int32_t* toptr, - int64_t tooffset, - const float* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint64_fromfloat32(int64_t* toptr, - int64_t tooffset, - const float* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint8_fromfloat32(uint8_t* toptr, - int64_t tooffset, - const float* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint16_fromfloat32(uint16_t* toptr, - int64_t tooffset, - const float* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint32_fromfloat32(uint32_t* toptr, - int64_t tooffset, - const float* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint64_fromfloat32(uint64_t* toptr, - int64_t tooffset, - const float* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat32_fromfloat32(float* toptr, - int64_t tooffset, - const float* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat64_fromfloat32(double* toptr, - int64_t tooffset, - const float* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} - -ERROR -awkward_NumpyArray_fill_toint8_fromfloat64(int8_t* toptr, - int64_t tooffset, - const double* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint16_fromfloat64(int16_t* toptr, - int64_t tooffset, - const double* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint32_fromfloat64(int32_t* toptr, - int64_t tooffset, - const double* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_toint64_fromfloat64(int64_t* toptr, - int64_t tooffset, - const double* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint8_fromfloat64(uint8_t* toptr, - int64_t tooffset, - const double* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint16_fromfloat64(uint16_t* toptr, - int64_t tooffset, - const double* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint32_fromfloat64(uint32_t* toptr, - int64_t tooffset, - const double* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_touint64_fromfloat64(uint64_t* toptr, - int64_t tooffset, - const double* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat32_fromfloat64(float* toptr, - int64_t tooffset, - const double* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} -ERROR -awkward_NumpyArray_fill_tofloat64_fromfloat64(double* toptr, - int64_t tooffset, - const double* fromptr, - int64_t length) { - return awkward_NumpyArray_fill( - toptr, tooffset, fromptr, length); -} diff --git a/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_rearrange_shifted.cpp b/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_rearrange_shifted.cpp index 56aa61a709..c616baed9a 100644 --- a/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_rearrange_shifted.cpp +++ b/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_rearrange_shifted.cpp @@ -8,25 +8,25 @@ template ERROR awkward_NumpyArray_rearrange_shifted( TO* toptr, - const FROM* shifts, + const FROM* fromshifts, int64_t length, - const FROM* offsets, + const FROM* fromoffsets, int64_t offsetslength, - const FROM* parents, + const FROM* fromparents, int64_t /* parentslength */, // FIXME: these arguments are not needed - const FROM* starts, + const FROM* fromstarts, int64_t /* startslength */) { int64_t k = 0; for (int64_t i = 0; i < offsetslength - 1; i++) { - for (int64_t j = 0; j < offsets[i + 1] - offsets[i]; j++) { - toptr[k] = toptr[k] + offsets[i]; + for (int64_t j = 0; j < fromoffsets[i + 1] - fromoffsets[i]; j++) { + toptr[k] = toptr[k] + fromoffsets[i]; k++; } } for (int64_t i = 0; i < length; i++) { - int64_t parent = parents[i]; - int64_t start = starts[parent]; - toptr[i] = toptr[i] + shifts[toptr[i]] - start; + int64_t parent = fromparents[i]; + int64_t start = fromstarts[parent]; + toptr[i] = toptr[i] + fromshifts[toptr[i]] - start; } return success(); diff --git a/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_subrange_equal.cpp b/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_subrange_equal.cpp index c4abbba845..0cd0fd701d 100644 --- a/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_subrange_equal.cpp +++ b/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_subrange_equal.cpp @@ -36,20 +36,6 @@ ERROR awkward_NumpyArray_subrange_equal( return success(); } - -ERROR awkward_NumpyArray_subrange_equal_bool( - bool* tmpptr, - const int64_t* fromstarts, - const int64_t* fromstops, - int64_t length, - bool* toequal) { - return awkward_NumpyArray_subrange_equal( - tmpptr, - fromstarts, - fromstops, - length, - toequal); -} ERROR awkward_NumpyArray_subrange_equal_int8( int8_t* tmpptr, const int64_t* fromstarts, diff --git a/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_subrange_equal_bool.cpp b/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_subrange_equal_bool.cpp new file mode 100644 index 0000000000..7d08ab7f9d --- /dev/null +++ b/awkward-cpp/src/cpu-kernels/awkward_NumpyArray_subrange_equal_bool.cpp @@ -0,0 +1,37 @@ +// 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_subrange_equal_bool.cpp", line) + +#include "awkward/kernels.h" + +ERROR awkward_NumpyArray_subrange_equal_bool( + bool* tmpptr, + const int64_t* fromstarts, + const int64_t* fromstops, + int64_t length, + bool* toequal) { + + bool differ = true; + int64_t leftlen; + int64_t rightlen; + + for (int64_t i = 0; i < length - 1; i++) { + leftlen = fromstops[i] - fromstarts[i]; + for (int64_t ii = i + 1; ii < length - 1; ii++) { + rightlen = fromstops[ii] - fromstarts[ii]; + if (leftlen == rightlen) { + differ = false; + for (int64_t j = 0; j < leftlen; j++) { + if ((tmpptr[fromstarts[i] + j] != 0) != (tmpptr[fromstarts[ii] + j] != 0)) { + differ = true; + break; + } + } + } + } + } + + *toequal = !differ; + + return success(); +} diff --git a/awkward-cpp/src/cpu-kernels/awkward_unique.cpp b/awkward-cpp/src/cpu-kernels/awkward_unique.cpp deleted file mode 100644 index fc9730edaf..0000000000 --- a/awkward-cpp/src/cpu-kernels/awkward_unique.cpp +++ /dev/null @@ -1,121 +0,0 @@ -// 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_unique.cpp", line) - -#include "awkward/kernels.h" - -template -ERROR awkward_unique( - T* toptr, - int64_t length, - int64_t* tolength) { - - int64_t j = 0; - for (int64_t i = 1; i < length; i++) { - if (toptr[j] != toptr[i]) { - j++; - toptr[j] = toptr[i]; - } - } - *tolength = j + 1; - return success(); -} -ERROR awkward_unique_bool( - bool* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique( - toptr, - length, - tolength); -} -ERROR awkward_unique_int8( - int8_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique( - toptr, - length, - tolength); -} -ERROR awkward_unique_uint8( - uint8_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique( - toptr, - length, - tolength); -} -ERROR awkward_unique_int16( - int16_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique( - toptr, - length, - tolength); -} -ERROR awkward_unique_uint16( - uint16_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique( - toptr, - length, - tolength); -} -ERROR awkward_unique_int32( - int32_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique( - toptr, - length, - tolength); -} -ERROR awkward_unique_uint32( - uint32_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique( - toptr, - length, - tolength); -} -ERROR awkward_unique_int64( - int64_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique( - toptr, - length, - tolength); -} -ERROR awkward_unique_uint64( - uint64_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique( - toptr, - length, - tolength); -} -ERROR awkward_unique_float32( - float* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique( - toptr, - length, - tolength); -} -ERROR awkward_unique_float64( - double* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique( - toptr, - length, - tolength); -} diff --git a/awkward-cpp/src/cpu-kernels/awkward_unique_copy.cpp b/awkward-cpp/src/cpu-kernels/awkward_unique_copy.cpp deleted file mode 100644 index 0b0f0e04b6..0000000000 --- a/awkward-cpp/src/cpu-kernels/awkward_unique_copy.cpp +++ /dev/null @@ -1,145 +0,0 @@ -// 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_unique_copy.cpp", line) - -#include "awkward/kernels.h" - -template -ERROR awkward_unique_copy( - const T* fromptr, - T* toptr, - int64_t length, - int64_t* tolength) { - - int64_t j = 0; - toptr[0] = fromptr[0]; - for (int64_t i = 1; i < length; i++) { - if (toptr[j] != fromptr[i]) { - j++; - toptr[j] = fromptr[i]; - } - } - *tolength = j + 1; - return success(); -} -ERROR awkward_unique_copy_bool( - const bool* fromptr, - bool* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique_copy( - fromptr, - toptr, - length, - tolength); -} -ERROR awkward_unique_copy_int8( - const int8_t* fromptr, - int8_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique_copy( - fromptr, - toptr, - length, - tolength); -} -ERROR awkward_unique_copy_uint8( - const uint8_t* fromptr, - uint8_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique_copy( - fromptr, - toptr, - length, - tolength); -} -ERROR awkward_unique_copy_int16( - const int16_t* fromptr, - int16_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique_copy( - fromptr, - toptr, - length, - tolength); -} -ERROR awkward_unique_copy_uint16( - const uint16_t* fromptr, - uint16_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique_copy( - fromptr, - toptr, - length, - tolength); -} -ERROR awkward_unique_copy_int32( - const int32_t* fromptr, - int32_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique_copy( - fromptr, - toptr, - length, - tolength); -} -ERROR awkward_unique_copy_uint32( - const uint32_t* fromptr, - uint32_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique_copy( - fromptr, - toptr, - length, - tolength); -} -ERROR awkward_unique_copy_int64( - const int64_t* fromptr, - int64_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique_copy( - fromptr, - toptr, - length, - tolength); -} -ERROR awkward_unique_copy_uint64( - const uint64_t* fromptr, - uint64_t* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique_copy( - fromptr, - toptr, - length, - tolength); -} -ERROR awkward_unique_copy_float32( - const float* fromptr, - float* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique_copy( - fromptr, - toptr, - length, - tolength); -} -ERROR awkward_unique_copy_float64( - const double* fromptr, - double* toptr, - int64_t length, - int64_t* tolength) { - return awkward_unique_copy( - fromptr, - toptr, - length, - tolength); -} diff --git a/awkward-cpp/src/cpu-kernels/awkward_unique_ranges.cpp b/awkward-cpp/src/cpu-kernels/awkward_unique_ranges.cpp index 496f55279b..6b8fa08024 100644 --- a/awkward-cpp/src/cpu-kernels/awkward_unique_ranges.cpp +++ b/awkward-cpp/src/cpu-kernels/awkward_unique_ranges.cpp @@ -26,20 +26,6 @@ ERROR awkward_unique_ranges( return success(); } -ERROR awkward_unique_ranges_bool( - bool* toptr, - int64_t length, - const int64_t* fromoffsets, - int64_t offsetslength, - int64_t* tooffsets) { - return awkward_unique_ranges( - toptr, - length, - fromoffsets, - offsetslength, - tooffsets); -} - ERROR awkward_unique_ranges_int8( int8_t* toptr, int64_t length, diff --git a/awkward-cpp/src/cpu-kernels/awkward_unique_ranges_bool.cpp b/awkward-cpp/src/cpu-kernels/awkward_unique_ranges_bool.cpp new file mode 100644 index 0000000000..4636338041 --- /dev/null +++ b/awkward-cpp/src/cpu-kernels/awkward_unique_ranges_bool.cpp @@ -0,0 +1,26 @@ +// 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_unique_ranges_bool.cpp", line) + +#include "awkward/kernels.h" + +ERROR awkward_unique_ranges_bool( + bool* toptr, + int64_t /* length */, // FIXME: this argument is not needed + const int64_t* fromoffsets, + int64_t offsetslength, + int64_t* tooffsets) { + int64_t m = 0; + for (int64_t i = 0; i < offsetslength - 1; i++) { + tooffsets[i] = m; + toptr[m++] = toptr[fromoffsets[i]]; + for (int64_t k = fromoffsets[i]; k < fromoffsets[i + 1]; k++) { + if ((toptr[m - 1] != 0) != (toptr[k] != 0)) { + toptr[m++] = toptr[k]; + } + } + } + tooffsets[offsetslength - 1] = m; + + return success(); +} diff --git a/dev/generate-kernel-signatures.py b/dev/generate-kernel-signatures.py index cd185ebd2a..539b0ef39a 100644 --- a/dev/generate-kernel-signatures.py +++ b/dev/generate-kernel-signatures.py @@ -24,7 +24,6 @@ "awkward_IndexedArray_numnull", "awkward_IndexedArray_numnull_parents", "awkward_IndexedArray_numnull_unique_64", - "awkward_NumpyArray_fill", "awkward_ListArray_fill", "awkward_IndexedArray_fill", "awkward_IndexedArray_fill_count", @@ -69,6 +68,7 @@ "awkward_ListArray_localindex", "awkward_NumpyArray_pad_zero_to_length", "awkward_NumpyArray_reduce_adjust_starts_64", + "awkward_NumpyArray_rearrange_shifted", "awkward_NumpyArray_reduce_adjust_starts_shifts_64", "awkward_RegularArray_getitem_next_at", "awkward_BitMaskedArray_to_IndexedOptionArray", diff --git a/dev/generate-tests.py b/dev/generate-tests.py index 99a07f3d30..01c221e24d 100644 --- a/dev/generate-tests.py +++ b/dev/generate-tests.py @@ -519,7 +519,13 @@ def gencpukerneltests(specdict): elif count == 2: f.write( " " * 4 - + f"{arg} = ctypes.pointer(ctypes.cast((ctypes.c_{typename}*len({arg}[0]))(*{arg}[0]),ctypes.POINTER(ctypes.c_{typename})))\n" + + f"{arg}_ptr = (ctypes.POINTER(ctypes.c_{typename}) * len({arg}))()\n" + + " " * 4 + + f"for i in range(len({arg})):\n" + + " " * 8 + + f"{arg}_ptr[i] = (ctypes.c_{typename} * len({arg}[i]))(*{arg}[i])\n" + + " " * 4 + + f"{arg} = {arg}_ptr\n" ) f.write(" " * 4 + "funcC = getattr(lib, '" + spec.name + "')\n") args = "" @@ -634,7 +640,13 @@ def gencpuunittests(specdict): elif count == 2: f.write( " " * 4 - + f"{arg} = ctypes.pointer(ctypes.cast((ctypes.c_{typename}*len({arg}[0]))(*{arg}[0]),ctypes.POINTER(ctypes.c_{typename})))\n" + + f"{arg}_ptr = (ctypes.POINTER(ctypes.c_{typename}) * len({arg}))()\n" + + " " * 4 + + f"for i in range(len({arg})):\n" + + " " * 8 + + f"{arg}_ptr[i] = (ctypes.c_{typename} * len({arg}[i]))(*{arg}[i])\n" + + " " * 4 + + f"{arg} = {arg}_ptr\n" ) for arg, val in test["inputs"].items(): typename = remove_const( @@ -657,7 +669,13 @@ def gencpuunittests(specdict): elif count == 2: f.write( " " * 4 - + f"{arg} = ctypes.pointer(ctypes.cast((ctypes.c_{typename}*len({arg}[0]))(*{arg}[0]),ctypes.POINTER(ctypes.c_{typename})))\n" + + f"{arg}_ptr = (ctypes.POINTER(ctypes.c_{typename}) * len({arg}))()\n" + + " " * 4 + + f"for i in range(len({arg})):\n" + + " " * 8 + + f"{arg}_ptr[i] = (ctypes.c_{typename} * len({arg}[i]))(*{arg}[i])\n" + + " " * 4 + + f"{arg} = {arg}_ptr\n" ) f.write(" " * 4 + "funcC = getattr(lib, '" + spec.name + "')\n") @@ -704,7 +722,6 @@ def gencpuunittests(specdict): "awkward_IndexedArray_numnull", "awkward_IndexedArray_numnull_parents", "awkward_IndexedArray_numnull_unique_64", - "awkward_NumpyArray_fill", "awkward_ListArray_fill", "awkward_IndexedArray_fill", "awkward_IndexedArray_fill_count", @@ -749,6 +766,7 @@ def gencpuunittests(specdict): "awkward_ListArray_localindex", "awkward_NumpyArray_pad_zero_to_length", "awkward_NumpyArray_reduce_adjust_starts_64", + "awkward_NumpyArray_rearrange_shifted", "awkward_NumpyArray_reduce_adjust_starts_shifts_64", "awkward_RegularArray_getitem_next_at", "awkward_BitMaskedArray_to_IndexedOptionArray", diff --git a/kernel-specification.yml b/kernel-specification.yml index f2dcb1a7db..8d0e94b330 100644 --- a/kernel-specification.yml +++ b/kernel-specification.yml @@ -2581,615 +2581,6 @@ kernels: k = k + 1 automatic-tests: true - - name: awkward_NumpyArray_fill - specializations: - - name: awkward_NumpyArray_fill_toint8_fromint8 - args: - - {name: toptr, type: "List[int8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint8_fromint16 - args: - - {name: toptr, type: "List[int8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint8_fromint32 - args: - - {name: toptr, type: "List[int8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint8_fromint64 - args: - - {name: toptr, type: "List[int8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint8_fromuint8 - args: - - {name: toptr, type: "List[int8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint8_fromuint16 - args: - - {name: toptr, type: "List[int8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint8_fromuint32 - args: - - {name: toptr, type: "List[int8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint8_fromuint64 - args: - - {name: toptr, type: "List[int8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint8_fromfloat32 - args: - - {name: toptr, type: "List[int8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[float]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint8_fromfloat64 - args: - - {name: toptr, type: "List[int8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[double]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint16_fromint8 - args: - - {name: toptr, type: "List[int16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint16_fromint16 - args: - - {name: toptr, type: "List[int16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint16_fromint32 - args: - - {name: toptr, type: "List[int16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint16_fromint64 - args: - - {name: toptr, type: "List[int16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint16_fromuint8 - args: - - {name: toptr, type: "List[int16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint16_fromuint16 - args: - - {name: toptr, type: "List[int16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint16_fromuint32 - args: - - {name: toptr, type: "List[int16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint16_fromuint64 - args: - - {name: toptr, type: "List[int16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint16_fromfloat32 - args: - - {name: toptr, type: "List[int16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[float]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint16_fromfloat64 - args: - - {name: toptr, type: "List[int16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[double]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint32_fromint8 - args: - - {name: toptr, type: "List[int32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint32_fromint16 - args: - - {name: toptr, type: "List[int32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint32_fromint32 - args: - - {name: toptr, type: "List[int32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint32_fromint64 - args: - - {name: toptr, type: "List[int32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint32_fromuint8 - args: - - {name: toptr, type: "List[int32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint32_fromuint16 - args: - - {name: toptr, type: "List[int32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint32_fromuint32 - args: - - {name: toptr, type: "List[int32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint32_fromuint64 - args: - - {name: toptr, type: "List[int32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint32_fromfloat32 - args: - - {name: toptr, type: "List[int32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[float]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint32_fromfloat64 - args: - - {name: toptr, type: "List[int32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[double]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint64_fromint8 - args: - - {name: toptr, type: "List[int64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint64_fromint16 - args: - - {name: toptr, type: "List[int64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint64_fromint32 - args: - - {name: toptr, type: "List[int64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint64_fromint64 - args: - - {name: toptr, type: "List[int64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint64_fromuint8 - args: - - {name: toptr, type: "List[int64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint64_fromuint16 - args: - - {name: toptr, type: "List[int64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint64_fromuint32 - args: - - {name: toptr, type: "List[int64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint64_fromuint64 - args: - - {name: toptr, type: "List[int64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint64_fromfloat32 - args: - - {name: toptr, type: "List[int64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[float]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_toint64_fromfloat64 - args: - - {name: toptr, type: "List[int64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[double]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint8_fromint8 - args: - - {name: toptr, type: "List[uint8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint8_fromint16 - args: - - {name: toptr, type: "List[uint8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint8_fromint32 - args: - - {name: toptr, type: "List[uint8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint8_fromint64 - args: - - {name: toptr, type: "List[uint8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint8_fromuint8 - args: - - {name: toptr, type: "List[uint8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint8_fromuint16 - args: - - {name: toptr, type: "List[uint8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint8_fromuint32 - args: - - {name: toptr, type: "List[uint8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint8_fromuint64 - args: - - {name: toptr, type: "List[uint8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint8_fromfloat32 - args: - - {name: toptr, type: "List[uint8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[float]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint8_fromfloat64 - args: - - {name: toptr, type: "List[uint8_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[double]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint16_fromint8 - args: - - {name: toptr, type: "List[uint16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint16_fromint16 - args: - - {name: toptr, type: "List[uint16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint16_fromint32 - args: - - {name: toptr, type: "List[uint16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint16_fromint64 - args: - - {name: toptr, type: "List[uint16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint16_fromuint8 - args: - - {name: toptr, type: "List[uint16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint16_fromuint16 - args: - - {name: toptr, type: "List[uint16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint16_fromuint32 - args: - - {name: toptr, type: "List[uint16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint16_fromuint64 - args: - - {name: toptr, type: "List[uint16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint16_fromfloat32 - args: - - {name: toptr, type: "List[uint16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[float]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint16_fromfloat64 - args: - - {name: toptr, type: "List[uint16_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[double]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint32_fromint8 - args: - - {name: toptr, type: "List[uint32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint32_fromint16 - args: - - {name: toptr, type: "List[uint32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint32_fromint32 - args: - - {name: toptr, type: "List[uint32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint32_fromint64 - args: - - {name: toptr, type: "List[uint32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint32_fromuint8 - args: - - {name: toptr, type: "List[uint32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint32_fromuint16 - args: - - {name: toptr, type: "List[uint32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint32_fromuint32 - args: - - {name: toptr, type: "List[uint32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint32_fromuint64 - args: - - {name: toptr, type: "List[uint32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint32_fromfloat32 - args: - - {name: toptr, type: "List[uint32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[float]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint32_fromfloat64 - args: - - {name: toptr, type: "List[uint32_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[double]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint64_fromint8 - args: - - {name: toptr, type: "List[uint64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint64_fromint16 - args: - - {name: toptr, type: "List[uint64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint64_fromint32 - args: - - {name: toptr, type: "List[uint64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint64_fromint64 - args: - - {name: toptr, type: "List[uint64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint64_fromuint8 - args: - - {name: toptr, type: "List[uint64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint64_fromuint16 - args: - - {name: toptr, type: "List[uint64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint64_fromuint32 - args: - - {name: toptr, type: "List[uint64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint64_fromuint64 - args: - - {name: toptr, type: "List[uint64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint64_fromfloat32 - args: - - {name: toptr, type: "List[uint64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[float]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_touint64_fromfloat64 - args: - - {name: toptr, type: "List[uint64_t]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[double]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat32_fromint8 - args: - - {name: toptr, type: "List[float]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat32_fromint16 - args: - - {name: toptr, type: "List[float]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat32_fromint32 - args: - - {name: toptr, type: "List[float]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat32_fromint64 - args: - - {name: toptr, type: "List[float]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat32_fromuint8 - args: - - {name: toptr, type: "List[float]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat32_fromuint16 - args: - - {name: toptr, type: "List[float]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat32_fromuint32 - args: - - {name: toptr, type: "List[float]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat32_fromuint64 - args: - - {name: toptr, type: "List[float]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat32_fromfloat32 - args: - - {name: toptr, type: "List[float]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[float]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat32_fromfloat64 - args: - - {name: toptr, type: "List[float]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[double]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat64_fromint8 - args: - - {name: toptr, type: "List[double]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat64_fromint16 - args: - - {name: toptr, type: "List[double]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat64_fromint32 - args: - - {name: toptr, type: "List[double]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat64_fromint64 - args: - - {name: toptr, type: "List[double]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[int64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat64_fromuint8 - args: - - {name: toptr, type: "List[double]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint8_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat64_fromuint16 - args: - - {name: toptr, type: "List[double]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint16_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat64_fromuint32 - args: - - {name: toptr, type: "List[double]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint32_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat64_fromuint64 - args: - - {name: toptr, type: "List[double]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[uint64_t]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat64_fromfloat32 - args: - - {name: toptr, type: "List[double]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[float]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - - name: awkward_NumpyArray_fill_tofloat64_fromfloat64 - args: - - {name: toptr, type: "List[double]", dir: out} - - {name: tooffset, type: "int64_t", dir: in, role: IndexedArray-index-offset} - - {name: fromptr, type: "Const[List[double]]", dir: in, role: NumpyArray-ptr} - - {name: length, type: "int64_t", dir: in, role: default} - description: null - definition: | - def awkward_NumpyArray_fill(toptr, tooffset, fromptr, length): - for i in range(length): - toptr[tooffset + i] = fromptr[i] - automatic-tests: false - - name: awkward_NumpyArray_rearrange_shifted specializations: - name: awkward_NumpyArray_rearrange_shifted_toint64_fromint64 @@ -3205,10 +2596,10 @@ kernels: - {name: startslength, type: "int64_t", dir: in, role: default} description: null definition: | - def awkward_NumpyArray_rearrange_shifted(toptr, fromshifts, length, fromoffsets, offsetslength, fromparents, parentslength): + def awkward_NumpyArray_rearrange_shifted(toptr, fromshifts, length, fromoffsets, offsetslength, fromparents, parentslength, fromstarts, startslength): k = 0 for i in range(offsetslength - 1): - for j in rage(fromoffsets[i + 1] - fromoffsets[i]): + for j in range(fromoffsets[i + 1] - fromoffsets[i]): toptr[k] = toptr[k] + fromoffsets[i] k = k + 1 @@ -3381,7 +2772,7 @@ kernels: l_to_char += 1 automatic-tests: false - - name: awkward_NumpyArray_subrange_equal + - name: awkward_NumpyArray_subrange_equal_bool specializations: - name: awkward_NumpyArray_subrange_equal_bool args: @@ -3390,6 +2781,27 @@ kernels: - {name: fromstops, type: "Const[List[int64_t]]", dir: in, role: ListArray-stops} - {name: length, type: "int64_t", dir: in, role: default} - {name: toequal, type: "List[bool]", dir: out} + description: null + definition: | + def awkward_NumpyArray_subrange_equal_bool( + tmpptr, fromstarts, fromstops, length, toequal + ): + differ = True + for i in range(length - 1): + leftlen = fromstops[i] - fromstarts[i] + for ii in range(i + 1, length - 1): + rightlen = fromstops[ii] - fromstarts[ii] + if leftlen == rightlen: + differ = False + for j in range (leftlen): + if (tmpptr[fromstarts[i] + j] != 0) != (tmpptr[fromstarts[ii] + j] != 0): + differ = True + break + toequal[0] = not differ + automatic-tests: false + + - name: awkward_NumpyArray_subrange_equal + specializations: - name: awkward_NumpyArray_subrange_equal_int8 args: - {name: tmpptr, type: "List[int8_t]", dir: in, role: IndexedArray-index} @@ -5717,158 +5129,6 @@ kernels: Insert Python definition here automatic-tests: false - - name: awkward_unique - specializations: - - name: awkward_unique_bool - args: - - {name: toptr, type: "List[bool]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_int8 - args: - - {name: toptr, type: "List[int8_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_int16 - args: - - {name: toptr, type: "List[int16_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_int32 - args: - - {name: toptr, type: "List[int32_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_int64 - args: - - {name: toptr, type: "List[int64_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_uint8 - args: - - {name: toptr, type: "List[uint8_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_uint16 - args: - - {name: toptr, type: "List[uint16_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_uint32 - args: - - {name: toptr, type: "List[uint32_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_uint64 - args: - - {name: toptr, type: "List[uint64_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_float32 - args: - - {name: toptr, type: "List[float]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_float64 - args: - - {name: toptr, type: "List[double]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - description: null - definition: | - def awkward_unique( - toptr, length, tolength - ): - j = 0 - for i in range(1, length): - if toptr[j] != toptr[i]: - j += 1 - toptr[j] = toptr[i] - tolength[0] = j + 1 - automatic-tests: false - - - name: awkward_unique_copy - specializations: - - name: awkward_unique_copy_bool - args: - - {name: fromptr, type: "Const[List[bool]]", dir: in} - - {name: toptr, type: "List[bool]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_copy_int8 - args: - - {name: fromptr, type: "Const[List[int8_t]]", dir: in} - - {name: toptr, type: "List[int8_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_copy_int16 - args: - - {name: fromptr, type: "Const[List[int16_t]]", dir: in} - - {name: toptr, type: "List[int16_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_copy_int32 - args: - - {name: fromptr, type: "Const[List[int32_t]]", dir: in} - - {name: toptr, type: "List[int32_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_copy_int64 - args: - - {name: fromptr, type: "Const[List[int64_t]]", dir: in} - - {name: toptr, type: "List[int64_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_copy_uint8 - args: - - {name: fromptr, type: "Const[List[uint8_t]]", dir: in} - - {name: toptr, type: "List[uint8_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_copy_uint16 - args: - - {name: fromptr, type: "Const[List[uint16_t]]", dir: in} - - {name: toptr, type: "List[uint16_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_copy_uint32 - args: - - {name: fromptr, type: "Const[List[uint32_t]]", dir: in} - - {name: toptr, type: "List[uint32_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_copy_uint64 - args: - - {name: fromptr, type: "Const[List[uint64_t]]", dir: in} - - {name: toptr, type: "List[uint64_t]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_copy_float32 - args: - - {name: fromptr, type: "Const[List[float]]", dir: in} - - {name: toptr, type: "List[float]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - - name: awkward_unique_copy_float64 - args: - - {name: fromptr, type: "Const[List[double]]", dir: in} - - {name: toptr, type: "List[double]", dir: out} - - {name: length, type: "int64_t", dir: in, role: default} - - {name: tolength, type: "List[int64_t]", dir: out, role: default} - description: null - definition: | - def awkward_unique_copy( - fromptr, toptr, length, tolength - ): - j = 0 - toptr[0] = fromptr[0] - for i in range(1, length): - if toptr[j] != fromptr[i]: - j += 1 - toptr[j] = fromptr[i] - tolength[0] = j + 1 - automatic-tests: false - - name: awkward_unique_offsets specializations: - name: awkward_unique_offsets_int8 @@ -5911,10 +5171,11 @@ kernels: if starts[j] == starts[j + 1]: tooffsets[j + 1] = fromoffsets[i] j += 1 + j += 1 tooffsets[startslength] = fromoffsets[length - 1] automatic-tests: false - - name: awkward_unique_ranges + - name: awkward_unique_ranges_bool specializations: - name: awkward_unique_ranges_bool args: @@ -5923,6 +5184,25 @@ kernels: - {name: fromoffsets, type: "Const[List[int64_t]]", dir: in, role: default} - {name: offsetslength, type: "int64_t", dir: in, role: default} - {name: tooffsets, type: "List[int64_t]", dir: out, role: default} + description: null + definition: | + def awkward_unique_ranges_bool( + toptr, length, fromoffsets, offsetslength, tooffsets + ): + m = 0 + for i in range(offsetslength - 1): + tooffsets[i] = m + toptr[m] = toptr[fromoffsets[i]] + m += 1 + for k in range(fromoffsets[i], fromoffsets[i + 1]): + if ((toptr[m - 1] != 0) != (toptr[k] != 0)): + toptr[m] = toptr[k] + m += 1 + tooffsets[offsetslength - 1] = m + automatic-tests: false + + - name: awkward_unique_ranges + specializations: - name: awkward_unique_ranges_int8 args: - {name: toptr, type: "List[int8_t]", dir: out} @@ -5999,7 +5279,7 @@ kernels: toptr, length, fromoffsets, offsetslength, tooffsets ): m = 0 - for i in range(offsetslength): + for i in range(offsetslength - 1): tooffsets[i] = m toptr[m] = toptr[fromoffsets[i]] m += 1 diff --git a/kernel-test-data.json b/kernel-test-data.json index 1d2da28ee5..a49090d6a3 100644 --- a/kernel-test-data.json +++ b/kernel-test-data.json @@ -15911,6 +15911,33 @@ "name": "awkward_UnionArray_flatten_length", "status": false, "tests": [ + { + "error": false, + "message": "", + "inputs": { + "fromtags": [], + "fromindex": [], + "length": 0, + "offsetsraws": [[], []] + }, + "outputs": { + "total_length": [0] + } + }, + { + "error": false, + "message": "", + "inputs": { + "fromtags": [0, 1, 0, 1], + "fromindex": [0, 0, 1, 1], + "length": 4, + "offsetsraws": [[0, 2, 2, 3, 5], [2, 2, 3, 5, 6]] + }, + "outputs": { + "total_length": [3] + } + }, + { "error": false, "message": "", @@ -15926,6 +15953,57 @@ } ] }, + { + "name": "awkward_UnionArray_flatten_combine", + "status": false, + "tests": [ + { + "error": false, + "message": "", + "inputs": { + "fromtags": [], + "fromindex": [], + "length": 0, + "offsetsraws": [[], []] + }, + "outputs": { + "totags": [], + "toindex": [], + "tooffsets": [0] + } + }, + { + "error": false, + "message": "", + "inputs": { + "fromtags": [0, 1, 0, 1], + "fromindex": [0, 0, 1, 1], + "length": 4, + "offsetsraws": [[0, 2, 2, 3, 5], [2, 2, 3, 5, 6]] + }, + "outputs": { + "totags": [0, 0, 1], + "toindex": [0, 1, 2], + "tooffsets": [0, 2, 2, 2, 3] + } + }, + { + "error": false, + "message": "", + "inputs": { + "fromtags": [0, 0, 0, 0], + "fromindex": [0, 1, 2, 3], + "length": 4, + "offsetsraws": [[0, 1, 3, 5, 7], [1, 3, 5, 7, 9]] + }, + "outputs": { + "totags": [0, 0, 0, 0, 0, 0, 0], + "toindex": [0, 1, 2, 3, 4, 5, 6], + "tooffsets": [0, 1, 3, 5, 7] + } + } + ] + }, { "name": "awkward_UnionArray_validity", "status": true, @@ -20481,28 +20559,30 @@ ] }, { - "name": "awkward_NumpyArray_fill", + "name": "awkward_NumpyArray_pad_zero_to_length", "status": true, "tests": [ { "error": false, "message": "", "inputs": { - "tooffset": 0, + "fromoffsets": [0, 2, 3], "fromptr": [0, 1, 3], - "length": 3 + "offsetslength": 3, + "target": 2 }, "outputs": { - "toptr": [0, 1, 3] + "toptr": [0, 1, 3, 0] } }, { "error": false, "message": "", "inputs": { - "tooffset": 0, + "fromoffsets": [], "fromptr": [], - "length": 0 + "offsetslength": 0, + "target": 0 }, "outputs": { "toptr": [] @@ -20512,65 +20592,59 @@ "error": false, "message": "", "inputs": { - "tooffset": 0, + "fromoffsets": [0, 2, 2, 4], "fromptr": [1, 3, 3, 5], - "length": 4 + "offsetslength": 4, + "target": 4 }, "outputs": { - "toptr": [1, 3, 3, 5] + "toptr": [1, 3, 0, 0, 0, 0, 0, 0, 3, 5, 0, 0] } }, { "error": false, "message": "", "inputs": { - "tooffset": 0, + "fromoffsets": [0, 0], "fromptr": [3, 5], - "length": 2 + "offsetslength": 2, + "target": 4 }, "outputs": { - "toptr": [3, 5] + "toptr": [0, 0, 0, 0] } }, { "error": false, "message": "", "inputs": { - "tooffset": 0, + "fromoffsets": [0, 1, 2, 3, 5], "fromptr": [0, 3, 3, 5, 6], - "length": 5 + "offsetslength": 5, + "target": 4 }, "outputs": { - "toptr": [0, 3, 3, 5, 6] + "toptr": [0, 0, 0, 0, 3, 0, 0, 0, 3, 0, 0, 0, 5, 6, 0, 0] } } ] }, { - "name": "awkward_NumpyArray_pad_zero_to_length", + "name": "awkward_NumpyArray_rearrange_shifted", "status": true, "tests": [ { "error": false, "message": "", "inputs": { - "fromoffsets": [0, 2, 3], - "fromptr": [0, 1, 3], - "offsetslength": 3, - "target": 2 - }, - "outputs": { - "toptr": [0, 1, 3, 0] - } - }, - { - "error": false, - "message": "", - "inputs": { + "fromshifts": [], + "length": 0, "fromoffsets": [], - "fromptr": [], "offsetslength": 0, - "target": 0 + "fromparents": [], + "parentslength": 0, + "fromstarts": [], + "startslength": 0 }, "outputs": { "toptr": [] @@ -20580,39 +20654,36 @@ "error": false, "message": "", "inputs": { - "fromoffsets": [0, 2, 2, 4], - "fromptr": [1, 3, 3, 5], - "offsetslength": 4, - "target": 4 - }, - "outputs": { - "toptr": [1, 3, 0, 0, 0, 0, 0, 0, 3, 5, 0, 0] - } - }, - { - "error": false, - "message": "", - "inputs": { - "fromoffsets": [0, 0], - "fromptr": [3, 5], - "offsetslength": 2, - "target": 4 + "toptr": [0, 1, 2, 3, 4, 5, 6, 7, 8], + "fromshifts": [0, 1, 2, 3, 4, 5, 6], + "length": 4, + "fromoffsets": [0, 1, 3, 3, 5, 7, 9], + "offsetslength": 7, + "fromparents": [0, 1, 3, 6], + "parentslength": 4, + "fromstarts": [0, 1, 2, 3, 4, 5, 6], + "startslength": 7 }, "outputs": { - "toptr": [0, 0, 0, 0] + "toptr": [0, 3, 3, 6, 7, 10, 11, 14, 15] } }, { "error": false, "message": "", "inputs": { - "fromoffsets": [0, 1, 2, 3, 5], - "fromptr": [0, 3, 3, 5, 6], - "offsetslength": 5, - "target": 4 + "toptr": [0, 0, 0, 0, 0, 0, 0, 0], + "fromshifts": [0, 1, 2, 3, 4, 5, 6], + "length": 4, + "fromoffsets": [0, 2, 5, 8], + "offsetslength": 4, + "fromparents": [0, 1, 3, 6], + "parentslength": 4, + "fromstarts": [0, 1, 2, 3, 4, 5, 6], + "startslength": 7 }, "outputs": { - "toptr": [0, 0, 0, 0, 3, 0, 0, 0, 3, 0, 0, 0, 5, 6, 0, 0] + "toptr": [0, -1, 1, -2, 2, 5, 5, 5] } } ] @@ -26041,6 +26112,252 @@ } } ] + }, + { + "name": "awkward_NumpyArray_subrange_equal", + "status": true, + "tests": [ + { + "error": false, + "message": "", + "inputs": { + "tmpptr": [], + "fromstarts": [], + "fromstops": [], + "length": 0 + }, + "outputs": { + "toequal": [0] + } + }, + { + "error": false, + "message": "", + "inputs": { + "tmpptr": [0, 1, 2, 3, 4, 5, 6, 7], + "fromstarts": [0, 1, 3, 5], + "fromstops": [1, 3, 5, 7], + "length": 4 + }, + "outputs": { + "toequal": [0] + } + } + ] + }, + { + "name": "awkward_NumpyArray_subrange_equal_bool", + "status": true, + "tests": [ + { + "error": false, + "message": "", + "inputs": { + "tmpptr": [], + "fromstarts": [], + "fromstops": [], + "length": 0 + }, + "outputs": { + "toequal": [0] + } + }, + { + "error": false, + "message": "", + "inputs": { + "tmpptr": [0, 1, 2, 3, 4, 5, 6, 7], + "fromstarts": [0, 1, 3, 5], + "fromstops": [1, 3, 5, 7], + "length": 4 + }, + "outputs": { + "toequal": [1] + } + } + ] + }, + { + "name": "awkward_unique_offsets", + "status": true, + "tests": [ + { + "error": false, + "message": "", + "inputs": { + "fromoffsets": [0], + "starts": [], + "startslength": 0, + "length": 1 + }, + "outputs": { + "tooffsets": [0] + } + }, + { + "error": false, + "message": "", + "inputs": { + "fromoffsets": [0, 2, 2, 3, 4, 5], + "starts": [0, 2, 4, 6], + "startslength": 4, + "length": 6 + }, + "outputs": { + "tooffsets": [0, 2, 2, 3, 5, 5] + } + }, + { + "error": false, + "message": "", + "inputs": { + "fromoffsets": [0, 1, 2, 2, 3], + "starts": [0, 1, 3, 4], + "startslength": 4, + "length": 5 + }, + "outputs": { + "tooffsets": [0, 1, 2, 2, 3] + } + }, + { + "error": false, + "message": "", + "inputs": { + "fromoffsets": [0, 1, 2, 2, 3], + "starts": [0, 1, 2, 3], + "startslength": 4, + "length": 5 + }, + "outputs": { + "tooffsets": [0, 1, 2, 2, 3] + } + } + ] + }, + { + "name": "awkward_unique_ranges", + "status": true, + "tests": [ + { + "error": false, + "message": "", + "inputs": { + "toptr": [], + "fromoffsets": [0], + "offsetslength": 1, + "length": 0 + }, + "outputs": { + "toptr": [], + "tooffsets": [0] + } + }, + { + "error": false, + "message": "", + "inputs": { + "toptr": [1, 2], + "fromoffsets": [0, 2], + "offsetslength": 2, + "length": 2 + }, + "outputs": { + "toptr": [1, 2], + "tooffsets": [0, 2] + } + }, + { + "error": false, + "message": "", + "inputs": { + "toptr": [0, 1, 2, 3, 4, 5], + "fromoffsets": [0, 3, 5, 6], + "offsetslength": 4, + "length": 6 + }, + "outputs": { + "toptr": [0, 1, 2, 3, 4, 5], + "tooffsets": [0, 3, 5, 6] + } + }, + { + "error": false, + "message": "", + "inputs": { + "toptr": [3, 2, 1, 0], + "fromoffsets": [0, 0, 3, 3], + "offsetslength": 4, + "length": 4 + }, + "outputs": { + "toptr": [3, 3, 1, 0], + "tooffsets": [0, 1, 3, 4] + } + } + ] + }, + { + "name": "awkward_unique_ranges_bool", + "status": true, + "tests": [ + { + "error": false, + "message": "", + "inputs": { + "toptr": [], + "fromoffsets": [0], + "offsetslength": 1, + "length": 1 + }, + "outputs": { + "toptr": [], + "tooffsets": [0] + } + }, + { + "error": false, + "message": "", + "inputs": { + "toptr": [1, 1], + "fromoffsets": [0, 2], + "offsetslength": 2, + "length": 1 + }, + "outputs": { + "toptr": [1, 1], + "tooffsets": [0, 1] + } + }, + { + "error": false, + "message": "", + "inputs": { + "toptr": [0, 1, 1, 1, 1, 1], + "fromoffsets": [0, 3, 5, 6], + "offsetslength": 4, + "length": 6 + }, + "outputs": { + "toptr": [0, 1, 1, 1, 1, 1], + "tooffsets": [0, 2, 3, 4] + } + }, + { + "error": false, + "message": "", + "inputs": { + "toptr": [1, 1, 1, 0], + "fromoffsets": [0, 0, 3, 3], + "offsetslength": 4, + "length": 4 + }, + "outputs": { + "toptr": [1, 1, 0, 0], + "tooffsets": [0, 1, 2, 3] + } + } + ] } ] } diff --git a/src/awkward/_connect/cuda/__init__.py b/src/awkward/_connect/cuda/__init__.py index 680ff80568..499bafa485 100644 --- a/src/awkward/_connect/cuda/__init__.py +++ b/src/awkward/_connect/cuda/__init__.py @@ -107,8 +107,7 @@ def fetch_template_specializations(kernel_dict): "awkward_ListOffsetArray_reduce_nonlocal_maxcount_offsetscopy_64", "awkward_ListOffsetArray_rpad_length_axis1", "awkward_MaskedArray_getitem_next_jagged_project", - "awkward_RegularArray_reduce_local_nextparents", - "awkward_RegularArray_reduce_nonlocal_preparenext", + "awkward_NumpyArray_rearrange_shifted", "awkward_UnionArray_project", "awkward_reduce_count_64", "awkward_reduce_sum", diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_fill.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_fill.cu deleted file mode 100644 index e0ef17c78f..0000000000 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_fill.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_fill( - 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] = (T)fromptr[thread_id]; - } - } -} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_rearrange_shifted.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_rearrange_shifted.cu new file mode 100644 index 0000000000..a7df7a5e55 --- /dev/null +++ b/src/awkward/_connect/cuda/cuda_kernels/awkward_NumpyArray_rearrange_shifted.cu @@ -0,0 +1,62 @@ +// BSD 3-Clause License; see https://github.com/scikit-hep/awkward-1.0/blob/main/LICENSE + +// BEGIN PYTHON +// def f(grid, block, args): +// (toptr, fromshifts, length, fromoffsets, offsetslength, fromparents, parentslength, fromstarts, startslength, invocation_index, err_code) = args +// cuda_kernel_templates.get_function(fetch_specialization(["awkward_NumpyArray_rearrange_shifted_a", toptr.dtype, fromshifts.dtype, fromoffsets.dtype, fromparents.dtype, fromstarts.dtype]))(grid, block, (toptr, fromshifts, length, fromoffsets, offsetslength, fromparents, parentslength, fromstarts, startslength, invocation_index, err_code)) +// cuda_kernel_templates.get_function(fetch_specialization(["awkward_NumpyArray_rearrange_shifted_b", toptr.dtype, fromshifts.dtype, fromoffsets.dtype, fromparents.dtype, fromstarts.dtype]))(grid, block, (toptr, fromshifts, length, fromoffsets, offsetslength, fromparents, parentslength, fromstarts, startslength, invocation_index, err_code)) +// out["awkward_NumpyArray_rearrange_shifted_a", {dtype_specializations}] = None +// out["awkward_NumpyArray_rearrange_shifted_b", {dtype_specializations}] = None +// END PYTHON + + +template +__global__ void +awkward_NumpyArray_rearrange_shifted_a( + T* toptr, + C* fromshifts, + int64_t length, + U* fromoffsets, + int64_t offsetslength, + V* fromparents, + int64_t /* parentslength */, + W* fromstarts, + int64_t /* startslength */, + 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 < offsetslength - 1) { + for (int64_t j = 0; j < fromoffsets[thread_id + 1] - fromoffsets[thread_id]; j++) { + int64_t idx = fromoffsets[thread_id] + j; + toptr[idx] = toptr[idx] + fromoffsets[thread_id]; + } + } + } +} + +template +__global__ void +awkward_NumpyArray_rearrange_shifted_b( + T* toptr, + C* fromshifts, + int64_t length, + U* fromoffsets, + int64_t offsetslength, + V* fromparents, + int64_t /* parentslength */, + W* fromstarts, + int64_t /* startslength */, + 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) { + int64_t parent = fromparents[thread_id]; + int64_t start = fromstarts[parent]; + toptr[thread_id] = toptr[thread_id] + fromshifts[toptr[thread_id]] - start; + } + } +} diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_RegularArray_reduce_local_nextparents.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_RegularArray_reduce_local_nextparents.cu index e879afcc5f..1d66dc53ef 100644 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_RegularArray_reduce_local_nextparents.cu +++ b/src/awkward/_connect/cuda/cuda_kernels/awkward_RegularArray_reduce_local_nextparents.cu @@ -3,35 +3,15 @@ // BEGIN PYTHON // def f(grid, block, args): // (nextparents, size, length, invocation_index, err_code) = args -// scan_in_array = cupy.zeros(length * size, dtype=cupy.int64) -// cuda_kernel_templates.get_function(fetch_specialization(['awkward_RegularArray_reduce_local_nextparents_a', nextparents.dtype]))(grid, block, (nextparents, size, length, scan_in_array, invocation_index, err_code)) +// scan_in_array = cupy.ones(length * size, dtype=cupy.int64) // scan_in_array = cupy.cumsum(scan_in_array) -// cuda_kernel_templates.get_function(fetch_specialization(['awkward_RegularArray_reduce_local_nextparents_b', nextparents.dtype]))(grid, block, (nextparents, size, length, scan_in_array, invocation_index, err_code)) -// out["awkward_RegularArray_reduce_local_nextparents_a", {dtype_specializations}] = None -// out["awkward_RegularArray_reduce_local_nextparents_b", {dtype_specializations}] = None +// cuda_kernel_templates.get_function(fetch_specialization(['awkward_RegularArray_reduce_local_nextparents', nextparents.dtype]))(grid, block, (nextparents, size, length, scan_in_array, invocation_index, err_code)) +// out["awkward_RegularArray_reduce_local_nextparents", {dtype_specializations}] = None // END PYTHON template __global__ void -awkward_RegularArray_reduce_local_nextparents_a( - T* nextparents, - int64_t size, - int64_t length, - 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; - int64_t len = length * size; - if (thread_id < len) { - scan_in_array[thread_id] = 1; - } - } -} - -template -__global__ void -awkward_RegularArray_reduce_local_nextparents_b( +awkward_RegularArray_reduce_local_nextparents( T* nextparents, int64_t size, int64_t length, diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_RegularArray_reduce_nonlocal_preparenext.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_RegularArray_reduce_nonlocal_preparenext.cu index 043446fcfa..20f9d50e8d 100644 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_RegularArray_reduce_nonlocal_preparenext.cu +++ b/src/awkward/_connect/cuda/cuda_kernels/awkward_RegularArray_reduce_nonlocal_preparenext.cu @@ -3,37 +3,15 @@ // BEGIN PYTHON // def f(grid, block, args): // (nextcarry, nextparents, parents, size, length, invocation_index, err_code) = args -// scan_in_array = cupy.zeros(length * size, dtype=cupy.int64) -// cuda_kernel_templates.get_function(fetch_specialization(['awkward_RegularArray_reduce_nonlocal_preparenext_a', nextcarry.dtype, nextparents.dtype, parents.dtype]))(grid, block, (nextcarry, nextparents, parents, size, length, scan_in_array, invocation_index, err_code)) +// scan_in_array = cupy.ones(length * size, dtype=cupy.int64) // scan_in_array = cupy.cumsum(scan_in_array) -// cuda_kernel_templates.get_function(fetch_specialization(['awkward_RegularArray_reduce_nonlocal_preparenext_b', nextcarry.dtype, nextparents.dtype, parents.dtype]))(grid, block, (nextcarry, nextparents, parents, size, length, scan_in_array, invocation_index, err_code)) -// out["awkward_RegularArray_reduce_nonlocal_preparenext_a", {dtype_specializations}] = None -// out["awkward_RegularArray_reduce_nonlocal_preparenext_b", {dtype_specializations}] = None +// cuda_kernel_templates.get_function(fetch_specialization(['awkward_RegularArray_reduce_nonlocal_preparenext', nextcarry.dtype, nextparents.dtype, parents.dtype]))(grid, block, (nextcarry, nextparents, parents, size, length, scan_in_array, invocation_index, err_code)) +// out["awkward_RegularArray_reduce_nonlocal_preparenext", {dtype_specializations}] = None // END PYTHON template __global__ void -awkward_RegularArray_reduce_nonlocal_preparenext_a( - T* nextcarry, - C* nextparents, - const U* parents, - int64_t size, - int64_t length, - int64_t* scan_in_array, - uint64_t invocation_index, - uint64_t* err_code) { -if (err_code[0] == NO_ERROR) { - int64_t thready_id = blockIdx.x * blockDim.x + threadIdx.x; - int64_t len = length * size; - if (thready_id < len) { - scan_in_array[thready_id] = 1; - } - } -} - -template -__global__ void -awkward_RegularArray_reduce_nonlocal_preparenext_b( +awkward_RegularArray_reduce_nonlocal_preparenext( T* nextcarry, C* nextparents, const U* parents, diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_sorting_ranges.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_sorting_ranges.cu index e8bfbf33dd..b05e0dfc4d 100644 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_sorting_ranges.cu +++ b/src/awkward/_connect/cuda/cuda_kernels/awkward_sorting_ranges.cu @@ -3,7 +3,7 @@ // BEGIN PYTHON // def f(grid, block, args): // (toindex, tolength, parents, parentslength, invocation_index, err_code) = args -// scan_in_array_k = cupy.zeros(parentslength, dtype=cupy.int64) +// scan_in_array_k = cupy.ones(parentslength, dtype=cupy.int64) // scan_in_array_j = cupy.zeros(parentslength, dtype=cupy.int64) // cuda_kernel_templates.get_function(fetch_specialization(['awkward_sorting_ranges_a', toindex.dtype, parents.dtype]))(grid, block, (toindex, tolength, parents, parentslength, scan_in_array_k, scan_in_array_j, invocation_index, err_code)) // scan_in_array_k = cupy.cumsum(scan_in_array_k) @@ -26,11 +26,10 @@ awkward_sorting_ranges_a( uint64_t* err_code) { if (err_code[0] == NO_ERROR) { int64_t thread_id = blockIdx.x * blockDim.x + threadIdx.x; - if (thread_id >= 1 && thread_id < parentslength) { + if (thread_id > 0 && thread_id < parentslength) { if (parents[thread_id - 1] != parents[thread_id]) { scan_in_array_j[thread_id] = 1; } - scan_in_array_k[thread_id] = 1; } } } @@ -51,7 +50,7 @@ awkward_sorting_ranges_b( toindex[0] = 0; if (thread_id > 0 && thread_id < parentslength) { if (parents[thread_id - 1] != parents[thread_id]) { - toindex[scan_in_array_j[thread_id]] = scan_in_array_k[thread_id]; + toindex[scan_in_array_j[thread_id]] = scan_in_array_k[thread_id - 1]; } } toindex[tolength - 1] = parentslength; diff --git a/src/awkward/_nplikes/array_module.py b/src/awkward/_nplikes/array_module.py index 916fe708ae..b474b6571c 100644 --- a/src/awkward/_nplikes/array_module.py +++ b/src/awkward/_nplikes/array_module.py @@ -2,6 +2,7 @@ from __future__ import annotations +import inspect import math from functools import lru_cache @@ -390,14 +391,26 @@ def where( def unique_values(self, x: ArrayLikeT) -> ArrayLikeT: assert not isinstance(x, PlaceholderArray) - return self._module.unique( - x, - return_counts=False, - return_index=False, - return_inverse=False, - equal_nan=False, + np_unique_accepts_equal_nan = ( + "equal_nan" in inspect.signature(self._module.unique).parameters ) + if np_unique_accepts_equal_nan: + return self._module.unique( + x, + return_counts=False, + return_index=False, + return_inverse=False, + equal_nan=False, + ) + else: + return self._module.unique( + x, + return_counts=False, + return_index=False, + return_inverse=False, + ) + def unique_all(self, x: ArrayLikeT) -> UniqueAllResult: assert not isinstance(x, PlaceholderArray) values, indices, inverse_indices, counts = self._module.unique( diff --git a/src/awkward/contents/indexedarray.py b/src/awkward/contents/indexedarray.py index 3a80fc095e..c5e6524733 100644 --- a/src/awkward/contents/indexedarray.py +++ b/src/awkward/contents/indexedarray.py @@ -750,37 +750,17 @@ def _unique_index(self, index, sorted=True): ) ) - assert ( - next.nplike is self._backend.index_nplike - and length.nplike is self._backend.index_nplike - ) - self._backend.maybe_kernel_error( - self._backend["awkward_unique", next.dtype.type, length.dtype.type]( - next.data, - self._index.length, - length.data, - ) - ) + assert ( + self._index.nplike is self._backend.index_nplike + and next.nplike is self._backend.index_nplike + and length.nplike is self._backend.index_nplike + ) - else: - assert ( - self._index.nplike is self._backend.index_nplike - and next.nplike is self._backend.index_nplike - and length.nplike is self._backend.index_nplike - ) - self._backend.maybe_kernel_error( - self._backend[ - "awkward_unique_copy", - self._index.dtype.type, - next.dtype.type, - length.dtype.type, - ]( - self._index.data, - next.data, - self._index.length, - length.data, - ) - ) + next = ak.index.Index64( + self._backend.index_nplike.unique_values(self._index), + nplike=self._backend.index_nplike, + ) + length[0] = next.data.size return next[0 : length[0]] diff --git a/src/awkward/contents/numpyarray.py b/src/awkward/contents/numpyarray.py index 35e9522c18..2428e617a8 100644 --- a/src/awkward/contents/numpyarray.py +++ b/src/awkward/contents/numpyarray.py @@ -565,72 +565,46 @@ def is_contiguous(self) -> bool: def _subranges_equal(self, starts, stops, length, sorted=True): is_equal = ak.index.Index64.zeros(1, nplike=self._backend.nplike) - tmp = self._backend.nplike.empty(length, dtype=self.dtype) - self._backend.maybe_kernel_error( - self._backend[ - "awkward_NumpyArray_fill", - self.dtype.type, - self._data.dtype.type, - ]( - tmp, - 0, - self._data, - length, - ) + assert ( + starts.nplike is self._backend.index_nplike + and stops.nplike is self._backend.index_nplike ) - - if not sorted: - tmp_beg_ptr = ak.index.Index64.empty( - ak._util.kMaxLevels, nplike=self._backend.index_nplike - ) - tmp_end_ptr = ak.index.Index64.empty( - ak._util.kMaxLevels, nplike=self._backend.index_nplike - ) - - assert ( - tmp_beg_ptr.nplike is self._backend.index_nplike - and tmp_end_ptr.nplike is self._backend.index_nplike - and starts.nplike is self._backend.index_nplike - and stops.nplike is self._backend.index_nplike - ) + if self.dtype == np.bool_: self._backend.maybe_kernel_error( self._backend[ - "awkward_quick_sort", + "awkward_NumpyArray_subrange_equal_bool", self.dtype.type, - tmp_beg_ptr.dtype.type, - tmp_end_ptr.dtype.type, starts.dtype.type, stops.dtype.type, + np.bool_, ]( - tmp, - tmp_beg_ptr.data, - tmp_end_ptr.data, + self._backend.nplike.astype( + self._data, dtype=self.dtype, copy=True + ), starts.data, stops.data, - True, starts.length, - ak._util.kMaxLevels, + is_equal.data, ) ) - assert ( - starts.nplike is self._backend.index_nplike - and stops.nplike is self._backend.index_nplike - ) - self._backend.maybe_kernel_error( - self._backend[ - "awkward_NumpyArray_subrange_equal", - self.dtype.type, - starts.dtype.type, - stops.dtype.type, - np.bool_, - ]( - tmp, - starts.data, - stops.data, - starts.length, - is_equal.data, + else: + self._backend.maybe_kernel_error( + self._backend[ + "awkward_NumpyArray_subrange_equal", + self.dtype.type, + starts.dtype.type, + stops.dtype.type, + np.bool_, + ]( + self._backend.nplike.astype( + self._data, dtype=self.dtype, copy=True + ), + starts.data, + stops.data, + starts.length, + is_equal.data, + ) ) - ) return True if is_equal[0] == 1 else False @@ -763,20 +737,10 @@ def _unique(self, negaxis, starts, parents, outlength): False, ) ) - nextlength = ak.index.Index64.empty(1, self._backend.index_nplike) assert nextlength.nplike is self._backend.index_nplike - self._backend.maybe_kernel_error( - self._backend[ - "awkward_unique", - out.dtype.type, - nextlength.dtype.type, - ]( - out, - out.shape[0], - nextlength.data, - ) - ) + out = self._backend.index_nplike.unique_values(out) + nextlength[0] = out.size return ak.contents.NumpyArray( self._backend.nplike.asarray(out[: nextlength[0]], dtype=self.dtype), @@ -858,20 +822,36 @@ def _unique(self, negaxis, starts, parents, outlength): offsets.nplike is self._backend.index_nplike and nextoffsets.nplike is self._backend.index_nplike ) - self._backend.maybe_kernel_error( - self._backend[ - "awkward_unique_ranges", - out.dtype.type, - offsets.dtype.type, - nextoffsets.dtype.type, - ]( - out, - out.shape[0], - offsets.data, - offsets.length, - nextoffsets.data, + if out.dtype == np.bool_: + self._backend.maybe_kernel_error( + self._backend[ + "awkward_unique_ranges_bool", + out.dtype.type, + offsets.dtype.type, + nextoffsets.dtype.type, + ]( + out, + out.shape[0], + offsets.data, + offsets.length, + nextoffsets.data, + ) + ) + else: + self._backend.maybe_kernel_error( + self._backend[ + "awkward_unique_ranges", + out.dtype.type, + offsets.dtype.type, + nextoffsets.dtype.type, + ]( + out, + out.shape[0], + offsets.data, + offsets.length, + nextoffsets.data, + ) ) - ) outoffsets = ak.index.Index64.empty( starts.length + 1, self._backend.index_nplike diff --git a/studies/tests-and-kernels-map/kernels_to_tests.json b/studies/tests-and-kernels-map/kernels_to_tests.json index a19974b169..b848f1ee71 100644 --- a/studies/tests-and-kernels-map/kernels_to_tests.json +++ b/studies/tests-and-kernels-map/kernels_to_tests.json @@ -1150,12 +1150,6 @@ "test_0395_complex_type_arrays.py::test_astype_complex", "test_0652_tests_of_complex_numbers.py::test_reducers" ], - "awkward_unique_copy": [ - "test_0404_array_validity_check.py::test_IndexedArray", - "test_0404_array_validity_check.py::test_same_categories", - "test_0404_array_validity_check.py::test_categorical", - "test_2616_use_pyarrow_for_strings.py::test_split_whitespace - AssertionError" - ], "awkward_ListArray_getitem_next_range_spreadadvanced": [ "test_0011_listarray.py::test_listarray_array_slice", "test_0011_listarray.py::test_listoffsetarray_array_slice", @@ -2445,26 +2439,6 @@ "test_2860_enforce_concatenated_form.py::test_symmetric[right14-left11]", "test_2860_enforce_concatenated_form.py::test_symmetric[right14-left14]" ], - "awkward_unique": [ - "test_0404_array_validity_check.py::test_NumpyArray", - "test_0404_array_validity_check.py::test_ByteMaskedArray_2", - "test_0404_array_validity_check.py::test_ByteMaskedArray_1", - "test_0404_array_validity_check.py::test_BitMaskedArray", - "test_0404_array_validity_check.py::test_UnmaskedArray", - "test_0404_array_validity_check.py::test_RecordArray", - "test_0404_array_validity_check.py::test_ListOffsetArray", - "test_0404_array_validity_check.py::test_IndexedArray", - "test_0404_array_validity_check.py::test_2d_in_axis", - "test_1149_datetime_sort.py::test_date_time_sort_argsort_unique", - "test_1149_datetime_sort.py::test_time_delta_sort_argsort_unique", - "test_0404_array_validity_check.py::test_3d_non_unique", - "test_0404_array_validity_check.py::test_ByteMaskedArray_0", - "test_0404_array_validity_check.py::test_RegularArray", - "test_0404_array_validity_check.py::test_3d", - "test_0404_array_validity_check.py::test_2d", - "test_0404_array_validity_check.py::test_UnionArray", - "test_2550_validity_error_recursive.py::test" - ], "awkward_IndexedArray_reduce_next_nonlocal_nextshifts_64": [ "test_0074_argsort_and_sort.py::test_ByteMaskedArray", "test_0074_argsort_and_sort.py::test_IndexedOptionArray", @@ -4405,9 +4379,6 @@ "test_0404_array_validity_check.py::test_RegularArray", "test_0404_array_validity_check.py::test_ByteMaskedArray_1" ], - "awkward_NumpyArray_fill": [ - "test_0404_array_validity_check.py::test_subranges_equal" - ], "awkward_reduce_max_complex": [ "test_0395_complex_type_arrays.py::test_astype_complex", "test_0395_complex_type_arrays.py::test_count_max_complex", diff --git a/studies/tests-and-kernels-map/tests_to_kernels.json b/studies/tests-and-kernels-map/tests_to_kernels.json index bf17cc73b3..1548ca3551 100644 --- a/studies/tests-and-kernels-map/tests_to_kernels.json +++ b/studies/tests-and-kernels-map/tests_to_kernels.json @@ -3698,7 +3698,6 @@ "awkward_ListOffsetArray_reduce_local_nextparents_64", "awkward_sorting_ranges_length", "awkward_sort", - "awkward_unique", "awkward_sorting_ranges", "awkward_unique_ranges", "awkward_unique_offsets" @@ -3711,7 +3710,6 @@ "awkward_ListOffsetArray_reduce_local_nextparents_64", "awkward_sorting_ranges_length", "awkward_sort", - "awkward_unique", "awkward_sorting_ranges", "awkward_unique_ranges", "awkward_unique_offsets" @@ -3724,7 +3722,6 @@ "awkward_ListOffsetArray_reduce_local_nextparents_64", "awkward_sorting_ranges_length", "awkward_sort", - "awkward_unique", "awkward_sorting_ranges", "awkward_unique_ranges", "awkward_unique_offsets" @@ -3737,7 +3734,6 @@ "awkward_ListOffsetArray_reduce_local_nextparents_64", "awkward_sorting_ranges_length", "awkward_sort", - "awkward_unique", "awkward_sorting_ranges", "awkward_unique_ranges", "awkward_unique_offsets" @@ -3753,7 +3749,6 @@ "awkward_IndexedArray_reduce_next_64", "awkward_sort", "awkward_IndexedArray_unique_next_index_and_offsets_64", - "awkward_unique", "awkward_IndexedArray_numnull_parents", "awkward_IndexedArray_numnull", "awkward_IndexedArray_numnull_unique", @@ -3771,7 +3766,6 @@ "awkward_sorting_ranges_length", "awkward_IndexedArray_reduce_next_64", "awkward_sort", - "awkward_unique", "awkward_IndexedArray_numnull", "awkward_ListArray_compact_offsets", "awkward_IndexedArray_numnull_unique", @@ -3792,7 +3786,6 @@ "awkward_sorting_ranges_length", "awkward_IndexedArray_reduce_next_64", "awkward_sort", - "awkward_unique", "awkward_IndexedArray_numnull", "awkward_ListArray_compact_offsets", "awkward_IndexedArray_numnull_unique", @@ -3813,7 +3806,6 @@ "awkward_sorting_ranges_length", "awkward_IndexedArray_reduce_next_64", "awkward_sort", - "awkward_unique", "awkward_IndexedArray_numnull", "awkward_ListArray_compact_offsets", "awkward_IndexedArray_numnull_unique", @@ -3832,10 +3824,8 @@ "kernels": [ "awkward_ListOffsetArray_reduce_local_nextparents_64", "awkward_sorting_ranges_length", - "awkward_unique_copy", "awkward_IndexedArray_reduce_next_64", "awkward_sort", - "awkward_unique", "awkward_IndexedArray_numnull", "awkward_ListArray_compact_offsets", "awkward_sorting_ranges", @@ -3855,7 +3845,6 @@ "awkward_NumpyArray_sort_asstrings_uint8", "awkward_sorting_ranges_length", "awkward_sort", - "awkward_unique", "awkward_ListArray_compact_offsets", "awkward_sorting_ranges", "awkward_NumpyArray_unique_strings_uint8", @@ -3869,16 +3858,14 @@ "test_0404_array_validity_check.py::test_NumpyArray": { "count": 2, "kernels": [ - "awkward_sort", - "awkward_unique" + "awkward_sort" ], "status": false }, "test_0404_array_validity_check.py::test_RecordArray": { "count": 2, "kernels": [ - "awkward_sort", - "awkward_unique" + "awkward_sort" ], "status": false }, @@ -3890,7 +3877,6 @@ "awkward_IndexedArray_reduce_next_64", "awkward_sort", "awkward_IndexedArray_unique_next_index_and_offsets_64", - "awkward_unique", "awkward_IndexedArray_numnull_parents", "awkward_IndexedArray_numnull", "awkward_IndexedArray_numnull_unique", @@ -3909,7 +3895,6 @@ "awkward_UnionArray_simplify_one", "awkward_sort", "awkward_ListArray_fill", - "awkward_unique", "awkward_ListArray_compact_offsets", "awkward_sorting_ranges", "awkward_ListArray_broadcast_tooffsets", @@ -3921,8 +3906,7 @@ "test_0404_array_validity_check.py::test_UnmaskedArray": { "count": 2, "kernels": [ - "awkward_sort", - "awkward_unique" + "awkward_sort" ], "status": false }, @@ -3930,7 +3914,6 @@ "count": 6, "kernels": [ "awkward_NumpyArray_sort_asstrings_uint8", - "awkward_unique_copy", "awkward_IndexedArray_simplify", "awkward_ListArray_validity", "awkward_NumpyArray_unique_strings_uint8", @@ -3943,7 +3926,6 @@ "kernels": [ "awkward_ListOffsetArray_reduce_local_nextparents_64", "awkward_NumpyArray_sort_asstrings_uint8", - "awkward_unique_copy", "awkward_ListOffsetArray_reduce_local_outoffsets_64", "awkward_reduce_count_64", "awkward_ListArray_compact_offsets", @@ -3960,7 +3942,6 @@ "awkward_ListOffsetArray_reduce_local_nextparents_64", "awkward_sorting_ranges_length", "awkward_sort", - "awkward_NumpyArray_fill", "awkward_sorting_ranges", "awkward_NumpyArray_subrange_equal" ], @@ -7346,7 +7327,6 @@ "kernels": [ "awkward_sorting_ranges_length", "awkward_sort", - "awkward_unique", "awkward_sorting_ranges", "awkward_argsort" ], @@ -7357,7 +7337,6 @@ "kernels": [ "awkward_sorting_ranges_length", "awkward_sort", - "awkward_unique", "awkward_sorting_ranges", "awkward_argsort" ], @@ -10933,7 +10912,6 @@ "count": 3, "kernels": [ "awkward_sort", - "awkward_unique", "awkward_ListArray_validity" ], "status": false @@ -11071,7 +11049,6 @@ "kernels": [ "awkward_UnionArray_simplify_one", "awkward_reduce_prod_complex", - "awkward_unique_copy", "awkward_reduce_countnonzero_complex", "awkward_ByteMaskedArray_reduce_next_64", "awkward_IndexedArray_fill", diff --git a/tests/test_3059_boolean_kernels.py b/tests/test_3059_boolean_kernels.py new file mode 100644 index 0000000000..e299dd054a --- /dev/null +++ b/tests/test_3059_boolean_kernels.py @@ -0,0 +1,117 @@ +# BSD 3-Clause License; see https://github.com/scikit-hep/awkward/blob/master/LICENSE + +from __future__ import annotations + +import numpy as np + +import awkward as ak + +to_list = ak.operations.to_list + + +def test_bool_RecordArray(): + array = ak.highlevel.Array( + [ + {"x": True, "y": [True]}, + {"x": False, "y": [True, False]}, + ] + ) + + assert ak._do.is_unique(array.layout) is False + assert ak._do.is_unique(array["x"].layout) is True + assert ak._do.is_unique(array["y"].layout) is False + + +def test_bool_UnionArray(): + content1 = ak.operations.from_iter([[], [False], [True, True]], highlevel=False) + content2 = ak.operations.from_iter([[False, False], [True], []], highlevel=False) + tags = ak.index.Index8(np.array([0, 1, 0, 1, 0, 1], dtype=np.int8)) + index = ak.index.Index64(np.array([0, 0, 1, 1, 2, 2], dtype=np.int64)) + array = ak.contents.UnionArray.simplified(tags, index, [content1, content2]) + + assert to_list(array) == [ + [], + [False, False], + [False], + [True], + [True, True], + [], + ] + assert ak._do.is_unique(array) is False + assert to_list(ak._do.unique(array, axis=None)) == [False, True] + assert to_list(ak._do.unique(array, axis=-1)) == [ + [], + [False], + [False], + [True], + [True], + [], + ] + + +def test_bool_IndexedArray(): + content = ak.from_iter([[True], [False]], highlevel=False) + index = ak.index.Index64(np.array([1, 0], dtype=np.int64)) + indexedarray = ak.contents.IndexedArray(index, content) + + assert ak._do.is_unique(indexedarray) is True + + listoffsetarray = ak.operations.from_iter([[True], [False]], highlevel=False) + + index = ak.index.Index64(np.array([1, 0], dtype=np.int64)) + indexedarray = ak.contents.IndexedArray(index, listoffsetarray) + assert to_list(indexedarray) == [ + [False], + [True], + ] + assert ak._do.is_unique(indexedarray) is True + + assert ak._do.is_unique(indexedarray) is True + assert to_list(ak._do.unique(indexedarray)) == [ + False, + True, + ] + assert to_list(ak._do.unique(indexedarray, axis=-1)) == [ + [False], + [True], + ] + + +def test_bool_subranges_equal(): + array = ak.contents.NumpyArray( + np.array( + [ + [True, False, True, True, False], + [False, False, True, True, True], + [False, True, False, True, True], + ] + ) + ) + + starts = ak.index.Index64(np.array([0, 5, 10])) + stops = ak.index.Index64(np.array([5, 10, 15])) + + result = ak.sort(array, axis=-1, highlevel=False).content._subranges_equal( + starts, stops, 15 + ) + assert result is True + + starts = ak.index.Index64(np.array([0, 7])) + stops = ak.index.Index64(np.array([7, 15])) + + assert ( + ak.sort(array, axis=-1, highlevel=False).content._subranges_equal( + starts, stops, 15 + ) + is False + ) + + starts = ak.index.Index64(np.array([0])) + stops = ak.index.Index64(np.array([15])) + + assert ( + ak.sort(array, axis=-1, highlevel=False).content._subranges_equal( + starts, stops, 15 + ) + is False + )