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