Checkpoint

device/cuda/include/traccc/cuda/seeding2/kernels/seed_finding_kernel.hpp

@@ -9,9 +9,9 @@
 
 #include <traccc/cuda/seeding2/types/internal_sp.hpp>
 #include <traccc/cuda/seeding2/types/kd_tree.hpp>
+#include <traccc/edm/alt_seed.hpp>
 #include <traccc/edm/internal_spacepoint.hpp>
 #include <traccc/edm/spacepoint.hpp>
-#include <traccc/edm/alt_seed.hpp>
 #include <traccc/seeding/detail/seeding_config.hpp>
 
 namespace traccc::cuda {

device/cuda/include/traccc/cuda/utils/device_traits.hpp

@@ -19,7 +19,8 @@
 #define MAX_RESIDENT_THREADS_PER_SM 1536
 #elif __CUDA_ARCH__ <= 900
 #define MAX_RESIDENT_THREADS_PER_SM 2048
-#warning "Unknown CUDA architecture, setting maximum resident threads per block to 1024."
+#warning \
+    "Unknown CUDA architecture, setting maximum resident threads per block to 1024."
 #define MAX_RESIDENT_THREADS_PER_SM 1024
 #endif
 

device/cuda/include/traccc/cuda/utils/sort.hpp

@@ -0,0 +1,72 @@
+/** TRACCC library, part of the ACTS project (R&D line)
+ *
+ * (c) 2021 CERN for the benefit of the ACTS project
+ *
+ * Mozilla Public License Version 2.0
+ */
+
+#pragma once
+
+namespace traccc::cuda {
+template <typename T>
+__device__ void swap(T& a, T& b) {
+    T t = a;
+    a = b;
+    b = t;
+}
+
+template <typename K, typename C>
+__device__ void blockOddEvenKeySort(K* keys, uint32_t num_keys,
+                                    C&& comparison) {
+    bool sorted;
+
+    do {
+        sorted = true;
+
+        for (uint32_t j = 2 * threadIdx.x + 1; j < num_keys - 1;
+             j += 2 * blockDim.x) {
+            if (comparison(keys[j + 1], keys[j])) {
+                swap(keys[j + 1], keys[j]);
+                sorted = false;
+            }
+        }
+
+        __syncthreads();
+
+        for (uint32_t j = 2 * threadIdx.x; j < num_keys - 1;
+             j += 2 * blockDim.x) {
+            if (comparison(keys[j + 1], keys[j])) {
+                swap(keys[j + 1], keys[j]);
+                sorted = false;
+            }
+        }
+    } while (__syncthreads_or(!sorted));
+}
+
+template <typename K, typename C>
+__device__ void warpOddEvenKeySort(K* keys, uint32_t num_keys, C&& comparison) {
+    bool sorted;
+
+    do {
+        sorted = true;
+
+        for (uint32_t j = 2 * (threadIdx.x % WARP_SIZE) + 1; j < num_keys - 1;
+             j += 2 * WARP_SIZE) {
+            if (comparison(keys[j + 1], keys[j])) {
+                swap(keys[j + 1], keys[j]);
+                sorted = false;
+            }
+        }
+
+        __syncwarp(__activemask());
+
+        for (uint32_t j = 2 * (threadIdx.x % WARP_SIZE); j < num_keys - 1;
+             j += 2 * WARP_SIZE) {
+            if (comparison(keys[j + 1], keys[j])) {
+                swap(keys[j + 1], keys[j]);
+                sorted = false;
+            }
+        }
+    } while (__any_sync(__activemask(), !sorted));
+}
+}  // namespace traccc::cuda