dnrops
/
clang
mirror of https://github.com/microsoft/clang.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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

device-side compilation.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@236029 91177308-0d34-0410-b5e6-96231b3b80d8
This commit is contained in:
Artem Belevich 2015-04-28 20:31:49 +00:00
parent ce80aa27f0
commit 1a5e6db211
2 changed files with 20 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
lib/Sema/SemaDecl.cpp
View File

@ -5769,12 +5769,16 @@ Sema::ActOnVariableDeclarator(Scope *S, Declarator &D, DeclContext *DC,
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

14
test/SemaCUDA/qualifiers.cu
View File

@ -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) {}