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
|
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -fsyntax-only -verify=host -xhip %s
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsyntax-only -fcuda-is-device -verify=dev -xhip %s
// dev-no-diagnostics
#include "Inputs/cuda.h"
__global__ void kernel() {
__attribute__((address_space(0))) void *mem_ptr;
(void)__builtin_amdgcn_is_shared(mem_ptr);
}
template<typename T>
__global__ void template_kernel(T *p) {
__attribute__((address_space(0))) void *mem_ptr;
(void)__builtin_amdgcn_is_shared(mem_ptr);
}
void hfun() {
__attribute__((address_space(0))) void *mem_ptr;
(void)__builtin_amdgcn_is_shared(mem_ptr); // host-error {{reference to __device__ function '__builtin_amdgcn_is_shared' in __host__ function}}
}
template<typename T>
void template_hfun(T *p) {
__attribute__((address_space(0))) void *mem_ptr;
(void)__builtin_amdgcn_is_shared(mem_ptr); // host-error {{reference to __device__ function '__builtin_amdgcn_is_shared' in __host__ function}}
}
int main() {
int *p;
kernel<<<1,1>>>();
template_kernel<<<1,1>>>(p);
template_hfun(p); // host-note {{called by 'main'}}
}
|
