61 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_20 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}4 5define float @foo(<2 x float> %a) {6; CHECK-LABEL: foo(7; CHECK: {8; CHECK-NEXT: .reg .b32 %r<6>;9; CHECK-EMPTY:10; CHECK-NEXT: // %bb.0:11; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [foo_param_0];12; CHECK-NEXT: mul.rn.f32 %r3, %r2, %r2;13; CHECK-NEXT: mul.rn.f32 %r4, %r1, %r1;14; CHECK-NEXT: add.rn.f32 %r5, %r4, %r3;15; CHECK-NEXT: st.param.b32 [func_retval0], %r5;16; CHECK-NEXT: ret;17 %t1 = fmul <2 x float> %a, %a18 %t2 = extractelement <2 x float> %t1, i32 019 %t3 = extractelement <2 x float> %t1, i32 120 %t4 = fadd float %t2, %t321 ret float %t422}23 24 25define float @bar(<4 x float> %a) {26; CHECK-LABEL: bar(27; CHECK: {28; CHECK-NEXT: .reg .b32 %r<8>;29; CHECK-EMPTY:30; CHECK-NEXT: // %bb.0:31; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [bar_param_0];32; CHECK-NEXT: mul.rn.f32 %r5, %r2, %r2;33; CHECK-NEXT: mul.rn.f32 %r6, %r1, %r1;34; CHECK-NEXT: add.rn.f32 %r7, %r6, %r5;35; CHECK-NEXT: st.param.b32 [func_retval0], %r7;36; CHECK-NEXT: ret;37 %t1 = fmul <4 x float> %a, %a38 %t2 = extractelement <4 x float> %t1, i32 039 %t3 = extractelement <4 x float> %t1, i32 140 %t4 = fadd float %t2, %t341 ret float %t442}43 44 45define <4 x float> @baz(<4 x float> %a) {46; CHECK-LABEL: baz(47; CHECK: {48; CHECK-NEXT: .reg .b32 %r<9>;49; CHECK-EMPTY:50; CHECK-NEXT: // %bb.0:51; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [baz_param_0];52; CHECK-NEXT: mul.rn.f32 %r5, %r4, %r4;53; CHECK-NEXT: mul.rn.f32 %r6, %r3, %r3;54; CHECK-NEXT: mul.rn.f32 %r7, %r2, %r2;55; CHECK-NEXT: mul.rn.f32 %r8, %r1, %r1;56; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};57; CHECK-NEXT: ret;58 %t1 = fmul <4 x float> %a, %a59 ret <4 x float> %t160}61