device/cuda/mask_store_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 "mask_store_cuda_kernel.hpp"
0013
0014 namespace detray {
0015
0016 /// test kernel function to fill the output vector with is_inside function
0017 /// return values
0018 __global__ void mask_test_kernel(
0019     typename host_store_type::view_type store_data,
0020     vecmem::data::vector_view<point3> input_point3_data,
0021     vecmem::data::jagged_vector_view<int> output_data) {
0022   // get mask store
0023   device_store_type store(store_data);
0024
0025   // get mask objects
0026   vecmem::device_vector<point3> input_point3(input_point3_data);
0027   vecmem::jagged_device_vector<int> output_device(output_data);
0028
0029   const auto& rectangle_mask = store.get<mask_id::e_rectangle2D>()[0];
0030   const auto& trapezoid_mask = store.get<mask_id::e_trapezoid2D>()[0];
0031   const auto& ring_mask = store.get<mask_id::e_ring2D>()[0];
0032   const auto& cylinder_mask = store.get<mask_id::e_cylinder2D>()[0];
0033   const auto& annulus_mask = store.get<mask_id::e_annulus2D>()[0];
0034
0035   // get device results from is_inside function
0036   for (int i = 0; i < n_points; i++) {
0037     output_device[0].push_back(rectangle_mask.is_inside(input_point3[i]));
0038     output_device[1].push_back(trapezoid_mask.is_inside(input_point3[i]));
0039     output_device[2].push_back(ring_mask.is_inside(input_point3[i]));
0040     output_device[3].push_back(cylinder_mask.is_inside(input_point3[i]));
0041     output_device[4].push_back(annulus_mask.is_inside(input_point3[i]));
0042   }
0043 }
0044
0045 void mask_test(typename host_store_type::view_type store_data,
0046                vecmem::data::vector_view<point3> input_point3_data,
0047                vecmem::data::jagged_vector_view<int> output_data) {
0048   int block_dim = 1;
0049   int thread_dim = 1;
0050
0051   // run the test kernel
0052   mask_test_kernel<<<block_dim, thread_dim>>>(store_data, input_point3_data,
0053                                               output_data);
0054
0055   // cuda error check
0056   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0057   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0058 }
0059
0060 }  // namespace detray