Page MenuHomePhabricator

[libomptarget][nfc] Make interface.h target independent
ClosedPublic

Authored by JonChesterfield on Oct 7 2019, 5:51 PM.

Details

Summary

[libomptarget][nfc] Make interface.h target independent

Move interface.h under a top level include directory.
Remove #includes to avoid the interface depending on the implementation.

Diff Detail

Event Timeline

JonChesterfield created this revision.Oct 7 2019, 5:51 PM
Herald added a project: Restricted Project. · View Herald TranscriptOct 7 2019, 5:51 PM
JonChesterfield marked 3 inline comments as done.Oct 7 2019, 6:06 PM
JonChesterfield added inline comments.
openmp/libomptarget/deviceRTLs/interface.h
20

__CUDACC__ seems the preferred way of spelling this constraint. The amdgcn instantiation would add an equivalent few lines to this header and otherwise leave it unchanged:

#ifdef __AMDGCN__
#define EXTERN extern "C" __attribute__((device))
typedef uint64_t __kmpc_impl_lanemask_t;
#endif
openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
37–41

Google suggests this is idiomatic cmake for 'find parent directory'

jdoerfert added a comment.EditedOct 7 2019, 6:27 PM

I like that we start to extract common code out of .../nvptx but I'm confused about the how. As I mention below, I expected "all generic code" to include target_impl.h in order to avoid target specific code (here the __device__ stuff) under guards. Maybe I miss something so please let me know what you think.


I was expecting something like:
include/{interface.h, option.h, ...} all contain an #include "target_impl.h" at the beginning (see below on two options for that).
Code like the __kmpc_impl_lanemask_t typedef and defines like NOINLINE are in this setup part of target_impl.h.

openmp/libomptarget/deviceRTLs/interface.h
24

I would have thought we include target_impl.h here to get the target specific stuff.
Either via

#ifdef __CUDACC__
#include "nvptx/target_impl.h"
#endif

or via cmake include path magic.

JonChesterfield added a comment.EditedOct 7 2019, 7:05 PM

I'm expecting most code moved under common to need #include "target_impl.h", located as you describe. loop.cu etc fit that category, a couple of files may not need any target stuff at all.

The library API is a flat description of everything the user (well, mostly the compiler) can find within the compiled library. It's currently called interface.h (I'd prefer libomptarget/deviceRTLs/deviceRTL.h). Including options.h or target_impl.h in the interface (as we have now) pulls in a bunch of macros and symbols that happen to be nearby.

Equivalent to this change would be creating a file nvptx/src/target_interface.h (or maybe nvptx_interface.h) containing the #define and typedef, then changing interface.h to:

#ifdef __CUDACC__
#include "../nvptx/src/target_interface.h"
#endif

Or s$src$include$, and optionally replace the relative include with more cmake.

I'd like to be able to #include "deviceRTL/interface.h" from test code and have no idea which instantiation it runs against, but it's not a disaster if the interface.h also needs a bunch of cmake wiring to locate the rest of said interface. My main goal here is to get the implementation symbols out of the library interface.

  • Move interface details into separate file

Sketch two. Thanks for the feedback - I like the separate file more. It provides an obvious place to put architecture specific extensions to the API. Avoided relative includes by dropping the include dir.

Harbormaster completed remote builds in B39144: Diff 223722.

I like this better. I'll review it tomorrow (I hope).

Ping. I'm hoping this is uncontentious.

jdoerfert accepted this revision.Oct 14 2019, 4:19 PM

Ping. I'm hoping this is uncontentious.

LGTM.

This revision is now accepted and ready to land.Oct 14 2019, 4:19 PM
This revision was automatically updated to reflect the committed changes.