Skip to content

Commit 76be42d

Browse files
AlexMacleanAlexMaclean
committed
test updates
1 parent 124bcf8 commit 76be42d

File tree

84 files changed

+9024
-7177
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

84 files changed

+9024
-7177
lines changed

‎llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll

Lines changed: 53 additions & 53 deletions
Original file line numberDiff line numberDiff line change
@@ -45,36 +45,36 @@ define half @fh(ptr %p) {
4545
; ENABLED-LABEL: fh(
4646
; ENABLED: {
4747
; ENABLED-NEXT: .reg .b16 %rs<10>;
48-
; ENABLED-NEXT: .reg .b32 %f<13>;
48+
; ENABLED-NEXT: .reg .b32 %r<13>;
4949
; ENABLED-NEXT: .reg .b64 %rd<2>;
5050
; ENABLED-EMPTY:
5151
; ENABLED-NEXT: // %bb.0:
5252
; ENABLED-NEXT: ld.param.b64 %rd1, [fh_param_0];
5353
; ENABLED-NEXT: ld.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [%rd1];
5454
; ENABLED-NEXT: ld.b16 %rs5, [%rd1+8];
55-
; ENABLED-NEXT: cvt.f32.f16 %f1, %rs2;
56-
; ENABLED-NEXT: cvt.f32.f16 %f2, %rs1;
57-
; ENABLED-NEXT: add.rn.f32 %f3, %f2, %f1;
58-
; ENABLED-NEXT: cvt.rn.f16.f32 %rs6, %f3;
59-
; ENABLED-NEXT: cvt.f32.f16 %f4, %rs4;
60-
; ENABLED-NEXT: cvt.f32.f16 %f5, %rs3;
61-
; ENABLED-NEXT: add.rn.f32 %f6, %f5, %f4;
62-
; ENABLED-NEXT: cvt.rn.f16.f32 %rs7, %f6;
63-
; ENABLED-NEXT: cvt.f32.f16 %f7, %rs7;
64-
; ENABLED-NEXT: cvt.f32.f16 %f8, %rs6;
65-
; ENABLED-NEXT: add.rn.f32 %f9, %f8, %f7;
66-
; ENABLED-NEXT: cvt.rn.f16.f32 %rs8, %f9;
67-
; ENABLED-NEXT: cvt.f32.f16 %f10, %rs8;
68-
; ENABLED-NEXT: cvt.f32.f16 %f11, %rs5;
69-
; ENABLED-NEXT: add.rn.f32 %f12, %f10, %f11;
70-
; ENABLED-NEXT: cvt.rn.f16.f32 %rs9, %f12;
55+
; ENABLED-NEXT: cvt.f32.f16 %r1, %rs2;
56+
; ENABLED-NEXT: cvt.f32.f16 %r2, %rs1;
57+
; ENABLED-NEXT: add.rn.f32 %r3, %r2, %r1;
58+
; ENABLED-NEXT: cvt.rn.f16.f32 %rs6, %r3;
59+
; ENABLED-NEXT: cvt.f32.f16 %r4, %rs4;
60+
; ENABLED-NEXT: cvt.f32.f16 %r5, %rs3;
61+
; ENABLED-NEXT: add.rn.f32 %r6, %r5, %r4;
62+
; ENABLED-NEXT: cvt.rn.f16.f32 %rs7, %r6;
63+
; ENABLED-NEXT: cvt.f32.f16 %r7, %rs7;
64+
; ENABLED-NEXT: cvt.f32.f16 %r8, %rs6;
65+
; ENABLED-NEXT: add.rn.f32 %r9, %r8, %r7;
66+
; ENABLED-NEXT: cvt.rn.f16.f32 %rs8, %r9;
67+
; ENABLED-NEXT: cvt.f32.f16 %r10, %rs8;
68+
; ENABLED-NEXT: cvt.f32.f16 %r11, %rs5;
69+
; ENABLED-NEXT: add.rn.f32 %r12, %r10, %r11;
70+
; ENABLED-NEXT: cvt.rn.f16.f32 %rs9, %r12;
7171
; ENABLED-NEXT: st.param.b16 [func_retval0], %rs9;
7272
; ENABLED-NEXT: ret;
7373
;
7474
; DISABLED-LABEL: fh(
7575
; DISABLED: {
7676
; DISABLED-NEXT: .reg .b16 %rs<10>;
77-
; DISABLED-NEXT: .reg .b32 %f<13>;
77+
; DISABLED-NEXT: .reg .b32 %r<13>;
7878
; DISABLED-NEXT: .reg .b64 %rd<2>;
7979
; DISABLED-EMPTY:
8080
; DISABLED-NEXT: // %bb.0:
@@ -84,22 +84,22 @@ define half @fh(ptr %p) {
8484
; DISABLED-NEXT: ld.b16 %rs3, [%rd1+4];
8585
; DISABLED-NEXT: ld.b16 %rs4, [%rd1+6];
8686
; DISABLED-NEXT: ld.b16 %rs5, [%rd1+8];
87-
; DISABLED-NEXT: cvt.f32.f16 %f1, %rs2;
88-
; DISABLED-NEXT: cvt.f32.f16 %f2, %rs1;
89-
; DISABLED-NEXT: add.rn.f32 %f3, %f2, %f1;
90-
; DISABLED-NEXT: cvt.rn.f16.f32 %rs6, %f3;
91-
; DISABLED-NEXT: cvt.f32.f16 %f4, %rs4;
92-
; DISABLED-NEXT: cvt.f32.f16 %f5, %rs3;
93-
; DISABLED-NEXT: add.rn.f32 %f6, %f5, %f4;
94-
; DISABLED-NEXT: cvt.rn.f16.f32 %rs7, %f6;
95-
; DISABLED-NEXT: cvt.f32.f16 %f7, %rs7;
96-
; DISABLED-NEXT: cvt.f32.f16 %f8, %rs6;
97-
; DISABLED-NEXT: add.rn.f32 %f9, %f8, %f7;
98-
; DISABLED-NEXT: cvt.rn.f16.f32 %rs8, %f9;
99-
; DISABLED-NEXT: cvt.f32.f16 %f10, %rs8;
100-
; DISABLED-NEXT: cvt.f32.f16 %f11, %rs5;
101-
; DISABLED-NEXT: add.rn.f32 %f12, %f10, %f11;
102-
; DISABLED-NEXT: cvt.rn.f16.f32 %rs9, %f12;
87+
; DISABLED-NEXT: cvt.f32.f16 %r1, %rs2;
88+
; DISABLED-NEXT: cvt.f32.f16 %r2, %rs1;
89+
; DISABLED-NEXT: add.rn.f32 %r3, %r2, %r1;
90+
; DISABLED-NEXT: cvt.rn.f16.f32 %rs6, %r3;
91+
; DISABLED-NEXT: cvt.f32.f16 %r4, %rs4;
92+
; DISABLED-NEXT: cvt.f32.f16 %r5, %rs3;
93+
; DISABLED-NEXT: add.rn.f32 %r6, %r5, %r4;
94+
; DISABLED-NEXT: cvt.rn.f16.f32 %rs7, %r6;
95+
; DISABLED-NEXT: cvt.f32.f16 %r7, %rs7;
96+
; DISABLED-NEXT: cvt.f32.f16 %r8, %rs6;
97+
; DISABLED-NEXT: add.rn.f32 %r9, %r8, %r7;
98+
; DISABLED-NEXT: cvt.rn.f16.f32 %rs8, %r9;
99+
; DISABLED-NEXT: cvt.f32.f16 %r10, %rs8;
100+
; DISABLED-NEXT: cvt.f32.f16 %r11, %rs5;
101+
; DISABLED-NEXT: add.rn.f32 %r12, %r10, %r11;
102+
; DISABLED-NEXT: cvt.rn.f16.f32 %rs9, %r12;
103103
; DISABLED-NEXT: st.param.b16 [func_retval0], %rs9;
104104
; DISABLED-NEXT: ret;
105105
%p.1 = getelementptr half, ptr %p, i32 1
@@ -121,37 +121,37 @@ define half @fh(ptr %p) {
121121
define float @ff(ptr %p) {
122122
; ENABLED-LABEL: ff(
123123
; ENABLED: {
124-
; ENABLED-NEXT: .reg .b32 %f<10>;
124+
; ENABLED-NEXT: .reg .b32 %r<10>;
125125
; ENABLED-NEXT: .reg .b64 %rd<2>;
126126
; ENABLED-EMPTY:
127127
; ENABLED-NEXT: // %bb.0:
128128
; ENABLED-NEXT: ld.param.b64 %rd1, [ff_param_0];
129-
; ENABLED-NEXT: ld.v4.b32 {%f1, %f2, %f3, %f4}, [%rd1];
130-
; ENABLED-NEXT: ld.b32 %f5, [%rd1+16];
131-
; ENABLED-NEXT: add.rn.f32 %f6, %f1, %f2;
132-
; ENABLED-NEXT: add.rn.f32 %f7, %f3, %f4;
133-
; ENABLED-NEXT: add.rn.f32 %f8, %f6, %f7;
134-
; ENABLED-NEXT: add.rn.f32 %f9, %f8, %f5;
135-
; ENABLED-NEXT: st.param.b32 [func_retval0], %f9;
129+
; ENABLED-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
130+
; ENABLED-NEXT: ld.b32 %r5, [%rd1+16];
131+
; ENABLED-NEXT: add.rn.f32 %r6, %r1, %r2;
132+
; ENABLED-NEXT: add.rn.f32 %r7, %r3, %r4;
133+
; ENABLED-NEXT: add.rn.f32 %r8, %r6, %r7;
134+
; ENABLED-NEXT: add.rn.f32 %r9, %r8, %r5;
135+
; ENABLED-NEXT: st.param.b32 [func_retval0], %r9;
136136
; ENABLED-NEXT: ret;
137137
;
138138
; DISABLED-LABEL: ff(
139139
; DISABLED: {
140-
; DISABLED-NEXT: .reg .b32 %f<10>;
140+
; DISABLED-NEXT: .reg .b32 %r<10>;
141141
; DISABLED-NEXT: .reg .b64 %rd<2>;
142142
; DISABLED-EMPTY:
143143
; DISABLED-NEXT: // %bb.0:
144144
; DISABLED-NEXT: ld.param.b64 %rd1, [ff_param_0];
145-
; DISABLED-NEXT: ld.b32 %f1, [%rd1];
146-
; DISABLED-NEXT: ld.b32 %f2, [%rd1+4];
147-
; DISABLED-NEXT: ld.b32 %f3, [%rd1+8];
148-
; DISABLED-NEXT: ld.b32 %f4, [%rd1+12];
149-
; DISABLED-NEXT: ld.b32 %f5, [%rd1+16];
150-
; DISABLED-NEXT: add.rn.f32 %f6, %f1, %f2;
151-
; DISABLED-NEXT: add.rn.f32 %f7, %f3, %f4;
152-
; DISABLED-NEXT: add.rn.f32 %f8, %f6, %f7;
153-
; DISABLED-NEXT: add.rn.f32 %f9, %f8, %f5;
154-
; DISABLED-NEXT: st.param.b32 [func_retval0], %f9;
145+
; DISABLED-NEXT: ld.b32 %r1, [%rd1];
146+
; DISABLED-NEXT: ld.b32 %r2, [%rd1+4];
147+
; DISABLED-NEXT: ld.b32 %r3, [%rd1+8];
148+
; DISABLED-NEXT: ld.b32 %r4, [%rd1+12];
149+
; DISABLED-NEXT: ld.b32 %r5, [%rd1+16];
150+
; DISABLED-NEXT: add.rn.f32 %r6, %r1, %r2;
151+
; DISABLED-NEXT: add.rn.f32 %r7, %r3, %r4;
152+
; DISABLED-NEXT: add.rn.f32 %r8, %r6, %r7;
153+
; DISABLED-NEXT: add.rn.f32 %r9, %r8, %r5;
154+
; DISABLED-NEXT: st.param.b32 [func_retval0], %r9;
155155
; DISABLED-NEXT: ret;
156156
%p.1 = getelementptr float, ptr %p, i32 1
157157
%p.2 = getelementptr float, ptr %p, i32 2

‎llvm/test/CodeGen/NVPTX/access-non-generic.ll

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -23,10 +23,10 @@ define void @ld_st_shared_f32(i32 %i, float %v) {
2323
; load cast
2424
%1 = load float, ptr addrspacecast (ptr addrspace(3) @scalar to ptr), align 4
2525
call void @use(float %1)
26-
; PTX: ld.shared.b32 %f{{[0-9]+}}, [scalar];
26+
; PTX: ld.shared.b32 %r{{[0-9]+}}, [scalar];
2727
; store cast
2828
store float %v, ptr addrspacecast (ptr addrspace(3) @scalar to ptr), align 4
29-
; PTX: st.shared.b32 [scalar], %f{{[0-9]+}};
29+
; PTX: st.shared.b32 [scalar], %r{{[0-9]+}};
3030
; use syncthreads to disable optimizations across components
3131
call void @llvm.nvvm.barrier0()
3232
; PTX: bar.sync 0;
@@ -35,31 +35,31 @@ define void @ld_st_shared_f32(i32 %i, float %v) {
3535
%2 = addrspacecast ptr addrspace(3) @scalar to ptr
3636
%3 = load float, ptr %2, align 4
3737
call void @use(float %3)
38-
; PTX: ld.shared.b32 %f{{[0-9]+}}, [scalar];
38+
; PTX: ld.shared.b32 %r{{[0-9]+}}, [scalar];
3939
; cast; store
4040
store float %v, ptr %2, align 4
41-
; PTX: st.shared.b32 [scalar], %f{{[0-9]+}};
41+
; PTX: st.shared.b32 [scalar], %r{{[0-9]+}};
4242
call void @llvm.nvvm.barrier0()
4343
; PTX: bar.sync 0;
4444

4545
; load gep cast
4646
%4 = load float, ptr getelementptr inbounds ([10 x float], ptr addrspacecast (ptr addrspace(3) @array to ptr), i32 0, i32 5), align 4
4747
call void @use(float %4)
48-
; PTX: ld.shared.b32 %f{{[0-9]+}}, [array+20];
48+
; PTX: ld.shared.b32 %r{{[0-9]+}}, [array+20];
4949
; store gep cast
5050
store float %v, ptr getelementptr inbounds ([10 x float], ptr addrspacecast (ptr addrspace(3) @array to ptr), i32 0, i32 5), align 4
51-
; PTX: st.shared.b32 [array+20], %f{{[0-9]+}};
51+
; PTX: st.shared.b32 [array+20], %r{{[0-9]+}};
5252
call void @llvm.nvvm.barrier0()
5353
; PTX: bar.sync 0;
5454

5555
; gep cast; load
5656
%5 = getelementptr inbounds [10 x float], ptr addrspacecast (ptr addrspace(3) @array to ptr), i32 0, i32 5
5757
%6 = load float, ptr %5, align 4
5858
call void @use(float %6)
59-
; PTX: ld.shared.b32 %f{{[0-9]+}}, [array+20];
59+
; PTX: ld.shared.b32 %r{{[0-9]+}}, [array+20];
6060
; gep cast; store
6161
store float %v, ptr %5, align 4
62-
; PTX: st.shared.b32 [array+20], %f{{[0-9]+}};
62+
; PTX: st.shared.b32 [array+20], %r{{[0-9]+}};
6363
call void @llvm.nvvm.barrier0()
6464
; PTX: bar.sync 0;
6565

@@ -68,10 +68,10 @@ define void @ld_st_shared_f32(i32 %i, float %v) {
6868
%8 = getelementptr inbounds [10 x float], ptr %7, i32 0, i32 %i
6969
%9 = load float, ptr %8, align 4
7070
call void @use(float %9)
71-
; PTX: ld.shared.b32 %f{{[0-9]+}}, [%{{(r|rl|rd)[0-9]+}}];
71+
; PTX: ld.shared.b32 %r{{[0-9]+}}, [%{{(r|rl|rd)[0-9]+}}];
7272
; cast; gep; store
7373
store float %v, ptr %8, align 4
74-
; PTX: st.shared.b32 [%{{(r|rl|rd)[0-9]+}}], %f{{[0-9]+}};
74+
; PTX: st.shared.b32 [%{{(r|rl|rd)[0-9]+}}], %r{{[0-9]+}};
7575
call void @llvm.nvvm.barrier0()
7676
; PTX: bar.sync 0;
7777

‎llvm/test/CodeGen/NVPTX/aggregate-return.ll

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -10,7 +10,7 @@ define void @test_v2f32(<2 x float> %input, ptr %output) {
1010
; CHECK-LABEL: @test_v2f32
1111
%call = tail call <2 x float> @barv(<2 x float> %input)
1212
; CHECK: .param .align 8 .b8 retval0[8];
13-
; CHECK: ld.param.v2.b32 {[[E0:%f[0-9]+]], [[E1:%f[0-9]+]]}, [retval0];
13+
; CHECK: ld.param.v2.b32 {[[E0:%r[0-9]+]], [[E1:%r[0-9]+]]}, [retval0];
1414
store <2 x float> %call, ptr %output, align 8
1515
; CHECK: st.v2.b32 [{{%rd[0-9]+}}], {[[E0]], [[E1]]}
1616
ret void
@@ -21,10 +21,10 @@ define void @test_v3f32(<3 x float> %input, ptr %output) {
2121
;
2222
%call = tail call <3 x float> @barv3(<3 x float> %input)
2323
; CHECK: .param .align 16 .b8 retval0[16];
24-
; CHECK-DAG: ld.param.v2.b32 {[[E0:%f[0-9]+]], [[E1:%f[0-9]+]]}, [retval0];
25-
; CHECK-DAG: ld.param.b32 [[E2:%f[0-9]+]], [retval0+8];
24+
; CHECK-DAG: ld.param.v2.b32 {[[E0:%r[0-9]+]], [[E1:%r[0-9]+]]}, [retval0];
25+
; CHECK-DAG: ld.param.b32 [[E2:%r[0-9]+]], [retval0+8];
2626
; Make sure we don't load more values than than we need to.
27-
; CHECK-NOT: ld.param.b32 [[E3:%f[0-9]+]], [retval0+12];
27+
; CHECK-NOT: ld.param.b32 [[E3:%r[0-9]+]], [retval0+12];
2828
store <3 x float> %call, ptr %output, align 8
2929
; CHECK-DAG: st.b32 [{{%rd[0-9]}}+8],
3030
; -- This is suboptimal. We should do st.v2.f32 instead
@@ -38,8 +38,8 @@ define void @test_a2f32([2 x float] %input, ptr %output) {
3838
; CHECK-LABEL: @test_a2f32
3939
%call = tail call [2 x float] @bara([2 x float] %input)
4040
; CHECK: .param .align 4 .b8 retval0[8];
41-
; CHECK-DAG: ld.param.b32 [[ELEMA1:%f[0-9]+]], [retval0];
42-
; CHECK-DAG: ld.param.b32 [[ELEMA2:%f[0-9]+]], [retval0+4];
41+
; CHECK-DAG: ld.param.b32 [[ELEMA1:%r[0-9]+]], [retval0];
42+
; CHECK-DAG: ld.param.b32 [[ELEMA2:%r[0-9]+]], [retval0+4];
4343
store [2 x float] %call, ptr %output, align 4
4444
; CHECK: }
4545
; CHECK-DAG: st.b32 [{{%rd[0-9]+}}], [[ELEMA1]]
@@ -52,8 +52,8 @@ define void @test_s2f32({float, float} %input, ptr %output) {
5252
; CHECK-LABEL: @test_s2f32
5353
%call = tail call {float, float} @bars({float, float} %input)
5454
; CHECK: .param .align 4 .b8 retval0[8];
55-
; CHECK-DAG: ld.param.b32 [[ELEMS1:%f[0-9]+]], [retval0];
56-
; CHECK-DAG: ld.param.b32 [[ELEMS2:%f[0-9]+]], [retval0+4];
55+
; CHECK-DAG: ld.param.b32 [[ELEMS1:%r[0-9]+]], [retval0];
56+
; CHECK-DAG: ld.param.b32 [[ELEMS2:%r[0-9]+]], [retval0+4];
5757
store {float, float} %call, ptr %output, align 4
5858
; CHECK: }
5959
; CHECK-DAG: st.b32 [{{%rd[0-9]+}}], [[ELEMS1]]

‎llvm/test/CodeGen/NVPTX/and-or-setcc.ll

Lines changed: 12 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -8,15 +8,14 @@ define i1 @and_ord(float %a, float %b) {
88
; CHECK-LABEL: and_ord(
99
; CHECK: {
1010
; CHECK-NEXT: .reg .pred %p<2>;
11-
; CHECK-NEXT: .reg .b32 %r<2>;
12-
; CHECK-NEXT: .reg .b32 %f<3>;
11+
; CHECK-NEXT: .reg .b32 %r<4>;
1312
; CHECK-EMPTY:
1413
; CHECK-NEXT: // %bb.0:
15-
; CHECK-NEXT: ld.param.b32 %f1, [and_ord_param_0];
16-
; CHECK-NEXT: ld.param.b32 %f2, [and_ord_param_1];
17-
; CHECK-NEXT: setp.num.f32 %p1, %f1, %f2;
18-
; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;
19-
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
14+
; CHECK-NEXT: ld.param.b32 %r1, [and_ord_param_0];
15+
; CHECK-NEXT: ld.param.b32 %r2, [and_ord_param_1];
16+
; CHECK-NEXT: setp.num.f32 %p1, %r1, %r2;
17+
; CHECK-NEXT: selp.b32 %r3, 1, 0, %p1;
18+
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
2019
; CHECK-NEXT: ret;
2120
%c = fcmp ord float %a, 0.0
2221
%d = fcmp ord float %b, 0.0
@@ -28,15 +27,14 @@ define i1 @or_uno(float %a, float %b) {
2827
; CHECK-LABEL: or_uno(
2928
; CHECK: {
3029
; CHECK-NEXT: .reg .pred %p<2>;
31-
; CHECK-NEXT: .reg .b32 %r<2>;
32-
; CHECK-NEXT: .reg .b32 %f<3>;
30+
; CHECK-NEXT: .reg .b32 %r<4>;
3331
; CHECK-EMPTY:
3432
; CHECK-NEXT: // %bb.0:
35-
; CHECK-NEXT: ld.param.b32 %f1, [or_uno_param_0];
36-
; CHECK-NEXT: ld.param.b32 %f2, [or_uno_param_1];
37-
; CHECK-NEXT: setp.nan.f32 %p1, %f1, %f2;
38-
; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;
39-
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
33+
; CHECK-NEXT: ld.param.b32 %r1, [or_uno_param_0];
34+
; CHECK-NEXT: ld.param.b32 %r2, [or_uno_param_1];
35+
; CHECK-NEXT: setp.nan.f32 %p1, %r1, %r2;
36+
; CHECK-NEXT: selp.b32 %r3, 1, 0, %p1;
37+
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
4038
; CHECK-NEXT: ret;
4139
%c = fcmp uno float %a, 0.0
4240
%d = fcmp uno float %b, 0.0

‎llvm/test/CodeGen/NVPTX/arithmetic-fp-sm20.ll

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -11,28 +11,28 @@
1111
;;; f64
1212

1313
define double @fadd_f64(double %a, double %b) {
14-
; CHECK: add.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}
14+
; CHECK: add.f64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}}
1515
; CHECK: ret
1616
%ret = fadd double %a, %b
1717
ret double %ret
1818
}
1919

2020
define double @fsub_f64(double %a, double %b) {
21-
; CHECK: sub.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}
21+
; CHECK: sub.f64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}}
2222
; CHECK: ret
2323
%ret = fsub double %a, %b
2424
ret double %ret
2525
}
2626

2727
define double @fmul_f64(double %a, double %b) {
28-
; CHECK: mul.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}
28+
; CHECK: mul.f64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}}
2929
; CHECK: ret
3030
%ret = fmul double %a, %b
3131
ret double %ret
3232
}
3333

3434
define double @fdiv_f64(double %a, double %b) {
35-
; CHECK: div.rn.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}
35+
; CHECK: div.rn.f64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, %rd{{[0-9]+}}
3636
; CHECK: ret
3737
%ret = fdiv double %a, %b
3838
ret double %ret
@@ -44,28 +44,28 @@ define double @fdiv_f64(double %a, double %b) {
4444
;;; f32
4545

4646
define float @fadd_f32(float %a, float %b) {
47-
; CHECK: add.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}
47+
; CHECK: add.f32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}}
4848
; CHECK: ret
4949
%ret = fadd float %a, %b
5050
ret float %ret
5151
}
5252

5353
define float @fsub_f32(float %a, float %b) {
54-
; CHECK: sub.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}
54+
; CHECK: sub.f32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}}
5555
; CHECK: ret
5656
%ret = fsub float %a, %b
5757
ret float %ret
5858
}
5959

6060
define float @fmul_f32(float %a, float %b) {
61-
; CHECK: mul.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}
61+
; CHECK: mul.f32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}}
6262
; CHECK: ret
6363
%ret = fmul float %a, %b
6464
ret float %ret
6565
}
6666

6767
define float @fdiv_f32(float %a, float %b) {
68-
; CHECK: div.rn.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}
68+
; CHECK: div.rn.f32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}}
6969
; CHECK: ret
7070
%ret = fdiv float %a, %b
7171
ret float %ret

‎llvm/test/CodeGen/NVPTX/atomics-with-scope.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -121,11 +121,11 @@ define void @test_atomics_scope_imm(ptr %fp, float %f,
121121
; CHECK: atom.cta.add.u64{{.*}}, 2;
122122
%tmp2i = tail call i64 @llvm.nvvm.atomic.add.gen.i.cta.i64.p0(ptr %llp, i64 2)
123123

124-
; CHECK: atom.cta.add.f32{{.*}}, %f{{[0-9]+}};
124+
; CHECK: atom.cta.add.f32{{.*}}, %r{{[0-9]+}};
125125
%tmp3r = tail call float @llvm.nvvm.atomic.add.gen.f.cta.f32.p0(ptr %fp, float %f)
126126
; CHECK: atom.cta.add.f32{{.*}}, 0f40400000;
127127
%tmp3i = tail call float @llvm.nvvm.atomic.add.gen.f.cta.f32.p0(ptr %fp, float 3.0)
128-
; CHECK: atom.cta.add.f64{{.*}}, %fd{{[0-9]+}};
128+
; CHECK: atom.cta.add.f64{{.*}}, %rd{{[0-9]+}};
129129
%tmp4r = tail call double @llvm.nvvm.atomic.add.gen.f.cta.f64.p0(ptr %dfp, double %df)
130130
; CHECK: atom.cta.add.f64{{.*}}, 0d4010000000000000;
131131
%tmp4i = tail call double @llvm.nvvm.atomic.add.gen.f.cta.f64.p0(ptr %dfp, double 4.0)

0 commit comments

Comments
 (0)