device/cuda/container_cuda_kernel.cu

0001 // This file is part of the ACTS project.
0002 //
0003 // Copyright (C) 2016 CERN for the benefit of the ACTS project
0004 //
0005 // This Source Code Form is subject to the terms of the Mozilla Public
0006 // License, v. 2.0. If a copy of the MPL was not distributed with this
0007 // file, You can obtain one at https://mozilla.org/MPL/2.0/.
0008
0009 #include "detray/definitions/detail/cuda_definitions.hpp"
0010
0011 // Detray test include(s)
0012 #include "container_cuda_kernel.hpp"
0013
0014 namespace detray {
0015
0016 __global__ void test_single_store_kernel(
0017     single_store_t::view_type store_view,
0018     vecmem::data::vector_view<double> sum_data) {
0019   single_store_dev_t store(store_view);
0020   vecmem::device_vector<double> sum(sum_data);
0021
0022   for (const double e : store) {
0023     sum[0] += e;
0024   }
0025 }
0026
0027 __global__ void test_tuple_cont_kernel(
0028     typename tuple_cont_t::view_type store_view,
0029     vecmem::data::vector_view<double> sum_data) {
0030   tuple_cont_dev_t store(store_view);
0031   vecmem::device_vector<double> sum(sum_data);
0032
0033   const auto& vec0 = store.get<0>();
0034   const auto& vec1 = store.get<1>();
0035
0036   for (const auto e : vec0) {
0037     sum[0] += e;
0038   }
0039   for (const auto e : vec1) {
0040     sum[0] += e;
0041   }
0042 }
0043
0044 __global__ void test_reg_multi_store_kernel(
0045     reg_multi_store_t::view_type store_view,
0046     vecmem::data::vector_view<double> sum_data) {
0047   using enum reg_type_ids;
0048
0049   reg_multi_store_dev_t store(store_view);
0050   vecmem::device_vector<double> sum(sum_data);
0051
0052   const auto& vec0 = store.get<e_size>();
0053   const auto& vec1 = store.get<e_float>();
0054   const auto& vec2 = store.get<e_double>();
0055
0056   for (const auto e : vec0) {
0057     sum[0] += e;
0058   }
0059   for (const auto e : vec1) {
0060     sum[0] += e;
0061   }
0062   for (const auto e : vec2) {
0063     sum[0] += e;
0064   }
0065 }
0066
0067 __global__ void test_multi_store_kernel(
0068     multi_store_t::view_type store_view,
0069     vecmem::data::vector_view<double> sum_data) {
0070   using enum type_ids;
0071
0072   multi_store_dev_t store(store_view);
0073   vecmem::device_vector<double> sum(sum_data);
0074
0075   const auto& vec0 = store.get<e_float>();
0076   const auto& vec1 = store.get<e_test_class>().first;
0077   const auto& vec2 = store.get<e_test_class>().second;
0078
0079   for (const auto e : vec0) {
0080     sum[0] += e;
0081   }
0082   for (const auto e : vec1) {
0083     sum[0] += e;
0084   }
0085   for (const auto e : vec2) {
0086     sum[0] += e;
0087   }
0088 }
0089
0090 void test_single_store(single_store_t::view_type store_view,
0091                        vecmem::data::vector_view<double> sum_data) {
0092   // run the test kernel
0093   test_single_store_kernel<<<1, 1>>>(store_view, sum_data);
0094
0095   // cuda error check
0096   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0097   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0098 }
0099
0100 void test_tuple_container(tuple_cont_t::view_type store_view,
0101                           vecmem::data::vector_view<double> sum_data) {
0102   // run the test kernel
0103   test_tuple_cont_kernel<<<1, 1>>>(store_view, sum_data);
0104
0105   // cuda error check
0106   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0107   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0108 }
0109
0110 void test_reg_multi_store(reg_multi_store_t::view_type store_view,
0111                           vecmem::data::vector_view<double> sum_data) {
0112   // run the test kernel
0113   test_reg_multi_store_kernel<<<1, 1>>>(store_view, sum_data);
0114
0115   // cuda error check
0116   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0117   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0118 }
0119
0120 void test_multi_store(typename multi_store_t::view_type store_view,
0121                       vecmem::data::vector_view<double> sum_data) {
0122   // run the test kernel
0123   test_multi_store_kernel<<<1, 1>>>(store_view, sum_data);
0124
0125   // cuda error check
0126   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0127   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0128 }
0129
0130 }  // namespace detray