use unsigned int instead of size_t for seed finding

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.

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<bin_information, internal_spacepoint<spacepoint> >;
 
-inline size_t find_vector_id_from_global_id(
-    size_t global_bin, vecmem::vector<bin_information>& headers) {
+inline unsigned int find_vector_id_from_global_id(
+    unsigned int global_bin, vecmem::vector<bin_information>& headers) {
     auto iterator =
         std::find_if(headers.begin(), headers.end(),
                      [&global_bin](const bin_information& bin_info) {

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,

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(

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];
 

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,19 +155,19 @@ __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();
-             ++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
@@ -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()) {

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;

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;
     }

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;

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