diff --git a/tutorials/CMakeLists.txt b/tutorials/CMakeLists.txt index 83731b4d4..b714631c3 100644 --- a/tutorials/CMakeLists.txt +++ b/tutorials/CMakeLists.txt @@ -6,6 +6,7 @@ # Set the common C++ flags. include( detray-compiler-options-cpp ) +include_directories( SYSTEM $ ) # Set up a tutorial library, which showcases how to use detray components. file( GLOB _detray_tutorial_headers diff --git a/tutorials/src/device/cuda/CMakeLists.txt b/tutorials/src/device/cuda/CMakeLists.txt index 91849c0b1..08b787d64 100644 --- a/tutorials/src/device/cuda/CMakeLists.txt +++ b/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) diff --git a/tutorials/src/device/cuda/detector_construction.cpp b/tutorials/src/device/cuda/detector_construction.cpp index 23324e1bc..af853dae7 100644 --- a/tutorials/src/device/cuda/detector_construction.cpp +++ b/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) diff --git a/tutorials/src/device/cuda/propagation.cpp b/tutorials/src/device/cuda/propagation.cpp index 5fa882c06..35e23ceb3 100644 --- a/tutorials/src/device/cuda/propagation.cpp +++ b/tutorials/src/device/cuda/propagation.cpp @@ -14,8 +14,7 @@ // Vecmem include(s) #include #include - -#include "vecmem/utils/cuda/copy.hpp" +#include // 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::T, - 0. * detray::unit::T, - 2. * detray::unit::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 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); } diff --git a/tutorials/src/device/cuda/propagation.hpp b/tutorials/src/device/cuda/propagation.hpp index 6c8e1c07f..bced2dea8 100644 --- a/tutorials/src/device/cuda/propagation.hpp +++ b/tutorials/src/device/cuda/propagation.hpp @@ -23,8 +23,22 @@ #include "detray/tracks/tracks.hpp" #include "detray/tutorial/types.hpp" +// Covfie include(s) +#include + namespace detray::tutorial { +namespace bfield::cuda { + +// Inhomogeneous field (cuda) +using inhom_bknd_t = covfie::backend::affine< + covfie::backend::nearest_neighbour>>>>; + +} // namespace bfield::cuda + // Detector using detector_host_t = detector; using detector_device_t = @@ -35,8 +49,11 @@ using navigator_t = navigator; using intersection_t = navigator_t::intersection_type; // Stepper -using field_t = detray::bfield::const_field_t; -using stepper_t = rk_stepper; +using host_field_t = covfie::field; +using device_field_t = + covfie::field; +using stepper_t = + rk_stepper; // Actors using actor_chain_t = @@ -51,7 +68,7 @@ using propagator_t = propagator; /// 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> tracks_data, diff --git a/tutorials/src/device/cuda/propagation_kernel.cu b/tutorials/src/device/cuda/propagation_kernel.cu index d832925c1..d5dbafe53 100644 --- a/tutorials/src/device/cuda/propagation_kernel.cu +++ b/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> 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> tracks_data, diff --git a/utils/include/detray/detectors/bfield.hpp b/utils/include/detray/detectors/bfield.hpp index 0216913eb..2415226cd 100644 --- a/utils/include/detray/detectors/bfield.hpp +++ b/utils/include/detray/detectors/bfield.hpp @@ -22,8 +22,8 @@ namespace detray::bfield { /// Constant bfield (host and device) using const_bknd_t = - covfie::backend::constant, - covfie::vector::vector_d>; + covfie::backend::constant, + covfie::vector::vector_d>; using const_field_t = covfie::field; @@ -31,7 +31,7 @@ using const_field_t = covfie::field; using inhom_bknd_t = covfie::backend::affine< covfie::backend::nearest_neighbour>>>>; + covfie::backend::array>>>>; using inhom_field_t = covfie::field;