| 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 | } |
|---|
| 39 | |
|---|