[cfe-dev] [CUDA] Why "Disallow 'extern __shared__' variables"?

[Justin Lebar via cfe-dev]

Clang will accept

  "extern __shared__ int x[];"

The CUDA programming guide is not precise, but as I read it, we're
doing the right thing here:

  https://docs.nvidia.com/cuda/cuda-c-programming-guide/#shared

Given that "extern __shared__" means "get me a pointer to the
dynamically-allocated shared memory for this kernel," using a
non-array / non-pointer type would be...odd?

nvcc may accept it without the "[]", but this is a pretty low bar to
set for semantic correctness.  :)  clang already rejects plenty of
semantically-questionable code that nvcc accepts.

I agree the error should be better, though.  I'd be happy to review a patch?

> I was using it successfully, and I believe correctly, in both nvcc and clang-cuda to build multiple relocatable objects with extern shared memory references that linked to these globals with nvlink.

I am impressed and a little alarmed that you've gotten this to work
with clang.  There is code in clang that's explicitly incompatible
with multiple translation unit device-side compilation.  Off the top
of my head: We mark all device functions as internal, which we need in
order to get some key optimizations in llvm.  But this means that if
you declare a function it won't be emitted unless it's used.

This is something we've wanted to fix, but it's going to be a bit
tricky to do it in a way that doesn't regress performance for people
doing single-TU compilation.  We'll probably need to add a flag.  If
you're interested in making this work, I'd be happy to provide
guidance patching.

-Justin

On Fri, Oct 28, 2016 at 6:32 AM, Arpith C Jacob <acjacob at us.ibm.com> wrote:
> Hi,
>
> I was wondering why Clang-cuda now disallows adding 'extern' on __shared__
> variables with this patch:
> https://reviews.llvm.org/D25125?id=73133
>
> I was using it successfully, and I believe correctly, in both nvcc and
> clang-cuda to build multiple relocatable objects with extern shared memory
> references that linked to these globals with nvlink. Please let me know if
> you would like specific examples and I'll be happy to describe it.
>
> Thanks,
> Arpith