94 lines · cpp
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 -verify3 4// Note: These would happen implicitly, within the implementation of the5// 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-diagnostics11 12#define __device__ __attribute__((device))13#define __global__ __attribute__((global))14 15extern "C" void host_fn() {}16 17struct Dummy {};18 19struct S {20 S() {}21 ~S() { host_fn(); }22 23 int x;24};25 26struct T {27 __device__ void hd() { host_fn(); }28 29 __device__ void hd3();30 31 void h() {}32 33 void operator+();34 void operator-(const T&) {}35 36 operator Dummy() { return Dummy(); }37};38 39__device__ void T::hd3() { host_fn(); }40 41template <typename T> __device__ void hd2() { host_fn(); }42 43__global__ void kernel() { hd2<int>(); }44 45__device__ void hd() { host_fn(); }46 47template <typename T> __device__ void hd3() { host_fn(); }48__device__ void device_fn() { hd3<int>(); }49 50__device__ void local_var() {51 S s;52}53 54__device__ void explicit_destructor(S *s) {55 s->~S();56}57 58__device__ void hd_member_fn() {59 T t;60 61 t.hd();62}63 64__device__ void h_member_fn() {65 T t;66 t.h();67}68 69__device__ void unaryOp() {70 T t;71 (void) +t;72}73 74__device__ void binaryOp() {75 T t;76 (void) (t - t);77}78 79__device__ void implicitConversion() {80 T t;81 Dummy d = t;82}83 84template <typename T>85struct TmplStruct {86 template <typename U> __device__ void fn() {}87};88 89template <>90template <>91__device__ void TmplStruct<int>::fn<int>() { host_fn(); }92 93__device__ void double_specialization() { TmplStruct<int>().fn<int>(); }94