204 lines · plain
1// REQUIRES: amdgpu-registered-target2// 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 %s8// 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_RETURN15// 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_RETURN22// 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 %s28// 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 %s35// 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 %s42// 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 %s50// 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 %s57 58// expected-no-diagnostics59 60// Check handling of overriden, implicitly __host__ dtor (should emit as a61// 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 875 76// Check support for pure and deleted virtual functions77struct 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 894// 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 895// 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 >= 201103L104// 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_fma114// CXX14: call contract noundef half @llvm.fma.f16115__device__ double test_fma(_Float16 h, int i) {116 return fma(h, h, i);117}118 119#endif120 121// CHECK-LABEL: amdgpu_kernel void @_Z4kernPff122__global__ void kern(float *x, float y) {123 *x = sin(y);124}125 126// CHECK-LABEL: define{{.*}} i64 @_Z11test_size_tv127// CHECK: ret i64 8128__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_floorv135// CHECK: call {{.*}}double @llvm.floor.f64(double136__device__ float test_floor() {137 return floor(5);138}139 140// CHECK-LABEL: define{{.*}}@_Z8test_maxv141// CHECK: call {{.*}}double @llvm.maxnum.f64(double {{.*}}, double142__device__ float test_max() {143 return max(5, 6.0);144}145 146// CHECK-LABEL: define{{.*}}@_Z10test_isnanv147__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_malloc166// MALLOC: call {{.*}}ptr @malloc(i64167// MALLOC: call {{.*}}ptr @malloc(i64168// MALLOC-LABEL: define weak {{.*}}ptr @malloc(i64169// MALLOC: call i64 @__ockl_dm_alloc170// NOMALLOC: call void @llvm.trap171// MALLOC-ASAN-LABEL: define weak {{.*}}ptr @malloc(i64172// 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_free180// MALLOC: call {{.*}}void @free(ptr181// MALLOC: call {{.*}}void @free(ptr182// MALLOC-LABEL: define weak {{.*}}void @free(ptr183// MALLOC: call void @__ockl_dm_dealloc184// NOMALLOC: call void @llvm.trap185// MALLOC-ASAN-LABEL: define weak {{.*}}void @free(ptr186// 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_host194// MALLOC-HOST: call {{.*}}ptr @_Z6mallocm(i64195// MALLOC-HOST: call {{.*}}void @_Z4freePv(ptr196// MALLOC-HOST: call {{.*}}ptr @_Z6mallocm(i64197// MALLOC-HOST: call {{.*}}void @_Z4freePv(ptr198void test_malloc_host(void *a) {199 a = malloc(42);200 free(a);201 a = std::malloc(42);202 std::free(a);203}204