clang/test/Headers/hip-header.hip

   1 // REQUIRES: amdgpu-registered-target
   2 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
   3 // RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
   4 // RUN:   -internal-isystem %S/Inputs/include \
   5 // RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
   6 // RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
   7 // RUN:   -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,NOMALLOC %s
   8 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
   9 // RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
  10 // RUN:   -internal-isystem %S/Inputs/include \
  11 // RUN:   -include cmath \
  12 // RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
  13 // RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
  14 // RUN:   -D__HIPCC_RTC__ | FileCheck %s  -check-prefixes=AMD_BOOL_RETURN
  15 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
  16 // RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
  17 // RUN:   -internal-isystem %S/Inputs/include \
  18 // RUN:   -include cmath \
  19 // RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
  20 // RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
  21 // RUN:   -D__HIPCC_RTC__ -DUSE_ISNAN_WITH_INT_RETURN | FileCheck %s -check-prefixes=AMD_INT_RETURN
  22 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
  23 // RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
  24 // RUN:   -internal-isystem %S/Inputs/include \
  25 // RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
  26 // RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
  27 // RUN:   -D__HIPCC_RTC__ -std=c++14 | FileCheck -check-prefixes=CHECK,CXX14 %s
  28 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
  29 // RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
  30 // RUN:   -internal-isystem %S/Inputs/include \
  31 // RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
  32 // RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
  33 // RUN:   -DHIP_VERSION_MAJOR=4 -DHIP_VERSION_MINOR=5 \
  34 // RUN:   -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,MALLOC %s
  35 // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
  36 // RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
  37 // RUN:   -internal-isystem %S/Inputs/include \
  38 // RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
  39 // RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
  40 // RUN:   -DHIP_VERSION_MAJOR=4 -DHIP_VERSION_MINOR=5 \
  41 // RUN:   -fsanitize=address -disable-llvm-passes -D__HIPCC_RTC__ \
  42 // RUN:   | FileCheck -check-prefixes=MALLOC-ASAN %s
  43
  44 // expected-no-diagnostics
  45
  46 // Check handling of overriden, implicitly __host__ dtor (should emit as a
  47 // nullptr to global)
  48
  49 struct vbase {
  50     virtual ~vbase();
  51 };
  52
  53 template<typename T>
  54 struct vderived : public vbase {
  55     ~vderived();
  56 };
  57
  58 template struct vderived<void>;
  59
  60 // CHECK: @_ZTV8vderivedIvE = weak_odr unnamed_addr addrspace(1) constant { [4 x ptr addrspace(1)] } zeroinitializer, comdat, align 8
  61
  62 // Check support for pure and deleted virtual functions
  63 struct base {
  64   __host__
  65   __device__
  66   virtual void pv() = 0;
  67   __host__
  68   __device__
  69   virtual void dv() = delete;
  70 };
  71 struct derived:base {
  72   __host__
  73   __device__
  74   virtual void pv() override {};
  75 };
  76 __device__ void test_vf() {
  77     derived d;
  78 }
  79 // CHECK: @_ZTV7derived = linkonce_odr unnamed_addr addrspace(1) constant { [4 x ptr addrspace(1)] } { [4 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) addrspacecast (ptr @_ZN7derived2pvEv to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @__cxa_deleted_virtual to ptr addrspace(1))] }, comdat, align 8
  80 // CHECK: @_ZTV4base = linkonce_odr unnamed_addr addrspace(1) constant { [4 x ptr addrspace(1)] } { [4 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) addrspacecast (ptr @__cxa_pure_virtual to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @__cxa_deleted_virtual to ptr addrspace(1))] }, comdat, align 8
  81 // CHECK: define{{.*}}void @__cxa_pure_virtual()
  82 // CHECK: define{{.*}}void @__cxa_deleted_virtual()
  83
  84 struct Number {
  85   __device__ Number(float _x) : x(_x) {}
  86   float x;
  87 };
  88
  89 #if __cplusplus >= 201103L
  90 // Check __hip::__numeric_type can be used with a class without default ctor.
  91 __device__ void test_numeric_type() {
  92   int x = __hip::__numeric_type<Number>::value;
  93 }
  94
  95 // ToDo: Fix __clang_hip_cmake.h to specialize __hip::is_arithmetic<_Float16>
  96 // to resolve fma(_Float16, _Float16, int) to fma(double, double, double)
  97 // instead of fma(_Float16, _Float16, _Float16).
  98
  99 // CXX14-LABEL: define{{.*}}@_Z8test_fma
 100 // CXX14: call contract half @llvm.fma.f16
 101 __device__ double test_fma(_Float16 h, int i) {
 102   return fma(h, h, i);
 103 }
 104
 105 #endif
 106
 107 // CHECK-LABEL: amdgpu_kernel void @_Z4kernPff
 108 __global__ void kern(float *x, float y) {
 109   *x = sin(y);
 110 }
 111
 112 // CHECK-LABEL: define{{.*}} i64 @_Z11test_size_tv
 113 // CHECK: ret i64 8
 114 __device__ size_t test_size_t() {
 115   return sizeof(size_t);
 116 }
 117
 118 // Check there is no ambiguity when calling overloaded math functions.
 119
 120 // CHECK-LABEL: define{{.*}}@_Z10test_floorv
 121 // CHECK: call {{.*}}double @llvm.floor.f64(double
 122 __device__ float test_floor() {
 123   return floor(5);
 124 }
 125
 126 // CHECK-LABEL: define{{.*}}@_Z8test_maxv
 127 // CHECK: call {{.*}}double @llvm.maxnum.f64(double {{.*}}, double
 128 __device__ float test_max() {
 129   return max(5, 6.0);
 130 }
 131
 132 // CHECK-LABEL: define{{.*}}@_Z10test_isnanv
 133 __device__ double test_isnan() {
 134   double r = 0;
 135   double d = 5.0;
 136   float f = 5.0;
 137
 138   // AMD_INT_RETURN: call i1 @llvm.is.fpclass.f32(float {{.*}}, i32 3)
 139   // AMD_BOOL_RETURN: call i1 @llvm.is.fpclass.f32(float {{.*}}, i32 3)
 140   r += isnan(f);
 141
 142   // AMD_INT_RETURN: call i1 @llvm.is.fpclass.f64(double {{.*}}, i32 3)
 143   // AMD_BOOL_RETURN: call i1 @llvm.is.fpclass.f64(double {{.*}}, i32 3)
 144   r += isnan(d);
 145
 146   return r ;
 147 }
 148
 149 // Check that device malloc and free do not conflict with std headers.
 150 #include <cstdlib>
 151 // CHECK-LABEL: define{{.*}}@_Z11test_malloc
 152 // CHECK: call {{.*}}ptr @malloc(i64
 153 // CHECK-LABEL: define weak {{.*}}ptr @malloc(i64
 154 // MALLOC:  call i64 @__ockl_dm_alloc
 155 // NOMALLOC:  call void @llvm.trap
 156 // MALLOC-ASAN-LABEL: define weak {{.*}}ptr @malloc(i64
 157 // MALLOC-ASAN:  call ptr @llvm.returnaddress(i32 0)
 158 // MALLOC-ASAN:  call i64 @__asan_malloc_impl(i64 {{.*}}, i64 {{.*}})
 159 __device__ void test_malloc(void *a) {
 160   a = malloc(42);
 161 }
 162
 163 // CHECK-LABEL: define{{.*}}@_Z9test_free
 164 // CHECK: call {{.*}}void @free(ptr
 165 // CHECK-LABEL: define weak {{.*}}void @free(ptr
 166 // MALLOC:  call void @__ockl_dm_dealloc
 167 // NOMALLOC: call void @llvm.trap
 168 // MALLOC-ASAN-LABEL: define weak {{.*}}void @free(ptr
 169 // MALLOC-ASAN:  call ptr @llvm.returnaddress(i32 0)
 170 // MALLOC-ASAN:  call void @__asan_free_impl(i64 {{.*}}, i64 {{.*}})
 171 __device__ void test_free(void *a) {
 172   free(a);
 173 }