|
150
|
1 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
|
|
|
2 // RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -verify -fcuda-is-device %s
|
|
|
3 //
|
|
|
4 // We run clang_cc1 with 'not' because source file contains
|
|
|
5 // intentional errors. CC1 failure is expected and must be ignored
|
|
|
6 // here. We're interested in what ends up in AST and that's what
|
|
|
7 // FileCheck verifies.
|
|
|
8 // RUN: not %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -ast-dump %s \
|
|
|
9 // RUN: | FileCheck %s --check-prefix=CHECK-ALL --check-prefix=CHECK-HOST
|
|
|
10 // RUN: not %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -ast-dump -fcuda-is-device %s \
|
|
|
11 // RUN: | FileCheck %s --check-prefix=CHECK-ALL --check-prefix=CHECK-DEVICE
|
|
|
12
|
|
|
13 #include "Inputs/cuda.h"
|
|
|
14
|
|
|
15 // Host (x86) supports TLS and device-side compilation should ignore
|
|
|
16 // host variables. No errors in either case.
|
|
|
17 int __thread host_tls_var;
|
|
|
18 // CHECK-ALL: host_tls_var 'int' tls
|
|
|
19
|
|
|
20 #if defined(__CUDA_ARCH__)
|
|
|
21 // NVPTX does not support TLS
|
|
|
22 __device__ int __thread device_tls_var; // expected-error {{thread-local storage is not supported for the current target}}
|
|
|
23 // CHECK-DEVICE: device_tls_var 'int' tls
|
|
|
24 __shared__ int __thread shared_tls_var; // expected-error {{thread-local storage is not supported for the current target}}
|
|
|
25 // CHECK-DEVICE: shared_tls_var 'int' tls
|
|
|
26 #else
|
|
|
27 // Device-side vars should not produce any errors during host-side
|
|
|
28 // compilation.
|
|
|
29 __device__ int __thread device_tls_var;
|
|
|
30 // CHECK-HOST: device_tls_var 'int' tls
|
|
|
31 __shared__ int __thread shared_tls_var;
|
|
|
32 // CHECK-HOST: shared_tls_var 'int' tls
|
|
|
33 #endif
|
|
|
34
|
|
|
35 __global__ void g1(int x) {}
|
|
|
36 __global__ int g2(int x) { // expected-error {{must have void return type}}
|
|
|
37 return 1;
|
|
|
38 }
|
