[PATCH] D15305: [CUDA] Do not allow dynamic initialization of global device side variables.

[Justin Lebar via cfe-commits]

jlebar added a comment.

tra asked me to check for coverage.  Looks pretty good in that respect.


================
Comment at: include/clang/Basic/DiagnosticSemaKinds.td:6419
@@ +6418,3 @@
+    "dynamic initialization is not supported for "
+    "__device__, __constant__ and __shared__ variables.">;
+def err_shared_var_init : Error<
----------------
Nit, but, since we're all language nerds here, suggest adding an Oxford comma.

================
Comment at: lib/Sema/SemaCUDA.cpp:436
@@ +435,3 @@
+  if (CD->isTrivial())
+    return true;
+
----------------
The test passes if I comment out this if statement.  I'm not sure if that's expected; this may or may not be entirely covered below.

================
Comment at: lib/Sema/SemaCUDA.cpp:442
@@ +441,3 @@
+  // and the function body is an empty compound statement.
+  if (!(CD->isDefined() && CD->getNumParams() == 0 && CD->hasTrivialBody()))
+    return false;
----------------
Tests pass if I comment out the isDefined check.

================
Comment at: lib/Sema/SemaDecl.cpp:10183
@@ +10182,3 @@
+  // 7.5).  We also allow constant initializers for __constant__ and
+  // __device__ variables.
+  const Expr *Init = VD->getInit();
----------------
> We also allow constant initializers for __constant__ and __device__ variables.

Consider rephrasing this -- it sounds like this is a clang extension, but I just checked and it does not appear to be.

================
Comment at: lib/Sema/SemaDecl.cpp:10186
@@ +10185,3 @@
+  const bool IsGlobal = VD->hasGlobalStorage() && !VD->isStaticLocal();
+  if (Init && IsGlobal && getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+      (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
----------------
Test passes if I comment out IsGlobal or CUDAIsDevice.  (I'm not sure if you care to test the latter, but the former seems important.)


http://reviews.llvm.org/D15305