Merge pull request #568 from niermann999/fix-tutorial

fix: warnings in propagation tutorial and switch tutorial to inhomogeneous field

tutorials/CMakeLists.txt

@@ -6,6 +6,7 @@
 
 # Set the common C++ flags.
 include( detray-compiler-options-cpp )
+include_directories( SYSTEM $<TARGET_PROPERTY:covfie::core,INTERFACE_INCLUDE_DIRECTORIES> )
 
 # Set up a tutorial library, which showcases how to use detray components.
 file( GLOB _detray_tutorial_headers

tutorials/src/device/cuda/CMakeLists.txt

@@ -14,18 +14,18 @@ enable_language(CUDA)
 include(detray-compiler-options-cuda)
 
 # Examples using the array algebra plugin.
-#detray_add_tutorial( detector_construction_cuda
-#   "detector_construction.cpp"
-#   "detector_construction.hpp"
-#   "detector_construction.cu"
-#   LINK_LIBRARIES vecmem::cuda detray::core_array
-#                  detray::tutorial detray::utils)
-#target_compile_definitions(detray_tutorial_detector_construction_cuda PRIVATE array=array)
+detray_add_tutorial( detector_construction_cuda
+   "detector_construction.cpp"
+   "detector_construction.hpp"
+   "detector_construction.cu"
+   LINK_LIBRARIES vecmem::cuda detray::core_array detray::tutorial
+                  detray::utils)
+target_compile_definitions(detray_tutorial_detector_construction_cuda PRIVATE array=array)
 
-#detray_add_tutorial( propagation_cuda
-#   "propagation.cpp"
-#   "propagation.hpp"
-#   "propagation_kernel.cu"
-#   LINK_LIBRARIES vecmem::cuda detray::core_array covfie::core covfie::cuda
-#                  detray::tutorial detray::utils)
-#target_compile_definitions(detray_tutorial_propagation_cuda PRIVATE array=array)
+detray_add_tutorial( propagation_cuda
+   "propagation.cpp"
+   "propagation.hpp"
+   "propagation_kernel.cu"
+   LINK_LIBRARIES vecmem::cuda detray::core_array covfie::core covfie::cuda
+                  detray::tutorial detray::utils)
+target_compile_definitions(detray_tutorial_propagation_cuda PRIVATE array=array)

tutorials/src/device/cuda/detector_construction.cpp

@@ -8,7 +8,6 @@
 // Project include(s)
 #include "detector_construction.hpp"
 
-#include "detray/detectors/bfield.hpp"
 #include "detray/detectors/create_toy_geometry.hpp"
 
 // Vecmem include(s)

tutorials/src/device/cuda/propagation.cpp

@@ -14,8 +14,7 @@
 // Vecmem include(s)
 #include <vecmem/memory/cuda/device_memory_resource.hpp>
 #include <vecmem/memory/cuda/managed_memory_resource.hpp>
-
-#include "vecmem/utils/cuda/copy.hpp"
+#include <vecmem/utils/cuda/copy.hpp>
 
 // System
 
@@ -24,13 +23,8 @@ int main() {
     // VecMem memory resource(s)
     vecmem::cuda::managed_memory_resource mng_mr;
 
-    // Create the bfield
-    const auto B =
-        detray::tutorial::vector3{0. * detray::unit<detray::scalar>::T,
-                                  0. * detray::unit<detray::scalar>::T,
-                                  2. * detray::unit<detray::scalar>::T};
-
-    detray::tutorial::field_t bfield = detray::bfield::create_const_field(B);
+    // Create the host bfield
+    auto bfield = detray::bfield::create_inhom_field();
 
     // Create the toy geometry
     auto [det, names] = detray::create_toy_geometry(mng_mr);
@@ -60,6 +54,8 @@ int main() {
 
     // Get data for device
     auto det_data = detray::get_data(det);
+    covfie::field<detray::tutorial::bfield::cuda::inhom_bknd_t> device_bfield(
+        bfield);
     auto tracks_data = detray::get_data(tracks);
 
     // Create navigator candidates buffer
@@ -69,6 +65,6 @@ int main() {
     copy.setup(candidates_buffer);
 
     // Run the propagator test for GPU device
-    detray::tutorial::propagation(det_data, bfield, tracks_data,
+    detray::tutorial::propagation(det_data, device_bfield, tracks_data,
                                   candidates_buffer);
 }

tutorials/src/device/cuda/propagation.hpp

@@ -23,8 +23,22 @@
 #include "detray/tracks/tracks.hpp"
 #include "detray/tutorial/types.hpp"
 
+// Covfie include(s)
+#include <covfie/cuda/backend/primitive/cuda_device_array.hpp>
+
 namespace detray::tutorial {
 
+namespace bfield::cuda {
+
+// Inhomogeneous field (cuda)
+using inhom_bknd_t = covfie::backend::affine<
+    covfie::backend::nearest_neighbour<covfie::backend::strided<
+        covfie::vector::ulong3,
+        covfie::backend::cuda_device_array<
+            covfie::vector::vector_d<detray::scalar, 3>>>>>;
+
+}  // namespace bfield::cuda
+
 // Detector
 using detector_host_t = detector<detray::toy_metadata, host_container_types>;
 using detector_device_t =
@@ -35,8 +49,11 @@ using navigator_t = navigator<detector_device_t>;
 using intersection_t = navigator_t::intersection_type;
 
 // Stepper
-using field_t = detray::bfield::const_field_t;
-using stepper_t = rk_stepper<field_t::view_t, detray::tutorial::transform3>;
+using host_field_t = covfie::field<detray::bfield::inhom_bknd_t>;
+using device_field_t =
+    covfie::field<detray::tutorial::bfield::cuda::inhom_bknd_t>;
+using stepper_t =
+    rk_stepper<device_field_t::view_t, detray::tutorial::transform3>;
 
 // Actors
 using actor_chain_t =
@@ -51,7 +68,7 @@ using propagator_t = propagator<stepper_t, navigator_t, actor_chain_t>;
 /// Propagation tutorial function
 void propagation(
     typename detector_host_t::view_type det_data,
-    typename field_t::view_t field_data,
+    typename device_field_t::view_t field_data,
     const vecmem::data::vector_view<
         free_track_parameters<detray::tutorial::transform3>>
         tracks_data,

tutorials/src/device/cuda/propagation_kernel.cu

@@ -17,7 +17,7 @@ inline constexpr detray::scalar path_limit{2.f *
 /// Kernel that runs the entire propagation loop
 __global__ void propagation_kernel(
     typename detray::tutorial::detector_host_t::view_type det_data,
-    typename detray::tutorial::field_t::view_t field_data,
+    typename detray::tutorial::device_field_t::view_t field_data,
     const vecmem::data::vector_view<
         detray::free_track_parameters<detray::tutorial::transform3>>
         tracks_data,
@@ -67,7 +67,7 @@ __global__ void propagation_kernel(
 
 void propagation(
     typename detray::tutorial::detector_host_t::view_type det_data,
-    typename detray::tutorial::field_t::view_t field_data,
+    typename detray::tutorial::device_field_t::view_t field_data,
     const vecmem::data::vector_view<
         detray::free_track_parameters<detray::tutorial::transform3>>
         tracks_data,

utils/include/detray/detectors/bfield.hpp

@@ -22,16 +22,16 @@ namespace detray::bfield {
 
 /// Constant bfield (host and device)
 using const_bknd_t =
-    covfie::backend::constant<covfie::vector::vector_d<scalar, 3>,
-                              covfie::vector::vector_d<scalar, 3>>;
+    covfie::backend::constant<covfie::vector::vector_d<detray::scalar, 3>,
+                              covfie::vector::vector_d<detray::scalar, 3>>;
 
 using const_field_t = covfie::field<const_bknd_t>;
 
 /// Inhomogeneous field (host)
 using inhom_bknd_t = covfie::backend::affine<
     covfie::backend::nearest_neighbour<covfie::backend::strided<
         covfie::vector::ulong3,
-        covfie::backend::array<covfie::vector::vector_d<scalar, 3>>>>>;
+        covfie::backend::array<covfie::vector::vector_d<detray::scalar, 3>>>>>;
 
 using inhom_field_t = covfie::field<inhom_bknd_t>;