[Openmp-commits] [PATCH] D66855: [libomptarget] Refactor syncthreads macro to inline function

[Alexey Bataev via Phabricator via Openmp-commits]

ABataev added inline comments.


================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu:556
     }
     // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
+    __kmpc_impl_syncthreads();
----------------
Remove FIXME


================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu:571
   if (isSPMDExecutionMode) {
     // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
+    __kmpc_impl_syncthreads();
----------------
And here too


================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu:129
   }
   // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
+  __kmpc_impl_syncthreads();
----------------
Remove FIXME


================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu:172
 
   // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
+  __kmpc_impl_syncthreads();
----------------
Same here


================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu:78
   PRINT0(LD_SYNC, "call kmpc_barrier_simple_spmd\n");
   // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
+  __kmpc_impl_syncthreads();
----------------
Same


================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h:51
+  asm volatile("bar.sync %0;" : : "r"(0) : "memory");
+#endif
+}
----------------
`#endif // __clang__`


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D66855/new/

https://reviews.llvm.org/D66855