[CUDA][HIP] Warn shared var initialization

[yaxunl]

Summary

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 a.cu

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


__global__ void foo() {
  __shared__ A a;
}

$nvcc -c a.cu
a.cu(10): 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.

[tra]

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]

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.

[jlebar]

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.

[yaxunl]

We need to revive this review since we found it is very inconvenient without this.

Also, nvcc only emits a warning for it

https://cuda.godbolt.org/z/AzrDTt

So I think maybe we even want to emit it as warning by default.

[tra]

NVCC is not always right.

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

I don't think I've seen the answer to my question above.

However, user can turn it into a warning by -Wno-error=cuda-shared-init.

I believe that producing the error by default is the correct thing to do for correctness sake. Your example on godbolt illustrates it quite clearly -- there is no initialization done whatsoever, which is contrary to what the written code intended to do. I am strongly against turning this into a warning by default.

If there's specific *need* for making it a warning in some cases, we can consider adding an *option* to make it a warning.

include/clang/Basic/DiagnosticSemaKinds.td

 def note_cuda_conflicting_device_function_declared_here : Note<
   "conflicting __device__ function declared here">;
 def err_cuda_device_exceptions : Error<
   "cannot use '%0' in "
   "%select{__device__|__global__|__host__|__host__ __device__}1 function">;
 def err_dynamic_var_init : Error<
     "dynamic initialization is not supported for "
     "__device__, __constant__, and __shared__ variables.">;
-def err_shared_var_init : Error<
-    "initialization is not supported for __shared__ variables.">;
+def warn_shared_var_init : Warning<
+    "initialization is not supported for __shared__ variables.">,
+    InGroup<DiagGroup<"cuda-shared-init">>, DefaultError;
 def err_device_static_local_var : Error<
     "within a %select{__device__|__global__|__host__|__host__ __device__}0 "
     "function, only __shared__ variables or const variables without device "
     "memory qualifier may be marked 'static'">;
 def err_cuda_vla : Error<
     "cannot use variable-length arrays in "
     "%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
 def err_cuda_extern_shared : Error<"__shared__ variable %0 cannot be 'extern'">;

lib/Sema/SemaCUDA.cpp

if (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
     // Also make sure that destructor, if there is one, is empty.
     if (AllowedInit)
       if (CXXRecordDecl *RD = VD->getType()->getAsCXXRecordDecl())
         AllowedInit =
             isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
 
     if (!AllowedInit) {
       Diag(VD->getLocation(), VD->hasAttr<CUDASharedAttr>()
-                                  ? diag::err_shared_var_init
+                                  ? diag::warn_shared_var_init
                                   : diag::err_dynamic_var_init)
           << Init->getSourceRange();
       VD->setInvalidDecl();
     }
   } else {
     // This is a host-side global variable.  Check that the initializer is
     // callable from the host side.
     const FunctionDecl *InitFn = nullptr;