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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
