1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
|
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
define float @foo(<2 x float> %a) {
; CHECK-LABEL: foo(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<6>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [foo_param_0];
; CHECK-NEXT: mul.rn.f32 %r3, %r2, %r2;
; CHECK-NEXT: mul.rn.f32 %r4, %r1, %r1;
; CHECK-NEXT: add.rn.f32 %r5, %r4, %r3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r5;
; CHECK-NEXT: ret;
%t1 = fmul <2 x float> %a, %a
%t2 = extractelement <2 x float> %t1, i32 0
%t3 = extractelement <2 x float> %t1, i32 1
%t4 = fadd float %t2, %t3
ret float %t4
}
define float @bar(<4 x float> %a) {
; CHECK-LABEL: bar(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<8>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [bar_param_0];
; CHECK-NEXT: mul.rn.f32 %r5, %r2, %r2;
; CHECK-NEXT: mul.rn.f32 %r6, %r1, %r1;
; CHECK-NEXT: add.rn.f32 %r7, %r6, %r5;
; CHECK-NEXT: st.param.b32 [func_retval0], %r7;
; CHECK-NEXT: ret;
%t1 = fmul <4 x float> %a, %a
%t2 = extractelement <4 x float> %t1, i32 0
%t3 = extractelement <4 x float> %t1, i32 1
%t4 = fadd float %t2, %t3
ret float %t4
}
define <4 x float> @baz(<4 x float> %a) {
; CHECK-LABEL: baz(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<9>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [baz_param_0];
; CHECK-NEXT: mul.rn.f32 %r5, %r4, %r4;
; CHECK-NEXT: mul.rn.f32 %r6, %r3, %r3;
; CHECK-NEXT: mul.rn.f32 %r7, %r2, %r2;
; CHECK-NEXT: mul.rn.f32 %r8, %r1, %r1;
; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};
; CHECK-NEXT: ret;
%t1 = fmul <4 x float> %a, %a
ret <4 x float> %t1
}
|
