Warning, /acts/Detray/tests/unit_tests/device/hip/detector_hip_kernel.hip is written in an unsupported language. File is not indexed.
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/hip_definitions.hpp"
0010 #include "detray/geometry/tracking_surface.hpp"
0011
0012 // Detray test include(s)
0013 #include "detector_hip_kernel.hpp"
0014
0015 namespace detray {
0016
0017 __global__ void detector_test_kernel(
0018
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
0027 // convert toy detector_data into detector w/ device vectors
0028 detector_device_t det_device(det_data);
0029
0030 // convert subdetector data objects into objects w/ device vectors
0031 vecmem::device_vector<det_volume_t> volumes_device(volumes_data);
0032 vecmem::device_vector<det_surface_t> surfaces_device(surfaces_data);
0033 vecmem::device_vector<transform_t> transforms_device(transforms_data);
0034 vecmem::device_vector<rectangle_t> rectangles_device(rectangles_data);
0035 vecmem::device_vector<disc_t> discs_device(discs_data);
0036 vecmem::device_vector<cylinder_t> cylinders_device(cylinders_data);
0037
0038 // copy objects - volume
0039 for (unsigned int i = 0u; i < det_device.volumes().size(); i++) {
0040 volumes_device[i] = det_device.volumes()[i];
0041 }
0042
0043 // copy objects - surfaces
0044 for (unsigned int i = 0u; i < det_device.surfaces().size(); i++) {
0045 surfaces_device[i] = det_device.surfaces()[i];
0046 }
0047
0048 // copy objects - transforms
0049 auto& trfs = det_device.transform_store();
0050 auto ctx = typename detector_host_t::geometry_context{};
0051 for (unsigned int i = 0u; i < trfs.size(ctx); i++) {
0052 transforms_device[i] = trfs.at(i, ctx);
0053 }
0054
0055 // copy objects - masks
0056 auto& masks = det_device.mask_store();
0057 auto& rectangles =
0058 masks.template get<detector_host_t::masks::id::e_rectangle2D>();
0059 for (unsigned int i = 0u; i < rectangles.size(); i++) {
0060 rectangles_device[i] = rectangles[i];
0061 }
0062
0063 auto& discs = masks.template get<detector_host_t::masks::id::e_ring2D>();
0064 for (unsigned int i = 0u; i < discs.size(); i++) {
0065 discs_device[i] = discs[i];
0066 }
0067
0068 auto& cylinders = masks.template get<
0069 detector_host_t::masks::id::e_concentric_cylinder2D>();
0070 for (unsigned int i = 0u; i < cylinders.size(); i++) {
0071 cylinders_device[i] = cylinders[i];
0072 }
0073 }
0074
0075 /// implementation of the test function for detector
0076 void detector_test(typename detector_host_t::view_type det_data,
0077 vecmem::data::vector_view<det_volume_t> volumes_data,
0078 vecmem::data::vector_view<det_surface_t> surfaces_data,
0079 vecmem::data::vector_view<transform_t> transforms_data,
0080 vecmem::data::vector_view<rectangle_t> rectangles_data,
0081 vecmem::data::vector_view<disc_t> discs_data,
0082 vecmem::data::vector_view<cylinder_t> cylinders_data) {
0083
0084 constexpr int block_dim = 1u;
0085 constexpr int thread_dim = 1u;
0086 // run the test kernel
0087 hipLaunchKernelGGL(detector_test_kernel, dim3(block_dim), dim3(thread_dim),
0088 0, 0, det_data, volumes_data, surfaces_data,
0089 transforms_data, rectangles_data, discs_data,
0090 cylinders_data);
0091 // HIP error check
0092 DETRAY_HIP_ERROR_CHECK(hipGetLastError());
0093 DETRAY_HIP_ERROR_CHECK(hipDeviceSynchronize());
0094 }
0095 } // namespace detray
