| 1 | // Tests that host and target builtins can be used in the same TU, |
|---|
| 2 | // have appropriate host/device attributes and that CUDA call |
|---|
| 3 | // restrictions are enforced. Also verify that non-target builtins can |
|---|
| 4 | // be used from both host and device functions. |
|---|
| 5 | // |
|---|
| 6 | // REQUIRES: x86-registered-target |
|---|
| 7 | // REQUIRES: nvptx-registered-target |
|---|
| 8 | // RUN: %clang_cc1 -triple x86_64-unknown-unknown \ |
|---|
| 9 | // RUN: -aux-triple nvptx64-unknown-cuda \ |
|---|
| 10 | // RUN: -fsyntax-only -verify %s |
|---|
| 11 | // RUN: %clang_cc1 -triple nvptx64-unknown-cuda -fcuda-is-device \ |
|---|
| 12 | // RUN: -aux-triple x86_64-unknown-unknown \ |
|---|
| 13 | // RUN: -fsyntax-only -verify %s |
|---|
| 14 | |
|---|
| 15 | #if !(defined(__amd64__) && defined(__PTX__)) |
|---|
| 16 | #error "Expected to see preprocessor macros from both sides of compilation." |
|---|
| 17 | #endif |
|---|
| 18 | |
|---|
| 19 | void hf() { |
|---|
| 20 | int x = __builtin_ia32_rdtsc(); |
|---|
| 21 | int y = __nvvm_read_ptx_sreg_tid_x(); // expected-note {{'__nvvm_read_ptx_sreg_tid_x' declared here}} |
|---|
| 22 | // expected-error@-1 {{reference to __device__ function '__nvvm_read_ptx_sreg_tid_x' in __host__ function}} |
|---|
| 23 | x = __builtin_abs(1); |
|---|
| 24 | } |
|---|
| 25 | |
|---|
| 26 | __attribute__((device)) void df() { |
|---|
| 27 | int x = __nvvm_read_ptx_sreg_tid_x(); |
|---|
| 28 | int y = __builtin_ia32_rdtsc(); // expected-error {{reference to __host__ function '__builtin_ia32_rdtsc' in __device__ function}} |
|---|
| 29 | // expected-note@20 {{'__builtin_ia32_rdtsc' declared here}} |
|---|
| 30 | x = __builtin_abs(1); |
|---|
| 31 | } |
|---|
| 32 | |
|---|