37 lines · plain
1// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -fsyntax-only -verify=host -xhip %s2// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsyntax-only -fcuda-is-device -verify=dev -xhip %s3 4// dev-no-diagnostics5 6#include "Inputs/cuda.h"7 8__global__ void kernel() { 9 __attribute__((address_space(0))) void *mem_ptr;10 (void)__builtin_amdgcn_is_shared(mem_ptr);11}12 13template<typename T>14__global__ void template_kernel(T *p) { 15 __attribute__((address_space(0))) void *mem_ptr;16 (void)__builtin_amdgcn_is_shared(mem_ptr);17}18 19void hfun() {20 __attribute__((address_space(0))) void *mem_ptr;21 (void)__builtin_amdgcn_is_shared(mem_ptr); // host-error {{reference to __device__ function '__builtin_amdgcn_is_shared' in __host__ function}}22}23 24template<typename T>25void template_hfun(T *p) {26 __attribute__((address_space(0))) void *mem_ptr;27 (void)__builtin_amdgcn_is_shared(mem_ptr); // host-error {{reference to __device__ function '__builtin_amdgcn_is_shared' in __host__ function}}28}29 30 31int main() {32 int *p;33 kernel<<<1,1>>>();34 template_kernel<<<1,1>>>(p);35 template_hfun(p); // host-note {{called by 'main'}}36}37