• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
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 
hf()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 
df()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