GPU: Make GPUCommonAlgorithm::sortInBlock deterministic with GPUCA_DETERMINISTIC_MODE

Author: davidrohr

GPU/Common/GPUCommonAlgorithm.h

@@ -283,21 +283,29 @@ GPUdi() void GPUCommonAlgorithm::sortInBlock(T* begin, T* end, const S& comp)
 #ifndef GPUCA_GPUCODE
   GPUCommonAlgorithm::sort(begin, end, comp);
 #else
-  int32_t n = end - begin;
-  for (int32_t i = 0; i < n; i++) {
-    for (int32_t tIdx = get_local_id(0); tIdx < n; tIdx += get_local_size(0)) {
-      int32_t offset = i % 2;
-      int32_t curPos = 2 * tIdx + offset;
-      int32_t nextPos = curPos + 1;
-
-      if (nextPos < n) {
-        if (!comp(begin[curPos], begin[nextPos])) {
-          IterSwap(&begin[curPos], &begin[nextPos]);
+  GPUCA_DETERMINISTIC_CODE( // clang-format off
+    GPUbarrier();
+    if (get_local_id(0) == 0) {
+      GPUCommonAlgorithm::sort(begin, end, comp);
+    }
+    GPUbarrier();
+  , // !GPUCA_DETERMINISTIC_CODE
+    int32_t n = end - begin;
+    for (int32_t i = 0; i < n; i++) {
+      for (int32_t tIdx = get_local_id(0); tIdx < n; tIdx += get_local_size(0)) {
+        int32_t offset = i % 2;
+        int32_t curPos = 2 * tIdx + offset;
+        int32_t nextPos = curPos + 1;
+
+        if (nextPos < n) {
+          if (!comp(begin[curPos], begin[nextPos])) {
+            IterSwap(&begin[curPos], &begin[nextPos]);
+          }
         }
       }
+      GPUbarrier();
     }
-    GPUbarrier();
-  }
+  ) // clang-format on
 #endif
 }