[Openmp-commits] [PATCH] D14254: [OpenMP] Initial implementation of OpenMP offloading library - libomptarget device RTLs.

[Guansong Zhang via Phabricator via Openmp-commits]

guansong added a comment.

Here is a general comment, we should follow Alex's suggestion to make a generic GPU device and extend that to NVPTX and AMDGCN. My current idea is:

1. think all the current spell of nvptx as generic_gpu
2. add template to some functions, and instantiate them with nvptx and amdgcn required types
3. for example the warp size in nvptx is 32 while in amdgcn is 64

Thoughts?



================
Comment at: libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h:100
+};
+extern __device__ __shared__ DataSharingStateTy DataSharingState;
+
----------------
grokos wrote:
> arpith-jacob wrote:
> > Hahnfeld wrote:
> > > arpith-jacob wrote:
> > > > grokos wrote:
> > > > > Hahnfeld wrote:
> > > > > > This won't compile with current Clang 5.0.0:
> > > > > > ```
> > > > > > error: __shared__ variable 'DataSharingState' cannot be 'extern'
> > > > > > ```
> > > > > This looks like a clang problem. Shared variables can be extern. This code can be compiled with clang 4.0, maybe we should submit a bug report for clang 5.0?
> > > > Clang doesn't support this attribute and it requires additional investigation.  For now, I would disable building the bclib version of libomptarget-nvptx so that the runtime library is built using nvcc.
> > > This was explicitly introduced in https://reviews.llvm.org/D25125. The documentation suggests that the error is correct:
> > > ```
> > > __device__, __shared__, and __constant__ variables cannot be defined as external using the extern keyword. The only exception is for dynamically allocated __shared__ variables as described in __shared__.
> > > ```
> > > (http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#qualifiers)
> > That restriction only applies in whole program compilation mode, not separate compilation mode.  We need to teach Clang/llvm to support separate compilation mode.
> > 
> > ```
> > When compiling in the whole program compilation mode (see the nvcc user manual for a description of this mode), __device__, __shared__, and __constant__ variables cannot be defined as external using the extern keyword. The only exception is for dynamically allocated __shared__ variables as described in __shared__.
> > 
> > When compiling in the separate compilation mode (see the nvcc user manual for a description of this mode), __device__, __shared__, and __constant__ variables can be defined as external using the extern keyword. nvlink will generate an error when it cannot find a definition for an external variable (unless it is a dynamically allocated __shared__ variable).
> > ```
> > 
> > 
> OK, for the time being the default is set to disabled so as not to create any problems while building the runtime.
I know this issue. We do have a patch on clang 6.0 to allow this. But I am not confident enough to submit it I guess. Who should we talk to?


Repository:
  rL LLVM

https://reviews.llvm.org/D14254