[PATCH] D41485: [OpenMP][libomptarget] Add data sharing support in libomptarget

[Jonas Hahnfeld via Phabricator via llvm-commits]

Hahnfeld added inline comments.


================
Comment at: libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu:35
 extern __device__ __shared__ uint32_t execution_param;
+__device__ __shared__ omptarget_nvptx_SharedArgs omptarget_nvptx_sharedArgs;
 
----------------
gtbercea wrote:
> Hahnfeld wrote:
> > Hahnfeld wrote:
> > > Hahnfeld wrote:
> > > > Hahnfeld wrote:
> > > > > Hahnfeld wrote:
> > > > > > I'm not sure if this needs to be //defined// here... Hopefully this lightens up after rebasing ;-)
> > > > > Yeah, I think this needs be be declared `extern` here and defined in `omp_data.cu`.
> > > > Actually, all of these `extern` declarations should go away because they already exist in `omptarget-nvptx.h`! We don't need to reiterate them here...
> > > I still think this won't work: If there are multiple translation units with `target` regions, the linker should complain about duplicate symbol definitions... I think that's why there is a separate file `omp_data.cu`.
> > (obviously only when building and using the bclib which we don't yet have in trunk; but I think we shouldn't introduce the problem if we can easily avoid it)
> Do you have a failing example for this?
> 
> Also I'm not sure what you want me to move where. The comments keep going back and forth.
Compiling and linking multiple source files with `target` regions should fail when using the bclib.

Just move the definition (ie without `extern`) to `omp_data.cu`. (Terminology taken from https://stackoverflow.com/a/1410632)


Repository:
  rOMP OpenMP

https://reviews.llvm.org/D41485