Skip to content

Commit

Permalink
ref: Use angle from surface class instead of intersector (#557)
Browse files Browse the repository at this point in the history 
While debugging the detray update I found that in the seed_generator the cosine of the incidence angle on the intersection is quite often invalid, because the mask tolerance and overstepping tolerance that are given to the intersection_update call do not match what the navigator used to find the track parameters on surface. This leads to essentially a zero path through the material (thickness gets divided by the cosine). Because putting the navigation configuration in there to set all of the tolerances is a bit awkward, I used the angle calculation of the surface directly.
  • Loading branch information
@niermann999
niermann999 authored Apr 26, 2024
1 parent 2150b3f commit 74772b9
Show file tree
Hide file tree
Showing 5 changed files with 15 additions and 70 deletions.
19 changes: 3 additions & 16 deletions core/include/traccc/finding/finding_algorithm.ipp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -11,9 +11,6 @@
// detray include(s).
#include "detray/geometry/barcode.hpp"
#include "detray/geometry/surface.hpp"
#include "detray/navigation/detail/ray.hpp"
#include "detray/navigation/intersection/ray_intersector.hpp"
#include "detray/navigation/intersection_kernel.hpp"

// System include
#include <algorithm>
Expand Down Expand Up @@ -129,27 +126,17 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
*************************/

// Get intersection at surface
const detray::surface<detector_type> sf{det,
in_param.surface_link()};
const detray::surface sf{det, in_param.surface_link()};

const cxt_t ctx{};
const auto free_vec =
sf.bound_to_free_vector(ctx, in_param.vector());
intersection_type sfi;

const auto sf_desc = det.surface(in_param.surface_link());
sfi.sf_desc = sf_desc;
sf.template visit_mask<
detray::intersection_update<detray::ray_intersector>>(
detray::detail::ray<transform3_type>(free_vec), sfi,
det.transform_store());

// Apply interactor
typename interactor_type::state interactor_state;
interactor_type{}.update(
in_param, interactor_state,
static_cast<int>(detray::navigation::direction::e_forward), sf,
sfi.cos_incidence_angle);
std::abs(
sf.cos_angle(ctx, in_param.dir(), in_param.bound_local())));

