tests/cuda/test_barrier.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/barrier.hpp"
0014
0015 __global__ void testBarrierAnd(bool* out) {
0016   traccc::cuda::barrier bar;
0017
0018   bool v;
0019
0020   v = bar.blockAnd(false);
0021   if (threadIdx.x == 0) {
0022     out[0] = v;
0023   }
0024
0025   v = bar.blockAnd(true);
0026   if (threadIdx.x == 0) {
0027     out[1] = v;
0028   }
0029
0030   v = bar.blockAnd(threadIdx.x % 2 == 0);
0031   if (threadIdx.x == 0) {
0032     out[2] = v;
0033   }
0034
0035   v = bar.blockAnd(threadIdx.x < 32);
0036   if (threadIdx.x == 0) {
0037     out[3] = v;
0038   }
0039 }
0040
0041 TEST(CUDABarrier, BarrierAnd) {
0042   vecmem::cuda::managed_memory_resource mr;
0043   constexpr std::size_t n_bools = 4;
0044
0045   vecmem::unique_alloc_ptr<bool[]> out =
0046       vecmem::make_unique_alloc<bool[]>(mr, n_bools);
0047
0048   testBarrierAnd<<<1, 1024>>>(out.get());
0049
0050   ASSERT_EQ(cudaGetLastError(), cudaSuccess);
0051   ASSERT_EQ(cudaDeviceSynchronize(), cudaSuccess);
0052
0053   EXPECT_FALSE(out.get()[0]);
0054   EXPECT_TRUE(out.get()[1]);
0055   EXPECT_FALSE(out.get()[2]);
0056   EXPECT_FALSE(out.get()[3]);
0057 }
0058
0059 __global__ void testBarrierOr(bool* out) {
0060   traccc::cuda::barrier bar;
0061
0062   bool v;
0063
0064   v = bar.blockOr(false);
0065   if (threadIdx.x == 0) {
0066     out[0] = v;
0067   }
0068
0069   v = bar.blockOr(true);
0070   if (threadIdx.x == 0) {
0071     out[1] = v;
0072   }
0073
0074   v = bar.blockOr(threadIdx.x % 2 == 0);
0075   if (threadIdx.x == 0) {
0076     out[2] = v;
0077   }
0078
0079   v = bar.blockOr(threadIdx.x < 32);
0080   if (threadIdx.x == 0) {
0081     out[3] = v;
0082   }
0083 }
0084
0085 TEST(CUDABarrier, BarrierOr) {
0086   vecmem::cuda::managed_memory_resource mr;
0087   constexpr std::size_t n_bools = 4;
0088
0089   vecmem::unique_alloc_ptr<bool[]> out =
0090       vecmem::make_unique_alloc<bool[]>(mr, n_bools);
0091
0092   testBarrierOr<<<1, 1024>>>(out.get());
0093
0094   ASSERT_EQ(cudaGetLastError(), cudaSuccess);
0095   ASSERT_EQ(cudaDeviceSynchronize(), cudaSuccess);
0096
0097   EXPECT_FALSE(out.get()[0]);
0098   EXPECT_TRUE(out.get()[1]);
0099   EXPECT_TRUE(out.get()[2]);
0100   EXPECT_TRUE(out.get()[3]);
0101 }
0102
0103 __global__ void testBarrierCount(int* out) {
0104   traccc::cuda::barrier bar;
0105
0106   int v;
0107
0108   v = bar.blockCount(false);
0109   if (threadIdx.x == 0) {
0110     out[0] = v;
0111   }
0112
0113   v = bar.blockCount(true);
0114   if (threadIdx.x == 0) {
0115     out[1] = v;
0116   }
0117
0118   v = bar.blockCount(threadIdx.x % 2 == 0);
0119   if (threadIdx.x == 0) {
0120     out[2] = v;
0121   }
0122
0123   v = bar.blockCount(threadIdx.x < 32);
0124   if (threadIdx.x == 0) {
0125     out[3] = v;
0126   }
0127 }
0128
0129 TEST(CUDABarrier, BarrierCount) {
0130   vecmem::cuda::managed_memory_resource mr;
0131   constexpr std::size_t n_ints = 4;
0132
0133   vecmem::unique_alloc_ptr<int[]> out =
0134       vecmem::make_unique_alloc<int[]>(mr, n_ints);
0135
0136   testBarrierCount<<<1, 1024>>>(out.get());
0137
0138   ASSERT_EQ(cudaGetLastError(), cudaSuccess);
0139   ASSERT_EQ(cudaDeviceSynchronize(), cudaSuccess);
0140
0141   EXPECT_EQ(out.get()[0], 0);
0142   EXPECT_EQ(out.get()[1], 1024);
0143   EXPECT_EQ(out.get()[2], 512);
0144   EXPECT_EQ(out.get()[3], 32);
0145 }