[PATCH] D42923: [CUDA] Allow external variables in separate compilation

[Jonas Hahnfeld via Phabricator via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rL325136: [CUDA] Allow external variables in separate 
compilation (authored by Hahnfeld, committed by ).
Herald added a subscriber: llvm-commits.

Changed prior to commit:
  https://reviews.llvm.org/D42923?vs=132866&id=134231#toc

Repository:
  rL LLVM

https://reviews.llvm.org/D42923

Files:
  cfe/trunk/lib/Sema/SemaDeclAttr.cpp
  cfe/trunk/test/SemaCUDA/extern-shared.cu


Index: cfe/trunk/test/SemaCUDA/extern-shared.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/extern-shared.cu
+++ cfe/trunk/test/SemaCUDA/extern-shared.cu
@@ -1,6 +1,11 @@
 // RUN: %clang_cc1 -fsyntax-only -verify %s
 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
 
+// RUN: %clang_cc1 -fsyntax-only -fcuda-rdc -verify=rdc %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -fcuda-rdc -verify=rdc %s
+// These declarations are fine in separate compilation mode:
+// rdc-no-diagnostics
+
 #include "Inputs/cuda.h"
 
 __device__ void foo() {
Index: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp
@@ -4112,7 +4112,8 @@
   auto *VD = cast<VarDecl>(D);
   // extern __shared__ is only allowed on arrays with no length (e.g.
   // "int x[]").
-  if (VD->hasExternalStorage() && !isa<IncompleteArrayType>(VD->getType())) {
+  if (!S.getLangOpts().CUDARelocatableDeviceCode && VD->hasExternalStorage() &&
+      !isa<IncompleteArrayType>(VD->getType())) {
     S.Diag(Attr.getLoc(), diag::err_cuda_extern_shared) << VD;
     return;
   }



Index: cfe/trunk/test/SemaCUDA/extern-shared.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/extern-shared.cu
+++ cfe/trunk/test/SemaCUDA/extern-shared.cu
@@ -1,6 +1,11 @@
 // RUN: %clang_cc1 -fsyntax-only -verify %s
 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
 
+// RUN: %clang_cc1 -fsyntax-only -fcuda-rdc -verify=rdc %s
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -fcuda-rdc -verify=rdc %s
+// These declarations are fine in separate compilation mode:
+// rdc-no-diagnostics
+
 #include "Inputs/cuda.h"
 
 __device__ void foo() {
Index: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp
@@ -4112,7 +4112,8 @@
   auto *VD = cast<VarDecl>(D);
   // extern __shared__ is only allowed on arrays with no length (e.g.
   // "int x[]").
-  if (VD->hasExternalStorage() && !isa<IncompleteArrayType>(VD->getType())) {
+  if (!S.getLangOpts().CUDARelocatableDeviceCode && VD->hasExternalStorage() &&
+      !isa<IncompleteArrayType>(VD->getType())) {
     S.Diag(Attr.getLoc(), diag::err_cuda_extern_shared) << VD;
     return;
   }

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits