clang/test/SemaHipStdPar/device-can-call-host.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 -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
  15 extern "C" void host_fn() {}
  16
  17 struct Dummy {};
  18
  19 struct S {
  20   S() {}
  21   ~S() { host_fn(); }
  22
  23   int x;
  24 };
  25
  26 struct 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
  41 template <typename T> __device__ void hd2() { host_fn(); }
  42
  43 __global__ void kernel() { hd2<int>(); }
  44
  45 __device__ void hd() { host_fn(); }
  46
  47 template <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
  84 template <typename T>
  85 struct TmplStruct {
  86   template <typename U> __device__ void fn() {}
  87 };
  88
  89 template <>
  90 template <>
  91 __device__ void TmplStruct<int>::fn<int>() { host_fn(); }
  92
  93 __device__ void double_specialization() { TmplStruct<int>().fn<int>(); }