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

[Johannes Doerfert via Phabricator via Openmp-commits]

jdoerfert added a comment.

I'm fine with this, @ABataev ?



================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/loop.cu:392
   INLINE static uint64_t NextIter() {
-    unsigned int active = __ACTIVEMASK();
-    int leader = __ffs(active) - 1;
----------------
I would prefer `__SHFL_SYNC` `__ACTIVEMASK` etc. also to be function calls to `__kmpc_XXXX` functions but I won't require it.


================
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) {
----------------
JonChesterfield wrote:
> ABataev wrote:
> > JonChesterfield wrote:
> > > ABataev wrote:
> > > > Better to use original `INLINE` macro defined in the project rather than to define the new one.
> > > I'd prefer that too, but INLINE maps to `__inline__`, rather than `__forceinline__`, and that leaves calls to these functions in the bitcode library for nvptx.
> > Then better to fix original `INLINE` macro and replace `__inline__` with `__forceinline__`. I assume we'd like to inline all the functions.
> That works for me. I suspect everything marked INLINE was intended to be inlined. Diff at D65876.
We can even have different definitions of `INLINE` if that becomes necessary.



================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h:39
+
+INLINE int __kmpc_impl_popc(uint32_t x) { return __popc(x); }
+
----------------
`int` -> `(u)int32_t` ?


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D65836