52 lines · plain
1// Tests that host and target builtins can be used in the same TU,2// have appropriate host/device attributes and that CUDA call3// restrictions are enforced. Also verify that non-target builtins can4// be used from both host and device functions.5//6// REQUIRES: x86-registered-target7// REQUIRES: nvptx-registered-target8// RUN: %clang_cc1 -triple x86_64-unknown-unknown \9// RUN: -aux-triple nvptx64-unknown-cuda \10// RUN: -fsyntax-only -verify=host %s11// 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 %s15 16#if !(defined(__amd64__) && defined(__PTX__))17#error "Expected to see preprocessor macros from both sides of compilation."18#endif19 20void 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__ >= 80034__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#endif52