| Artem Belevich | fa62ad4 | 2015-04-27 19:37:53 +0000 | [diff] [blame^] | 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 |
| Peter Collingbourne | 6ab610c | 2010-12-01 03:15:31 +0000 | [diff] [blame] | 3 | |
| Eli Bendersky | 3468d9d | 2014-04-28 22:21:28 +0000 | [diff] [blame] | 4 | #include "Inputs/cuda.h" |
| Peter Collingbourne | 6ab610c | 2010-12-01 03:15:31 +0000 | [diff] [blame] | 5 | |
| Artem Belevich | fa62ad4 | 2015-04-27 19:37:53 +0000 | [diff] [blame^] | 6 | // Host (x86) supports TLS and device-side compilation should ignore |
| 7 | // host variables. No errors in either case. |
| 8 | int __thread host_tls_var; |
| 9 | |
| 10 | #if defined(__CUDA_ARCH__) |
| 11 | // NVPTX does not support TLS |
| 12 | __device__ int __thread device_tls_var; // expected-error {{thread-local storage is not supported for the current target}} |
| 13 | __shared__ int __thread shared_tls_var; // expected-error {{thread-local storage is not supported for the current target}} |
| 14 | #else |
| 15 | // Device-side vars should not produce any errors during host-side |
| 16 | // compilation. |
| 17 | __device__ int __thread device_tls_var; |
| 18 | __shared__ int __thread shared_tls_var; |
| 19 | #endif |
| 20 | |
| Peter Collingbourne | 6ab610c | 2010-12-01 03:15:31 +0000 | [diff] [blame] | 21 | __global__ void g1(int x) {} |
| Peter Collingbourne | e8cfaf4 | 2010-12-12 23:02:57 +0000 | [diff] [blame] | 22 | __global__ int g2(int x) { // expected-error {{must have void return type}} |
| 23 | return 1; |
| 24 | } |