From c6cb2d2852e1b1fcc7aa621c4b13ea71cea4301e Mon Sep 17 00:00:00 2001 From: beomki-yeo Date: Wed, 1 Sep 2021 11:56:09 -0700 Subject: [PATCH 1/2] use unsigned int instead of size_t --- core/include/edm/container.hpp | 2 +- core/include/edm/internal_spacepoint.hpp | 12 ++++++------ core/include/seeding/detail/singlet.hpp | 4 ++-- core/include/seeding/doublet_finding.hpp | 10 ++++++---- core/include/seeding/seed_finding.hpp | 6 +++--- device/cuda/src/seeding/doublet_finding.cu | 12 ++++++------ device/cuda/src/seeding/seed_selecting.cu | 12 ++++++------ device/cuda/src/seeding/triplet_counting.cu | 2 +- device/cuda/src/seeding/triplet_finding.cu | 4 ++-- device/cuda/src/seeding/weight_updating.cu | 8 ++++---- 10 files changed, 37 insertions(+), 35 deletions(-) diff --git a/core/include/edm/container.hpp b/core/include/edm/container.hpp index a1147408d9..a43d9d83af 100644 --- a/core/include/edm/container.hpp +++ b/core/include/edm/container.hpp @@ -162,7 +162,7 @@ class container { /** * @brief Default Constructor */ - TRACCC_HOST container() = default; + container() = default; /** * @brief Bounds-checking mutable element accessor. diff --git a/core/include/edm/internal_spacepoint.hpp b/core/include/edm/internal_spacepoint.hpp index 74037b7ba5..a186bde165 100644 --- a/core/include/edm/internal_spacepoint.hpp +++ b/core/include/edm/internal_spacepoint.hpp @@ -24,19 +24,19 @@ namespace traccc { struct neighbor_idx { - size_t counts; + unsigned int counts; /// global_indices: the global indices of neighbor bins provided by axis /// class size of 9 is from (3 neighbors on z-axis) x (3 neighbors on /// phi-axis) - size_t global_indices[9]; + unsigned int global_indices[9]; /// vector_indices: the actual indices of neighbor bins, which are used for /// navigating internal_spacepoint_container - size_t vector_indices[9]; + unsigned int vector_indices[9]; }; /// Header: bin information (global bin index, neighborhood bin indices) struct bin_information { - size_t global_index; + unsigned int global_index; neighbor_idx bottom_idx; neighbor_idx top_idx; }; @@ -154,8 +154,8 @@ using internal_spacepoint_container_buffer = using internal_spacepoint_container_view = container_view >; -inline size_t find_vector_id_from_global_id( - size_t global_bin, vecmem::vector& headers) { +inline unsigned int find_vector_id_from_global_id( + unsigned int global_bin, vecmem::vector& headers) { auto iterator = std::find_if(headers.begin(), headers.end(), [&global_bin](const bin_information& bin_info) { diff --git a/core/include/seeding/detail/singlet.hpp b/core/include/seeding/detail/singlet.hpp index a2fbb5577e..5260a6bf50 100644 --- a/core/include/seeding/detail/singlet.hpp +++ b/core/include/seeding/detail/singlet.hpp @@ -17,9 +17,9 @@ namespace traccc { /// location of spacepoint in internal spacepoint container struct sp_location { /// index of the bin of the spacepoint grid - std::size_t bin_idx; + unsigned int bin_idx; /// index of the spacepoint in the bin - std::size_t sp_idx; + unsigned int sp_idx; }; inline TRACCC_HOST_DEVICE bool operator==(const sp_location& lhs, diff --git a/core/include/seeding/doublet_finding.hpp b/core/include/seeding/doublet_finding.hpp index 6ff546263e..c872399af1 100644 --- a/core/include/seeding/doublet_finding.hpp +++ b/core/include/seeding/doublet_finding.hpp @@ -74,12 +74,13 @@ struct doublet_finding // for middle-bottom doublets if (bottom) { - for (size_t i = 0; i < counts; ++i) { + for (unsigned int i = 0; i < counts; ++i) { auto& bin_idx = bottom_bin_indices[i]; auto& spacepoints = isp_container.get_items()[bin_idx]; - for (size_t sp_idx = 0; sp_idx < spacepoints.size(); ++sp_idx) { + for (unsigned int sp_idx = 0; sp_idx < spacepoints.size(); + ++sp_idx) { auto& spB = spacepoints[sp_idx]; if (!doublet_finding_helper::isCompatible( @@ -102,11 +103,12 @@ struct doublet_finding auto& counts = bin_information.top_idx.counts; auto& top_bin_indices = bin_information.top_idx.vector_indices; - for (size_t i = 0; i < counts; ++i) { + for (unsigned int i = 0; i < counts; ++i) { auto& bin_idx = top_bin_indices[i]; auto& spacepoints = isp_container.get_items()[bin_idx]; - for (size_t sp_idx = 0; sp_idx < spacepoints.size(); ++sp_idx) { + for (unsigned int sp_idx = 0; sp_idx < spacepoints.size(); + ++sp_idx) { auto& spT = spacepoints[sp_idx]; if (!doublet_finding_helper::isCompatible( diff --git a/core/include/seeding/seed_finding.hpp b/core/include/seeding/seed_finding.hpp index 39d48bf5f3..cb2012efe5 100644 --- a/core/include/seeding/seed_finding.hpp +++ b/core/include/seeding/seed_finding.hpp @@ -59,7 +59,7 @@ struct seed_finding const bool top = false; // iterate over grid bins - for (size_t i = 0; i < isp_container.size(); ++i) { + for (unsigned int i = 0; i < isp_container.size(); ++i) { auto& bin_information = isp_container.get_headers()[i]; auto& spM_collection = isp_container.get_items()[i]; @@ -70,7 +70,7 @@ struct seed_finding */ /// iterate over middle spacepoints - for (size_t j = 0; j < spM_collection.size(); ++j) { + for (unsigned int j = 0; j < spM_collection.size(); ++j) { sp_location spM_location({i, j}); // middule-bottom doublet search @@ -91,7 +91,7 @@ struct seed_finding // triplet search from the combinations of two doublets which // share middle spacepoint - for (size_t k = 0; k < mid_bot.first.size(); ++k) { + for (unsigned int k = 0; k < mid_bot.first.size(); ++k) { auto& doublet_mb = mid_bot.first[k]; auto& lb = mid_bot.second[k]; diff --git a/device/cuda/src/seeding/doublet_finding.cu b/device/cuda/src/seeding/doublet_finding.cu index 5283a17e81..93ac2a4a80 100644 --- a/device/cuda/src/seeding/doublet_finding.cu +++ b/device/cuda/src/seeding/doublet_finding.cu @@ -54,7 +54,7 @@ void doublet_finding(const seedfinder_config& config, // N_i is the number of blocks for i-th bin, defined as // num_compatible_middle_sp_per_bin / num_threads + 1 unsigned int num_blocks = 0; - for (size_t i = 0; i < internal_sp_view.headers.size(); ++i) { + for (unsigned int i = 0; i < internal_sp_view.headers.size(); ++i) { num_blocks += doublet_counter_container.get_headers()[i] / num_threads + 1; } @@ -155,18 +155,18 @@ __global__ void doublet_finding_kernel( // spacepoints unsigned int mid_bot_start_idx = 0; unsigned int mid_top_start_idx = 0; - for (size_t i = 0; i < gid; i++) { + for (unsigned int i = 0; i < gid; i++) { mid_bot_start_idx += doublet_counter_per_bin[i].n_mid_bot; mid_top_start_idx += doublet_counter_per_bin[i].n_mid_top; } // Loop over (bottom and top) internal spacepoints in tje neighbor bins - for (size_t i_n = 0; i_n < bin_info.bottom_idx.counts; ++i_n) { + for (unsigned int i_n = 0; i_n < bin_info.bottom_idx.counts; ++i_n) { const auto& neigh_bin = bin_info.bottom_idx.vector_indices[i_n]; const auto& neigh_internal_sp_per_bin = internal_sp_device.get_items().at(neigh_bin); - for (size_t spB_idx = 0; spB_idx < neigh_internal_sp_per_bin.size(); + for (unsigned int spB_idx = 0; spB_idx < neigh_internal_sp_per_bin.size(); ++spB_idx) { const auto& neigh_isp = neigh_internal_sp_per_bin[spB_idx]; @@ -183,7 +183,7 @@ __global__ void doublet_finding_kernel( doublet_counter_per_bin[gid].n_mid_bot && num_mid_bot_doublets_per_bin < mid_bot_doublets_per_bin.size()) { - size_t pos = mid_bot_start_idx + n_mid_bot_per_spM; + unsigned int pos = mid_bot_start_idx + n_mid_bot_per_spM; // prevent overflow again if (pos >= mid_bot_doublets_per_bin.size()) { @@ -209,7 +209,7 @@ __global__ void doublet_finding_kernel( doublet_counter_per_bin[gid].n_mid_top && num_mid_top_doublets_per_bin < mid_top_doublets_per_bin.size()) { - size_t pos = mid_top_start_idx + n_mid_top_per_spM; + unsigned int pos = mid_top_start_idx + n_mid_top_per_spM; // prevent overflow again if (pos >= mid_top_doublets_per_bin.size()) { diff --git a/device/cuda/src/seeding/seed_selecting.cu b/device/cuda/src/seeding/seed_selecting.cu index c15b8ba5c5..2d79b920c7 100644 --- a/device/cuda/src/seeding/seed_selecting.cu +++ b/device/cuda/src/seeding/seed_selecting.cu @@ -83,7 +83,7 @@ void seed_selecting(const seedfilter_config& filter_config, // N_i is the number of blocks for i-th bin, defined as num_triplets_per_bin // / num_threads + 1 unsigned int num_blocks = 0; - for (size_t i = 0; i < internal_sp_view.headers.size(); ++i) { + for (unsigned int i = 0; i < internal_sp_view.headers.size(); ++i) { num_blocks += triplet_counter_container.get_headers()[i] / num_threads + 1; } @@ -170,13 +170,13 @@ __global__ void seed_selecting_kernel( auto& spM = internal_sp_per_bin[spM_idx]; // number of triplets per compatible middle spacepoint - size_t n_triplets_per_spM = 0; + unsigned int n_triplets_per_spM = 0; // the start index of triplets_per_spM - size_t stride = threadIdx.x * filter_config.max_triplets_per_spM; + unsigned int stride = threadIdx.x * filter_config.max_triplets_per_spM; // iterate over the triplets in the bin - for (size_t i = 0; i < num_triplets_per_bin; ++i) { + for (unsigned int i = 0; i < num_triplets_per_bin; ++i) { auto& aTriplet = triplets_per_bin[i]; auto& spB_loc = aTriplet.sp1; auto& spT_loc = aTriplet.sp3; @@ -256,10 +256,10 @@ __global__ void seed_selecting_kernel( triplet_weight_descending()); // the number of good seed per compatible middle spacepoint - size_t n_seeds_per_spM = 0; + unsigned int n_seeds_per_spM = 0; // iterate over the good triplets for final selection of seeds - for (size_t i = stride; i < stride + n_triplets_per_spM; ++i) { + for (unsigned int i = stride; i < stride + n_triplets_per_spM; ++i) { auto& aTriplet = triplets_per_spM[i]; auto& spB_loc = aTriplet.sp1; auto& spT_loc = aTriplet.sp3; diff --git a/device/cuda/src/seeding/triplet_counting.cu b/device/cuda/src/seeding/triplet_counting.cu index 955743bb93..c02a62c163 100644 --- a/device/cuda/src/seeding/triplet_counting.cu +++ b/device/cuda/src/seeding/triplet_counting.cu @@ -58,7 +58,7 @@ void triplet_counting(const seedfinder_config& config, // N_i is the number of blocks for i-th bin, defined as // num_mid_bot_doublet_per_bin / num_threads + 1 unsigned int num_blocks = 0; - for (size_t i = 0; i < internal_sp_view.headers.size(); ++i) { + for (unsigned int i = 0; i < internal_sp_view.headers.size(); ++i) { num_blocks += mid_bot_doublet_container.get_headers()[i] / num_threads + 1; } diff --git a/device/cuda/src/seeding/triplet_finding.cu b/device/cuda/src/seeding/triplet_finding.cu index 6125eff54c..c749758115 100644 --- a/device/cuda/src/seeding/triplet_finding.cu +++ b/device/cuda/src/seeding/triplet_finding.cu @@ -61,7 +61,7 @@ void triplet_finding(const seedfinder_config& config, // N_i is the number of blocks for i-th bin, defined as // num_compatible_mid_bot_doublets_per_bin / num_threads + 1 unsigned int num_blocks = 0; - for (size_t i = 0; i < internal_sp_view.headers.size(); ++i) { + for (unsigned int i = 0; i < internal_sp_view.headers.size(); ++i) { num_blocks += triplet_counter_container.get_headers()[i] / num_threads + 1; } @@ -245,7 +245,7 @@ __global__ void triplet_finding_kernel( if (triplet_finding_helper::isCompatible( spM, lb, lt, config, iSinTheta2, scatteringInRegion2, curvature, impact_parameter)) { - size_t pos = triplet_start_idx + n_triplets_per_mb; + unsigned int pos = triplet_start_idx + n_triplets_per_mb; // prevent the overflow if (pos >= triplets_per_bin.size()) { continue; diff --git a/device/cuda/src/seeding/weight_updating.cu b/device/cuda/src/seeding/weight_updating.cu index 4874acf5de..c2cb9a370f 100644 --- a/device/cuda/src/seeding/weight_updating.cu +++ b/device/cuda/src/seeding/weight_updating.cu @@ -48,7 +48,7 @@ void weight_updating(const seedfilter_config& filter_config, // N_i is the number of blocks for i-th bin, defined as num_triplets_per_bin // / num_threads + 1 unsigned int num_blocks = 0; - for (size_t i = 0; i < internal_sp_view.headers.size(); ++i) { + for (unsigned int i = 0; i < internal_sp_view.headers.size(); ++i) { num_blocks += triplet_container.get_headers()[i] / num_threads + 1; } @@ -117,8 +117,8 @@ __global__ void weight_updating_kernel( // find the reference index (start and end) of the triplet container item // vector - size_t start_idx = 0; - size_t end_idx = 0; + unsigned int start_idx = 0; + unsigned int end_idx = 0; for (auto triplet_counter : triplet_counter_per_bin) { end_idx += triplet_counter.n_triplets; @@ -187,7 +187,7 @@ __global__ void weight_updating_kernel( bool newCompSeed = true; - for (size_t i_s = 0; i_s < num_compat_seedR; ++i_s) { + for (unsigned int i_s = 0; i_s < num_compat_seedR; ++i_s) { scalar previousDiameter = compat_seedR[i_s]; // original ATLAS code uses higher min distance for 2nd found From 4f1c4292467361cc46ae1112d9cf830a71922165 Mon Sep 17 00:00:00 2001 From: beomki-yeo Date: Wed, 1 Sep 2021 14:54:06 -0700 Subject: [PATCH 2/2] clang format --- device/cuda/src/seeding/doublet_finding.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/device/cuda/src/seeding/doublet_finding.cu b/device/cuda/src/seeding/doublet_finding.cu index 93ac2a4a80..693878ea19 100644 --- a/device/cuda/src/seeding/doublet_finding.cu +++ b/device/cuda/src/seeding/doublet_finding.cu @@ -166,8 +166,8 @@ __global__ void doublet_finding_kernel( const auto& neigh_internal_sp_per_bin = internal_sp_device.get_items().at(neigh_bin); - for (unsigned int spB_idx = 0; spB_idx < neigh_internal_sp_per_bin.size(); - ++spB_idx) { + for (unsigned int spB_idx = 0; + spB_idx < neigh_internal_sp_per_bin.size(); ++spB_idx) { const auto& neigh_isp = neigh_internal_sp_per_bin[spB_idx]; // Check if middle and bottom sp can form a doublet