[Openmp-commits] [PATCH] D69718: [libomptarget] Implement target_impl for amdgcn

[Jon Chesterfield via Phabricator via Openmp-commits]

JonChesterfield added inline comments.


================
Comment at: openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h:23
+#define DEVICE __device__
+#define INLINE __inline__ DEVICE
+#define NOINLINE __noinline__ DEVICE
----------------
ABataev wrote:
> `__forceinline__`?
inline is sufficient when building with clang. It's possible the nvptx version should have an `#ifdef` that chooses between the two based on clang vs nvcc.


================
Comment at: openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h:81-82
+INLINE void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi) {
+  lo = (uint32_t)(val & 0x00000000FFFFFFFFL);
+  hi = (uint32_t)((val & 0xFFFFFFFF00000000L) >> 32);
+}
----------------
ABataev wrote:
> Maybe ULL suffix instead of just L?
Nice spot. Don't really want the sign conversions here, fixed. Although to use UINT64_C instead of the suffix.

Strictly speaking that should probably be 32u (or UINT32_C(32), but that feels like a step too far) as well. Thoughts?


================
Comment at: openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h:90
+static const __kmpc_impl_lanemask_t __kmpc_impl_all_lanes =
+    UINT64_C(0xffffffffffffffff);
+
----------------
ABataev wrote:
> 0xffffffffffffffffULL?
I'd prefer to use the UINT64_C notation. It's longer but it's harder to get wrong.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D69718