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