| Artem Belevich | 9674a64 | 2015-09-22 17:23:05 +0000 | [diff] [blame^] | 1 | // Tests that target-specific builtins have appropriate host/device |
| 2 | // attributes and that CUDA call restrictions are enforced. Also |
| 3 | // verify that non-target builtins can be used from both host and |
| 4 | // device functions. |
| 5 | // |
| 6 | // REQUIRES: x86-registered-target |
| 7 | // REQUIRES: nvptx-registered-target |
| 8 | // RUN: %clang_cc1 -triple x86_64-unknown-unknown \ |
| 9 | // RUN: -fcuda-target-overloads -fsyntax-only -verify %s |
| 10 | // RUN: %clang_cc1 -triple nvptx64-unknown-cuda -fcuda-is-device \ |
| 11 | // RUN: -fcuda-target-overloads -fsyntax-only -verify %s |
| 12 | |
| 13 | |
| 14 | #ifdef __CUDA_ARCH__ |
| 15 | // Device-side builtins are not allowed to be called from host functions. |
| 16 | void hf() { |
| 17 | int x = __builtin_ptx_read_tid_x(); // expected-note {{'__builtin_ptx_read_tid_x' declared here}} |
| 18 | // expected-error@-1 {{reference to __device__ function '__builtin_ptx_read_tid_x' in __host__ function}} |
| 19 | x = __builtin_abs(1); |
| 20 | } |
| 21 | __attribute__((device)) void df() { |
| 22 | int x = __builtin_ptx_read_tid_x(); |
| 23 | x = __builtin_abs(1); |
| 24 | } |
| 25 | #else |
| 26 | // Host-side builtins are not allowed to be called from device functions. |
| 27 | __attribute__((device)) void df() { |
| 28 | int x = __builtin_ia32_rdtsc(); // expected-note {{'__builtin_ia32_rdtsc' declared here}} |
| 29 | // expected-error@-1 {{reference to __host__ function '__builtin_ia32_rdtsc' in __device__ function}} |
| 30 | x = __builtin_abs(1); |
| 31 | } |
| 32 | void hf() { |
| 33 | int x = __builtin_ia32_rdtsc(); |
| 34 | x = __builtin_abs(1); |
| 35 | } |
| 36 | #endif |