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

[Jon Chesterfield via Phabricator via Openmp-commits]

JonChesterfield marked 4 inline comments as done.
JonChesterfield added a comment.

In D65836#1618858 <https://reviews.llvm.org/D65836#1618858>, @ABataev wrote:

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


Part of the motivation behind this change is that smaller diffs are easier to discuss. Hopefully this contributes to reaching said agreement. I think moving inline nvptx behind an interface is prerequisite for any movement towards sharing code between architectures.



================
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) {
----------------
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.


================
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));
----------------
ABataev wrote:
> Why pointers? Use references.
Habit. References look like pass by value at the call site so I tend to write out parameters as pointers. Changed.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D65836