diff --git a/device/UT/common/include/CompassUTStructs.cuh b/device/UT/common/include/CompassUTStructs.cuh index ea9beadd80bac8414547d788b8f46592ee37e7d4..58c8fe904617d7f6fe5e16ed9dd196e94b980042 100644 --- a/device/UT/common/include/CompassUTStructs.cuh +++ b/device/UT/common/include/CompassUTStructs.cuh @@ -15,7 +15,10 @@ #include "BackendCommon.h" namespace CompassUT::Structs { - using LayerHitRanges = std::array; + + constexpr static ushort invalid_hit = std::numeric_limits::max(); + + using LayerHitRanges = std::array; struct alignas(4) Candidate { uint16_t velo_index; @@ -25,7 +28,7 @@ namespace CompassUT::Structs { struct alignas(16) VeloUTTrack { unsigned velo_index; float score; - int16_t ut_hits[4]; + uint16_t ut_hits[4]; } __attribute__((packed)); struct FitResult { diff --git a/device/UT/compass_ut/src/CompassUTDefineCandidates.cu b/device/UT/compass_ut/src/CompassUTDefineCandidates.cu index 16abc6607d7ead946c1d4b3e1d4831956dcd48f6..6869962132901deb571b2672bb4e40819d045811 100644 --- a/device/UT/compass_ut/src/CompassUTDefineCandidates.cu +++ b/device/UT/compass_ut/src/CompassUTDefineCandidates.cu @@ -95,6 +95,9 @@ __global__ void compass_ut_define_candidates::compass_ut_hit_preselection( const unsigned event_number = parameters.dev_event_list[blockIdx.x]; const unsigned number_of_events = parameters.dev_number_of_events[0]; + // Alias + constexpr auto invalid_hit = CompassUT::Structs::invalid_hit; + // UT hits const unsigned total_number_of_hits = parameters.dev_ut_hit_offsets[number_of_events * UT::Constants::n_groups]; const UT::HitOffsets ut_hit_offsets {parameters.dev_ut_hit_offsets, event_number}; @@ -169,8 +172,10 @@ __global__ void compass_ut_define_candidates::compass_ut_hit_preselection( // Find hit ranges // unsigned num_ranges = 0; - CompassUT::Structs::LayerHitRanges hit_ranges = { - short2 {-1, -1}, short2 {-1, -1}, short2 {-1, -1}, short2 {-1, -1}}; + CompassUT::Structs::LayerHitRanges hit_ranges = {ushort2 {invalid_hit, invalid_hit}, + ushort2 {invalid_hit, invalid_hit}, + ushort2 {invalid_hit, invalid_hit}, + ushort2 {invalid_hit, invalid_hit}}; unsigned num_candidates = 0; const auto fired_sectors = dev_ut_layer_geometry->find_sectors(layer, expected_layer_y - yTolLayer, expected_layer_y + yTolLayer); @@ -199,7 +204,7 @@ __global__ void compass_ut_define_candidates::compass_ut_hit_preselection( if (hit_end <= hit_start) continue; // Fill result - hit_ranges[num_ranges] = short2 {static_cast(hit_start), static_cast(hit_end)}; + hit_ranges[num_ranges] = ushort2 {static_cast(hit_start), static_cast(hit_end)}; num_ranges++; num_candidates += hit_end - hit_start; } @@ -219,6 +224,9 @@ compass_ut_define_candidates::compass_ut_fill_cadidates(Parameters parameters, f const unsigned event_number = parameters.dev_event_list[blockIdx.x]; const unsigned number_of_events = parameters.dev_number_of_events[0]; + // Alias + constexpr auto invalid_hit = CompassUT::Structs::invalid_hit; + // UT hits const unsigned total_number_of_hits = parameters.dev_ut_hit_offsets[number_of_events * UT::Constants::n_groups]; const UT::HitOffsets ut_hit_offsets {parameters.dev_ut_hit_offsets, event_number}; @@ -270,7 +278,7 @@ compass_ut_define_candidates::compass_ut_fill_cadidates(Parameters parameters, f const auto hit_start = hit_ranges[sector_idx].x; const auto hit_end = hit_ranges[sector_idx].y; - if (hit_start == -1 || hit_end == -1) break; + if (hit_start == invalid_hit || hit_end == invalid_hit) break; for (int h = hit_start; h < hit_end; h++) { const auto zAtYEq0 = ut_hits.zAtYEq0(h); diff --git a/device/UT/compass_ut/src/CompassUTFindTracks.cu b/device/UT/compass_ut/src/CompassUTFindTracks.cu index abad1a8f284d3f307a31232591486837839833e3..b800084db68ad1f38b5cb1e04417a75fe3cd6b43 100644 --- a/device/UT/compass_ut/src/CompassUTFindTracks.cu +++ b/device/UT/compass_ut/src/CompassUTFindTracks.cu @@ -60,6 +60,9 @@ __global__ void compass_ut_find_tracks::compass_ut_find_tracks( const unsigned event_number = parameters.dev_event_list[blockIdx.x]; const unsigned number_of_events = parameters.dev_number_of_events[0]; + // Alias + constexpr auto invalid_hit = CompassUT::Structs::invalid_hit; + // UT hits const unsigned total_number_of_hits = parameters.dev_ut_hit_offsets[number_of_events * UT::Constants::n_groups]; const UT::HitOffsets ut_hit_offsets {parameters.dev_ut_hit_offsets, event_number}; @@ -99,7 +102,7 @@ __global__ void compass_ut_find_tracks::compass_ut_find_tracks( ((h0_z0 - vp_state.z()) * (h0_z0 - vp_state.z())); // Find best hits - int16_t best_hits[4] = {-1, -1, -1, -1}; + uint16_t best_hits[4] = {invalid_hit, invalid_hit, invalid_hit, invalid_hit}; best_hits[layers[case_idx][0]] = candidate.hit_idx; unsigned num_hits = 1; float score = 0.f; diff --git a/device/UT/compass_ut/src/CompassUTFitTracks.cu b/device/UT/compass_ut/src/CompassUTFitTracks.cu index f9760dd5e8030ad426a7ea7ebcc8128065860272..1205a39a767182012cd3f3b28edfd786a2d5589f 100644 --- a/device/UT/compass_ut/src/CompassUTFitTracks.cu +++ b/device/UT/compass_ut/src/CompassUTFitTracks.cu @@ -73,7 +73,7 @@ namespace { const float* bdl_table, UT::ConstHits& ut_hits, const MiniState& velo_state, - const int16_t* best_hits, + const uint16_t* best_hits, const float sigma_velo_slope) { // @@ -107,7 +107,7 @@ namespace { for (unsigned i = 0; i < UT::Constants::n_layers; ++i) { const auto hit_index = best_hits[i]; - if (hit_index >= 0) { + if (hit_index != CompassUT::Structs::invalid_hit) { const float wi = ut_hits.weight(hit_index); const float ci = ut_hits.cosT(hit_index); last_z = ut_hits.zAtYEq0(hit_index); @@ -143,7 +143,7 @@ namespace { for (unsigned i = 0; i < UT::Constants::n_layers; ++i) { const auto hit_index = best_hits[i]; - if (hit_index >= 0) { + if (hit_index != CompassUT::Structs::invalid_hit) { const float zd = ut_hits.zAtYEq0(hit_index); const float xd = xUTFit + xSlopeUTFit * (zd - UT::Constants::zMidUT); // x_pos_layer @@ -207,7 +207,7 @@ namespace { const float yyProto = velo_state.y() - velo_state.ty() * velo_state.z(); for (unsigned i = 0; i < UT::Constants::n_layers; ++i) { - if (best_hits[i] != -1) { + if (best_hits[i] != CompassUT::Structs::invalid_hit) { const auto hit = best_hits[i]; const float yy = yyProto + (velo_state.ty() * ut_hits.zAtYEq0(hit)); @@ -253,7 +253,7 @@ namespace { // // float chi2 = weight * (distX * distX * distCorrectionX2 + offsetY * offsetY / (1.0f + ty * ty)); // for (unsigned i = 0; i < UT::Constants::n_layers; ++i) { - // if (best_hits[i] != -1) { + // if (best_hits[i] != CompassUT::Structs::invalid_hit) { // const auto hit = best_hits[i]; // const float w = ut_hits.weight(hit); @@ -303,7 +303,7 @@ namespace { // Fit UT hits without VELO constraints __device__ bool - free_fit(UT::ConstHits& ut_hits, const MiniState& vp_state, const int16_t* found_hits, free_state_t& out) + free_fit(UT::ConstHits& ut_hits, const MiniState& vp_state, const uint16_t* found_hits, free_state_t& out) { using UT::Constants::n_layers; using UT::Constants::zMidUT; @@ -312,7 +312,7 @@ namespace { for (unsigned i = 0; i < n_layers; i++) { const auto hit = found_hits[i]; - if (hit == -1) continue; + if (hit == CompassUT::Structs::invalid_hit) continue; const auto w = ut_hits.weight(hit); const auto zAtYEq0 = ut_hits.zAtYEq0(hit); const auto dz = zAtYEq0 - zMidUT; @@ -393,6 +393,9 @@ __global__ void compass_ut_fit_tracks::compass_ut_fit_tracks( const unsigned event_number = parameters.dev_event_list[blockIdx.x]; const unsigned number_of_events = parameters.dev_number_of_events[0]; + // Alias + constexpr auto invalid_hit = CompassUT::Structs::invalid_hit; + // UT hits const unsigned total_number_of_hits = parameters.dev_ut_hit_offsets[number_of_events * UT::Constants::n_groups]; const UT::HitOffsets ut_hit_offsets {parameters.dev_ut_hit_offsets, event_number}; @@ -428,8 +431,8 @@ __global__ void compass_ut_fit_tracks::compass_ut_fit_tracks( if ((p < min_momentum_final) || (pt < min_pt_final)) continue; // Compute number of hits - const auto nHits = (candidate.ut_hits[0] >= 0) + (candidate.ut_hits[1] >= 0) + (candidate.ut_hits[2] >= 0) + - (candidate.ut_hits[3] >= 0); + const auto nHits = (candidate.ut_hits[0] != invalid_hit) + (candidate.ut_hits[1] != invalid_hit) + + (candidate.ut_hits[2] != invalid_hit) + (candidate.ut_hits[3] != invalid_hit); // Ghost killing const auto model = (nHits == 3) ? dev_velout_ghostkiller_3hits : dev_velout_ghostkiller_4hits; diff --git a/device/UT/compass_ut/src/CompassUTSelectTracks.cu b/device/UT/compass_ut/src/CompassUTSelectTracks.cu index 62e8453245476cc8b1bd7a8be908f563f1f43623..da15e342513c75887b115a99afafa0d28d3d0b70 100644 --- a/device/UT/compass_ut/src/CompassUTSelectTracks.cu +++ b/device/UT/compass_ut/src/CompassUTSelectTracks.cu @@ -107,7 +107,7 @@ __global__ void compass_ut_select_tracks::compass_ut_kill_clone_tracks( for (unsigned i = threadIdx.x; i < input_size; i += blockDim.x) { const auto candidate = parameters.dev_ut_track_output_tracks[input_offset + i]; const auto hit = candidate.ut_hits[layer]; - if (hit >= 0) { + if (hit != CompassUT::Structs::invalid_hit) { atomicMin(best_scores + get_idx(hit - hit_layer_offset), __float_as_int(candidate.score)); } } diff --git a/device/UT/consolidate/src/ConsolidateUT.cu b/device/UT/consolidate/src/ConsolidateUT.cu index eea7aeee334fc97341e588e2969f0924e437559a..d27df0ffdcd3ef3bab03a5029e41e18e1e7067da 100644 --- a/device/UT/consolidate/src/ConsolidateUT.cu +++ b/device/UT/consolidate/src/ConsolidateUT.cu @@ -90,7 +90,7 @@ __device__ void populate(const UT::TrackHits& track, const F& assign) int hit_number = 0; for (unsigned i = 0; i < UT::Constants::n_layers; ++i) { const auto hit_index = track.hits[i]; - if (hit_index != -1) { + if (hit_index != UT::TrackHits::invalid_hit) { assign(hit_number++, hit_index); } } @@ -102,7 +102,7 @@ __device__ void populate_plane_code(const UT::TrackHits& track, const F& assign) int hit_number = 0; for (unsigned i = 0; i < UT::Constants::n_layers; ++i) { const auto hit_index = track.hits[i]; - if (hit_index != -1) { + if (hit_index != UT::TrackHits::invalid_hit) { assign(hit_number++, i); } } diff --git a/device/downstream/common/DownstreamV2Structs.cuh b/device/downstream/common/DownstreamV2Structs.cuh index 22774defb66bf6b3dc0872407b72978c4aafd5c4..a86f6e51703a97ed87b19e974803149e84541b22 100644 --- a/device/downstream/common/DownstreamV2Structs.cuh +++ b/device/downstream/common/DownstreamV2Structs.cuh @@ -13,7 +13,10 @@ #include "BackendCommon.h" namespace Downstream::Structs { - using LayerHitRanges = std::array; + + constexpr static ushort invalid_hit = std::numeric_limits::max(); + + using LayerHitRanges = std::array; struct alignas(4) HitCombinations { uint16_t layers[4]; @@ -40,7 +43,7 @@ namespace Downstream::Structs { struct alignas(4) CompactTrack { unsigned scifi_idx; - int16_t ut_hits[4]; + uint16_t ut_hits[4]; } __attribute__((packed)); struct alignas(16) CompactState { diff --git a/device/downstream/tracking/src/DownstreamV2DefineSciFiCandidates.cu b/device/downstream/tracking/src/DownstreamV2DefineSciFiCandidates.cu index a17e5938ce4e27872d6b2cd9369d449305caf8cc..114cfc4bb7bb1a1d5295cf3b7b1335c84084ba8b 100644 --- a/device/downstream/tracking/src/DownstreamV2DefineSciFiCandidates.cu +++ b/device/downstream/tracking/src/DownstreamV2DefineSciFiCandidates.cu @@ -100,6 +100,9 @@ __global__ void downstream_v2_define_scifi_candidates::downstream_v2_find_hit_ra const auto event_number = parameters.dev_event_list[blockIdx.x]; const auto number_of_events = parameters.dev_number_of_events[0]; + // Alias + constexpr auto invalid_hit = Downstream::Structs::invalid_hit; + // Offsets const auto seed_offset = parameters.dev_offsets_seeding_tracks[event_number]; const auto num_seed = parameters.dev_offsets_seeding_tracks[event_number + 1] - seed_offset; @@ -167,8 +170,10 @@ __global__ void downstream_v2_define_scifi_candidates::downstream_v2_find_hit_ra // Get fired sectors const auto fired_sectors = dev_ut_layer_geometry->find_sectors(layer, layer_y - yTolLayer, layer_y + yTolLayer); unsigned num_ranges = 0; - Downstream::Structs::LayerHitRanges out_ranges = { - short2 {-1, -1}, short2 {-1, -1}, short2 {-1, -1}, short2 {-1, -1}}; + Downstream::Structs::LayerHitRanges out_ranges = {ushort2 {invalid_hit, invalid_hit}, + ushort2 {invalid_hit, invalid_hit}, + ushort2 {invalid_hit, invalid_hit}, + ushort2 {invalid_hit, invalid_hit}}; for (unsigned sector_idx = 0; sector_idx < 4; sector_idx++) { const auto sector = fired_sectors[sector_idx]; if (sector == -1) break; @@ -208,6 +213,9 @@ __global__ void downstream_v2_define_scifi_candidates::downstream_v2_compute_num // Basics const auto event_number = parameters.dev_event_list[blockIdx.x]; + // Alias + constexpr auto invalid_hit = Downstream::Structs::invalid_hit; + // Offsets const auto input_offset = parameters.dev_offsets_seeding_tracks[event_number]; const auto output_offset = parameters.dev_downstream_seed_offsets[event_number]; @@ -221,10 +229,10 @@ __global__ void downstream_v2_define_scifi_candidates::downstream_v2_compute_num for (unsigned i = threadIdx.x; i < input_size; i += blockDim.x) { const auto ranges = input_seeds[i].hits; - output_counters[i * 4 + 0] = (ranges[0].x == -1) ? 0 : ranges[0].y - ranges[0].x; - output_counters[i * 4 + 1] = (ranges[1].x == -1) ? 0 : ranges[1].y - ranges[1].x; - output_counters[i * 4 + 2] = (ranges[2].x == -1) ? 0 : ranges[2].y - ranges[2].x; - output_counters[i * 4 + 3] = (ranges[3].x == -1) ? 0 : ranges[3].y - ranges[3].x; + output_counters[i * 4 + 0] = (ranges[0].x == invalid_hit) ? 0 : ranges[0].y - ranges[0].x; + output_counters[i * 4 + 1] = (ranges[1].x == invalid_hit) ? 0 : ranges[1].y - ranges[1].x; + output_counters[i * 4 + 2] = (ranges[2].x == invalid_hit) ? 0 : ranges[2].y - ranges[2].x; + output_counters[i * 4 + 3] = (ranges[3].x == invalid_hit) ? 0 : ranges[3].y - ranges[3].x; } } @@ -247,7 +255,7 @@ __global__ void downstream_v2_define_scifi_candidates::downstream_v2_fill_pairs( const auto output_idx = output_offset + i; for (unsigned r = 0; r < 4; r++) { const auto range = seed.hits[r]; - if (range.x == -1) break; + if (range.x == Downstream::Structs::invalid_hit) break; const auto output_hit_offset = parameters.dev_downstream_seed_hits_offsets[output_idx * 4 + r]; const auto output_num_hits = diff --git a/device/downstream/tracking/src/DownstreamV2FindTracks.cu b/device/downstream/tracking/src/DownstreamV2FindTracks.cu index 8492e6458d768edc7726ce1f2f4810272313ca25..4715c71937517560e98c28ebec0a61c1ba7a3699 100644 --- a/device/downstream/tracking/src/DownstreamV2FindTracks.cu +++ b/device/downstream/tracking/src/DownstreamV2FindTracks.cu @@ -127,6 +127,9 @@ __global__ void downstream_v2_find_tracks::downstream_v2_find_hits_in_layer_0( const auto event_number = parameters.dev_event_list[blockIdx.x]; const auto number_of_events = parameters.dev_number_of_events[0]; + // Alias + constexpr auto invalid_hit = Downstream::Structs::invalid_hit; + // UT hits const auto total_number_of_ut_hits = parameters.dev_ut_hit_offsets[number_of_events * UT::Constants::n_groups]; UT::HitOffsets ut_hit_offsets {parameters.dev_ut_hit_offsets, event_number}; @@ -227,8 +230,8 @@ __global__ void downstream_v2_find_tracks::downstream_v2_find_hits_in_layer_0( Downstream::Structs::CompactTrack out; out.scifi_idx = scifi_offset + scifi_idx; out.ut_hits[0] = best_hit.best() + hit_cache.HitOffset(); - out.ut_hits[1] = -1; - out.ut_hits[2] = -1; + out.ut_hits[1] = invalid_hit; + out.ut_hits[2] = invalid_hit; out.ut_hits[3] = h3_idx; output_tracks[idx] = out; } @@ -378,6 +381,9 @@ __global__ void downstream_v2_find_tracks::downstream_v2_find_hits_in_layer_2( const auto event_number = parameters.dev_event_list[blockIdx.x]; const auto number_of_events = parameters.dev_number_of_events[0]; + // Alias + constexpr auto invalid_hit = Downstream::Structs::invalid_hit; + // UT hits const auto total_number_of_ut_hits = parameters.dev_ut_hit_offsets[number_of_events * UT::Constants::n_groups]; UT::HitOffsets ut_hit_offsets {parameters.dev_ut_hit_offsets, event_number}; @@ -435,11 +441,12 @@ __global__ void downstream_v2_find_tracks::downstream_v2_find_hits_in_layer_2( // Get candidates const auto dist_10 = (x0 + tx * ut_hits.zAtYEq0(h10_idx)) - (ut_hits.xAt(h10_idx, scifi_state.yAt(ut_hits.zAtYEq0(h10_idx)))); - const auto dist_11 = (h11_idx != -1) ? (x0 + tx * ut_hits.zAtYEq0(h11_idx)) - - (ut_hits.xAt(h11_idx, scifi_state.yAt(ut_hits.zAtYEq0(h11_idx)))) : - std::numeric_limits::infinity(); + const auto dist_11 = + (h11_idx != invalid_hit) ? + (x0 + tx * ut_hits.zAtYEq0(h11_idx)) - (ut_hits.xAt(h11_idx, scifi_state.yAt(ut_hits.zAtYEq0(h11_idx)))) : + std::numeric_limits::infinity(); - Downstream::Helpers::BestSelector best_combination; + Downstream::Helpers::BestSelector best_combination; // Compute the search range const auto layer_dxdy = dev_ut_layer_geometry->mean_dxdy[layer]; @@ -471,9 +478,9 @@ __global__ void downstream_v2_find_tracks::downstream_v2_find_hits_in_layer_2( } const float xdist = (x0 + tx * hit.zAtYEq0()) - hit.xAt(expected_hit_y); if (fabsf(xdist) < xTol) { - best_combination.add(xdist + dist_10, make_short2(h10_idx, idx)); - if (h11_idx != -1) { - best_combination.add(xdist + dist_11, make_short2(h11_idx, idx)); + best_combination.add(xdist + dist_10, make_ushort2(h10_idx, idx)); + if (h11_idx != invalid_hit) { + best_combination.add(xdist + dist_11, make_ushort2(h11_idx, idx)); } } }); diff --git a/device/event_model/UT/include/UTEventModel.cuh b/device/event_model/UT/include/UTEventModel.cuh index e9669ecb3b923d27132a2002f0ad861fd06e50f2..6ce92aa233d09d25219a61b97b7338292491104f 100644 --- a/device/event_model/UT/include/UTEventModel.cuh +++ b/device/event_model/UT/include/UTEventModel.cuh @@ -61,12 +61,14 @@ namespace UT { }; struct TrackHits { + constexpr static ushort invalid_hit = std::numeric_limits::max(); + float qop; float x, z; float tx; unsigned short hits_num = 0; unsigned short velo_track_index; - short hits[UT::Constants::max_track_size]; + unsigned short hits[UT::Constants::max_track_size]; friend std::ostream& operator<<(std::ostream& stream, const TrackHits& hit) {