From 9674a64cd9f83b6ad7b8b4d1dab82966ad7e050b Mon Sep 17 00:00:00 2001 From: Artem Belevich Date: Tue, 22 Sep 2015 17:23:05 +0000 Subject: [CUDA] Add appropriate host/device attribute to builtins. The changes are part of attribute-based CUDA function overloading (D12453) and as such are only enabled when it's in effect (-fcuda-target-overloads). Differential Revision: http://reviews.llvm.org/D12122 llvm-svn: 248296 --- clang/test/SemaCUDA/builtins.cu | 36 +++++++++++++++++++++++++++++++ clang/test/SemaCUDA/implicit-intrinsic.cu | 9 +++++--- 2 files changed, 42 insertions(+), 3 deletions(-) create mode 100644 clang/test/SemaCUDA/builtins.cu (limited to 'clang/test/SemaCUDA') diff --git a/clang/test/SemaCUDA/builtins.cu b/clang/test/SemaCUDA/builtins.cu new file mode 100644 index 00000000000..2c619b56339 --- /dev/null +++ b/clang/test/SemaCUDA/builtins.cu @@ -0,0 +1,36 @@ +// Tests that target-specific builtins have appropriate host/device +// attributes and that CUDA call restrictions are enforced. Also +// verify that non-target builtins can be used from both host and +// device functions. +// +// REQUIRES: x86-registered-target +// REQUIRES: nvptx-registered-target +// RUN: %clang_cc1 -triple x86_64-unknown-unknown \ +// RUN: -fcuda-target-overloads -fsyntax-only -verify %s +// RUN: %clang_cc1 -triple nvptx64-unknown-cuda -fcuda-is-device \ +// RUN: -fcuda-target-overloads -fsyntax-only -verify %s + + +#ifdef __CUDA_ARCH__ +// Device-side builtins are not allowed to be called from host functions. +void hf() { + int x = __builtin_ptx_read_tid_x(); // expected-note {{'__builtin_ptx_read_tid_x' declared here}} + // expected-error@-1 {{reference to __device__ function '__builtin_ptx_read_tid_x' in __host__ function}} + x = __builtin_abs(1); +} +__attribute__((device)) void df() { + int x = __builtin_ptx_read_tid_x(); + x = __builtin_abs(1); +} +#else +// Host-side builtins are not allowed to be called from device functions. +__attribute__((device)) void df() { + int x = __builtin_ia32_rdtsc(); // expected-note {{'__builtin_ia32_rdtsc' declared here}} + // expected-error@-1 {{reference to __host__ function '__builtin_ia32_rdtsc' in __device__ function}} + x = __builtin_abs(1); +} +void hf() { + int x = __builtin_ia32_rdtsc(); + x = __builtin_abs(1); +} +#endif diff --git a/clang/test/SemaCUDA/implicit-intrinsic.cu b/clang/test/SemaCUDA/implicit-intrinsic.cu index 3d24aa719e5..0793d64b101 100644 --- a/clang/test/SemaCUDA/implicit-intrinsic.cu +++ b/clang/test/SemaCUDA/implicit-intrinsic.cu @@ -1,10 +1,13 @@ -// RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fsyntax-only -verify %s +// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device \ +// RUN: -fsyntax-only -verify %s +// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device \ +// RUN: -fcuda-target-overloads -fsyntax-only -verify %s #include "Inputs/cuda.h" // expected-no-diagnostics __device__ void __threadfence_system() { - // This shouldn't produce an error, since __nvvm_membar_sys is inferred to - // be __host__ __device__ and thus callable from device code. + // This shouldn't produce an error, since __nvvm_membar_sys should be + // __device__ and thus callable from device code. __nvvm_membar_sys(); } -- cgit v1.2.3