[Openmp-commits] [PATCH] D14254: [OpenMP] Initial implementation of OpenMP offloading library - libomptarget device RTLs.

[Jonas Hahnfeld via Phabricator via Openmp-commits]

Hahnfeld added a comment.

I talked to @ABataev and Clang never generates calls to `__kmpc_atomic_` functions. I've never heard about `__array_atomic` but I guess they aren't used either as array reductions are performed element-wise.



================
Comment at: libomptarget/deviceRTLs/nvptx/src/interface.h:462
+
+// atomic
+// EXTERN void __array_atomicfixed4_add(kmp_Indent *loc, int tid, int32_t
----------------
You can remove the next ~820 lines.


================
Comment at: libomptarget/deviceRTLs/nvptx/src/reduction.cu:97-116
+  int globalThreadId = GetGlobalThreadId();
+  omptarget_nvptx_TaskDescr *currTaskDescr =
+      getMyTopTaskDescriptor(globalThreadId);
+  int numthread;
+  if (currTaskDescr->IsParallelConstruct()) {
+    numthread = omp_get_num_threads();
+  } else {
----------------
grokos wrote:
> Hahnfeld wrote:
> > Same as `__kmpc_reduce_gpu` (and `__kmpc_reduce_nowait41`)?
> `__kmpc_reduce_nowait41` has been removed from new diff.
It's still there...


================
Comment at: libomptarget/deviceRTLs/nvptx/src/reduction.cu:663-666
+template <>
+INLINE __device__ float Compute<float, omptarget_nvptx_add>(float a, float b) {
+  return a + b;
+}
----------------
grokos wrote:
> Hahnfeld wrote:
> > Could you add a comment explaining why this has to be specialized?
> A comment has been added in the new diff. The specialization is useful in order to avoid the switch at runtime.
`binop` is a template argument so the compiler will do that for you. Anyway, I think these functions can be removed altogether now!


================
Comment at: libomptarget/deviceRTLs/nvptx/src/reduction.cu:20-32
+// cannot implement atomic_start and atomic_end for GPU. Report runtime error
+EXTERN void __kmpc_atomic_start() {
+  printf("__kmpc_atomic_start not supported\n");
+  asm("trap;");
+  return;
+}
+
----------------
Remove


================
Comment at: libomptarget/deviceRTLs/nvptx/src/reduction.cu:110
+
+// implement different data type or operations  with atomicCAS
+#define omptarget_nvptx_add(x, y) ((x) + (y))
----------------
I think you can remove the next ~730 lines until `__kmpc_shuffle_int32`


================
Comment at: libomptarget/deviceRTLs/nvptx/src/reduction.cu:951-993
+//
+// runtime support for array reduction
+//
+
+#define ARRAYATOMIC_GENOP(_name, _dtype, _op)                                  \
+  EXTERN void __array_atomic_##_name##_##_op(                                  \
+      kmp_Indent *id_ref, int32_t gtid, _dtype *lhs, _dtype *rhs, int64_t n) { \
----------------
These macros are never used.


================
Comment at: libomptarget/deviceRTLs/nvptx/src/stdio.cu:18
+
+EXTERN int __kmpc_printf(const char *str) { return printf("%s", str); }
----------------
I think this function isn't used anymore?


Repository:
  rL LLVM

https://reviews.llvm.org/D14254