[Openmp-commits] [PATCH] D65836: Factor architecture dependent code out of loop.cu

[Alexey Bataev via Phabricator via Openmp-commits]

ABataev added a comment.

I would suggest at first to come to an agreement on the design of this reworked library at first.



================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h:17-18
+
+#define FORCEINLINE __forceinline__ __device__
+
+FORCEINLINE void __kmpc_impl_unpack(int64_t val, int32_t *lo, int32_t *hi) {
----------------
Better to use original `INLINE` macro defined in the project rather than to define the new one.


================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h:19
+
+FORCEINLINE void __kmpc_impl_unpack(int64_t val, int32_t *lo, int32_t *hi) {
+  asm volatile("mov.b64 {%0,%1}, %2;" : "=r"(*lo), "=r"(*hi) : "l"(val));
----------------
Why pointers? Use references.


================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h:32
+FORCEINLINE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt() {
+  uint32_t res;
+  asm("mov.u32 %0, %%lanemask_lt;" : "=r"(res));
----------------
`uitn32_t`->`__kmpc_impl_lanemask_t`?


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D65836