1 // RUN: %clang_cc1 %s -emit-llvm -o - -triple -amdgcn-amd-amdhsa \
2 // RUN: -target-cpu gfx906 -fcuda-is-device -x hip \
3 // RUN: | FileCheck --check-prefixes=COMMON,CRDIV %s
4 // RUN: %clang_cc1 %s -emit-llvm -o - -triple -amdgcn-amd-amdhsa \
5 // RUN: -target-cpu gfx906 -fcuda-is-device -x hip \
6 // RUN: -fno-hip-fp32-correctly-rounded-divide-sqrt \
7 // RUN: | FileCheck --check-prefixes=COMMON,NCRDIV %s
8
9 #include "Inputs/cuda.h"
10
11 typedef __attribute__(( ext_vector_type(4) )) float float4;
12
13 // COMMON-LABEL: @_Z11spscalardiv
14 // COMMON: fdiv{{.*}},
15 // NCRDIV: !fpmath ![[MD:[0-9]+]]
16 // CRDIV-NOT: !fpmath
spscalardiv(float a,float b)17 __device__ float spscalardiv(float a, float b) {
18 return a / b;
19 }
20
21 // COMMON-LABEL: @_Z11spvectordiv
22 // COMMON: fdiv{{.*}},
23 // NCRDIV: !fpmath ![[MD]]
24 // CRDIV-NOT: !fpmath
spvectordiv(float4 a,float4 b)25 __device__ float4 spvectordiv(float4 a, float4 b) {
26 return a / b;
27 }
28
29 // COMMON-LABEL: @_Z11dpscalardiv
30 // COMMON-NOT: !fpmath
dpscalardiv(double a,double b)31 __device__ double dpscalardiv(double a, double b) {
32 return a / b;
33 }
34
35 // COMMON-LABEL: @_Z12spscalarsqrt
36 // NCRDIV: call contract float @llvm.sqrt.f32(float %{{.+}}), !fpmath ![[MD:[0-9]+]]
37 // CRDIV: call contract float @llvm.sqrt.f32(float %{{.+}}){{$}}
spscalarsqrt(float a)38 __device__ float spscalarsqrt(float a) {
39 return __builtin_sqrtf(a);
40 }
41
42 // COMMON-LABEL: @_Z12dpscalarsqrt
43 // COMMON: call contract double @llvm.sqrt.f64(double %{{.+}}){{$}}
44 // COMMON-NOT: !fpmath
dpscalarsqrt(double a)45 __device__ double dpscalarsqrt(double a) {
46 return __builtin_sqrt(a);
47 }
48
49 // COMMON-LABEL: @_Z28test_builtin_elementwise_f32f
50 // NCRDIV: call contract float @llvm.sqrt.f32(float %{{.+}}), !fpmath ![[MD:[0-9]+]]
51 // CRDIV: call contract float @llvm.sqrt.f32(float %{{.+}}){{$}}
test_builtin_elementwise_f32(float a)52 __device__ float test_builtin_elementwise_f32(float a) {
53 return __builtin_elementwise_sqrt(a);
54 }
55
56 // COMMON-LABEL: @_Z28test_builtin_elementwise_f64d
57 // COMMON: call contract double @llvm.sqrt.f64(double %{{.+}}){{$}}
58 // COMMON-NOT: !fpmath
test_builtin_elementwise_f64(double a)59 __device__ double test_builtin_elementwise_f64(double a) {
60 return __builtin_elementwise_sqrt(a);
61 }
62
63 // NCRSQRT: ![[MD]] = !{float 2.500000e+00}
64