[Openmp-commits] [PATCH] D95298: [OpenMP][libomptarget][WIP] Refined the file structure of `deviceRTLs`

[Jon Chesterfield via Phabricator via Openmp-commits]

JonChesterfield added a comment.

I really like about half of this change.

Creating a target_interface.h containing the common parts of the two target_impl is good. Maybe better located under common/ but doesn't matter much.

Expanding nvptx_interface.h and amdgcn_interface.h to contain the extra bits of target_interface.h seems reasonable, but it means a lot of symbols and definitions are pasted into the top of interface.h. It is then no longer a clear description of the interface to the libary.

Previously, interface.h was not included in the source files. There are pros and cons to that, overall it's probably better to have it in scope for the files that are implementing the functions within it as it increases the chance of the prototype matching the implementation.

Minimising the number of includes is not as good. It means that dropping debug.h, or dropping the include of target_interface from it, will break a bunch of files. Better for each source that calls functions defined in a header to include that header directly. In practice, this probably means writing `#include "target_interface"`at the top of most of them. It's slightly more typing at the top of the source files in exchange for a build that doesn't break when removing apparently dead includes from headers.

I think my concerns amount to putting #include target_interface everywhere that currently has #include target_impl, and keeping stuff that is not part of the library interface out of interface.h.

Top of interface.h could change to:

  #ifndef _INTERFACES_H_
  #define _INTERFACES_H_
  
  #include <stddef.h>
  #include <stdint.h>
  
  #ifdef __AMDGCN__ // probably variant at some point, same effect
  #define EXTERN extern "C" __attribute__((device))
  typedef uint64_t __kmpc_impl_lanemask_t;
  typedef uint32_t omp_lock_t; /* arbitrary type of the right length */
  EXTERN uint32_t __kmpc_amdgcn_gpu_num_threads();
  #endif
  
  #ifdef __CUDACC__
  #define EXTERN extern "C" __device__
  typedef uint32_t __kmpc_impl_lanemask_t;
  typedef uint32_t omp_lock_t; /* arbitrary type of the right length */
  #endif

That has the advantage of clarity, and the disadvantage of putting a bunch of code behind a macro. It'll cause the rocm compiler some grief as there are a bunch of extensions under amdgcn_interface.h at present.



================
Comment at: openmp/libomptarget/deviceRTLs/common/debug.h:32
 #include "common/device_environment.h"
+#include "target_interface.h"
+
----------------
A lot of files only compile after this change because they include debug and debug happens to include target_interface. I'd much rather all the files that use functions declared in target_interface continue to include the target_interface header directly.


================
Comment at: openmp/libomptarget/deviceRTLs/common/debug.h:34
+
+#include <cassert>
 
----------------
The amdgcn implementation is freestanding. At some point I moved the libc and libc++ headers until target_impl to make that work. This means it can build without libc++ and without libc, which is important as neither exist on the device and including bits of the host libc seems hazardous.


================
Comment at: openmp/libomptarget/deviceRTLs/common/omptarget.h:17
 
-#include "target_impl.h"
-#include "common/debug.h"     // debug
-#include "interface.h" // interfaces with omp, compiler, and user
+#include "common/debug.h" // debug
 #include "common/state-queue.h"
----------------
lets lose the "debug.h" // debug comment


================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h:12
 
-#include <stdint.h>
+#define DEVICE __device__
+#define EXTERN extern "C" DEVICE
----------------
stdint was a good thing, means the header is self contained. I.e. doesn't have to be fortunately included after stdint.h at the include site.


================
Comment at: openmp/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h:19
+
+////////////////////////////////////////////////////////////////////////////////
+// Kernel options
----------------
if nvtpx wants <cassert> or <stdlib>, please put them here


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D95298