[PATCH] D20039: [CUDA] Restrict init of local __shared__ variables to empty constructors only.

Justin Lebar via cfe-commits cfe-commits at lists.llvm.org
Fri May 6 17:58:28 PDT 2016


jlebar added a comment.

While I think this is 100% the right thing to do, I am worried about breaking existing targets.  Maybe we need an escape valve, at least until we get that sorted out?  Unless you're pretty confident this isn't happening / will be easy enough to fix.


================
Comment at: lib/CodeGen/CGDecl.cpp:376
@@ +375,3 @@
+  // CUDA's local and local static __shared__ variables should not
+  // have any non-empty initializers which is ensured by Sema.
+  // Whatever initializer such variable may have when it gets here is
----------------
Please set off "which is ensured by Sema" somehow.  I'd probably say

> initializers.  (This is ensured by Sema.)

================
Comment at: lib/Sema/SemaDecl.cpp:10416
@@ -10415,2 +10415,3 @@
   // 7.5). CUDA also allows constant initializers for __constant__ and
-  // __device__ variables.
+  // __device__ variables.  We also must have the same checks applied
+  // to all __shared__ variables whether they are local or
----------------
s/have the same checks applied/apply the same checks


http://reviews.llvm.org/D20039





More information about the cfe-commits mailing list