tests/cuda/test_unique_lock.cu

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 #include <mutex>
0010
0011 #include <gtest/gtest.h>
0012 #include <vecmem/memory/cuda/managed_memory_resource.hpp>
0013 #include <vecmem/memory/unique_ptr.hpp>
0014
0015 #include "../../device/cuda/src/utils/cuda_error_handling.hpp"
0016 #include "traccc/device/mutex.hpp"
0017 #include "traccc/device/unique_lock.hpp"
0018
0019 __global__ void unique_lock_add_kernel_try_lock(uint32_t *out,
0020                                                 uint32_t *_lock) {
0021   traccc::device::mutex m(*_lock);
0022
0023   if (threadIdx.x == 0) {
0024     traccc::device::unique_lock lock(m, std::try_to_lock);
0025
0026     if (!lock.owns_lock()) {
0027       lock.lock();
0028     }
0029
0030     uint32_t tmp = *out;
0031     tmp += 1;
0032     *out = tmp;
0033   }
0034 }
0035
0036 __global__ void unique_lock_add_kernel_defer_lock(uint32_t *out,
0037                                                   uint32_t *_lock) {
0038   traccc::device::mutex m(*_lock);
0039   traccc::device::unique_lock lock(m, std::defer_lock);
0040
0041   if (threadIdx.x == 0) {
0042     lock.lock();
0043
0044     uint32_t tmp = *out;
0045     tmp += 1;
0046     *out = tmp;
0047   }
0048 }
0049
0050 __global__ void unique_lock_add_kernel_adopt_lock(uint32_t *out,
0051                                                   uint32_t *_lock) {
0052   traccc::device::mutex m(*_lock);
0053
0054   if (threadIdx.x == 0) {
0055     m.lock();
0056     traccc::device::unique_lock lock(m, std::adopt_lock);
0057
0058     uint32_t tmp = *out;
0059     tmp += 1;
0060     *out = tmp;
0061   }
0062 }
0063
0064 TEST(CUDAUniqueLock, MassAdditionKernelTryLock) {
0065   vecmem::cuda::managed_memory_resource mr;
0066
0067   vecmem::unique_alloc_ptr<uint32_t> out =
0068       vecmem::make_unique_alloc<uint32_t>(mr);
0069   vecmem::unique_alloc_ptr<uint32_t> lock =
0070       vecmem::make_unique_alloc<uint32_t>(mr);
0071
0072   TRACCC_CUDA_ERROR_CHECK(cudaMemset(lock.get(), 0, sizeof(uint32_t)));
0073   TRACCC_CUDA_ERROR_CHECK(cudaMemset(out.get(), 0, sizeof(uint32_t)));
0074
0075   uint32_t n_blocks = 262144;
0076   uint32_t n_threads = 32;
0077
0078   unique_lock_add_kernel_try_lock<<<n_blocks, n_threads>>>(out.get(),
0079                                                            lock.get());
0080
0081   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
0082   TRACCC_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0083
0084   EXPECT_EQ(n_blocks, *out.get());
0085 }
0086
0087 TEST(CUDAUniqueLock, MassAdditionKernelDeferLock) {
0088   vecmem::cuda::managed_memory_resource mr;
0089
0090   vecmem::unique_alloc_ptr<uint32_t> out =
0091       vecmem::make_unique_alloc<uint32_t>(mr);
0092   vecmem::unique_alloc_ptr<uint32_t> lock =
0093       vecmem::make_unique_alloc<uint32_t>(mr);
0094
0095   TRACCC_CUDA_ERROR_CHECK(cudaMemset(lock.get(), 0, sizeof(uint32_t)));
0096   TRACCC_CUDA_ERROR_CHECK(cudaMemset(out.get(), 0, sizeof(uint32_t)));
0097
0098   uint32_t n_blocks = 262144;
0099   uint32_t n_threads = 32;
0100
0101   unique_lock_add_kernel_defer_lock<<<n_blocks, n_threads>>>(out.get(),
0102                                                              lock.get());
0103
0104   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
0105   TRACCC_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0106
0107   EXPECT_EQ(n_blocks, *out.get());
0108 }
0109
0110 TEST(CUDAUniqueLock, MassAdditionKernelAdoptLock) {
0111   vecmem::cuda::managed_memory_resource mr;
0112
0113   vecmem::unique_alloc_ptr<uint32_t> out =
0114       vecmem::make_unique_alloc<uint32_t>(mr);
0115   vecmem::unique_alloc_ptr<uint32_t> lock =
0116       vecmem::make_unique_alloc<uint32_t>(mr);
0117
0118   TRACCC_CUDA_ERROR_CHECK(cudaMemset(lock.get(), 0, sizeof(uint32_t)));
0119   TRACCC_CUDA_ERROR_CHECK(cudaMemset(out.get(), 0, sizeof(uint32_t)));
0120
0121   uint32_t n_blocks = 262144;
0122   uint32_t n_threads = 32;
0123
0124   unique_lock_add_kernel_adopt_lock<<<n_blocks, n_threads>>>(out.get(),
0125                                                              lock.get());
0126
0127   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
0128   TRACCC_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0129
0130   EXPECT_EQ(n_blocks, *out.get());
0131 }