tests/sycl/test_mutex.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 "traccc/device/mutex.hpp"
0014
0015 // VecMem include(s).
0016 #include <vecmem/memory/sycl/device_memory_resource.hpp>
0017 #include <vecmem/memory/sycl/shared_memory_resource.hpp>
0018 #include <vecmem/memory/unique_ptr.hpp>
0019
0020 // GoogleTest include(s).
0021 #include <gtest/gtest.h>
0022
0023 TEST(SYCLMutex, MassAdditionKernel) {
0024   ::sycl::queue queue;
0025   vecmem::sycl::device_memory_resource device_mr{&queue};
0026   vecmem::sycl::shared_memory_resource shared_mr{&queue};
0027
0028   vecmem::unique_alloc_ptr<uint32_t> out =
0029       vecmem::make_unique_alloc<uint32_t>(shared_mr);
0030   vecmem::unique_alloc_ptr<uint32_t> lock =
0031       vecmem::make_unique_alloc<uint32_t>(device_mr);
0032
0033   queue.memset(lock.get(), 0, sizeof(uint32_t)).wait_and_throw();
0034   queue.memset(out.get(), 0, sizeof(uint32_t)).wait_and_throw();
0035
0036   uint32_t n_blocks = 262144;
0037   uint32_t n_threads = 32;
0038
0039   ::sycl::nd_range test_range(::sycl::range<1>(n_blocks * n_threads),
0040                               ::sycl::range<1>(n_threads));
0041
0042   queue
0043       .submit([&, out = out.get(), lock = lock.get()](::sycl::handler &h) {
0044         h.parallel_for<class MassAdditionTest>(
0045             test_range, [=](::sycl::nd_item<1> item) {
0046               traccc::device::mutex m(*lock);
0047
0048               if (item.get_local_id()[0] == 0) {
0049                 m.lock();
0050                 uint32_t tmp = *out;
0051                 tmp += 1;
0052                 *out = tmp;
0053                 m.unlock();
0054               }
0055             });
0056       })
0057       .wait_and_throw();
0058
0059   EXPECT_EQ(n_blocks, *out.get());
0060 }