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

[Artem Belevich via cfe-commits]

tra added inline comments.

================
Comment at: lib/Sema/SemaCUDA.cpp:429-430
@@ +428,4 @@
+                                  CXXConstructorDecl *CD) {
+  if (!CD->isDefined() && CD->isTemplateInstantiation())
+    InstantiateFunctionDefinition(VarLoc, CD->getFirstDecl());
+
----------------
rsmith wrote:
> The function might still not be defined after this (if the template is not defined); you should presumably return `false` here in that case.
I don't think it's needed. If it's still not definied, it will be caught by hasTrivialBody() check below.

================
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;
+
----------------
rsmith wrote:
> 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 &&`?
NVCC produces an error (probably because it does not support c++14):
zz.cu(1): error: statement may not appear in a constexpr constructor

clang w/ this patch indeed considers it to be a non-empty initializer and produces an error.

I agree that allowing constant initializer is the right thing to do. Your example requires c++14, so there's no direct comparison with nvcc, but I think allowing it is indeed the right thing to do here.



Repository:
  rL LLVM

http://reviews.llvm.org/D15305