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

[Richard Smith via cfe-commits]

rsmith accepted this revision.
rsmith added a comment.
This revision is now accepted and ready to land.

Some minor things, but feel free to commit after addressing them.

I agree that we should figure out what to do about the zero/undef initialization separately.


================
Comment at: lib/Sema/SemaCUDA.cpp:429-430
@@ +428,4 @@
+                                  CXXConstructorDecl *CD) {
+  if (!CD->isDefined() && CD->isTemplateInstantiation())
+    InstantiateFunctionDefinition(VarLoc, CD->getFirstDecl());
+
----------------
The function might still not be defined after this (if the template is not defined); you should presumably return `false` here in that case.

================
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;
----------------
Please do remove the `isDefined` check here. Including it makes a reader wonder what case it's trying to handle.

================
Comment at: lib/Sema/SemaCUDA.cpp:455-457
@@ +454,5 @@
+
+  // Its class has no virtual functions and no virtual base classes.
+  if (CD->getParent()->isDynamicClass())
+    return false;
+
----------------
Maybe reorder this before the `CXXCtorInitializer` check? It's a much cheaper test.

================
Comment at: lib/Sema/SemaDecl.cpp:10191-10198
@@ +10190,10 @@
+      bool AllowedInit = false;
+      if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
+        AllowedInit =
+            isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
+      else if ((VD->hasAttr<CUDADeviceAttr>() ||
+                VD->hasAttr<CUDAConstantAttr>()) &&
+               VD->getInit()->isConstantInitializer(
+                   Context, VD->getType()->isReferenceType()))
+        AllowedInit = true;
+
----------------
What should happen if the init is a constant initializer that is a `CXXConstructExpr`, but it uses a constructor that is not empty from CUDA's perspective? Such as:

  struct X { constexpr X() { int n = 0; } };
  __device__ X x;

I would assume this should be valid, but I think you'll reject it. Maybe change `else if (` to `if (!AllowedInit &&`?

================
Comment at: lib/Sema/SemaDecl.cpp:10196-10198
@@ +10195,5 @@
+                VD->hasAttr<CUDAConstantAttr>()) &&
+               VD->getInit()->isConstantInitializer(
+                   Context, VD->getType()->isReferenceType()))
+        AllowedInit = true;
+
----------------
Might be clearer as

  if (__device__ || __constant__)
    AllowedInit = isConstantInitializer(...)


http://reviews.llvm.org/D15305