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

[Artem Belevich via cfe-commits]

tra added a comment.

Richard, I've updated the patch as you've suggested -- it indeed simplifies things quite a bit and handles the corner cases you've mentioned.


================
Comment at: lib/CodeGen/CGDeclCXX.cpp:323-324
@@ +322,4 @@
+
+  // The constructor function has no parameters,
+  if (CD->getNumParams() != 0)
+    return false;
----------------
rsmith wrote:
> What if the constructor is a C-style varargs function:
> 
>   struct X { X(...) {} };
> 
> ?
CUDA does not support varargs on device side. nvcc fails with an error:

> error: a "device" function cannot have ellipsis

That's another thing I'll need to fix (as a separate patch) as clang currently accepts varargs everywhere.

This patch will ignore number of arguments passed to varargs constructor, but the checks for empty body still do apply.

================
Comment at: lib/CodeGen/CGDeclCXX.cpp:329
@@ +328,3 @@
+  for (const CXXCtorInitializer *CI: CD->inits())
+    if (CI->isAnyMemberInitializer() && CI->isWritten())
+      return false;
----------------
rsmith wrote:
> tra wrote:
> > @rsmith: is this a good way to find member initializer list items?
> > 
> > ```
> > struct S {
> >     int a,b,c;
> >     S() : a(1),b(2),c(3) {}
> > };
> > ```
> > I'm looking for a(),b(),c() which is what I think CUDA spec wants to check for, but CD->inits() appears to have other initializers on the list as well.
> You shouldn't need to check `isAnyMemberInitializer`: if there's any written inits, the constructor violates the rules.
As it turns out, the rules don't apply to all written initializers. For instance, nvcc allows empty constructor on init list:
```
struct A {  __device__ A(){}; };
struct B {  __device__ B(){}; };

struct C : A {
  B b;
  __device__ C() : A(), b() {}
};

__device__ C c;
```

I've simplified the patch so that in only checks for constructor's 'emptiness', but disregards how that constructor gets to be executed.

================
Comment at: lib/CodeGen/CGDeclCXX.cpp:333
@@ +332,3 @@
+  // and the function body is an empty compound statement.
+  // That does not always work.
+  if (!CD->hasTrivialBody())
----------------
rsmith wrote:
> What doesn't always work?
It was leftover from early patch variant that didn't defer emitting global vars.
If I don't defer and need to emit a global var before constructor definition is available, hasTrivialBody() returns false and triggers diagnostics. 

================
Comment at: lib/CodeGen/CGDeclCXX.cpp:347-367
@@ +346,23 @@
+
+  // The default constructors of all base classes of its class can be
+  // considered empty.
+  for (auto &Base : RD->bases())
+    if (hasNonEmptyDefaultConstructors(*this,
+                                       Base.getType()->getAsCXXRecordDecl()))
+      return false;
+
+  // For all the nonstatic data members of its class that are of class type
+  // (or array thereof), the default constructors can be considered empty.
+  for (const auto *I : RD->decls())
+    if (const FieldDecl *V = dyn_cast<FieldDecl>(I)) {
+      QualType T = V->getType();
+
+      if (const ArrayType *Ty = dyn_cast<ArrayType>(T))
+        while ((Ty = dyn_cast<ArrayType>(T)))
+          T = Ty->getElementType();
+
+      if (const CXXRecordDecl *R = T->getAsCXXRecordDecl())
+        if (hasNonEmptyDefaultConstructors(*this, R))
+          return false;
+    }
+
----------------
rsmith wrote:
> Rather than checking these properties this way, I'd suggest you check the initialization expression in each `CXXCtorInitializer` only contains `CXXConstructExpr`s for empty constructors (or any other whitelisted constructs). Your current approach will miss a couple of cases which the CUDA spec misses but presumably meant to exclude:
> 
> 1) Default member initializers
> 
>   int f();
>   struct X { int n = f(); X() {} };
> 
> 2) Cases where a constructor other than a default constructor is implicitly invoked
> 
>   struct A { template<typename ...T> A(T...); };
>   struct B : A { B() {} };
Nice. This has simplified the checks a lot.

================
Comment at: lib/CodeGen/CodeGenModule.cpp:1347-1351
@@ -1346,2 +1346,7 @@
     return false;
+  // Delay codegen for device-side CUDA variables. We need to have all
+  // constructor definitions available before we can determine whether
+  // we can skip them or produce an error.
+  if (LangOpts.CUDA && LangOpts.CUDAIsDevice && isa<VarDecl>(Global))
+    return false;
 
----------------
rsmith wrote:
> According to the quoted specification, you're supposed to check whether the constructor can be considered empty at the point in the translation unit where the definition of the variable occurs, so I don't think you need to delay anything.
I guess it's a bug in their guide as nvcc accepts following code with constructor definition appearing *after* the variable:

```
struct S {  S(); };
__device__ S s;
S::S() {}

```




http://reviews.llvm.org/D15305