xref: /llvm-project/llvm/test/CodeGen/NVPTX/f32-lg2.ll (revision cdb4da32dbc362dc03125e965bf9847604856b31)
13ba339b5SPrinceton Ferro; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
23ba339b5SPrinceton Ferro; RUN: llc < %s -mcpu=sm_20 -mattr=+ptx32 | FileCheck --check-prefixes=CHECK %s
33ba339b5SPrinceton Ferro; RUN: %if ptxas %{ llc < %s -mcpu=sm_20 -mattr=+ptx32 | %ptxas-verify %}
4*cdb4da32SJoseph Hubertarget triple = "nvptx64-nvidia-cuda"
53ba339b5SPrinceton Ferro
63ba339b5SPrinceton Ferrodeclare float @llvm.nvvm.lg2.approx.f(float)
73ba339b5SPrinceton Ferrodeclare float @llvm.nvvm.lg2.approx.ftz.f(float)
83ba339b5SPrinceton Ferro
93ba339b5SPrinceton Ferro; CHECK-LABEL: lg2_float
103ba339b5SPrinceton Ferrodefine float @lg2_float(float %0) {
113ba339b5SPrinceton Ferro; CHECK-LABEL: lg2_float(
123ba339b5SPrinceton Ferro; CHECK:       {
133ba339b5SPrinceton Ferro; CHECK-NEXT:    .reg .f32 %f<3>;
143ba339b5SPrinceton Ferro; CHECK-EMPTY:
153ba339b5SPrinceton Ferro; CHECK-NEXT:  // %bb.0:
163ba339b5SPrinceton Ferro; CHECK-NEXT:    ld.param.f32 %f1, [lg2_float_param_0];
173ba339b5SPrinceton Ferro; CHECK-NEXT:    lg2.approx.f32 %f2, %f1;
183ba339b5SPrinceton Ferro; CHECK-NEXT:    st.param.f32 [func_retval0], %f2;
193ba339b5SPrinceton Ferro; CHECK-NEXT:    ret;
203ba339b5SPrinceton Ferro  %res = call float @llvm.nvvm.lg2.approx.f(float %0)
213ba339b5SPrinceton Ferro  ret float %res
223ba339b5SPrinceton Ferro}
233ba339b5SPrinceton Ferro
243ba339b5SPrinceton Ferro; CHECK-LABEL: lg2_float_ftz
253ba339b5SPrinceton Ferrodefine float @lg2_float_ftz(float %0) {
263ba339b5SPrinceton Ferro; CHECK-LABEL: lg2_float_ftz(
273ba339b5SPrinceton Ferro; CHECK:       {
283ba339b5SPrinceton Ferro; CHECK-NEXT:    .reg .f32 %f<3>;
293ba339b5SPrinceton Ferro; CHECK-EMPTY:
303ba339b5SPrinceton Ferro; CHECK-NEXT:  // %bb.0:
313ba339b5SPrinceton Ferro; CHECK-NEXT:    ld.param.f32 %f1, [lg2_float_ftz_param_0];
323ba339b5SPrinceton Ferro; CHECK-NEXT:    lg2.approx.ftz.f32 %f2, %f1;
333ba339b5SPrinceton Ferro; CHECK-NEXT:    st.param.f32 [func_retval0], %f2;
343ba339b5SPrinceton Ferro; CHECK-NEXT:    ret;
353ba339b5SPrinceton Ferro  %res = call float @llvm.nvvm.lg2.approx.ftz.f(float %0)
363ba339b5SPrinceton Ferro  ret float %res
373ba339b5SPrinceton Ferro}
38