device/cuda/detector_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 #include "detray/geometry/tracking_surface.hpp"
0011
0012 // Detray test include(s)
0013 #include "detector_cuda_kernel.hpp"
0014
0015 namespace detray {
0016
0017 // cuda kernel to copy sub-detector objects
0018 __global__ void detector_test_kernel(
0019     typename detector_host_t::view_type det_data,
0020     vecmem::data::vector_view<det_volume_t> volumes_data,
0021     vecmem::data::vector_view<det_surface_t> surfaces_data,
0022     vecmem::data::vector_view<transform_t> transforms_data,
0023     vecmem::data::vector_view<rectangle_t> rectangles_data,
0024     vecmem::data::vector_view<disc_t> discs_data,
0025     vecmem::data::vector_view<cylinder_t> cylinders_data) {
0026   // convert toy detector_data into detector w/ device vectors
0027   detector_device_t det_device(det_data);
0028
0029   // convert subdetector data objects into objects w/ device vectors
0030   vecmem::device_vector<det_volume_t> volumes_device(volumes_data);
0031   vecmem::device_vector<det_surface_t> surfaces_device(surfaces_data);
0032   vecmem::device_vector<transform_t> transforms_device(transforms_data);
0033   vecmem::device_vector<rectangle_t> rectangles_device(rectangles_data);
0034   vecmem::device_vector<disc_t> discs_device(discs_data);
0035   vecmem::device_vector<cylinder_t> cylinders_device(cylinders_data);
0036
0037   // copy objects - volume
0038   for (unsigned int i = 0u; i < det_device.volumes().size(); i++) {
0039     volumes_device[i] = det_device.volumes()[i];
0040   }
0041
0042   // copy objects - surfaces
0043   for (unsigned int i = 0u; i < det_device.surfaces().size(); i++) {
0044     surfaces_device[i] = det_device.surfaces()[i];
0045   }
0046
0047   // copy objects - transforms
0048   auto& trfs = det_device.transform_store();
0049   auto ctx = typename detector_host_t::geometry_context{};
0050   for (unsigned int i = 0u; i < trfs.size(ctx); i++) {
0051     transforms_device[i] = trfs.at(i, ctx);
0052   }
0053
0054   // copy objects - masks
0055   auto& masks = det_device.mask_store();
0056   auto& rectangles =
0057       masks.template get<detector_host_t::masks::id::e_rectangle2D>();
0058   for (unsigned int i = 0u; i < rectangles.size(); i++) {
0059     rectangles_device[i] = rectangles[i];
0060   }
0061
0062   auto& discs = masks.template get<detector_host_t::masks::id::e_ring2D>();
0063   for (unsigned int i = 0u; i < discs.size(); i++) {
0064     discs_device[i] = discs[i];
0065   }
0066
0067   auto& cylinders =
0068       masks.template get<detector_host_t::masks::id::e_concentric_cylinder2D>();
0069   for (unsigned int i = 0u; i < cylinders.size(); i++) {
0070     cylinders_device[i] = cylinders[i];
0071   }
0072
0073   // print output test for surface finder
0074   /*auto& accel_device = det_device.accelerator_store();
0075   for (unsigned int i_s = 0u; i_s < accel_device.size(); i_s++) {
0076       auto& grid = accel_device[i_s];
0077       for (unsigned int i = 0u; i < grid.axis_p0().bins(); i++) {
0078           for (unsigned int j = 0u; j < grid.axis_p1().bins(); j++) {
0079               const auto& bin = grid.bin(i, j);
0080               for (auto& id : bin) {
0081                   // printf("%d \n", id);
0082               }
0083           }
0084       }
0085   }*/
0086 }
0087
0088 /// implementation of the test function for detector
0089 void detector_test(typename detector_host_t::view_type det_data,
0090                    vecmem::data::vector_view<det_volume_t> volumes_data,
0091                    vecmem::data::vector_view<det_surface_t> surfaces_data,
0092                    vecmem::data::vector_view<transform_t> transforms_data,
0093                    vecmem::data::vector_view<rectangle_t> rectangles_data,
0094                    vecmem::data::vector_view<disc_t> discs_data,
0095                    vecmem::data::vector_view<cylinder_t> cylinders_data) {
0096   constexpr int block_dim = 1u;
0097   constexpr int thread_dim = 1u;
0098
0099   // run the test kernel
0100   detector_test_kernel<<<block_dim, thread_dim>>>(
0101       det_data, volumes_data, surfaces_data, transforms_data, rectangles_data,
0102       discs_data, cylinders_data);
0103
0104   // cuda error check
0105   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0106   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0107 }
0108
0109 // cuda kernel to extract surface transforms from two detector views - static
0110 // and misaligned - and to copy them into vectors
0111 __global__ void detector_alignment_test_kernel(
0112     typename detector_host_t::view_type det_data_static,
0113     typename detector_host_t::view_type det_data_aligned,
0114     vecmem::data::vector_view<transform_t> surfacexf_data_static,
0115     vecmem::data::vector_view<transform_t> surfacexf_data_aligned) {
0116   auto ctx = typename detector_host_t::geometry_context{};
0117
0118   // two instances of device detectors
0119   detector_device_t det_device_static(det_data_static);
0120   detector_device_t det_device_aligned(det_data_aligned);
0121
0122   // device vectors of surface transforms
0123   vecmem::device_vector<transform_t> surfacexf_device_static(
0124       surfacexf_data_static);
0125   vecmem::device_vector<transform_t> surfacexf_device_aligned(
0126       surfacexf_data_aligned);
0127
0128   // copy surface transforms into relevant vectors
0129   for (unsigned int i = 0u; i < det_device_static.surfaces().size(); i++) {
0130     const auto sf =
0131         tracking_surface{det_device_static, det_device_static.surfaces()[i]};
0132     surfacexf_device_static[i] = sf.transform(ctx);
0133   }
0134
0135   for (unsigned int i = 0u; i < det_device_aligned.surfaces().size(); i++) {
0136     const auto sf =
0137         tracking_surface{det_device_aligned, det_device_aligned.surfaces()[i]};
0138     surfacexf_device_aligned[i] = sf.transform(ctx);
0139   }
0140 }
0141
0142 /// implementation of the alignment test function for detector
0143 void detector_alignment_test(
0144     typename detector_host_t::view_type det_data_static,
0145     typename detector_host_t::view_type det_data_aligned,
0146     vecmem::data::vector_view<transform_t> surfacexf_data_static,
0147     vecmem::data::vector_view<transform_t> surfacexf_data_aligned) {
0148   constexpr int block_dim = 1u;
0149   constexpr int thread_dim = 1u;
0150
0151   // run the test kernel
0152   detector_alignment_test_kernel<<<block_dim, thread_dim>>>(
0153       det_data_static, det_data_aligned, surfacexf_data_static,
0154       surfacexf_data_aligned);
0155
0156   // cuda error check
0157   DETRAY_CUDA_ERROR_CHECK(cudaGetLastError());
0158   DETRAY_CUDA_ERROR_CHECK(cudaDeviceSynchronize());
0159 }
0160
0161 }  // namespace detray