tests/sycl/test_barrier.sycl

0001 /**
0002  * traccc library, part of the ACTS project (R&D line)
0003  *
0004  * (c) 2024-2026 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 "../../device/sycl/src/utils/barrier.hpp"
0014
0015 // VecMem include(s).
0016 #include <vecmem/memory/sycl/shared_memory_resource.hpp>
0017 #include <vecmem/memory/unique_ptr.hpp>
0018
0019 // GoogleTest include(s).
0020 #include <gtest/gtest.h>
0021
0022 TEST(SYCLBarrier, BarrierAnd) {
0023   vecmem::sycl::shared_memory_resource mr;
0024
0025   ::sycl::queue queue;
0026   constexpr std::size_t n_bools = 4;
0027
0028   vecmem::unique_alloc_ptr<bool[]> out =
0029       vecmem::make_unique_alloc<bool[]>(mr, n_bools);
0030
0031   ::sycl::nd_range test_range(::sycl::range<1>(128), ::sycl::range<1>(128));
0032
0033   queue
0034       .submit([&, out = out.get()](::sycl::handler& h) {
0035         h.parallel_for<class BarrierAndTest>(
0036             test_range, [=](::sycl::nd_item<1> item) {
0037               const traccc::sycl::details::barrier bar{item};
0038
0039               bool v;
0040
0041               v = bar.blockAnd(false);
0042               if (item.get_local_id()[0] == 0) {
0043                 out[0] = v;
0044               }
0045
0046               v = bar.blockAnd(true);
0047               if (item.get_local_id()[0] == 0) {
0048                 out[1] = v;
0049               }
0050
0051               v = bar.blockAnd(item.get_local_id()[0] % 2 == 0);
0052               if (item.get_local_id()[0] == 0) {
0053                 out[2] = v;
0054               }
0055
0056               v = bar.blockAnd(item.get_local_id()[0] < 32);
0057               if (item.get_local_id()[0] == 0) {
0058                 out[3] = v;
0059               }
0060             });
0061       })
0062       .wait_and_throw();
0063
0064   EXPECT_FALSE(out.get()[0]);
0065   EXPECT_TRUE(out.get()[1]);
0066   EXPECT_FALSE(out.get()[2]);
0067   EXPECT_FALSE(out.get()[3]);
0068 }
0069
0070 TEST(SYCLBarrier, BarrierOr) {
0071   vecmem::sycl::shared_memory_resource mr;
0072
0073   ::sycl::queue queue;
0074   constexpr std::size_t n_bools = 4;
0075
0076   vecmem::unique_alloc_ptr<bool[]> out =
0077       vecmem::make_unique_alloc<bool[]>(mr, n_bools);
0078
0079   ::sycl::nd_range test_range(::sycl::range<1>(128), ::sycl::range<1>(128));
0080
0081   queue
0082       .submit([&, out = out.get()](::sycl::handler& h) {
0083         h.parallel_for<class BarrierOrTest>(
0084             test_range, [=](::sycl::nd_item<1> item) {
0085               const traccc::sycl::details::barrier bar{item};
0086
0087               bool v;
0088
0089               v = bar.blockOr(false);
0090               if (item.get_local_id()[0] == 0) {
0091                 out[0] = v;
0092               }
0093
0094               v = bar.blockOr(true);
0095               if (item.get_local_id()[0] == 0) {
0096                 out[1] = v;
0097               }
0098
0099               v = bar.blockOr(item.get_local_id()[0] % 2 == 0);
0100               if (item.get_local_id()[0] == 0) {
0101                 out[2] = v;
0102               }
0103
0104               v = bar.blockOr(item.get_local_id()[0] < 32);
0105               if (item.get_local_id()[0] == 0) {
0106                 out[3] = v;
0107               }
0108             });
0109       })
0110       .wait_and_throw();
0111
0112   EXPECT_FALSE(out.get()[0]);
0113   EXPECT_TRUE(out.get()[1]);
0114   EXPECT_TRUE(out.get()[2]);
0115   EXPECT_TRUE(out.get()[3]);
0116 }
0117
0118 TEST(SYCLBarrier, BarrierCount) {
0119   vecmem::sycl::shared_memory_resource mr;
0120
0121   ::sycl::queue queue;
0122   constexpr std::size_t n_ints = 4;
0123
0124   vecmem::unique_alloc_ptr<unsigned int[]> out =
0125       vecmem::make_unique_alloc<unsigned int[]>(mr, n_ints);
0126
0127   ::sycl::nd_range test_range(::sycl::range<1>(128), ::sycl::range<1>(128));
0128
0129   queue
0130       .submit([&, out = out.get()](::sycl::handler& h) {
0131         h.parallel_for<class BarrierCountTest>(
0132             test_range, [=](::sycl::nd_item<1> item) {
0133               const traccc::sycl::details::barrier bar{item};
0134
0135               unsigned int v;
0136
0137               v = bar.blockCount(false);
0138               if (item.get_local_id()[0] == 0) {
0139                 out[0] = v;
0140               }
0141
0142               v = bar.blockCount(true);
0143               if (item.get_local_id()[0] == 0) {
0144                 out[1] = v;
0145               }
0146
0147               v = bar.blockCount(item.get_local_id()[0] % 2 == 0);
0148               if (item.get_local_id()[0] == 0) {
0149                 out[2] = v;
0150               }
0151
0152               v = bar.blockCount(item.get_local_id()[0] < 32);
0153               if (item.get_local_id()[0] == 0) {
0154                 out[3] = v;
0155               }
0156             });
0157       })
0158       .wait_and_throw();
0159
0160   EXPECT_EQ(out.get()[0], 0);
0161   EXPECT_EQ(out.get()[1], 128);
0162   EXPECT_EQ(out.get()[2], 64);
0163   EXPECT_EQ(out.get()[3], 32);
0164 }