[cuda] Preserve TLS storage class of host variable during device-side compilation.

[tra]

Summary

Side effect of r235907 is that host-side TLS variables ended up in AST as non-TLS during device-side compilation.
While missing thread storage class spec does not affect generated code because the variable will be ignored
during code generation, we've lost information which is useful during source analysis.

This patch preserves host variable's TLS class during device-side compilation.

[eliben]

Is there any way to create a test that would catch this?

[tra]

In D9327#162710, @eliben wrote:

Is there any way to create a test that would catch this?

I can check what ends up in AST. Stay tuned.

[tra]

Added a test case to verify that TLS info is preserved in AST.

[eliben]

test/SemaCUDA/qualifiers.cu

3

'not' RUNs are sometimes hard to decipher, so a comment could help here - just to explain what the test is doing and what's the result it expects

[eliben]

LGTM

[with a small nit]

[tra]

Added a comment explaining why some test runs use 'not' to run cc1.

[jpienaar]

LGTM

[tra]

Added better explanation of the test.

[tra]

test/SemaCUDA/qualifiers.cu

3

Done.

[tra]

Committed : http://reviews.llvm.org/rL236029

lib/Sema/SemaDecl.cpp

if (DeclSpec::TSCS TSCS = D.getDeclSpec().getThreadStorageClassSpec()) {
     if (NewVD->hasLocalStorage() &&
         (SCSpec != DeclSpec::SCS_unspecified ||
          TSCS != DeclSpec::TSCS_thread_local ||
          !DC->isFunctionOrMethod()))
       Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
            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
       NewVD->setTSCSpec(TSCS);
   }
 
   // C99 6.7.4p3
   //   An inline definition of a function with external linkage shall

test/SemaCUDA/qualifiers.cu

 // 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) {}
 __global__ int g2(int x) { // expected-error {{must have void return type}}
   return 1;
 }