[Openmp-dev] [RFC] Make shared variables work in opencl

[Jon Chesterfield via Openmp-dev]

Hello OpenMP dev,

Our deviceRTL has a collection of per-workgroup variables in omp_data.
Mostly integers, a couple of classes. These are __shared__ and at program
scope. It's great that they're all in a single source file. I am seeking
suggestions on how to make this work on OpenCL.

- This works on nvptx. I'm quite impressed by that.
- On OpenCL, we can use __local, but that doesn't compile. Local variables
can't be at program scope, and they can't have constructors. Primitive
types only
- In HIP, the branch I use implements this by inlining everything. Trunk
HIP doesn't support file scope shared variables yet. I could be persuaded
it never should.

I'm therefore interested in moving these __shared__ variables out of
program scope. Possibly to a function local variable at the top of the
device side call tree, passed down to the functions that use it by pointer.
Better ideas are very welcome - this feels like a problem opencl devs will
have run into before.

Thanks!

Jon

Closely related, workgroup scope memory is fast and therefore scarce, so
longer term it would be nice to specialise kernels based on which parts of
this state they actually use. Without just inlining everything.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/openmp-dev/attachments/20200123/dc1e2cd0/attachment.html>