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
9 #include "Inputs/cuda.h"
11 typedef __attribute__(( ext_vector_type(4) )) float float4;
13 // COMMON-LABEL: @_Z11spscalardiv
14 // COMMON: fdiv{{.*}},
15 // NCRDIV: !fpmath ![[MD:[0-9]+]]
17 __device__ float spscalardiv(float a, float b) {
21 // COMMON-LABEL: @_Z11spvectordiv
22 // COMMON: fdiv{{.*}},
23 // NCRDIV: !fpmath ![[MD]]
25 __device__ float4 spvectordiv(float4 a, float4 b) {
29 // COMMON-LABEL: @_Z11dpscalardiv
30 // COMMON-NOT: !fpmath
31 __device__ double dpscalardiv(double a, double b) {
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);
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);
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);
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);
63 // NCRSQRT: ![[MD]] = !{float 2.500000e+00}