diff --git a/src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_max.cu b/src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_max.cu index 2941aa417e..0d1ecdebcd 100644 --- a/src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_max.cu +++ b/src/awkward/_connect/cuda/cuda_kernels/awkward_reduce_max.cu @@ -60,7 +60,7 @@ awkward_reduce_max_b( T val = identity; if (idx >= stride && thread_id < lenparents && parents[thread_id] == parents[thread_id - stride]) { - val = temp[idx - stride]; + val = temp[thread_id - stride]; } __syncthreads(); temp[thread_id] = val > temp[thread_id] ? val : temp[thread_id]; diff --git a/tests-cuda/test_3149_complex_reducers.py b/tests-cuda/test_3149_complex_reducers.py index bd53020721..abc921b30d 100644 --- a/tests-cuda/test_3149_complex_reducers.py +++ b/tests-cuda/test_3149_complex_reducers.py @@ -302,8 +302,8 @@ def test_0652_minmax(): def test_block_boundary_sum_complex(): - np.random.seed(42) - array = np.random.randint(6000, size=6000) + rng = np.random.default_rng(seed=42) + array = rng.integers(6000, size=6000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -323,7 +323,6 @@ def test_block_boundary_sum_complex(): def test_block_boundary_prod_complex1(): - np.random.seed(42) complex_array = np.vectorize(complex)(np.full(1000, 0), np.full(1000, 1)) content = ak.contents.NumpyArray(complex_array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -341,7 +340,6 @@ def test_block_boundary_prod_complex1(): def test_block_boundary_prod_complex2(): - np.random.seed(42) complex_array = np.vectorize(complex)(np.full(1001, 0), np.full(1001, 1)) content = ak.contents.NumpyArray(complex_array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -359,7 +357,6 @@ def test_block_boundary_prod_complex2(): def test_block_boundary_prod_complex3(): - np.random.seed(42) complex_array = np.vectorize(complex)(np.full(1002, 0), np.full(1002, 1)) content = ak.contents.NumpyArray(complex_array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -377,7 +374,6 @@ def test_block_boundary_prod_complex3(): def test_block_boundary_prod_complex4(): - np.random.seed(42) complex_array = np.vectorize(complex)(np.full(1000, 0), np.full(1000, 1.01)) content = ak.contents.NumpyArray(complex_array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -397,7 +393,6 @@ def test_block_boundary_prod_complex4(): def test_block_boundary_prod_complex5(): - np.random.seed(42) complex_array = np.vectorize(complex)(np.full(1001, 0), np.full(1001, 1.01)) content = ak.contents.NumpyArray(complex_array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -417,7 +412,6 @@ def test_block_boundary_prod_complex5(): def test_block_boundary_prod_complex6(): - np.random.seed(42) complex_array = np.vectorize(complex)(np.full(1002, 0), np.full(1002, 1.01)) content = ak.contents.NumpyArray(complex_array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -437,7 +431,6 @@ def test_block_boundary_prod_complex6(): def test_block_boundary_prod_complex7(): - np.random.seed(42) complex_array = np.vectorize(complex)(np.full(1000, 0), np.full(1000, 0.99)) content = ak.contents.NumpyArray(complex_array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -457,7 +450,6 @@ def test_block_boundary_prod_complex7(): def test_block_boundary_prod_complex8(): - np.random.seed(42) complex_array = np.vectorize(complex)(np.full(1001, 0), np.full(1001, 0.99)) content = ak.contents.NumpyArray(complex_array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -477,7 +469,6 @@ def test_block_boundary_prod_complex8(): def test_block_boundary_prod_complex9(): - np.random.seed(42) complex_array = np.vectorize(complex)(np.full(1002, 0), np.full(1002, 0.99)) content = ak.contents.NumpyArray(complex_array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -497,7 +488,6 @@ def test_block_boundary_prod_complex9(): def test_block_boundary_prod_complex10(): - np.random.seed(42) complex_array = np.vectorize(complex)(np.full(1000, 0), np.full(1000, 1.1)) content = ak.contents.NumpyArray(complex_array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -517,7 +507,6 @@ def test_block_boundary_prod_complex10(): def test_block_boundary_prod_complex11(): - np.random.seed(42) complex_array = np.vectorize(complex)(np.full(1001, 0), np.full(1001, 1.1)) content = ak.contents.NumpyArray(complex_array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -537,7 +526,6 @@ def test_block_boundary_prod_complex11(): def test_block_boundary_prod_complex12(): - np.random.seed(42) complex_array = np.vectorize(complex)(np.full(1002, 0), np.full(1002, 1.1)) content = ak.contents.NumpyArray(complex_array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -557,8 +545,8 @@ def test_block_boundary_prod_complex12(): def test_block_boundary_prod_complex13(): - np.random.seed(42) - array = np.random.randint(50, size=1000) + rng = np.random.default_rng(seed=42) + array = rng.integers(50, size=1000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -580,8 +568,8 @@ def test_block_boundary_prod_complex13(): def test_block_boundary_any_complex(): - np.random.seed(42) - array = np.random.randint(6000, size=6000) + rng = np.random.default_rng(seed=42) + array = rng.integers(6000, size=6000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -601,8 +589,8 @@ def test_block_boundary_any_complex(): def test_block_boundary_all_complex(): - np.random.seed(42) - array = np.random.randint(6000, size=6000) + rng = np.random.default_rng(seed=42) + array = rng.integers(6000, size=6000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -622,8 +610,8 @@ def test_block_boundary_all_complex(): def test_block_boundary_min_complex1(): - np.random.seed(42) - array = np.random.randint(5, size=6000) + rng = np.random.default_rng(seed=42) + array = rng.integers(5, size=6000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -643,8 +631,8 @@ def test_block_boundary_min_complex1(): def test_block_boundary_min_complex2(): - np.random.seed(42) - array = np.random.randint(6000, size=6000) + rng = np.random.default_rng(seed=42) + array = rng.integers(6000, size=6000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -664,8 +652,8 @@ def test_block_boundary_min_complex2(): def test_block_boundary_max_complex1(): - np.random.seed(42) - array = np.random.randint(5, size=6000) + rng = np.random.default_rng(seed=42) + array = rng.integers(5, size=6000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -685,8 +673,8 @@ def test_block_boundary_max_complex1(): def test_block_boundary_max_complex2(): - np.random.seed(42) - array = np.random.randint(6000, size=6000) + rng = np.random.default_rng(seed=42) + array = rng.integers(6000, size=6000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -706,8 +694,8 @@ def test_block_boundary_max_complex2(): def test_block_boundary_sum_bool_complex(): - np.random.seed(42) - array = np.random.randint(2, size=6000, dtype=np.bool_) + rng = np.random.default_rng(seed=42) + array = rng.integers(2, size=6000, dtype=np.bool_) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -727,8 +715,8 @@ def test_block_boundary_sum_bool_complex(): def test_block_boundary_countnonzero_complex_1(): - np.random.seed(42) - array = np.random.randint(6000, size=6000) + rng = np.random.default_rng(seed=42) + array = rng.integers(6000, size=6000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -749,8 +737,8 @@ def test_block_boundary_countnonzero_complex_1(): def test_block_boundary_countnonzero_complex_2(): - np.random.seed(42) - array = np.random.randint(2, size=6000) + rng = np.random.default_rng(seed=42) + array = rng.integers(2, size=6000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -772,8 +760,8 @@ def test_block_boundary_countnonzero_complex_2(): @pytest.mark.skip(reason="awkward_reduce_argmax_complex is not implemented") def test_block_boundary_argmax_complex1(): - np.random.seed(42) - array = np.random.randint(5, size=6000) + rng = np.random.default_rng(seed=42) + array = rng.integers(5, size=6000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -794,8 +782,8 @@ def test_block_boundary_argmax_complex1(): @pytest.mark.skip(reason="awkward_reduce_argmax_complex is not implemented") def test_block_boundary_argmax_complex2(): - np.random.seed(42) - array = np.random.randint(6000, size=6000) + rng = np.random.default_rng(seed=42) + array = rng.integers(6000, size=6000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -816,8 +804,8 @@ def test_block_boundary_argmax_complex2(): @pytest.mark.skip(reason="awkward_reduce_argmin_complex is not implemented") def test_block_boundary_argmin_complex1(): - np.random.seed(42) - array = np.random.randint(5, size=6000) + rng = np.random.default_rng(seed=42) + array = rng.integers(5, size=6000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) @@ -838,8 +826,8 @@ def test_block_boundary_argmin_complex1(): @pytest.mark.skip(reason="awkward_reduce_argmin_complex is not implemented") def test_block_boundary_argmin_complex2(): - np.random.seed(42) - array = np.random.randint(6000, size=6000) + rng = np.random.default_rng(seed=42) + array = rng.integers(6000, size=6000) complex_array = np.vectorize(complex)( array[0 : len(array) : 2], array[1 : len(array) : 2] ) diff --git a/tests-cuda/test_3150_combinations_n_equal_2.py b/tests-cuda/test_3150_combinations_n_equal_2.py index d65ef8416a..801d83599a 100644 --- a/tests-cuda/test_3150_combinations_n_equal_2.py +++ b/tests-cuda/test_3150_combinations_n_equal_2.py @@ -1188,7 +1188,6 @@ def test_1074_combinations_UnmaskedArray(): def test_block_boundary_combinations(): - np.random.seed(42) content = ak.contents.NumpyArray(np.arange(300)) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -1219,7 +1218,6 @@ def test_block_boundary_combinations(): def test_block_boundary_argcombinations(): - np.random.seed(42) content = ak.contents.NumpyArray(np.arange(300)) cuda_content = ak.to_backend(content, "cuda", highlevel=False) diff --git a/tests-cuda/test_3162_block_boundary_reducers.py b/tests-cuda/test_3162_block_boundary_reducers.py index cd0b57a0c8..deb52da002 100644 --- a/tests-cuda/test_3162_block_boundary_reducers.py +++ b/tests-cuda/test_3162_block_boundary_reducers.py @@ -17,8 +17,9 @@ def cleanup_cuda(): def test_block_boundary_sum(): - np.random.seed(42) - content = ak.contents.NumpyArray(np.random.randint(3000, size=3000)) + rng = np.random.default_rng(seed=42) + array = rng.integers(3000, size=3000) + content = ak.contents.NumpyArray(array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) assert ak.sum(cuda_content, -1, highlevel=False) == ak.sum( content, -1, highlevel=False @@ -34,8 +35,9 @@ def test_block_boundary_sum(): def test_block_boundary_any(): - np.random.seed(42) - content = ak.contents.NumpyArray(np.random.randint(3000, size=3000)) + rng = np.random.default_rng(seed=42) + array = rng.integers(3000, size=3000) + content = ak.contents.NumpyArray(array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) assert ak.any(cuda_content, -1, highlevel=False) == ak.any( content, -1, highlevel=False @@ -51,8 +53,9 @@ def test_block_boundary_any(): def test_block_boundary_all(): - np.random.seed(42) - content = ak.contents.NumpyArray(np.random.randint(3000, size=3000)) + rng = np.random.default_rng(seed=42) + array = rng.integers(3000, size=3000) + content = ak.contents.NumpyArray(array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) assert ak.all(cuda_content, -1, highlevel=False) == ak.all( content, -1, highlevel=False @@ -68,8 +71,9 @@ def test_block_boundary_all(): def test_block_boundary_sum_bool(): - np.random.seed(42) - content = ak.contents.NumpyArray(np.random.randint(2, size=3000, dtype=np.bool_)) + rng = np.random.default_rng(seed=42) + array = rng.integers(2, size=3000, dtype=np.bool_) + content = ak.contents.NumpyArray(array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) assert ak.sum(cuda_content, -1, highlevel=False) == ak.sum( content, -1, highlevel=False @@ -85,9 +89,13 @@ def test_block_boundary_sum_bool(): def test_block_boundary_max(): - np.random.seed(42) - content = ak.contents.NumpyArray(np.random.randint(3000, size=3000)) + rng = np.random.default_rng(seed=42) + array = rng.integers(3000, size=3000) + print(array) + content = ak.contents.NumpyArray(array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) + print(ak.max(content, -1, highlevel=False)) + print(ak.max(cuda_content, -1, highlevel=False)) assert ak.max(cuda_content, -1, highlevel=False) == ak.max( content, -1, highlevel=False ) @@ -102,8 +110,27 @@ def test_block_boundary_max(): def test_block_boundary_min(): - np.random.seed(42) - content = ak.contents.NumpyArray(np.random.randint(3000, size=3000)) + rng = np.random.default_rng(seed=42) + array = rng.integers(3000, size=3000) + content = ak.contents.NumpyArray(array) + cuda_content = ak.to_backend(content, "cuda", highlevel=False) + assert ak.min(cuda_content, -1, highlevel=False) == ak.min( + content, -1, highlevel=False + ) + + offsets = ak.index.Index64(np.array([0, 1, 2998, 3000], dtype=np.int64)) + depth1 = ak.contents.ListOffsetArray(offsets, content) + cuda_depth1 = ak.to_backend(depth1, "cuda", highlevel=False) + assert to_list(ak.min(cuda_depth1, -1, highlevel=False)) == to_list( + ak.min(depth1, -1, highlevel=False) + ) + del cuda_content, cuda_depth1 + + +def test_block_boundary_negative_min(): + rng = np.random.default_rng(seed=42) + array = rng.integers(3000, size=3000) * -1 + content = ak.contents.NumpyArray(array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) assert ak.min(cuda_content, -1, highlevel=False) == ak.min( content, -1, highlevel=False @@ -120,8 +147,9 @@ def test_block_boundary_min(): @pytest.mark.skip(reason="awkward_reduce_argmin is not implemented") def test_block_boundary_argmin(): - np.random.seed(42) - content = ak.contents.NumpyArray(np.random.randint(3000, size=3000)) + rng = np.random.default_rng(seed=42) + array = rng.integers(3000, size=3000) + content = ak.contents.NumpyArray(array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) assert ak.argmin(cuda_content, -1, highlevel=False) == ak.argmin( content, -1, highlevel=False @@ -138,8 +166,9 @@ def test_block_boundary_argmin(): @pytest.mark.skip(reason="awkward_reduce_argmax is not implemented") def test_block_boundary_argmax(): - np.random.seed(42) - content = ak.contents.NumpyArray(np.random.randint(3000, size=3000)) + rng = np.random.default_rng(seed=42) + array = rng.integers(3000, size=3000) + content = ak.contents.NumpyArray(array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) assert ak.argmax(cuda_content, -1, highlevel=False) == ak.argmax( content, -1, highlevel=False @@ -155,8 +184,9 @@ def test_block_boundary_argmax(): def test_block_boundary_count(): - np.random.seed(42) - content = ak.contents.NumpyArray(np.random.randint(3000, size=3000)) + rng = np.random.default_rng(seed=42) + array = rng.integers(3000, size=3000) + content = ak.contents.NumpyArray(array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) assert ak.count(cuda_content, -1, highlevel=False) == ak.count( content, -1, highlevel=False @@ -172,8 +202,9 @@ def test_block_boundary_count(): def test_block_boundary_count_nonzero(): - np.random.seed(42) - content = ak.contents.NumpyArray(np.random.randint(2, size=3000)) + rng = np.random.default_rng(seed=42) + array = rng.integers(2, size=3000) + content = ak.contents.NumpyArray(array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) assert ak.count_nonzero(cuda_content, -1, highlevel=False) == ak.count_nonzero( content, -1, highlevel=False @@ -189,7 +220,6 @@ def test_block_boundary_count_nonzero(): def test_block_boundary_prod(): - np.random.seed(42) primes = [x for x in range(2, 30000) if all(x % n != 0 for n in range(2, x))] content = ak.contents.NumpyArray(primes) cuda_content = ak.to_backend(content, "cuda", highlevel=False) @@ -207,8 +237,9 @@ def test_block_boundary_prod(): def test_block_boundary_prod_bool(): - np.random.seed(42) - content = ak.contents.NumpyArray(np.random.randint(2, size=3000, dtype=np.bool_)) + rng = np.random.default_rng(seed=42) + array = rng.integers(2, size=3000, dtype=np.bool_) + content = ak.contents.NumpyArray(array) cuda_content = ak.to_backend(content, "cuda", highlevel=False) assert ak.prod(cuda_content, -1, highlevel=False) == ak.prod( content, -1, highlevel=False