tests/cuda/test_mutex.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 <gtest/gtest.h>
0010 #include <vecmem/memory/cuda/managed_memory_resource.hpp>
0011 #include <vecmem/memory/unique_ptr.hpp>
0012
0013 #include "../../device/cuda/src/utils/cuda_error_handling.hpp"
0014 #include "traccc/device/mutex.hpp"
0015
0016 __global__ void mutex_add_kernel(uint32_t *out, uint32_t *lock) {
0017   traccc::device::mutex m(*lock);
0018
0019   if (threadIdx.x == 0) {
0020     m.lock();
0021     uint32_t tmp = *out;
0022     tmp += 1;
0023     *out = tmp;
0024     m.unlock();
0025   }
0026 }
0027
0028 TEST(CUDAMutex, MassAdditionKernel) {
0029   vecmem::cuda::managed_memory_resource mr;
0030
0031   vecmem::unique_alloc_ptr<uint32_t> out =
0032       vecmem::make_unique_alloc<uint32_t>(mr);
0033   vecmem::unique_alloc_ptr<uint32_t> lock =
0034       vecmem::make_unique_alloc<uint32_t>(mr);
0035
0036   TRACCC_CUDA_ERROR_CHECK(cudaMemset(lock.get(), 0, sizeof(uint32_t)));
0037   TRACCC_CUDA_ERROR_CHECK(cudaMemset(out.get(), 0, sizeof(uint32_t)));
0038
0039   uint32_t n_blocks = 262144;
0040   uint32_t n_threads = 32;
0041
0042   mutex_add_kernel<<<n_blocks, n_threads>>>(out.get(), lock.get());
0043
0044   TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());
0045   TRACCC_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0046
0047   EXPECT_EQ(n_blocks, *out.get());
0048 }