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 %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 29// expected-no-diagnostics 30 31// Check support for pure and deleted virtual functions 32struct base { 33 __host__ 34 __device__ 35 virtual void pv() = 0; 36 __host__ 37 __device__ 38 virtual void dv() = delete; 39}; 40struct derived:base { 41 __host__ 42 __device__ 43 virtual void pv() override {}; 44}; 45__device__ void test_vf() { 46 derived d; 47} 48// CHECK: @_ZTV7derived = linkonce_odr unnamed_addr addrspace(1) constant { [4 x i8*] } { [4 x i8*] [i8* null, i8* null, i8* bitcast (void (%struct.derived*)* @_ZN7derived2pvEv to i8*), i8* bitcast (void ()* @__cxa_deleted_virtual to i8*)] }, comdat, align 8 49// CHECK: @_ZTV4base = linkonce_odr unnamed_addr addrspace(1) constant { [4 x i8*] } { [4 x i8*] [i8* null, i8* null, i8* bitcast (void ()* @__cxa_pure_virtual to i8*), i8* bitcast (void ()* @__cxa_deleted_virtual to i8*)] }, comdat, align 8 50 51// CHECK: define{{.*}}void @__cxa_pure_virtual() 52// CHECK: define{{.*}}void @__cxa_deleted_virtual() 53 54struct Number { 55 __device__ Number(float _x) : x(_x) {} 56 float x; 57}; 58 59#if __cplusplus >= 201103L 60// Check __hip::__numeric_type can be used with a class without default ctor. 61__device__ void test_numeric_type() { 62 int x = __hip::__numeric_type<Number>::value; 63} 64 65// ToDo: Fix __clang_hip_cmake.h to specialize __hip::is_arithmetic<_Float16> 66// to resolve fma(_Float16, _Float16, int) to fma(double, double, double) 67// instead of fma(_Float16, _Float16, _Float16). 68 69// CXX14-LABEL: define{{.*}}@_Z8test_fma 70// CXX14: call {{.*}}@__ocml_fma_f16 71__device__ double test_fma(_Float16 h, int i) { 72 return fma(h, h, i); 73} 74 75#endif 76 77// CHECK-LABEL: amdgpu_kernel void @_Z4kernPff 78__global__ void kern(float *x, float y) { 79 *x = sin(y); 80} 81 82// CHECK-LABEL: define{{.*}} i64 @_Z11test_size_tv 83// CHECK: ret i64 8 84__device__ size_t test_size_t() { 85 return sizeof(size_t); 86} 87 88// Check there is no ambiguity when calling overloaded math functions. 89 90// CHECK-LABEL: define{{.*}}@_Z10test_floorv 91// CHECK: call {{.*}}double @__ocml_floor_f64(double 92__device__ float test_floor() { 93 return floor(5); 94} 95 96// CHECK-LABEL: define{{.*}}@_Z8test_maxv 97// CHECK: call {{.*}}double @__ocml_fmax_f64(double {{.*}}, double 98__device__ float test_max() { 99 return max(5, 6.0); 100} 101 102// CHECK-LABEL: define{{.*}}@_Z10test_isnanv 103__device__ double test_isnan() { 104 double r = 0; 105 double d = 5.0; 106 float f = 5.0; 107 108 // AMD_INT_RETURN: call i32 @__ocml_isnan_f32(float 109 // AMD_BOOL_RETURN: call i32 @__ocml_isnan_f32(float 110 r += isnan(f); 111 112 // AMD_INT_RETURN: call i32 @__ocml_isnan_f64(double 113 // AMD_BOOL_RETURN: call i32 @__ocml_isnan_f64(double 114 r += isnan(d); 115 116 return r ; 117} 118 119// Check that device malloc and free do not conflict with std headers. 120#include <cstdlib> 121// CHECK-LABEL: define{{.*}}@_Z11test_malloc 122// CHECK: call {{.*}}i8* @malloc(i64 123// CHECK: define weak {{.*}}i8* @malloc(i64 124__device__ void test_malloc(void *a) { 125 a = malloc(42); 126} 127 128// CHECK-LABEL: define{{.*}}@_Z9test_free 129// CHECK: call {{.*}}i8* @free(i8* 130// CHECK: define weak {{.*}}i8* @free(i8* 131__device__ void test_free(void *a) { 132 free(a); 133} 134