Page MenuHomePhabricator

[CUDA][HIP] Warn shared var initialization
Needs RevisionPublic

Authored by yaxunl on Mar 21 2019, 8:39 AM.



In many cases the default constructor of a class contains initializer of data
members, which allows concise code. The class may be instantiated as
global or automatic variables in device code, which is totally legal. However
when such a class is instantiated as a shared variable in device code, clang
will emit an error saying shared variables cannot be initialized.

Usually, users would like to just add explicit initialization for the shared variable,
instead of remove the initializer of data members from the default constructor,
since that would requires adding explicit initialization to all instances of the class,
even as global or automatic variables.

This requires the diagnostic of initialization of shared variable to be a warning,
instead of an error.

nvcc emits an warning for such situation, e.g.

$ cat

struct A {
  int a;
  __device__ A():a(0){}

__global__ void foo() {
  __shared__ A a;

$nvcc -c warning: __shared__ memory variable with non-empty constructor or destructor (potential race between threads)

This patch turns the diagnostic of initialization of shared varibles into a warning.
By default it is still treated as error, therefore no behavior change of clang. However,
user can turn it into a warning by -Wno-error=cuda-shared-init.

Diff Detail

rC Clang

Event Timeline

yaxunl created this revision.Mar 21 2019, 8:39 AM
Herald added a project: Restricted Project. · View Herald TranscriptMar 21 2019, 8:39 AM
tra added a comment.Mar 21 2019, 11:14 AM

This looks like one of the things we should *not* do as it affects correctness -- non-trivial constructor may be arbitrarily complex and the per-TU flag to enable this behavior is way too coarse, IMO.
On the other hand, I can believe that someone somewhere did write the code and relies to NVCC accepting it.

Is there a specific use case for this, other than matching nvcc bug-for-bug?

I've added Justin and Richard as reviewers for the language-lawyering input.

jlebar requested changes to this revision.Mar 21 2019, 11:17 AM

I agree with Art. The fact that nvcc allows this is broken.

If you want a flag that makes this error a warning, that might work for me. The flag should probably say "unsafe" or "I promise I will not complain when this breaks me" or something to that effect.

This revision now requires changes to proceed.Mar 21 2019, 11:17 AM

By default it is still treated as error, therefore no behavior change of clang.

Oh, I see, you already did what I'd suggested. :)

That's better. I think this needs to be made *much scarier* though. "Maybe race condition" doesn't capture the danger here -- you can very quickly get UB.

Maybe Richard has thoughts on whether we should allow broken things like this.