tests/cuda/test_sort.cu

0001 /**
0002  * traccc library, part of the ACTS project (R&D line)
0003  *
0004  * (c) 2024-2025 CERN for the benefit of the ACTS project
0005  *
0006  * Mozilla Public License Version 2.0
0007  */
0008
0009 #include <gtest/gtest.h>
0010 #include <vecmem/memory/cuda/managed_memory_resource.hpp>
0011 #include <vecmem/memory/unique_ptr.hpp>
0012
0013 #include "../../cuda/src/utils/barrier.hpp"
0014 #include "../../cuda/src/utils/thread_id.hpp"
0015 #include "traccc/device/sort.hpp"
0016
0017 __global__ void testBlockSortKernel(uint32_t *keys, uint32_t n_keys) {
0018   traccc::cuda::details::thread_id1 thread_id;
0019   traccc::cuda::barrier barrier;
0020   traccc::device::blockOddEvenSort(thread_id, barrier, keys, n_keys,
0021                                    std::less<uint32_t>());
0022 }
0023
0024 TEST(CUDASort, BlockOddEvenSort) {
0025   vecmem::cuda::managed_memory_resource mr;
0026
0027   uint32_t n = 2803;
0028   vecmem::unique_alloc_ptr<uint32_t[]> arr =
0029       vecmem::make_unique_alloc<uint32_t[]>(mr, n);
0030
0031   // As long as 13 and n_keys are coprime, this will generate a big,
0032   // non-sorted array containing every element.
0033   for (uint32_t i = 0; i < n; i++) {
0034     arr[i] = (13 * 500 * i) % n;
0035   }
0036
0037   testBlockSortKernel<<<1, 1024u>>>(arr.get(), n);
0038
0039   ASSERT_EQ(cudaPeekAtLastError(), cudaSuccess);
0040   ASSERT_EQ(cudaDeviceSynchronize(), cudaSuccess);
0041
0042   for (uint32_t i = 0; i < n; ++i) {
0043     ASSERT_EQ(arr[i], i);
0044   }
0045 }