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