tests/cuda/test_array_wrapper.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 #include <traccc/utils/array_wrapper.hpp>
0011 #include <vecmem/memory/cuda/device_memory_resource.hpp>
0012
0013 template <template <typename> typename F>
0014 struct vector {
0015   using tuple_t = std::tuple<uint32_t, uint32_t, uint32_t>;
0016   F<uint32_t> x, y, z;
0017 };
0018
0019 template <template <typename...> typename Layout>
0020 __global__ void testArrayWrapperKernel(
0021     const typename traccc::array_wrapper<Layout, vector>::handle h,
0022     uint32_t *total) {
0023   __shared__ uint32_t block_total;
0024
0025   if (threadIdx.x == 0) {
0026     block_total = 0;
0027   }
0028
0029   __syncthreads();
0030
0031   int tid = blockIdx.x * blockDim.x + threadIdx.x;
0032
0033   uint32_t warp_total = h[tid].x;
0034
0035   for (int i = 16; i >= 1; i /= 2) {
0036     warp_total += __shfl_xor_sync(0xffffffff, warp_total, i, 32);
0037   }
0038
0039   if (threadIdx.x % 32 == 0) {
0040     atomicAdd(&block_total, warp_total);
0041   }
0042
0043   __syncthreads();
0044
0045   if (threadIdx.x == 0) {
0046     atomicAdd(total, block_total);
0047   }
0048 }
0049
0050 template <template <typename...> typename Layout>
0051 __global__ void fillWrapperKernel(
0052     typename traccc::array_wrapper<Layout, vector>::handle h) {
0053   int tid = blockIdx.x * blockDim.x + threadIdx.x;
0054
0055   if (tid < h.size()) {
0056     h[tid].x = 1u;
0057     h[tid].y = 2u;
0058     h[tid].z = 3u;
0059   }
0060 }
0061
0062 TEST(CUDAArrayWrapper, SoALayout) {
0063   vecmem::cuda::device_memory_resource mr;
0064
0065   uint32_t n = 1024u * 1024u;
0066
0067   traccc::array_wrapper<traccc::soa, vector>::owner o(mr, n);
0068
0069   fillWrapperKernel<traccc::soa><<<n / 1024u, 1024u>>>(
0070       typename traccc::array_wrapper<traccc::soa, vector>::handle(o));
0071
0072   ASSERT_EQ(cudaPeekAtLastError(), cudaSuccess);
0073   ASSERT_EQ(cudaDeviceSynchronize(), cudaSuccess);
0074
0075   uint32_t host_result;
0076   uint32_t *dev_result = nullptr;
0077
0078   ASSERT_EQ(cudaMalloc(&dev_result, sizeof(uint32_t)), cudaSuccess);
0079   ASSERT_EQ(cudaMemset(&dev_result, sizeof(uint32_t), 0), cudaSuccess);
0080
0081   testArrayWrapperKernel<traccc::soa><<<n / 1024u, 1024u>>>(
0082       typename traccc::array_wrapper<traccc::soa, vector>::handle(o),
0083       dev_result);
0084
0085   ASSERT_EQ(cudaPeekAtLastError(), cudaSuccess);
0086   ASSERT_EQ(cudaDeviceSynchronize(), cudaSuccess);
0087
0088   ASSERT_EQ(cudaMemcpy(&host_result, dev_result, sizeof(uint32_t),
0089                        cudaMemcpyDeviceToHost),
0090             cudaSuccess);
0091
0092   ASSERT_EQ(host_result, n);
0093
0094   ASSERT_EQ(cudaFree(dev_result), cudaSuccess);
0095 }
0096
0097 TEST(CUDAArrayWrapper, AoSLayout) {
0098   vecmem::cuda::device_memory_resource mr;
0099
0100   uint32_t n = 1024u * 1024u;
0101
0102   traccc::array_wrapper<traccc::aos, vector>::owner o(mr, n);
0103
0104   fillWrapperKernel<traccc::aos><<<n / 1024u, 1024u>>>(
0105       typename traccc::array_wrapper<traccc::aos, vector>::handle(o));
0106
0107   ASSERT_EQ(cudaPeekAtLastError(), cudaSuccess);
0108   ASSERT_EQ(cudaDeviceSynchronize(), cudaSuccess);
0109
0110   uint32_t host_result;
0111   uint32_t *dev_result = nullptr;
0112
0113   ASSERT_EQ(cudaMalloc(&dev_result, sizeof(uint32_t)), cudaSuccess);
0114   ASSERT_EQ(cudaMemset(&dev_result, sizeof(uint32_t), 0), cudaSuccess);
0115
0116   testArrayWrapperKernel<traccc::aos><<<n / 1024u, 1024u>>>(
0117       typename traccc::array_wrapper<traccc::aos, vector>::handle(o),
0118       dev_result);
0119
0120   ASSERT_EQ(cudaPeekAtLastError(), cudaSuccess);
0121   ASSERT_EQ(cudaDeviceSynchronize(), cudaSuccess);
0122
0123   ASSERT_EQ(cudaMemcpy(&host_result, dev_result, sizeof(uint32_t),
0124                        cudaMemcpyDeviceToHost),
0125             cudaSuccess);
0126
0127   ASSERT_EQ(host_result, n);
0128
0129   ASSERT_EQ(cudaFree(dev_result), cudaSuccess);
0130 }