From 4efb7495e3be6721a61effe57ce8314fba9a9305 Mon Sep 17 00:00:00 2001
From: Attila Krasznahorkay <Attila.Krasznahorkay@cern.ch>
Date: Wed, 15 May 2024 12:39:01 +0200
Subject: [PATCH] Switched all device code to using std::fabs / sycl::fabs.
This is needed for oneAPI 2024.1.0 compatibility, as it now
envforces the SYCL standard that sycl::abs can only be used
on integer values. (Though this update could help CUDA code
a little as well...)
---
core/CMakeLists.txt | 1 +
core/include/traccc/definitions/math.hpp | 27 +++++++++++++++++++
core/include/traccc/edm/cell.hpp | 5 ++--
core/include/traccc/edm/measurement.hpp | 13 ++++-----
core/include/traccc/edm/spacepoint.hpp | 11 ++++----
.../traccc/seeding/doublet_finding_helper.hpp | 8 +++---
.../seeding/spacepoint_binning_helper.hpp | 5 ++--
.../track_params_estimation_helper.hpp | 8 +++---
.../traccc/seeding/triplet_finding_helper.hpp | 8 +++---
.../finding/device/impl/apply_interaction.ipp | 5 +++-
.../device/impl/update_triplet_weights.ipp | 13 +++++----
11 files changed, 74 insertions(+), 30 deletions(-)
create mode 100644 core/include/traccc/definitions/math.hpp
diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt
index 881ae11726..b188c37555 100644
--- a/core/CMakeLists.txt
+++ b/core/CMakeLists.txt
@@ -11,6 +11,7 @@ include( traccc-compiler-options-cpp )
traccc_add_library( traccc_core core TYPE SHARED
# Common definitions.
"include/traccc/definitions/track_parametrization.hpp"
+ "include/traccc/definitions/math.hpp"
"include/traccc/definitions/primitives.hpp"
"include/traccc/definitions/common.hpp"
"include/traccc/definitions/qualifiers.hpp"
diff --git a/core/include/traccc/definitions/math.hpp b/core/include/traccc/definitions/math.hpp
new file mode 100644
index 0000000000..718ece9ae4
--- /dev/null
+++ b/core/include/traccc/definitions/math.hpp
@@ -0,0 +1,27 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2024 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+// SYCL include(s).
+#if defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
+#include <CL/sycl.hpp>
+#endif
+
+// System include(s).
+#include <cmath>
+
+namespace traccc {
+
+/// Namespace to pick up math functions from
+#if defined(CL_SYCL_LANGUAGE_VERSION) || defined(SYCL_LANGUAGE_VERSION)
+namespace math = cl::sycl;
+#else
+namespace math = std;
+#endif // SYCL
+
+} // namespace traccc
diff --git a/core/include/traccc/edm/cell.hpp b/core/include/traccc/edm/cell.hpp
index 75d89c27e0..307125352c 100644
--- a/core/include/traccc/edm/cell.hpp
+++ b/core/include/traccc/edm/cell.hpp
@@ -9,6 +9,7 @@
// traccc include(s).
#include "traccc/definitions/common.hpp"
+#include "traccc/definitions/math.hpp"
#include "traccc/definitions/primitives.hpp"
#include "traccc/edm/container.hpp"
#include "traccc/geometry/pixel_data.hpp"
@@ -88,8 +89,8 @@ inline bool operator==(const cell& lhs, const cell& rhs) {
return ((lhs.module_link == rhs.module_link) &&
(lhs.channel0 == rhs.channel0) && (lhs.channel1 == rhs.channel1) &&
- (std::abs(lhs.activation - rhs.activation) < float_epsilon) &&
- (std::abs(lhs.time - rhs.time) < float_epsilon));
+ (math::fabs(lhs.activation - rhs.activation) < float_epsilon) &&
+ (math::fabs(lhs.time - rhs.time) < float_epsilon));
}
} // namespace traccc
diff --git a/core/include/traccc/edm/measurement.hpp b/core/include/traccc/edm/measurement.hpp
index 7005d8a194..e75ef7ae99 100644
--- a/core/include/traccc/edm/measurement.hpp
+++ b/core/include/traccc/edm/measurement.hpp
@@ -1,6 +1,6 @@
/** TRACCC library, part of the ACTS project (R&D line)
*
- * (c) 2022 CERN for the benefit of the ACTS project
+ * (c) 2022-2024 CERN for the benefit of the ACTS project
*
* Mozilla Public License Version 2.0
*/
@@ -8,6 +8,7 @@
#pragma once
// Project include(s).
+#include "traccc/definitions/math.hpp"
#include "traccc/definitions/primitives.hpp"
#include "traccc/definitions/qualifiers.hpp"
#include "traccc/definitions/track_parametrization.hpp"
@@ -62,7 +63,7 @@ TRACCC_HOST_DEVICE inline bool operator<(const measurement& lhs,
if (lhs.surface_link != rhs.surface_link) {
return lhs.surface_link < rhs.surface_link;
- } else if (std::abs(lhs.local[0] - rhs.local[0]) > float_epsilon) {
+ } else if (math::fabs(lhs.local[0] - rhs.local[0]) > float_epsilon) {
return (lhs.local[0] < rhs.local[0]);
} else {
return (lhs.local[1] < rhs.local[1]);
@@ -74,10 +75,10 @@ TRACCC_HOST_DEVICE
inline bool operator==(const measurement& lhs, const measurement& rhs) {
return ((lhs.surface_link == rhs.surface_link) &&
- (std::abs(lhs.local[0] - rhs.local[0]) < float_epsilon) &&
- (std::abs(lhs.local[1] - rhs.local[1]) < float_epsilon) &&
- (std::abs(lhs.variance[0] - rhs.variance[0]) < float_epsilon) &&
- (std::abs(lhs.variance[1] - rhs.variance[1]) < float_epsilon));
+ (math::fabs(lhs.local[0] - rhs.local[0]) < float_epsilon) &&
+ (math::fabs(lhs.local[1] - rhs.local[1]) < float_epsilon) &&
+ (math::fabs(lhs.variance[0] - rhs.variance[0]) < float_epsilon) &&
+ (math::fabs(lhs.variance[1] - rhs.variance[1]) < float_epsilon));
}
/// Comparator based on detray barcode value
diff --git a/core/include/traccc/edm/spacepoint.hpp b/core/include/traccc/edm/spacepoint.hpp
index c70493cf84..4814629248 100644
--- a/core/include/traccc/edm/spacepoint.hpp
+++ b/core/include/traccc/edm/spacepoint.hpp
@@ -9,6 +9,7 @@
// Project include(s).
#include "traccc/definitions/common.hpp"
+#include "traccc/definitions/math.hpp"
#include "traccc/definitions/primitives.hpp"
#include "traccc/definitions/qualifiers.hpp"
#include "traccc/edm/container.hpp"
@@ -47,9 +48,9 @@ struct spacepoint {
TRACCC_HOST_DEVICE
inline bool operator<(const spacepoint& lhs, const spacepoint& rhs) {
- if (std::abs(lhs.x() - rhs.x()) > float_epsilon) {
+ if (math::fabs(lhs.x() - rhs.x()) > float_epsilon) {
return (lhs.x() < rhs.x());
- } else if (std::abs(lhs.y() - rhs.y()) > float_epsilon) {
+ } else if (math::fabs(lhs.y() - rhs.y()) > float_epsilon) {
return (lhs.y() < rhs.y());
} else {
return (lhs.z() < rhs.z());
@@ -60,9 +61,9 @@ inline bool operator<(const spacepoint& lhs, const spacepoint& rhs) {
TRACCC_HOST_DEVICE
inline bool operator==(const spacepoint& lhs, const spacepoint& rhs) {
- return ((std::abs(lhs.x() - rhs.x()) < float_epsilon) &&
- (std::abs(lhs.y() - rhs.y()) < float_epsilon) &&
- (std::abs(lhs.z() - rhs.z()) < float_epsilon) &&
+ return ((math::fabs(lhs.x() - rhs.x()) < float_epsilon) &&
+ (math::fabs(lhs.y() - rhs.y()) < float_epsilon) &&
+ (math::fabs(lhs.z() - rhs.z()) < float_epsilon) &&
(lhs.meas == rhs.meas));
}
diff --git a/core/include/traccc/seeding/doublet_finding_helper.hpp b/core/include/traccc/seeding/doublet_finding_helper.hpp
index 427c7e0bb6..224b13da03 100644
--- a/core/include/traccc/seeding/doublet_finding_helper.hpp
+++ b/core/include/traccc/seeding/doublet_finding_helper.hpp
@@ -1,12 +1,14 @@
/** TRACCC library, part of the ACTS project (R&D line)
*
- * (c) 2021-2022 CERN for the benefit of the ACTS project
+ * (c) 2021-2024 CERN for the benefit of the ACTS project
*
* Mozilla Public License Version 2.0
*/
#pragma once
+// Local include(s).
+#include "traccc/definitions/math.hpp"
#include "traccc/edm/internal_spacepoint.hpp"
#include "traccc/seeding/detail/doublet.hpp"
#include "traccc/seeding/detail/lin_circle.hpp"
@@ -63,7 +65,7 @@ doublet_finding_helper::isCompatible(const internal_spacepoint<spacepoint>& sp1,
// actually zOrigin * deltaR to avoid division by 0 statements
scalar zOrigin = sp1.z() * deltaR - sp1.radius() * cotTheta;
if (deltaR > config.deltaRMax || deltaR < config.deltaRMin ||
- std::fabs(cotTheta) > config.cotThetaMax * deltaR ||
+ math::fabs(cotTheta) > config.cotThetaMax * deltaR ||
zOrigin < config.collisionRegionMin * deltaR ||
zOrigin > config.collisionRegionMax * deltaR) {
return false;
@@ -76,7 +78,7 @@ doublet_finding_helper::isCompatible(const internal_spacepoint<spacepoint>& sp1,
// actually zOrigin * deltaR to avoid division by 0 statements
scalar zOrigin = sp1.z() * deltaR - sp1.radius() * cotTheta;
if (deltaR > config.deltaRMax || deltaR < config.deltaRMin ||
- std::fabs(cotTheta) > config.cotThetaMax * deltaR ||
+ math::fabs(cotTheta) > config.cotThetaMax * deltaR ||
zOrigin < config.collisionRegionMin * deltaR ||
zOrigin > config.collisionRegionMax * deltaR) {
return false;
diff --git a/core/include/traccc/seeding/spacepoint_binning_helper.hpp b/core/include/traccc/seeding/spacepoint_binning_helper.hpp
index 81698fdeae..c027a34df0 100644
--- a/core/include/traccc/seeding/spacepoint_binning_helper.hpp
+++ b/core/include/traccc/seeding/spacepoint_binning_helper.hpp
@@ -8,6 +8,7 @@
#pragma once
// Project include(s).
+#include "traccc/definitions/math.hpp"
#include "traccc/edm/internal_spacepoint.hpp"
#include "traccc/edm/spacepoint.hpp"
#include "traccc/seeding/detail/seeding_config.hpp"
@@ -57,8 +58,8 @@ inline std::pair<detray::axis2::circular<>, detray::axis2::regular<>> get_axes(
// evaluating the azimutal deflection including the maximum impact
// parameter
scalar deltaAngleWithMaxD0 =
- std::abs(std::asin(grid_config.impactMax / (rMin)) -
- std::asin(grid_config.impactMax / grid_config.rMax));
+ math::fabs(std::asin(grid_config.impactMax / (rMin)) -
+ std::asin(grid_config.impactMax / grid_config.rMax));
// evaluating delta Phi based on the inner and outer angle, and the
// azimutal deflection including the maximum impact parameter Divide by
diff --git a/core/include/traccc/seeding/track_params_estimation_helper.hpp b/core/include/traccc/seeding/track_params_estimation_helper.hpp
index 5e235557bd..e1f9fc2b00 100644
--- a/core/include/traccc/seeding/track_params_estimation_helper.hpp
+++ b/core/include/traccc/seeding/track_params_estimation_helper.hpp
@@ -1,6 +1,6 @@
/** TRACCC library, part of the ACTS project (R&D line)
*
- * (c) 2021-2022 CERN for the benefit of the ACTS project
+ * (c) 2021-2024 CERN for the benefit of the ACTS project
*
* Mozilla Public License Version 2.0
*/
@@ -8,6 +8,7 @@
#pragma once
// Library include(s).
+#include "traccc/definitions/math.hpp"
#include "traccc/edm/seed.hpp"
#include "traccc/edm/spacepoint.hpp"
#include "traccc/edm/track_parameters.hpp"
@@ -117,8 +118,9 @@ inline TRACCC_HOST_DEVICE bound_vector seed_to_bound_vector(
// The estimated momentum, and its projection along the magnetic
// field diretion
- scalar pInGeV = std::abs(1.0f / getter::element(params, e_bound_qoverp, 0));
- scalar pzInGeV = 1.0f / std::abs(qOverPt) * invTanTheta;
+ scalar pInGeV =
+ math::fabs(1.0f / getter::element(params, e_bound_qoverp, 0));
+ scalar pzInGeV = 1.0f / math::fabs(qOverPt) * invTanTheta;
scalar massInGeV = mass / unit<scalar>::GeV;
// The estimated velocity, and its projection along the magnetic
diff --git a/core/include/traccc/seeding/triplet_finding_helper.hpp b/core/include/traccc/seeding/triplet_finding_helper.hpp
index 2323c780cf..855a5ddcd9 100644
--- a/core/include/traccc/seeding/triplet_finding_helper.hpp
+++ b/core/include/traccc/seeding/triplet_finding_helper.hpp
@@ -1,12 +1,14 @@
/** TRACCC library, part of the ACTS project (R&D line)
*
- * (c) 2021-2022 CERN for the benefit of the ACTS project
+ * (c) 2021-2024 CERN for the benefit of the ACTS project
*
* Mozilla Public License Version 2.0
*/
#pragma once
+// Local include(s).
+#include "traccc/definitions/math.hpp"
#include "traccc/edm/internal_spacepoint.hpp"
#include "traccc/seeding/detail/doublet.hpp"
#include "traccc/seeding/detail/lin_circle.hpp"
@@ -58,7 +60,7 @@ bool TRACCC_HOST_DEVICE triplet_finding_helper::isCompatible(
// if the error is larger than the difference in theta, no need to
// compare with scattering
if (deltaCotTheta2 - error2 > 0) {
- deltaCotTheta = std::abs(deltaCotTheta);
+ deltaCotTheta = math::fabs(deltaCotTheta);
// if deltaTheta larger than the scattering for the lower pT cut, skip
error = std::sqrt(error2);
dCotThetaMinusError2 = deltaCotTheta2 + error2 -
@@ -119,7 +121,7 @@ bool TRACCC_HOST_DEVICE triplet_finding_helper::isCompatible(
// A and B allow calculation of impact params in U/V plane with linear
// function
// (in contrast to having to solve a quadratic function in x/y plane)
- impact_parameter = std::abs((A - B * spM.radius()) * spM.radius());
+ impact_parameter = math::fabs((A - B * spM.radius()) * spM.radius());
if (impact_parameter > config.impactMax) {
return false;
diff --git a/device/common/include/traccc/finding/device/impl/apply_interaction.ipp b/device/common/include/traccc/finding/device/impl/apply_interaction.ipp
index 8af1945d30..3c8c973d9c 100644
--- a/device/common/include/traccc/finding/device/impl/apply_interaction.ipp
+++ b/device/common/include/traccc/finding/device/impl/apply_interaction.ipp
@@ -7,6 +7,9 @@
#pragma once
+// Project include(s).
+#include "traccc/definitions/math.hpp"
+
// Detray include(s).
#include "detray/geometry/surface.hpp"
@@ -43,7 +46,7 @@ TRACCC_DEVICE inline void apply_interaction(
interactor_type{}.update(
bound_param, interactor_state,
static_cast<int>(detray::navigation::direction::e_forward), sf,
- std::abs(
+ math::fabs(
sf.cos_angle(ctx, bound_param.dir(), bound_param.bound_local())));
}
diff --git a/device/common/include/traccc/seeding/device/impl/update_triplet_weights.ipp b/device/common/include/traccc/seeding/device/impl/update_triplet_weights.ipp
index a55f15def6..11eab9f946 100644
--- a/device/common/include/traccc/seeding/device/impl/update_triplet_weights.ipp
+++ b/device/common/include/traccc/seeding/device/impl/update_triplet_weights.ipp
@@ -1,15 +1,18 @@
/** TRACCC library, part of the ACTS project (R&D line)
*
- * (c) 2021-2023 CERN for the benefit of the ACTS project
+ * (c) 2021-2024 CERN for the benefit of the ACTS project
*
* Mozilla Public License Version 2.0
*/
+#pragma once
+
+// Project include(s).
+#include "traccc/definitions/math.hpp"
+
// System include(s).
#include <cassert>
-#pragma once
-
namespace traccc::device {
TRACCC_HOST_DEVICE
@@ -82,7 +85,7 @@ inline void update_triplet_weights(
// compared top SP should have at least deltaRMin distance
const scalar otherTop_r = other_spT.radius();
const scalar deltaR = currentTop_r - otherTop_r;
- if (std::abs(deltaR) < filter_config.deltaRMin) {
+ if (math::fabs(deltaR) < filter_config.deltaRMin) {
continue;
}
@@ -106,7 +109,7 @@ inline void update_triplet_weights(
// compatible seed (20mm instead of 5mm) add new compatible seed
// only if distance larger than rmin to all other compatible
// seeds
- if (std::abs(previousDiameter - otherTop_r) <
+ if (math::fabs(previousDiameter - otherTop_r) <
filter_config.deltaRMin) {
newCompSeed = false;
break;