[PATCH] D59647: [CUDA][HIP] Warn shared var initialization

Yaxun Liu via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Thu Mar 21 08:39:38 PDT 2019


yaxunl created this revision.
yaxunl added a reviewer: tra.
Herald added a project: clang.

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.


Repository:
  rC Clang

https://reviews.llvm.org/D59647

Files:
  include/clang/Basic/DiagnosticSemaKinds.td
  lib/Sema/SemaCUDA.cpp


Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -499,7 +499,7 @@
 
     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();
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -7194,8 +7194,9 @@
 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 "


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D59647.191702.patch
Type: text/x-patch
Size: 1360 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20190321/b9b6098d/attachment-0001.bin>


More information about the cfe-commits mailing list