1 // RUN: %clang_cc1 -x hip %s --hipstdpar -triple amdgcn-amd-amdhsa --std=c++17 \ 2 // RUN: -fcuda-is-device -emit-llvm -o /dev/null -verify 3 4 // Note: These would happen implicitly, within the implementation of the 5 // accelerator specific algorithm library, and not from user code. 6 7 // Calls from the accelerator side to implicitly host (i.e. unannotated) 8 // functions are fine. 9 10 // expected-no-diagnostics 11 12 #define __device__ __attribute__((device)) 13 #define __global__ __attribute__((global)) 14 host_fn()15extern "C" void host_fn() {} 16 17 struct Dummy {}; 18 19 struct S { SS20 S() {} ~SS21 ~S() { host_fn(); } 22 23 int x; 24 }; 25 26 struct T { hdT27 __device__ void hd() { host_fn(); } 28 29 __device__ void hd3(); 30 hT31 void h() {} 32 33 void operator+(); operator -T34 void operator-(const T&) {} 35 operator DummyT36 operator Dummy() { return Dummy(); } 37 }; 38 hd3()39__device__ void T::hd3() { host_fn(); } 40 hd2()41template <typename T> __device__ void hd2() { host_fn(); } 42 kernel()43__global__ void kernel() { hd2<int>(); } 44 hd()45__device__ void hd() { host_fn(); } 46 hd3()47template <typename T> __device__ void hd3() { host_fn(); } device_fn()48__device__ void device_fn() { hd3<int>(); } 49 local_var()50__device__ void local_var() { 51 S s; 52 } 53 explicit_destructor(S * s)54__device__ void explicit_destructor(S *s) { 55 s->~S(); 56 } 57 hd_member_fn()58__device__ void hd_member_fn() { 59 T t; 60 61 t.hd(); 62 } 63 h_member_fn()64__device__ void h_member_fn() { 65 T t; 66 t.h(); 67 } 68 unaryOp()69__device__ void unaryOp() { 70 T t; 71 (void) +t; 72 } 73 binaryOp()74__device__ void binaryOp() { 75 T t; 76 (void) (t - t); 77 } 78 implicitConversion()79__device__ void implicitConversion() { 80 T t; 81 Dummy d = t; 82 } 83 84 template <typename T> 85 struct TmplStruct { fnTmplStruct86 template <typename U> __device__ void fn() {} 87 }; 88 89 template <> 90 template <> fn()91__device__ void TmplStruct<int>::fn<int>() { host_fn(); } 92 double_specialization()93__device__ void double_specialization() { TmplStruct<int>().fn<int>(); } 94