From 5bc489425d413d21014d05bacdc32ae00a49b97d Mon Sep 17 00:00:00 2001
From: Da Yu Tou <dtou@n4050101.lbdaq.cern.ch>
Date: Mon, 8 Jan 2024 13:43:42 +0100
Subject: [PATCH 1/5] Fixed redundant copy ut track loops and removed unused
variables in UT search window.
---
device/UT/compassUT/src/SearchWindows.cu | 5 ++---
device/UT/consolidate/src/UTCopyTrackHitNumber.cu | 2 +-
2 files changed, 3 insertions(+), 4 deletions(-)
diff --git a/device/UT/compassUT/src/SearchWindows.cu b/device/UT/compassUT/src/SearchWindows.cu
index 44087e2fd23..8e94b22a8f0 100644
--- a/device/UT/compassUT/src/SearchWindows.cu
+++ b/device/UT/compassUT/src/SearchWindows.cu
@@ -176,8 +176,7 @@ __device__ std::tuple<int, int, int, int, int, int, int, int, int, int> calculat
if (3 + 4 * index >= UTMagnetTool::N_dxLay_vals) {
index = (UTMagnetTool::N_dxLay_vals - 3) / 4 - 1;
}
- const float normFact[4] {
- fudge_factors[4 * index], fudge_factors[1 + 4 * index], fudge_factors[2 + 4 * index], fudge_factors[3 + 4 * index]};
+ const float normFact = fudge_factors[4 * index + layer];
// -- this 500 seems a little odd...
// to do: change back!
@@ -190,7 +189,7 @@ __device__ std::tuple<int, int, int, int, int, int, int, int, int, int> calculat
const float z_at_layer = ut_hits.zAtYEq0(layer_offset);
const float y_track = velo_state.y + velo_state.ty * (z_at_layer - velo_state.z);
const float x_track = velo_state.x + velo_state.tx * (z_at_layer - velo_state.z);
- const float invNormFact = 1.0f / normFact[layer];
+ const float invNormFact = 1.0f / normFact;
const float xTolNormFact = xTol * invNormFact;
// Find sector group for lowerBoundX and upperBoundX
diff --git a/device/UT/consolidate/src/UTCopyTrackHitNumber.cu b/device/UT/consolidate/src/UTCopyTrackHitNumber.cu
index eeca0fe7fa8..6eb42239f0b 100644
--- a/device/UT/consolidate/src/UTCopyTrackHitNumber.cu
+++ b/device/UT/consolidate/src/UTCopyTrackHitNumber.cu
@@ -37,7 +37,7 @@ __global__ void ut_copy_track_hit_number::ut_copy_track_hit_number(ut_copy_track
unsigned* ut_track_hit_number = parameters.dev_ut_track_hit_number + accumulated_tracks;
// Loop over tracks.
- for (unsigned element = threadIdx.x; element < number_of_tracks; ++element) {
+ for (unsigned element = threadIdx.x; element < number_of_tracks; element += blockDim.x) {
ut_track_hit_number[element] = event_tracks[element].hits_num;
}
}
--
GitLab
From be1e303941a1c03b7cc63f2ae50a35b046e2586e Mon Sep 17 00:00:00 2001
From: dtou <toudayu@gmail.com>
Date: Wed, 10 Jan 2024 12:13:52 +0800
Subject: [PATCH 2/5] Updated fudge index logic.
---
device/UT/compassUT/src/SearchWindows.cu | 12 +++++-------
1 file changed, 5 insertions(+), 7 deletions(-)
diff --git a/device/UT/compassUT/src/SearchWindows.cu b/device/UT/compassUT/src/SearchWindows.cu
index 8e94b22a8f0..0b6838e4a83 100644
--- a/device/UT/compassUT/src/SearchWindows.cu
+++ b/device/UT/compassUT/src/SearchWindows.cu
@@ -167,16 +167,14 @@ __device__ std::tuple<int, int, int, int, int, int, int, int, int, int> calculat
const float min_pt,
const float min_momentum)
{
- // TODO: Understand and fix this logic
+ // TODO: Understand and fix this logic (especially ty_index)
// -- This is hardcoded, so faster
// -- If you ever change the Table in the magnet tool, this will be wrong
const float absSlopeY = fabsf(velo_state.ty);
- int index = static_cast<int>(absSlopeY * 100 + 0.5f);
- // assert(3 + 4 * index < UTMagnetTool::N_dxLay_vals);
- if (3 + 4 * index >= UTMagnetTool::N_dxLay_vals) {
- index = (UTMagnetTool::N_dxLay_vals - 3) / 4 - 1;
- }
- const float normFact = fudge_factors[4 * index + layer];
+ const int ty_index = static_cast<int>(absSlopeY * 100 + 0.5f);
+ const int fudge_index = min( UT::Constants::n_layers * index + layer,
+ UTMagnetTool::N_dxLay_vals - UT::Constants::n_layers + layer );
+ const float normFact = fudge_factors[fudge_index];
// -- this 500 seems a little odd...
// to do: change back!
--
GitLab
From 8b17cd87ca4910688a5184d2cdfc414d0a031416 Mon Sep 17 00:00:00 2001
From: Gitlab CI <noreply@cern.ch>
Date: Wed, 10 Jan 2024 04:14:46 +0000
Subject: [PATCH 3/5] Fixed formatting
patch generated by https://gitlab.cern.ch/lhcb/Allen/-/jobs/35141302
---
device/UT/compassUT/src/SearchWindows.cu | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/device/UT/compassUT/src/SearchWindows.cu b/device/UT/compassUT/src/SearchWindows.cu
index 0b6838e4a83..a2a6f64f84c 100644
--- a/device/UT/compassUT/src/SearchWindows.cu
+++ b/device/UT/compassUT/src/SearchWindows.cu
@@ -172,8 +172,8 @@ __device__ std::tuple<int, int, int, int, int, int, int, int, int, int> calculat
// -- If you ever change the Table in the magnet tool, this will be wrong
const float absSlopeY = fabsf(velo_state.ty);
const int ty_index = static_cast<int>(absSlopeY * 100 + 0.5f);
- const int fudge_index = min( UT::Constants::n_layers * index + layer,
- UTMagnetTool::N_dxLay_vals - UT::Constants::n_layers + layer );
+ const int fudge_index =
+ min(UT::Constants::n_layers * index + layer, UTMagnetTool::N_dxLay_vals - UT::Constants::n_layers + layer);
const float normFact = fudge_factors[fudge_index];
// -- this 500 seems a little odd...
--
GitLab
From e4cd5a480260f09e58918f88e4abb8ca5095455c Mon Sep 17 00:00:00 2001
From: dtou <toudayu@gmail.com>
Date: Wed, 10 Jan 2024 12:28:19 +0800
Subject: [PATCH 4/5] Updated UT Magnet fudge factor index logic.
---
device/UT/compassUT/src/SearchWindows.cu | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/device/UT/compassUT/src/SearchWindows.cu b/device/UT/compassUT/src/SearchWindows.cu
index a2a6f64f84c..98fbaee1cc9 100644
--- a/device/UT/compassUT/src/SearchWindows.cu
+++ b/device/UT/compassUT/src/SearchWindows.cu
@@ -173,7 +173,7 @@ __device__ std::tuple<int, int, int, int, int, int, int, int, int, int> calculat
const float absSlopeY = fabsf(velo_state.ty);
const int ty_index = static_cast<int>(absSlopeY * 100 + 0.5f);
const int fudge_index =
- min(UT::Constants::n_layers * index + layer, UTMagnetTool::N_dxLay_vals - UT::Constants::n_layers + layer);
+ min(UT::Constants::n_layers * ty_index + layer, UTMagnetTool::N_dxLay_vals - UT::Constants::n_layers + layer);
const float normFact = fudge_factors[fudge_index];
// -- this 500 seems a little odd...
--
GitLab
From 622888961238f388e67a0f67e1ae0d38010fa77d Mon Sep 17 00:00:00 2001
From: dtou <toudayu@gmail.com>
Date: Wed, 10 Jan 2024 13:45:15 +0800
Subject: [PATCH 5/5] Added TODO comments on ty->UT magnet table index
conversion.
---
device/UT/compassUT/src/SearchWindows.cu | 1 +
1 file changed, 1 insertion(+)
diff --git a/device/UT/compassUT/src/SearchWindows.cu b/device/UT/compassUT/src/SearchWindows.cu
index 98fbaee1cc9..4ecfcfb6f5a 100644
--- a/device/UT/compassUT/src/SearchWindows.cu
+++ b/device/UT/compassUT/src/SearchWindows.cu
@@ -170,6 +170,7 @@ __device__ std::tuple<int, int, int, int, int, int, int, int, int, int> calculat
// TODO: Understand and fix this logic (especially ty_index)
// -- This is hardcoded, so faster
// -- If you ever change the Table in the magnet tool, this will be wrong
+ // -- Need to understand where 100 and 0.5 comes from when converting ty->ty_index
const float absSlopeY = fabsf(velo_state.ty);
const int ty_index = static_cast<int>(absSlopeY * 100 + 0.5f);
const int fudge_index =
--
GitLab