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