|
221
|
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
|
|
|
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
|
|
|
25 __device__ float4 spvectordiv(float4 a, float4 b) {
|
|
|
26 return a / b;
|
|
|
27 }
|
|
|
28
|
|
|
29 // COMMON-LABEL: @_Z11dpscalardiv
|
|
|
30 // COMMON-NOT: !fpmath
|
|
|
31 __device__ double dpscalardiv(double a, double b) {
|
|
|
32 return a / b;
|
|
|
33 }
|
|
|
34
|
|
252
|
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 %{{.+}}){{$}}
|
|
|
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
|
|
|
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 %{{.+}}){{$}}
|
|
|
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
|
|
|
59 __device__ double test_builtin_elementwise_f64(double a) {
|
|
|
60 return __builtin_elementwise_sqrt(a);
|
|
|
61 }
|
|
|
62
|
|
|
63 // NCRSQRT: ![[MD]] = !{float 2.500000e+00}
|
