38 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mcpu=sm_50 -mattr=+ptx32 | FileCheck --check-prefixes=CHECK %s3; RUN: %if ptxas-sm_50 && ptxas-isa-3.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_50 -mattr=+ptx32 | %ptxas-verify -arch=sm_50 %}4target triple = "nvptx-nvidia-cuda"5 6declare float @llvm.nvvm.ex2.approx.f32(float)7declare float @llvm.nvvm.ex2.approx.ftz.f32(float)8 9; CHECK-LABEL: ex2_float10define float @ex2_float(float %0) {11; CHECK-LABEL: ex2_float(12; CHECK: {13; CHECK-NEXT: .reg .b32 %r<3>;14; CHECK-EMPTY:15; CHECK-NEXT: // %bb.0:16; CHECK-NEXT: ld.param.b32 %r1, [ex2_float_param_0];17; CHECK-NEXT: ex2.approx.f32 %r2, %r1;18; CHECK-NEXT: st.param.b32 [func_retval0], %r2;19; CHECK-NEXT: ret;20 %res = call float @llvm.nvvm.ex2.approx.f32(float %0)21 ret float %res22}23 24; CHECK-LABEL: ex2_float_ftz25define float @ex2_float_ftz(float %0) {26; CHECK-LABEL: ex2_float_ftz(27; CHECK: {28; CHECK-NEXT: .reg .b32 %r<3>;29; CHECK-EMPTY:30; CHECK-NEXT: // %bb.0:31; CHECK-NEXT: ld.param.b32 %r1, [ex2_float_ftz_param_0];32; CHECK-NEXT: ex2.approx.ftz.f32 %r2, %r1;33; CHECK-NEXT: st.param.b32 [func_retval0], %r2;34; CHECK-NEXT: ret;35 %res = call float @llvm.nvvm.ex2.approx.ftz.f32(float %0)36 ret float %res37}38