clang/test/SemaHipStdPar/device-can-call-host.cpp

// RUN: %clang_cc1 -x hip %s --hipstdpar -triple amdgcn-amd-amdhsa --std=c++17 \
// RUN:   -fcuda-is-device -emit-llvm -o /dev/null -verify

// Note: These would happen implicitly, within the implementation of the
//       accelerator specific algorithm library, and not from user code.

// Calls from the accelerator side to implicitly host (i.e. unannotated)
// functions are fine.

// expected-no-diagnostics

#define __device__ __attribute__((device))
#define __global__ __attribute__((global))

extern "C" void host_fn() {}

struct Dummy {};

struct S {
  S() {}
  ~S() { host_fn(); }

  int x;
};

struct T {
  __device__ void hd() { host_fn(); }

  __device__ void hd3();

  void h() {}

  void operator+();
  void operator-(const T&) {}

  operator Dummy() { return Dummy(); }
};

__device__ void T::hd3() { host_fn(); }

template <typename T> __device__ void hd2() { host_fn(); }

__global__ void kernel() { hd2<int>(); }

__device__ void hd() { host_fn(); }

template <typename T> __device__ void hd3() { host_fn(); }
__device__ void device_fn() { hd3<int>(); }

__device__ void local_var() {
  S s;
}

__device__ void explicit_destructor(S *s) {
  s->~S();
}

__device__ void hd_member_fn() {
  T t;

  t.hd();
}

__device__ void h_member_fn() {
  T t;
  t.h();
}

__device__ void unaryOp() {
  T t;
  (void) +t;
}

__device__ void binaryOp() {
  T t;
  (void) (t - t);
}

__device__ void implicitConversion() {
  T t;
  Dummy d = t;
}

template <typename T>
struct TmplStruct {
  template <typename U> __device__ void fn() {}
};

template <>
template <>
__device__ void TmplStruct<int>::fn<int>() { host_fn(); }

__device__ void double_specialization() { TmplStruct<int>().fn<int>(); }