device/cuda/utils_ranges_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 "utils_ranges_cuda_kernel.hpp"
0013
0014 namespace detray {
0015
0016 //
0017 // single
0018 //
0019 __global__ void single_kernel(const dindex value, dindex* result) {
0020   // single view should only add the value 'i' once
0021   for (auto i : detray::views::single(value)) {
0022     *result += i;
0023   }
0024 }
0025
0026 void test_single(const dindex value, dindex& check) {
0027   dindex* result{nullptr};
0028   cudaMallocManaged(&result, sizeof(dindex));
0029   *result = 0u;
0030
0031   // run the kernel
0032   single_kernel<<<1, 1>>>(value, result);
0033
0034   // cuda error check
0035   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0036   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0037
0038   check = *result;
0039   cudaFree(result);
0040 }
0041
0042 //
0043 // pointer
0044 //
0045 __global__ void pointer_kernel(const dindex value, dindex* result) {
0046   // pointer view should only add the value 'i' once
0047   for (auto i : detray::views::pointer(value)) {
0048     *result += i;
0049   }
0050 }
0051
0052 void test_pointer(const dindex value, dindex& check) {
0053   dindex* result{nullptr};
0054   cudaMallocManaged(&result, sizeof(dindex));
0055   *result = 0u;
0056
0057   // run the kernel
0058   pointer_kernel<<<1, 1>>>(value, result);
0059
0060   // cuda error check
0061   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0062   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0063
0064   check = *result;
0065   cudaFree(result);
0066 }
0067
0068 //
0069 // iota
0070 //
0071 __global__ void iota_kernel(const darray<dindex, 2> range,
0072                             vecmem::data::vector_view<dindex> check_data) {
0073   vecmem::device_vector<dindex> check(check_data);
0074
0075   for (auto i : detray::views::iota(range)) {
0076     check.push_back(i);
0077   }
0078 }
0079
0080 void test_iota(const darray<dindex, 2> range,
0081                vecmem::data::vector_view<dindex> check_data) {
0082   // run the kernel
0083   iota_kernel<<<1, 1>>>(range, check_data);
0084
0085   // cuda error check
0086   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0087   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0088 }
0089
0090 //
0091 // cartesian product
0092 //
0093 __global__ void cartesian_product_kernel(
0094     const darray<dindex, 2> range1, const darray<dindex, 2> range2,
0095     const darray<dindex, 2> range3,
0096     vecmem::data::vector_view<std::tuple<dindex, dindex, dindex>> check_data) {
0097   vecmem::device_vector<std::tuple<dindex, dindex, dindex>> check(check_data);
0098
0099   auto seq1 = detray::views::iota(range1);
0100   auto seq2 = detray::views::iota(range2);
0101   auto seq3 = detray::views::iota(range3);
0102
0103   for (const auto [i, j, k] : detray::views::cartesian_product(
0104            std::move(seq1), std::move(seq2), std::move(seq3))) {
0105     check.emplace_back(i, j, k);
0106   }
0107 }
0108
0109 void test_cartesian_product(
0110     const darray<dindex, 2> range1, const darray<dindex, 2> range2,
0111     const darray<dindex, 2> range3,
0112     vecmem::data::vector_view<std::tuple<dindex, dindex, dindex>> check_data) {
0113   // run the kernel
0114   cartesian_product_kernel<<<1, 1>>>(range1, range2, range3, check_data);
0115
0116   // cuda error check
0117   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0118   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0119 }
0120
0121 //
0122 // enumerate
0123 //
0124 __global__ void enumerate_kernel(
0125     vecmem::data::vector_view<uint_holder> seq_data,
0126     vecmem::data::vector_view<dindex> check_idx_data,
0127     vecmem::data::vector_view<dindex> check_value_data) {
0128   vecmem::device_vector<uint_holder> seq(seq_data);
0129   vecmem::device_vector<dindex> check_idx(check_idx_data);
0130   vecmem::device_vector<dindex> check_value(check_value_data);
0131
0132   for (auto [i, v] : detray::views::enumerate(seq)) {
0133     check_idx.push_back(i);
0134     check_value.push_back(v.ui);
0135   }
0136 }
0137
0138 void test_enumerate(vecmem::data::vector_view<uint_holder> seq_data,
0139                     vecmem::data::vector_view<dindex> check_idx_data,
0140                     vecmem::data::vector_view<dindex> check_value_data) {
0141   // run the kernel
0142   enumerate_kernel<<<1, 1>>>(seq_data, check_idx_data, check_value_data);
0143
0144   // cuda error check
0145   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0146   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0147 }
0148
0149 //
0150 // pick
0151 //
0152 __global__ void pick_kernel(
0153     vecmem::data::vector_view<uint_holder> seq_data,
0154     vecmem::data::vector_view<dindex> idx_data,
0155     vecmem::data::vector_view<dindex> check_idx_data,
0156     vecmem::data::vector_view<dindex> check_value_data) {
0157   vecmem::device_vector<uint_holder> seq(seq_data);
0158   vecmem::device_vector<dindex> idx(idx_data);
0159   vecmem::device_vector<dindex> check_idx(check_idx_data);
0160   vecmem::device_vector<dindex> check_value(check_value_data);
0161
0162   for (auto [i, v] : detray::views::pick(seq, idx)) {
0163     check_idx.push_back(i);
0164     check_value.push_back(v.ui);
0165   }
0166 }
0167
0168 void test_pick(vecmem::data::vector_view<uint_holder> seq_data,
0169                vecmem::data::vector_view<dindex> idx_data,
0170                vecmem::data::vector_view<dindex> check_idx_data,
0171                vecmem::data::vector_view<dindex> check_value_data) {
0172   // run the kernel
0173   pick_kernel<<<1, 1>>>(seq_data, idx_data, check_idx_data, check_value_data);
0174
0175   // cuda error check
0176   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0177   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0178 }
0179
0180 //
0181 // join
0182 //
0183 __global__ void join_kernel(
0184     vecmem::data::vector_view<uint_holder> seq_data_1,
0185     vecmem::data::vector_view<uint_holder> seq_data_2,
0186     vecmem::data::vector_view<dindex> check_value_data) {
0187   vecmem::device_vector<uint_holder> seq_1(seq_data_1);
0188   vecmem::device_vector<uint_holder> seq_2(seq_data_2);
0189   vecmem::device_vector<dindex> check_value(check_value_data);
0190   std::array<vecmem::device_vector<uint_holder>, 2> vectors{seq_1, seq_2};
0191
0192   for (auto v : detray::views::join(vectors)) {
0193     check_value.push_back(v.ui);
0194   }
0195 }
0196
0197 void test_join(vecmem::data::vector_view<uint_holder> seq_data_1,
0198                vecmem::data::vector_view<uint_holder> seq_data_2,
0199                vecmem::data::vector_view<dindex> check_value_data) {
0200   // run the kernel
0201   join_kernel<<<1, 1>>>(seq_data_1, seq_data_2, check_value_data);
0202
0203   // cuda error check
0204   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0205   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0206 }
0207
0208 //
0209 // static_join
0210 //
0211 __global__ void static_join_kernel(
0212     vecmem::data::vector_view<uint_holder> seq_data_1,
0213     vecmem::data::vector_view<uint_holder> seq_data_2,
0214     vecmem::data::vector_view<dindex> check_value_data) {
0215   vecmem::device_vector<uint_holder> seq_1(seq_data_1);
0216   vecmem::device_vector<uint_holder> seq_2(seq_data_2);
0217   vecmem::device_vector<dindex> check_value(check_value_data);
0218
0219   for (auto v : detray::views::static_join(seq_1, seq_2)) {
0220     check_value.push_back(v.ui);
0221   }
0222 }
0223
0224 void test_static_join(vecmem::data::vector_view<uint_holder> seq_data_1,
0225                       vecmem::data::vector_view<uint_holder> seq_data_2,
0226                       vecmem::data::vector_view<dindex> check_value_data) {
0227   // run the kernel
0228   static_join_kernel<<<1, 1>>>(seq_data_1, seq_data_2, check_value_data);
0229
0230   // cuda error check
0231   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0232   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0233 }
0234
0235 //
0236 // subrange
0237 //
0238 __global__ void subrange_kernel(vecmem::data::vector_view<int> seq_data,
0239                                 vecmem::data::vector_view<int> check_value_data,
0240                                 const std::size_t begin,
0241                                 const std::size_t end) {
0242   vecmem::device_vector<int> seq(seq_data);
0243   vecmem::device_vector<int> check(check_value_data);
0244
0245   for (const auto& v :
0246        detray::ranges::subrange(seq, std::array<std::size_t, 2>{begin, end})) {
0247     check.push_back(v);
0248   }
0249 }
0250
0251 void test_subrange(vecmem::data::vector_view<int> seq_data,
0252                    vecmem::data::vector_view<int> check_value_data,
0253                    const std::size_t begin, const std::size_t end) {
0254   // run the kernel
0255   subrange_kernel<<<1, 1>>>(seq_data, check_value_data, begin, end);
0256
0257   // cuda error check
0258   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0259   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0260 }
0261
0262 }  // namespace detray