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

[Artem Belevich via cfe-commits]

tra marked 3 inline comments as done.

================
Comment at: lib/Sema/SemaCUDA.cpp:436
@@ +435,3 @@
+  if (CD->isTrivial())
+    return true;
+
----------------
jlebar wrote:
> 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.
According to [[ http://en.cppreference.com/w/cpp/language/default_constructor#Trivial_default_constructor | CPP reference ]] trivial constructor will pass all other checks 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;
----------------
jlebar wrote:
> Tests pass if I comment out the isDefined check.
hasTrivialBody() would only return true if we have a body which only happens if function is defined. isDefined() is mostly for readability here.

================
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>() ||
----------------
jlebar wrote:
> 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.)
IsGlobal -- all test cases were using either global or local variables. I've added a static __shared__ variable in the device function. Now IsGlobal check (or, rather !isStaticLocal() part of it) is required in order for the tests to succeed.

CUDAIsDevice is not triggered because all test cases are run with -fcuda-is-device.
It's hard to run host-side test with -verify here because I'd have to put #ifdef around every  'expected-error'



http://reviews.llvm.org/D15305