65 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78| FileCheck --check-prefixes=CHECK %s3; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78| %ptxas-verify -arch=sm_90 %}4 5declare i32 @llvm.nvvm.f2tf32.rn(float %f1)6declare i32 @llvm.nvvm.f2tf32.rn.relu(float %f1)7declare i32 @llvm.nvvm.f2tf32.rz(float %f1)8declare i32 @llvm.nvvm.f2tf32.rz.relu(float %f1)9 10define i32 @cvt_rn_tf32_f32(float %f1) {11; CHECK-LABEL: cvt_rn_tf32_f32(12; CHECK: {13; CHECK-NEXT: .reg .b32 %r<3>;14; CHECK-EMPTY:15; CHECK-NEXT: // %bb.0:16; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_tf32_f32_param_0];17; CHECK-NEXT: cvt.rn.tf32.f32 %r2, %r1;18; CHECK-NEXT: st.param.b32 [func_retval0], %r2;19; CHECK-NEXT: ret;20 %val = call i32 @llvm.nvvm.f2tf32.rn(float %f1)21 ret i32 %val22}23 24define i32 @cvt_rn_relu_tf32_f32(float %f1) {25; CHECK-LABEL: cvt_rn_relu_tf32_f32(26; CHECK: {27; CHECK-NEXT: .reg .b32 %r<3>;28; CHECK-EMPTY:29; CHECK-NEXT: // %bb.0:30; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_relu_tf32_f32_param_0];31; CHECK-NEXT: cvt.rn.relu.tf32.f32 %r2, %r1;32; CHECK-NEXT: st.param.b32 [func_retval0], %r2;33; CHECK-NEXT: ret;34 %val = call i32 @llvm.nvvm.f2tf32.rn.relu(float %f1)35 ret i32 %val36}37 38define i32 @cvt_rz_tf32_f32(float %f1) {39; CHECK-LABEL: cvt_rz_tf32_f32(40; CHECK: {41; CHECK-NEXT: .reg .b32 %r<3>;42; CHECK-EMPTY:43; CHECK-NEXT: // %bb.0:44; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_tf32_f32_param_0];45; CHECK-NEXT: cvt.rz.tf32.f32 %r2, %r1;46; CHECK-NEXT: st.param.b32 [func_retval0], %r2;47; CHECK-NEXT: ret;48 %val = call i32 @llvm.nvvm.f2tf32.rz(float %f1)49 ret i32 %val50}51 52define i32 @cvt_rz_relu_tf32_f32(float %f1) {53; CHECK-LABEL: cvt_rz_relu_tf32_f32(54; CHECK: {55; CHECK-NEXT: .reg .b32 %r<3>;56; CHECK-EMPTY:57; CHECK-NEXT: // %bb.0:58; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_relu_tf32_f32_param_0];59; CHECK-NEXT: cvt.rz.relu.tf32.f32 %r2, %r1;60; CHECK-NEXT: st.param.b32 [func_retval0], %r2;61; CHECK-NEXT: ret;62 %val = call i32 @llvm.nvvm.f2tf32.rz.relu(float %f1)63 ret i32 %val64}65