r236029 - [cuda] Preserve TLS storage class of host variable even if it's a

[Artem Belevich]

Author: tra
Date: Tue Apr 28 15:31:49 2015
New Revision: 236029

URL: http://llvm.org/viewvc/llvm-project?rev=236029&view=rev
Log:
[cuda] Preserve TLS storage class of host variable even if it's a
device-side compilation.

Modified:
    cfe/trunk/lib/Sema/SemaDecl.cpp
    cfe/trunk/test/SemaCUDA/qualifiers.cu

Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=236029&r1=236028&r2=236029&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Tue Apr 28 15:31:49 2015
@@ -5769,12 +5769,16 @@ Sema::ActOnVariableDeclarator(Scope *S,
            diag::err_thread_non_global)
         << DeclSpec::getSpecifierName(TSCS);
     else if (!Context.getTargetInfo().isTLSSupported()) {
-      if (getLangOpts().CUDA)
+      if (getLangOpts().CUDA) {
         // Postpone error emission until we've collected attributes required to
         // figure out whether it's a host or device variable and whether the
         // error should be ignored.
         EmitTLSUnsupportedError = true;
-      else
+        // We still need to mark the variable as TLS so it shows up in AST with
+        // proper storage class for other tools to use even if we're not going
+        // to emit any code for it.
+        NewVD->setTSCSpec(TSCS);
+      } else
         Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
              diag::err_thread_unsupported);
     } else

Modified: cfe/trunk/test/SemaCUDA/qualifiers.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/qualifiers.cu?rev=236029&r1=236028&r2=236029&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/qualifiers.cu (original)
+++ cfe/trunk/test/SemaCUDA/qualifiers.cu Tue Apr 28 15:31:49 2015
@@ -1,21 +1,35 @@
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -verify -fcuda-is-device %s
+//
+// We run clang_cc1 with 'not' because source file contains
+// intentional errors. CC1 failure is expected and must be ignored
+// here. We're interested in what ends up in AST and that's what
+// FileCheck verifies.
+// RUN: not %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -ast-dump %s \
+// RUN:   | FileCheck %s --check-prefix=CHECK-ALL --check-prefix=CHECK-HOST
+// RUN: not %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -ast-dump -fcuda-is-device %s \
+// RUN:   | FileCheck %s --check-prefix=CHECK-ALL --check-prefix=CHECK-DEVICE
 
 #include "Inputs/cuda.h"
 
 // Host (x86) supports TLS and device-side compilation should ignore
 // host variables. No errors in either case.
 int __thread host_tls_var;
+// CHECK-ALL: host_tls_var 'int' tls
 
 #if defined(__CUDA_ARCH__)
 // NVPTX does not support TLS
 __device__ int __thread device_tls_var; // expected-error {{thread-local storage is not supported for the current target}}
+// CHECK-DEVICE: device_tls_var 'int' tls
 __shared__ int __thread shared_tls_var; // expected-error {{thread-local storage is not supported for the current target}}
+// CHECK-DEVICE: shared_tls_var 'int' tls
 #else
 // Device-side vars should not produce any errors during host-side
 // compilation.
 __device__ int __thread device_tls_var;
+// CHECK-HOST: device_tls_var 'int' tls
 __shared__ int __thread shared_tls_var;
+// CHECK-HOST: shared_tls_var 'int' tls
 #endif
 
 __global__ void g1(int x) {}