tests/sycl/test_sort.sycl

0001 /**
0002  * traccc library, part of the ACTS project (R&D line)
0003  *
0004  * (c) 2024 CERN for the benefit of the ACTS project
0005  *
0006  * Mozilla Public License Version 2.0
0007  */
0008
0009 // SYCL include(s).
0010 #include <sycl/sycl.hpp>
0011
0012 // Project include(s).
0013 #include "../../sycl/src/utils/barrier.hpp"
0014 #include "../../sycl/src/utils/thread_id.hpp"
0015 #include "traccc/device/sort.hpp"
0016
0017 // VecMem include(s).
0018 #include <vecmem/memory/sycl/shared_memory_resource.hpp>
0019 #include <vecmem/memory/unique_ptr.hpp>
0020
0021 // GTest include(s).
0022 #include <gtest/gtest.h>
0023
0024 TEST(SYCLSort, BlockOddEvenSort) {
0025   vecmem::sycl::shared_memory_resource mr;
0026   ::sycl::queue queue;
0027
0028   uint32_t n = 2803;
0029   vecmem::unique_alloc_ptr<uint32_t[]> arr =
0030       vecmem::make_unique_alloc<uint32_t[]>(mr, n);
0031
0032   // As long as 13 and n_keys are coprime, this will generate a big,
0033   // non-sorted array containing every element.
0034   for (uint32_t i = 0; i < n; i++) {
0035     arr[i] = (13 * 500 * i) % n;
0036   }
0037
0038   ::sycl::nd_range test_range(::sycl::range<1>(128), ::sycl::range<1>(128));
0039
0040   queue
0041       .submit([&, keys = arr.get()](::sycl::handler& h) {
0042         h.parallel_for<class BlockOddEvenSortKernel>(
0043             test_range, [=](::sycl::nd_item<1> item) {
0044               const traccc::sycl::details::thread_id thread_id{item};
0045               const traccc::sycl::details::barrier barrier{item};
0046               traccc::device::blockOddEvenSort(thread_id, barrier, keys, n,
0047                                                std::less<uint32_t>());
0048             });
0049       })
0050       .wait_and_throw();
0051
0052   for (uint32_t i = 0; i < n; ++i) {
0053     ASSERT_EQ(arr[i], i);
0054   }
0055 }