Dorothea Vom Bruch
--- a/cuda/utils/include/Sorting.cuh 0 → 100644
+++ b/cuda/utils/include/Sorting.cuh 0 → 100644
+ * @brief Sort by var stored in sorting_vars, store index in hit_permutations
+ */
+template<class T>
+__host__ __device__
+void findPermutation(
+  const T* sorting_vars,
+  const uint hit_start,
+  uint* hit_permutations,
+  const uint n_hits
+){
+#ifdef __CUDA_ARCH__
+  for (unsigned int i = 0; i < (n_hits + blockDim.x - 1); ++i) {
+    const unsigned int hit_rel_index = i*blockDim.x + threadIdx.x;
+    if ( hit_rel_index < n_hits ) {
+      const int hit_index = hit_start + hit_rel_index;
+      const T var = sorting_vars[hit_index];
--- a/cuda/utils/include/Sorting.cuh 0 → 100644
+++ b/cuda/utils/include/Sorting.cuh 0 → 100644
+      for (unsigned int j = 0; j < n_hits; ++j) {
+        const int other_hit_index = hit_start + j;
+        const T other_var = sorting_vars[other_hit_index];
+        // Stable sorting
+        position += var > other_var || ( var == other_var && hit_rel_index > j );
+      }
+      assert(position < n_hits);
+      
+      // Store it in hit_permutations 
+      hit_permutations[hit_start + position] = hit_index; 
+    }
+  }
+#else
+  for (unsigned int i = 0; i < n_hits; ++i) {
+    const int hit_index = hit_start + i;
+    const T var = sorting_vars[hit_index];
--- a/cuda/velo/calculate_phi_and_sort/src/SortByPhi.cu
+++ b/cuda/velo/calculate_phi_and_sort/src/SortByPhi.cu
  uint* new_hit_IDs = (uint*) hit_Zs;
  
  // Apply permutation across all arrays
-  apply_permutation(hit_permutations, event_hit_start, event_number_of_hits, hit_Xs, new_hit_Xs);
+  applyPermutation(hit_permutations, event_hit_start, event_number_of_hits, hit_Xs, new_hit_Xs);
--- a/stream/sequence/src/StreamSequence.cu
+++ b/stream/sequence/src/StreamSequence.cu
      cudaCheck(cudaMemcpyAsync(host_velo_states, argen.generate<arg::dev_velo_states>(argument_offsets), argen.size<arg::dev_velo_states>(host_number_of_reconstructed_velo_tracks[0]), cudaMemcpyDeviceToHost, stream)); 
    }

-    // VeloUT tracking
+    // UT hit sorting by x
    argument_sizes[arg::dev_ut_hits] = argen.size<arg::dev_ut_hits>(number_of_events);
+    argument_sizes[arg::dev_ut_hits_sorted] = argen.size<arg::dev_ut_hits_sorted>(number_of_events);
+    argument_sizes[arg::dev_ut_hit_permutations] = argen.size<arg::dev_ut_hit_permutations>(number_of_events * VeloUTTracking::max_numhits_per_event);
+    scheduler.setup_next(argument_sizes, argument_offsets, sequence_step++);
+    cudaCheck(cudaMemcpyAsync(argen.generate<arg::dev_ut_hits>(argument_offsets), host_ut_hits_events, number_of_events * sizeof(VeloUTTracking::HitsSoA), cudaMemcpyHostToDevice, stream ));