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

[Arpith Jacob via Phabricator via Openmp-commits]

arpith-jacob added inline comments.


================
Comment at: libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h:100
+};
+extern __device__ __shared__ DataSharingStateTy DataSharingState;
+
----------------
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).
```




Repository:
  rL LLVM

https://reviews.llvm.org/D14254