tests/cuda/test_sync.cu

0001 /**
0002  * TRACCC library, part of the ACTS project (R&D line)
0003  *
0004  * (c) 2021 CERN for the benefit of the ACTS project
0005  *
0006  * Mozilla Public License Version 2.0
0007  */
0008
0009 #include <gtest/gtest.h>
0010
0011 #include "../../device/cuda/src/utils/sync.cuh"
0012
0013 __global__ void testWarpIndexedBallotSyncBasicKernel(uint32_t *vts,
0014                                                      uint32_t *vis) {
0015   auto [vt, vi] = traccc::cuda::warp_indexed_ballot_sync(threadIdx.x % 2 == 0);
0016
0017   vts[threadIdx.x] = vt;
0018   vis[threadIdx.x] = vi;
0019 }
0020
0021 __global__ void testWarpIndexedBallotSyncWithExitKernel(uint32_t *vts,
0022                                                         uint32_t *vis) {
0023   if (threadIdx.x < 16) {
0024     return;
0025   }
0026
0027   auto [vt, vi] = traccc::cuda::warp_indexed_ballot_sync(threadIdx.x % 2 == 0);
0028
0029   vts[threadIdx.x] = vt;
0030   vis[threadIdx.x] = vi;
0031 }
0032
0033 TEST(CUDASync, WarpIndexedBallotSyncBasic) {
0034   uint32_t *dev_vt = nullptr, *dev_vi = nullptr;
0035   uint32_t host_vt[32], host_vi[32];
0036
0037   ASSERT_EQ(cudaMalloc(&dev_vt, 32u * sizeof(uint32_t)), cudaSuccess);
0038   ASSERT_EQ(cudaMalloc(&dev_vi, 32u * sizeof(uint32_t)), cudaSuccess);
0039   ASSERT_NE(dev_vt, nullptr);
0040   ASSERT_NE(dev_vi, nullptr);
0041
0042   testWarpIndexedBallotSyncBasicKernel<<<1, 32u>>>(dev_vt, dev_vi);
0043
0044   ASSERT_EQ(cudaPeekAtLastError(), cudaSuccess);
0045
0046   ASSERT_EQ(cudaMemcpy(host_vt, dev_vt, 32u * sizeof(uint32_t),
0047                        cudaMemcpyDeviceToHost),
0048             cudaSuccess);
0049   ASSERT_EQ(cudaMemcpy(host_vi, dev_vi, 32u * sizeof(uint32_t),
0050                        cudaMemcpyDeviceToHost),
0051             cudaSuccess);
0052
0053   for (uint32_t i = 0; i < 32u; ++i) {
0054     ASSERT_EQ(host_vt[i], 16u);
0055   }
0056
0057   for (uint32_t i = 0; i < 16u; ++i) {
0058     ASSERT_EQ(host_vi[i * 2], i);
0059   }
0060
0061   ASSERT_EQ(cudaFree(dev_vt), cudaSuccess);
0062   ASSERT_EQ(cudaFree(dev_vi), cudaSuccess);
0063 }
0064
0065 TEST(CUDASync, WarpIndexedBallotSyncWithExit) {
0066   uint32_t *dev_vt = nullptr, *dev_vi = nullptr;
0067   uint32_t host_vt[32], host_vi[32];
0068
0069   ASSERT_EQ(cudaMalloc(&dev_vt, 32u * sizeof(uint32_t)), cudaSuccess);
0070   ASSERT_EQ(cudaMalloc(&dev_vi, 32u * sizeof(uint32_t)), cudaSuccess);
0071   ASSERT_NE(dev_vt, nullptr);
0072   ASSERT_NE(dev_vi, nullptr);
0073
0074   testWarpIndexedBallotSyncWithExitKernel<<<1, 32u>>>(dev_vt, dev_vi);
0075
0076   ASSERT_EQ(cudaPeekAtLastError(), cudaSuccess);
0077
0078   ASSERT_EQ(cudaMemcpy(host_vt, dev_vt, 32u * sizeof(uint32_t),
0079                        cudaMemcpyDeviceToHost),
0080             cudaSuccess);
0081   ASSERT_EQ(cudaMemcpy(host_vi, dev_vi, 32u * sizeof(uint32_t),
0082                        cudaMemcpyDeviceToHost),
0083             cudaSuccess);
0084
0085   for (uint32_t i = 16; i < 32u; ++i) {
0086     ASSERT_EQ(host_vt[i], 8u);
0087   }
0088
0089   for (uint32_t i = 0; i < 8u; ++i) {
0090     ASSERT_EQ(host_vi[i * 2 + 16], i);
0091   }
0092
0093   ASSERT_EQ(cudaFree(dev_vt), cudaSuccess);
0094   ASSERT_EQ(cudaFree(dev_vi), cudaSuccess);
0095 }