39 lines · plain
1// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s2// RUN: %clang_cc1 -triple nvptx-unknown-cuda -fsyntax-only -verify -fcuda-is-device %s3//4// We run clang_cc1 with 'not' because source file contains5// intentional errors. CC1 failure is expected and must be ignored6// here. We're interested in what ends up in AST and that's what7// FileCheck verifies.8// RUN: not %clang_cc1 -triple x86_64-unknown-linux-gnu -ast-dump %s \9// RUN: | FileCheck %s --check-prefix=CHECK-ALL --check-prefix=CHECK-HOST10// RUN: not %clang_cc1 -triple nvptx-unknown-cuda -ast-dump -fcuda-is-device %s \11// RUN: | FileCheck %s --check-prefix=CHECK-ALL --check-prefix=CHECK-DEVICE12 13#include "Inputs/cuda.h"14 15// Host (x86) supports TLS and device-side compilation should ignore16// host variables. No errors in either case.17int __thread host_tls_var;18// CHECK-ALL: host_tls_var 'int' tls19 20#if defined(__CUDA_ARCH__)21// NVPTX does not support TLS22__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' tls24__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' tls26#else27// Device-side vars should not produce any errors during host-side28// compilation.29__device__ int __thread device_tls_var;30// CHECK-HOST: device_tls_var 'int' tls31__shared__ int __thread shared_tls_var;32// CHECK-HOST: shared_tls_var 'int' tls33#endif34 35__global__ void g1(int x) {}36__global__ int g2(int x) { // expected-error {{must have void return type}}37 return 1;38}39