/*************************
* CKF
Expand Down
17 changes: 3 additions & 14 deletions core/include/traccc/utils/seed_generator.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -13,8 +13,6 @@
// detray include(s).
#include "detray/geometry/barcode.hpp"
#include "detray/geometry/surface.hpp"
#include "detray/navigation/intersection/ray_intersector.hpp"
#include "detray/navigation/intersection_kernel.hpp"
#include "detray/propagator/actor_chain.hpp"
#include "detray/propagator/actors/aborters.hpp"
#include "detray/propagator/actors/parameter_resetter.hpp"
Expand Down Expand Up @@ -53,7 +51,7 @@ struct seed_generator {
const free_track_parameters& free_param) {

// Get bound parameter
const detray::surface<detector_t> sf{m_detector, surface_link};
const detray::surface sf{m_detector, surface_link};

const cxt_t ctx{};
auto bound_vec = sf.free_to_bound_vector(ctx, free_param.vector());
Expand All @@ -65,26 +63,17 @@ struct seed_generator {

// Type definitions
using transform3_type = typename detector_t::transform3;
using intersection_type =
detray::intersection2D<typename detector_t::surface_type,
transform3_type>;
using interactor_type =
detray::pointwise_material_interactor<transform3_type>;

intersection_type sfi;
sfi.sf_desc = m_detector.surface(surface_link);
sf.template visit_mask<
detray::intersection_update<detray::ray_intersector>>(
detray::detail::ray<transform3_type>(free_param.vector()), sfi,
m_detector.transform_store());

// Apply interactor
typename interactor_type::state interactor_state;
interactor_state.do_multiple_scattering = false;
interactor_type{}.update(
bound_param, interactor_state,
static_cast<int>(detray::navigation::direction::e_backward), sf,
sfi.cos_incidence_angle);
std::abs(sf.cos_angle(ctx, bound_param.dir(),
bound_param.bound_local())));

for (std::size_t i = 0; i < e_bound_size; i++) {

Expand Down
6 changes: 1 addition & 5 deletions device/common/include/traccc/finding/device/apply_interaction.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -17,20 +17,16 @@ namespace traccc::device {
///
/// @param[in] globalIndex The index of the current thread
/// @param[in] det_data Detector view object
/// @param[in] nav_candidates_buffer Buffer for navigation candidates
/// @param[in] n_params The number of parameters (or tracks)
/// @param[out] params_view Collection of output bound track_parameters
///
template <typename detector_t>
TRACCC_DEVICE inline void apply_interaction(
std::size_t globalIndex, typename detector_t::view_type det_data,
vecmem::data::jagged_vector_view<detray::intersection2D<
typename detector_t::surface_type, typename detector_t::transform3>>
nav_candidates_buffer,
const int n_params,
bound_track_parameters_collection_types::view params_view);

} // namespace traccc::device

// Include the implementation.
#include "traccc/finding/device/impl/apply_interaction.ipp"
#include "traccc/finding/device/impl/apply_interaction.ipp"
30 changes: 4 additions & 26 deletions device/common/include/traccc/finding/device/impl/apply_interaction.ipp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -9,36 +9,23 @@

// Detray include(s).
#include "detray/geometry/surface.hpp"
#include "detray/navigation/intersection/intersection.hpp"
#include "detray/navigation/intersection/ray_intersector.hpp"
#include "detray/navigation/intersection_kernel.hpp"

namespace traccc::device {

template <typename detector_t>
TRACCC_DEVICE inline void apply_interaction(
std::size_t globalIndex, typename detector_t::view_type det_data,
vecmem::data::jagged_vector_view<detray::intersection2D<
typename detector_t::surface_type, typename detector_t::transform3>>
nav_candidates_buffer,
const int n_params,
bound_track_parameters_collection_types::view params_view) {

// Type definitions
using transform3_type = typename detector_t::transform3;
using intersection_type =
detray::intersection2D<typename detector_t::surface_type,
transform3_type>;
using interactor_type =
detray::pointwise_material_interactor<transform3_type>;

// Detector
detector_t det(det_data);

// Navigation candidate buffer
vecmem::jagged_device_vector<intersection_type> nav_candidates(
nav_candidates_buffer);

// in param
bound_track_parameters_collection_types::device params(params_view);

Expand All @@ -49,25 +36,16 @@ TRACCC_DEVICE inline void apply_interaction(
auto& bound_param = params.at(globalIndex);

// Get intersection at surface
const detray::surface<detector_t> sf{det, bound_param.surface_link()};
using cxt_t = typename detector_t::geometry_context;
const cxt_t ctx{};
const auto free_vec = sf.bound_to_free_vector(ctx, bound_param.vector());

const auto& mask_store = det.mask_store();
intersection_type sfi;
sfi.sf_desc = det.surface(bound_param.surface_link());
sf.template visit_mask<
detray::intersection_update<detray::ray_intersector>>(
detray::detail::ray<transform3_type>(free_vec), sfi,
det.transform_store());
const detray::surface sf{det, bound_param.surface_link()};
const typename detector_t::geometry_context ctx{};

// Apply interactor
typename interactor_type::state interactor_state;
interactor_type{}.update(
bound_param, interactor_state,
static_cast<int>(detray::navigation::direction::e_forward), sf,
sfi.cos_incidence_angle);
std::abs(
sf.cos_angle(ctx, bound_param.dir(), bound_param.bound_local())));
}

} // namespace traccc::device
13 changes: 4 additions & 9 deletions device/cuda/src/finding/finding_algorithm.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -60,17 +60,12 @@ __global__ void make_barcode_sequence(
/// CUDA kernel for running @c traccc::device::apply_interaction
template <typename detector_t>
__global__ void apply_interaction(
typename detector_t::view_type det_data,
vecmem::data::jagged_vector_view<detray::intersection2D<
typename detector_t::surface_type, typename detector_t::transform3>>
nav_candidates_buffer,
const int n_params,
typename detector_t::view_type det_data, const int n_params,
bound_track_parameters_collection_types::view params_view) {

int gid = threadIdx.x + blockIdx.x * blockDim.x;

device::apply_interaction<detector_t>(gid, det_data, nav_candidates_buffer,
n_params, params_view);
device::apply_interaction<detector_t>(gid, det_data, n_params, params_view);
}

/// CUDA kernel for running @c traccc::device::count_measurements
Expand Down Expand Up @@ -295,8 +290,8 @@ finding_algorithm<stepper_t, navigator_t>::operator()(
nThreads = m_warp_size * 2;
nBlocks = (n_in_params + nThreads - 1) / nThreads;
kernels::apply_interaction<detector_type>
<<<nBlocks, nThreads, 0, stream>>>(det_view, navigation_buffer,
n_in_params, in_params_buffer);
<<<nBlocks, nThreads, 0, stream>>>(det_view, n_in_params,
in_params_buffer);
TRACCC_CUDA_ERROR_CHECK(cudaGetLastError());

/*****************************************************************
Expand Down

0 comments on commit 74772b9

Please sign in to comment.