tests/sycl/test_unique_lock.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 #include "traccc/device/unique_lock.hpp"
0015
0016 // VecMem include(s).
0017 #include <vecmem/memory/sycl/device_memory_resource.hpp>
0018 #include <vecmem/memory/sycl/shared_memory_resource.hpp>
0019 #include <vecmem/memory/unique_ptr.hpp>
0020
0021 // GoogleTest include(s).
0022 #include <gtest/gtest.h>
0023
0024 TEST(SYCLUniqueLock, MassAdditionKernelTryLock) {
0025   ::sycl::queue queue;
0026   vecmem::sycl::device_memory_resource device_mr{&queue};
0027   vecmem::sycl::shared_memory_resource shared_mr{&queue};
0028
0029   vecmem::unique_alloc_ptr<uint32_t> out =
0030       vecmem::make_unique_alloc<uint32_t>(shared_mr);
0031   vecmem::unique_alloc_ptr<uint32_t> lock =
0032       vecmem::make_unique_alloc<uint32_t>(device_mr);
0033
0034   queue.memset(lock.get(), 0, sizeof(uint32_t)).wait_and_throw();
0035   queue.memset(out.get(), 0, sizeof(uint32_t)).wait_and_throw();
0036
0037   uint32_t n_blocks = 262144;
0038   uint32_t n_threads = 32;
0039
0040   ::sycl::nd_range test_range(::sycl::range<1>(n_blocks * n_threads),
0041                               ::sycl::range<1>(n_threads));
0042
0043   queue
0044       .submit([&, out = out.get(), _lock = lock.get()](::sycl::handler& h) {
0045         h.parallel_for<class MassAdditionTryLockTest>(
0046             test_range, [=](::sycl::nd_item<1> item) {
0047               traccc::device::mutex m(*_lock);
0048
0049               if (item.get_local_id()[0] == 0) {
0050                 traccc::device::unique_lock lock(m, std::try_to_lock);
0051
0052                 if (!lock.owns_lock()) {
0053                   lock.lock();
0054                 }
0055
0056                 uint32_t tmp = *out;
0057                 tmp += 1;
0058                 *out = tmp;
0059               }
0060             });
0061       })
0062       .wait_and_throw();
0063
0064   EXPECT_EQ(n_blocks, *out.get());
0065 }
0066
0067 TEST(SYCLUniqueLock, MassAdditionKernelDeferLock) {
0068   ::sycl::queue queue;
0069   vecmem::sycl::device_memory_resource device_mr{&queue};
0070   vecmem::sycl::shared_memory_resource shared_mr{&queue};
0071
0072   vecmem::unique_alloc_ptr<uint32_t> out =
0073       vecmem::make_unique_alloc<uint32_t>(shared_mr);
0074   vecmem::unique_alloc_ptr<uint32_t> lock =
0075       vecmem::make_unique_alloc<uint32_t>(device_mr);
0076
0077   queue.memset(lock.get(), 0, sizeof(uint32_t)).wait_and_throw();
0078   queue.memset(out.get(), 0, sizeof(uint32_t)).wait_and_throw();
0079
0080   uint32_t n_blocks = 262144;
0081   uint32_t n_threads = 32;
0082
0083   ::sycl::nd_range test_range(::sycl::range<1>(n_blocks * n_threads),
0084                               ::sycl::range<1>(n_threads));
0085
0086   queue
0087       .submit([&, out = out.get(), _lock = lock.get()](::sycl::handler& h) {
0088         h.parallel_for<class MassAdditionDeferLockTest>(
0089             test_range, [=](::sycl::nd_item<1> item) {
0090               traccc::device::mutex m(*_lock);
0091               traccc::device::unique_lock lock(m, std::defer_lock);
0092
0093               if (item.get_local_id()[0] == 0) {
0094                 lock.lock();
0095
0096                 uint32_t tmp = *out;
0097                 tmp += 1;
0098                 *out = tmp;
0099               }
0100             });
0101       })
0102       .wait_and_throw();
0103
0104   EXPECT_EQ(n_blocks, *out.get());
0105 }
0106
0107 TEST(SYCLUniqueLock, MassAdditionKernelAdoptLock) {
0108   ::sycl::queue queue;
0109   vecmem::sycl::device_memory_resource device_mr{&queue};
0110   vecmem::sycl::shared_memory_resource shared_mr{&queue};
0111
0112   vecmem::unique_alloc_ptr<uint32_t> out =
0113       vecmem::make_unique_alloc<uint32_t>(shared_mr);
0114   vecmem::unique_alloc_ptr<uint32_t> lock =
0115       vecmem::make_unique_alloc<uint32_t>(device_mr);
0116
0117   queue.memset(lock.get(), 0, sizeof(uint32_t)).wait_and_throw();
0118   queue.memset(out.get(), 0, sizeof(uint32_t)).wait_and_throw();
0119
0120   uint32_t n_blocks = 262144;
0121   uint32_t n_threads = 32;
0122
0123   ::sycl::nd_range test_range(::sycl::range<1>(n_blocks * n_threads),
0124                               ::sycl::range<1>(n_threads));
0125
0126   queue
0127       .submit([&, out = out.get(), _lock = lock.get()](::sycl::handler& h) {
0128         h.parallel_for<class MassAdditionAdoptLockTest>(
0129             test_range, [=](::sycl::nd_item<1> item) {
0130               traccc::device::mutex m(*_lock);
0131
0132               if (item.get_local_id()[0] == 0) {
0133                 m.lock();
0134                 traccc::device::unique_lock lock(m, std::adopt_lock);
0135
0136                 uint32_t tmp = *out;
0137                 tmp += 1;
0138                 *out = tmp;
0139               }
0140             });
0141       })
0142       .wait_and_throw();
0143
0144   EXPECT_EQ(n_blocks, *out.get());
0145 }