143 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mcpu=sm_80 -mattr=+ptx70 | FileCheck --check-prefixes=CHECK %s3; RUN: %if ptxas-sm_80 && ptxas-isa-7.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %}4 5target triple = "nvptx-nvidia-cuda"6 7declare float @llvm.nvvm.fabs.f32(float)8declare float @llvm.nvvm.fabs.ftz.f32(float)9declare double @llvm.nvvm.fabs.f64(double)10declare half @llvm.nvvm.fabs.f16(half)11declare half @llvm.nvvm.fabs.ftz.f16(half)12declare <2 x half> @llvm.nvvm.fabs.v2f16(<2 x half>)13declare <2 x half> @llvm.nvvm.fabs.ftz.v2f16(<2 x half>)14declare bfloat @llvm.nvvm.fabs.bf16(bfloat)15declare <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat>)16 17 18define float @fabs_float(float %a) {19; CHECK-LABEL: fabs_float(20; CHECK: {21; CHECK-NEXT: .reg .b32 %r<3>;22; CHECK-EMPTY:23; CHECK-NEXT: // %bb.0:24; CHECK-NEXT: ld.param.b32 %r1, [fabs_float_param_0];25; CHECK-NEXT: abs.f32 %r2, %r1;26; CHECK-NEXT: st.param.b32 [func_retval0], %r2;27; CHECK-NEXT: ret;28 %ret = call float @llvm.nvvm.fabs.f32(float %a)29 ret float %ret30}31 32define float @fabs_float_ftz(float %a) {33; CHECK-LABEL: fabs_float_ftz(34; CHECK: {35; CHECK-NEXT: .reg .b32 %r<3>;36; CHECK-EMPTY:37; CHECK-NEXT: // %bb.0:38; CHECK-NEXT: ld.param.b32 %r1, [fabs_float_ftz_param_0];39; CHECK-NEXT: abs.ftz.f32 %r2, %r1;40; CHECK-NEXT: st.param.b32 [func_retval0], %r2;41; CHECK-NEXT: ret;42 %ret = call float @llvm.nvvm.fabs.ftz.f32(float %a)43 ret float %ret44}45 46define double @fabs_double(double %a) {47; CHECK-LABEL: fabs_double(48; CHECK: {49; CHECK-NEXT: .reg .b64 %rd<3>;50; CHECK-EMPTY:51; CHECK-NEXT: // %bb.0:52; CHECK-NEXT: ld.param.b64 %rd1, [fabs_double_param_0];53; CHECK-NEXT: abs.f64 %rd2, %rd1;54; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;55; CHECK-NEXT: ret;56 %ret = call double @llvm.nvvm.fabs.f64(double %a)57 ret double %ret58}59 60define half @fabs_half(half %a) {61; CHECK-LABEL: fabs_half(62; CHECK: {63; CHECK-NEXT: .reg .b16 %rs<3>;64; CHECK-EMPTY:65; CHECK-NEXT: // %bb.0:66; CHECK-NEXT: ld.param.b16 %rs1, [fabs_half_param_0];67; CHECK-NEXT: abs.f16 %rs2, %rs1;68; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;69; CHECK-NEXT: ret;70 %ret = call half @llvm.nvvm.fabs.f16(half %a)71 ret half %ret72}73 74define half @fabs_half_ftz(half %a) {75; CHECK-LABEL: fabs_half_ftz(76; CHECK: {77; CHECK-NEXT: .reg .b16 %rs<3>;78; CHECK-EMPTY:79; CHECK-NEXT: // %bb.0:80; CHECK-NEXT: ld.param.b16 %rs1, [fabs_half_ftz_param_0];81; CHECK-NEXT: abs.ftz.f16 %rs2, %rs1;82; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;83; CHECK-NEXT: ret;84 %ret = call half @llvm.nvvm.fabs.ftz.f16(half %a)85 ret half %ret86}87 88define <2 x half> @fabs_v2half(<2 x half> %a) {89; CHECK-LABEL: fabs_v2half(90; CHECK: {91; CHECK-NEXT: .reg .b32 %r<3>;92; CHECK-EMPTY:93; CHECK-NEXT: // %bb.0:94; CHECK-NEXT: ld.param.b32 %r1, [fabs_v2half_param_0];95; CHECK-NEXT: abs.f16x2 %r2, %r1;96; CHECK-NEXT: st.param.b32 [func_retval0], %r2;97; CHECK-NEXT: ret;98 %ret = call <2 x half> @llvm.nvvm.fabs.v2f16(<2 x half> %a)99 ret <2 x half> %ret100}101 102define <2 x half> @fabs_v2half_ftz(<2 x half> %a) {103; CHECK-LABEL: fabs_v2half_ftz(104; CHECK: {105; CHECK-NEXT: .reg .b32 %r<3>;106; CHECK-EMPTY:107; CHECK-NEXT: // %bb.0:108; CHECK-NEXT: ld.param.b32 %r1, [fabs_v2half_ftz_param_0];109; CHECK-NEXT: abs.ftz.f16x2 %r2, %r1;110; CHECK-NEXT: st.param.b32 [func_retval0], %r2;111; CHECK-NEXT: ret;112 %ret = call <2 x half> @llvm.nvvm.fabs.ftz.v2f16(<2 x half> %a)113 ret <2 x half> %ret114}115 116define bfloat @fabs_bf16(bfloat %a) {117; CHECK-LABEL: fabs_bf16(118; CHECK: {119; CHECK-NEXT: .reg .b16 %rs<3>;120; CHECK-EMPTY:121; CHECK-NEXT: // %bb.0:122; CHECK-NEXT: ld.param.b16 %rs1, [fabs_bf16_param_0];123; CHECK-NEXT: abs.bf16 %rs2, %rs1;124; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;125; CHECK-NEXT: ret;126 %ret = call bfloat @llvm.nvvm.fabs.bf16(bfloat %a)127 ret bfloat %ret128}129 130define <2 x bfloat> @fabs_v2bf16(<2 x bfloat> %a) {131; CHECK-LABEL: fabs_v2bf16(132; CHECK: {133; CHECK-NEXT: .reg .b32 %r<3>;134; CHECK-EMPTY:135; CHECK-NEXT: // %bb.0:136; CHECK-NEXT: ld.param.b32 %r1, [fabs_v2bf16_param_0];137; CHECK-NEXT: abs.bf16x2 %r2, %r1;138; CHECK-NEXT: st.param.b32 [func_retval0], %r2;139; CHECK-NEXT: ret;140 %ret = call <2 x bfloat> @llvm.nvvm.fabs.v2bf16(<2 x bfloat> %a)141 ret <2 x bfloat> %ret142}143