1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
|
// RUN: %clang_cc1 -triple x86_64 -verify=expected,dev \
// RUN: -verify-ignore-unexpected=note \
// RUN: -fopenmp -fopenmp-version=50 -fopenmp-targets=amdgcn-amd-amdhsa -o - %s
// RUN: %clang_cc1 -triple x86_64 -verify -verify-ignore-unexpected=note\
// RUN: -fopenmp -fopenmp-version=50 -fopenmp-targets=amdgcn-amd-amdhsa -o - -x c++ %s
// RUN: %clang_cc1 -triple x86_64 -verify=dev -verify-ignore-unexpected=note\
// RUN: -fcuda-is-device -o - %s
#if __CUDA__
#include "Inputs/cuda.h"
__device__ void cu_devf();
#endif
void bazz() {}
#pragma omp declare target to(bazz) device_type(nohost)
void bazzz() {bazz();}
#pragma omp declare target to(bazzz) device_type(nohost)
void any() {bazz();} // expected-error {{function with 'device_type(nohost)' is not available on host}}
void host1() {bazz();} // expected-error {{function with 'device_type(nohost)' is not available on host}}
#pragma omp declare target to(host1) device_type(host)
void host2() {bazz();} // expected-error {{function with 'device_type(nohost)' is not available on host}}
#pragma omp declare target to(host2)
void device() {host1();}
#pragma omp declare target to(device) device_type(nohost)
void host3() {host1();}
#pragma omp declare target to(host3)
#pragma omp declare target
void any1() {any();}
void any2() {host1();}
void any3() {device();} // expected-error {{function with 'device_type(nohost)' is not available on host}}
void any4() {any2();}
#pragma omp end declare target
void any5() {any();}
void any6() {host1();}
void any7() {device();} // expected-error {{function with 'device_type(nohost)' is not available on host}}
void any8() {any2();}
#if __CUDA__
void cu_hostf() { cu_devf(); } // dev-error {{no matching function for call to 'cu_devf'}}
__device__ void cu_devf2() { cu_hostf(); } // dev-error{{no matching function for call to 'cu_hostf'}}
#endif
|
