llvm-project

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=host %s
11
// RUN: %clang_cc1 -triple nvptx64-unknown-cuda -fcuda-is-device \
12
// RUN:     -aux-triple x86_64-unknown-unknown \
13
// RUN:     -target-cpu sm_80 -target-feature +ptx70 \
14
// RUN:     -fsyntax-only -verify=dev %s
15

16
#if !(defined(__amd64__) && defined(__PTX__))
17
#error "Expected to see preprocessor macros from both sides of compilation."
18
#endif
19

20
void hf() {
21
  int x = __builtin_ia32_rdtsc();
22
  int y = __nvvm_read_ptx_sreg_tid_x();
23
  // host-error@-1 {{reference to __device__ function '__nvvm_read_ptx_sreg_tid_x' in __host__ function}}
24
  x = __builtin_abs(1);
25
}
26

27
__attribute__((device)) void df() {
28
  int x = __nvvm_read_ptx_sreg_tid_x();
29
  int y = __builtin_ia32_rdtsc(); // dev-error {{reference to __host__ function '__builtin_ia32_rdtsc' in __device__ function}}
30
  x = __builtin_abs(1);
31
}
32

33
#if __CUDA_ARCH__ >= 800
34
__attribute__((device)) void nvvm_async_copy(__attribute__((address_space(3))) void* dst,
35
                                             __attribute__((address_space(1))) const void* src) {
36
  __nvvm_cp_async_ca_shared_global_4(dst, src);
37
  __nvvm_cp_async_ca_shared_global_8(dst, src);
38
  __nvvm_cp_async_ca_shared_global_16(dst, src);
39
  __nvvm_cp_async_cg_shared_global_16(dst, src);
40

41
  __nvvm_cp_async_ca_shared_global_4(dst, src, 2);
42
  __nvvm_cp_async_ca_shared_global_8(dst, src, 2);
43
  __nvvm_cp_async_ca_shared_global_16(dst, src, 2);
44
  __nvvm_cp_async_cg_shared_global_16(dst, src, 2);
45

46
  __nvvm_cp_async_ca_shared_global_4(dst, src, 2, 3); // dev-error {{too many arguments to function call}}
47
  __nvvm_cp_async_ca_shared_global_8(dst, src, 2, 3); // dev-error {{too many arguments to function call}}
48
  __nvvm_cp_async_ca_shared_global_16(dst, src, 2, 3); // dev-error {{too many arguments to function call}}
49
  __nvvm_cp_async_cg_shared_global_16(dst, src, 2, 3); // dev-error {{too many arguments to function call}}
50
}
51
#endif
52