Skip to content

Commit 4bb455f

Browse files
AlexMacleanAlexMaclean
committed
update more tests
1 parent 6f89025 commit 4bb455f

File tree

5 files changed

+41
-41
lines changed

5 files changed

+41
-41
lines changed

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

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -10,8 +10,8 @@ define i32 @bmsk_wrap(i32 %a, i32 %b) {
1010
; CHECK-NEXT: .reg .b32 %r<4>;
1111
; CHECK-EMPTY:
1212
; CHECK-NEXT: // %bb.0:
13-
; CHECK-NEXT: ld.param.u32 %r1, [bmsk_wrap_param_0];
14-
; CHECK-NEXT: ld.param.u32 %r2, [bmsk_wrap_param_1];
13+
; CHECK-NEXT: ld.param.b32 %r1, [bmsk_wrap_param_0];
14+
; CHECK-NEXT: ld.param.b32 %r2, [bmsk_wrap_param_1];
1515
; CHECK-NEXT: bmsk.wrap.b32 %r3, %r1, %r2;
1616
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
1717
; CHECK-NEXT: ret;
@@ -25,8 +25,8 @@ define i32 @bmsk_clamp(i32 %a, i32 %b) {
2525
; CHECK-NEXT: .reg .b32 %r<4>;
2626
; CHECK-EMPTY:
2727
; CHECK-NEXT: // %bb.0:
28-
; CHECK-NEXT: ld.param.u32 %r1, [bmsk_clamp_param_0];
29-
; CHECK-NEXT: ld.param.u32 %r2, [bmsk_clamp_param_1];
28+
; CHECK-NEXT: ld.param.b32 %r1, [bmsk_clamp_param_0];
29+
; CHECK-NEXT: ld.param.b32 %r2, [bmsk_clamp_param_1];
3030
; CHECK-NEXT: bmsk.clamp.b32 %r3, %r1, %r2;
3131
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
3232
; CHECK-NEXT: ret;
@@ -54,7 +54,7 @@ define i32 @bmsk_clamp_ir(i32 %a) {
5454
; CHECK-NEXT: .reg .b32 %r<3>;
5555
; CHECK-EMPTY:
5656
; CHECK-NEXT: // %bb.0:
57-
; CHECK-NEXT: ld.param.u32 %r1, [bmsk_clamp_ir_param_0];
57+
; CHECK-NEXT: ld.param.b32 %r1, [bmsk_clamp_ir_param_0];
5858
; CHECK-NEXT: bmsk.clamp.b32 %r2, %r1, 7;
5959
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
6060
; CHECK-NEXT: ret;
@@ -68,7 +68,7 @@ define i32 @bmsk_wrap_ri(i32 %a) {
6868
; CHECK-NEXT: .reg .b32 %r<3>;
6969
; CHECK-EMPTY:
7070
; CHECK-NEXT: // %bb.0:
71-
; CHECK-NEXT: ld.param.u32 %r1, [bmsk_wrap_ri_param_0];
71+
; CHECK-NEXT: ld.param.b32 %r1, [bmsk_wrap_ri_param_0];
7272
; CHECK-NEXT: bmsk.wrap.b32 %r2, 5, %r1;
7373
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
7474
; CHECK-NEXT: ret;

‎llvm/test/CodeGen/NVPTX/byval-const-global.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -13,8 +13,8 @@ define void @foo() {
1313
; CHECK-NEXT: .reg .b64 %rd<3>;
1414
; CHECK-EMPTY:
1515
; CHECK-NEXT: // %bb.0:
16-
; CHECK-NEXT: ld.global.u64 %rd1, [G];
17-
; CHECK-NEXT: ld.global.u64 %rd2, [G+8];
16+
; CHECK-NEXT: ld.global.b64 %rd1, [G];
17+
; CHECK-NEXT: ld.global.b64 %rd2, [G+8];
1818
; CHECK-NEXT: { // callseq 0, 0
1919
; CHECK-NEXT: .param .align 8 .b8 param0[16];
2020
; CHECK-NEXT: st.param.b64 [param0], %rd1;

‎llvm/test/CodeGen/NVPTX/i128-ld-st.ll

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -10,11 +10,11 @@ define i128 @foo(ptr %p, ptr %o) {
1010
; CHECK-NEXT: .reg .b64 %rd<5>;
1111
; CHECK-EMPTY:
1212
; CHECK-NEXT: // %bb.0:
13-
; CHECK-NEXT: ld.param.u64 %rd2, [foo_param_1];
14-
; CHECK-NEXT: ld.param.u64 %rd1, [foo_param_0];
15-
; CHECK-NEXT: ld.u8 %rd3, [%rd1];
13+
; CHECK-NEXT: ld.param.b64 %rd2, [foo_param_1];
14+
; CHECK-NEXT: ld.param.b64 %rd1, [foo_param_0];
15+
; CHECK-NEXT: ld.b8 %rd3, [%rd1];
1616
; CHECK-NEXT: mov.b64 %rd4, 0;
17-
; CHECK-NEXT: st.v2.u64 [%rd2], {%rd3, %rd4};
17+
; CHECK-NEXT: st.v2.b64 [%rd2], {%rd3, %rd4};
1818
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd3, %rd4};
1919
; CHECK-NEXT: ret;
2020
%c = load i8, ptr %p, align 1

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

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -10,8 +10,8 @@ define i32 @szext_wrap_u32(i32 %a, i32 %b) {
1010
; CHECK-NEXT: .reg .b32 %r<4>;
1111
; CHECK-EMPTY:
1212
; CHECK-NEXT: // %bb.0:
13-
; CHECK-NEXT: ld.param.u32 %r1, [szext_wrap_u32_param_0];
14-
; CHECK-NEXT: ld.param.u32 %r2, [szext_wrap_u32_param_1];
13+
; CHECK-NEXT: ld.param.b32 %r1, [szext_wrap_u32_param_0];
14+
; CHECK-NEXT: ld.param.b32 %r2, [szext_wrap_u32_param_1];
1515
; CHECK-NEXT: szext.wrap.u32 %r3, %r1, %r2;
1616
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
1717
; CHECK-NEXT: ret;
@@ -25,8 +25,8 @@ define i32 @szext_clamp_u32(i32 %a, i32 %b) {
2525
; CHECK-NEXT: .reg .b32 %r<4>;
2626
; CHECK-EMPTY:
2727
; CHECK-NEXT: // %bb.0:
28-
; CHECK-NEXT: ld.param.u32 %r1, [szext_clamp_u32_param_0];
29-
; CHECK-NEXT: ld.param.u32 %r2, [szext_clamp_u32_param_1];
28+
; CHECK-NEXT: ld.param.b32 %r1, [szext_clamp_u32_param_0];
29+
; CHECK-NEXT: ld.param.b32 %r2, [szext_clamp_u32_param_1];
3030
; CHECK-NEXT: szext.clamp.u32 %r3, %r1, %r2;
3131
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
3232
; CHECK-NEXT: ret;
@@ -40,8 +40,8 @@ define i32 @szext_wrap_s32(i32 %a, i32 %b) {
4040
; CHECK-NEXT: .reg .b32 %r<4>;
4141
; CHECK-EMPTY:
4242
; CHECK-NEXT: // %bb.0:
43-
; CHECK-NEXT: ld.param.u32 %r1, [szext_wrap_s32_param_0];
44-
; CHECK-NEXT: ld.param.u32 %r2, [szext_wrap_s32_param_1];
43+
; CHECK-NEXT: ld.param.b32 %r1, [szext_wrap_s32_param_0];
44+
; CHECK-NEXT: ld.param.b32 %r2, [szext_wrap_s32_param_1];
4545
; CHECK-NEXT: szext.wrap.s32 %r3, %r1, %r2;
4646
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
4747
; CHECK-NEXT: ret;
@@ -55,8 +55,8 @@ define i32 @szext_clamp_s32(i32 %a, i32 %b) {
5555
; CHECK-NEXT: .reg .b32 %r<4>;
5656
; CHECK-EMPTY:
5757
; CHECK-NEXT: // %bb.0:
58-
; CHECK-NEXT: ld.param.u32 %r1, [szext_clamp_s32_param_0];
59-
; CHECK-NEXT: ld.param.u32 %r2, [szext_clamp_s32_param_1];
58+
; CHECK-NEXT: ld.param.b32 %r1, [szext_clamp_s32_param_0];
59+
; CHECK-NEXT: ld.param.b32 %r2, [szext_clamp_s32_param_1];
6060
; CHECK-NEXT: szext.clamp.s32 %r3, %r1, %r2;
6161
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
6262
; CHECK-NEXT: ret;
@@ -84,7 +84,7 @@ define i32 @szext_wrap_s32_ir(i32 %a) {
8484
; CHECK-NEXT: .reg .b32 %r<3>;
8585
; CHECK-EMPTY:
8686
; CHECK-NEXT: // %bb.0:
87-
; CHECK-NEXT: ld.param.u32 %r1, [szext_wrap_s32_ir_param_0];
87+
; CHECK-NEXT: ld.param.b32 %r1, [szext_wrap_s32_ir_param_0];
8888
; CHECK-NEXT: szext.wrap.s32 %r2, 5, %r1;
8989
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
9090
; CHECK-NEXT: ret;
@@ -98,7 +98,7 @@ define i32 @szext_clamp_u32_ri(i32 %a) {
9898
; CHECK-NEXT: .reg .b32 %r<3>;
9999
; CHECK-EMPTY:
100100
; CHECK-NEXT: // %bb.0:
101-
; CHECK-NEXT: ld.param.u32 %r1, [szext_clamp_u32_ri_param_0];
101+
; CHECK-NEXT: ld.param.b32 %r1, [szext_clamp_u32_ri_param_0];
102102
; CHECK-NEXT: szext.clamp.u32 %r2, %r1, 7;
103103
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
104104
; CHECK-NEXT: ret;

‎llvm/test/CodeGen/NVPTX/tag-invariant-loads.ll

Lines changed: 19 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -18,12 +18,12 @@ define ptx_kernel void @basic(ptr noalias readonly %a, ptr %out) {
1818
; PTX-NEXT: .reg .b32 %f<2>;
1919
; PTX-EMPTY:
2020
; PTX-NEXT: // %bb.0:
21-
; PTX-NEXT: ld.param.u32 %r1, [basic_param_0];
21+
; PTX-NEXT: ld.param.b32 %r1, [basic_param_0];
2222
; PTX-NEXT: cvta.to.global.u32 %r2, %r1;
23-
; PTX-NEXT: ld.param.u32 %r3, [basic_param_1];
23+
; PTX-NEXT: ld.param.b32 %r3, [basic_param_1];
2424
; PTX-NEXT: cvta.to.global.u32 %r4, %r3;
25-
; PTX-NEXT: ld.global.nc.f32 %f1, [%r2];
26-
; PTX-NEXT: st.global.f32 [%r4], %f1;
25+
; PTX-NEXT: ld.global.nc.b32 %f1, [%r2];
26+
; PTX-NEXT: st.global.b32 [%r4], %f1;
2727
; PTX-NEXT: ret;
2828
%a_global = addrspacecast ptr %a to ptr addrspace(1)
2929
%val = load float, ptr addrspace(1) %a_global
@@ -47,18 +47,18 @@ define ptx_kernel void @select(ptr noalias readonly %a, ptr noalias readonly %b,
4747
; PTX-NEXT: .reg .b32 %r<9>;
4848
; PTX-EMPTY:
4949
; PTX-NEXT: // %bb.0:
50-
; PTX-NEXT: ld.param.u8 %rs1, [select_param_2];
50+
; PTX-NEXT: ld.param.b8 %rs1, [select_param_2];
5151
; PTX-NEXT: and.b16 %rs2, %rs1, 1;
5252
; PTX-NEXT: setp.ne.b16 %p1, %rs2, 0;
53-
; PTX-NEXT: ld.param.u32 %r1, [select_param_0];
53+
; PTX-NEXT: ld.param.b32 %r1, [select_param_0];
5454
; PTX-NEXT: cvta.to.global.u32 %r2, %r1;
55-
; PTX-NEXT: ld.param.u32 %r3, [select_param_1];
55+
; PTX-NEXT: ld.param.b32 %r3, [select_param_1];
5656
; PTX-NEXT: cvta.to.global.u32 %r4, %r3;
57-
; PTX-NEXT: ld.param.u32 %r5, [select_param_3];
57+
; PTX-NEXT: ld.param.b32 %r5, [select_param_3];
5858
; PTX-NEXT: cvta.to.global.u32 %r6, %r5;
5959
; PTX-NEXT: selp.b32 %r7, %r2, %r4, %p1;
60-
; PTX-NEXT: ld.global.nc.u32 %r8, [%r7];
61-
; PTX-NEXT: st.global.u32 [%r6], %r8;
60+
; PTX-NEXT: ld.global.nc.b32 %r8, [%r7];
61+
; PTX-NEXT: st.global.b32 [%r6], %r8;
6262
; PTX-NEXT: ret;
6363
%select = select i1 %c, ptr %a, ptr %b
6464
%select_global = addrspacecast ptr %select to ptr addrspace(1)
@@ -81,11 +81,11 @@ define void @not_kernel(ptr noalias readonly %a, ptr %out) {
8181
; PTX-NEXT: .reg .b32 %f<2>;
8282
; PTX-EMPTY:
8383
; PTX-NEXT: // %bb.0:
84-
; PTX-NEXT: ld.param.u32 %r1, [not_kernel_param_0];
84+
; PTX-NEXT: ld.param.b32 %r1, [not_kernel_param_0];
8585
; PTX-NEXT: cvta.to.global.u32 %r2, %r1;
86-
; PTX-NEXT: ld.param.u32 %r3, [not_kernel_param_1];
87-
; PTX-NEXT: ld.global.f32 %f1, [%r2];
88-
; PTX-NEXT: st.f32 [%r3], %f1;
86+
; PTX-NEXT: ld.param.b32 %r3, [not_kernel_param_1];
87+
; PTX-NEXT: ld.global.b32 %f1, [%r2];
88+
; PTX-NEXT: st.b32 [%r3], %f1;
8989
; PTX-NEXT: ret;
9090
%a_global = addrspacecast ptr %a to ptr addrspace(1)
9191
%val = load float, ptr addrspace(1) %a_global
@@ -114,17 +114,17 @@ define ptx_kernel void @global_load(ptr noalias readonly %a, i1 %c, ptr %out) {
114114
; PTX-NEXT: .reg .b64 %rd<2>;
115115
; PTX-EMPTY:
116116
; PTX-NEXT: // %bb.0:
117-
; PTX-NEXT: ld.param.u8 %rs1, [global_load_param_1];
117+
; PTX-NEXT: ld.param.b8 %rs1, [global_load_param_1];
118118
; PTX-NEXT: and.b16 %rs2, %rs1, 1;
119119
; PTX-NEXT: setp.ne.b16 %p1, %rs2, 0;
120-
; PTX-NEXT: ld.param.u32 %r1, [global_load_param_0];
120+
; PTX-NEXT: ld.param.b32 %r1, [global_load_param_0];
121121
; PTX-NEXT: cvta.to.global.u32 %r2, %r1;
122-
; PTX-NEXT: ld.param.u32 %r3, [global_load_param_2];
122+
; PTX-NEXT: ld.param.b32 %r3, [global_load_param_2];
123123
; PTX-NEXT: cvta.to.global.u32 %r4, %r3;
124124
; PTX-NEXT: mov.b32 %r5, G;
125125
; PTX-NEXT: selp.b32 %r6, %r5, %r2, %p1;
126-
; PTX-NEXT: ld.global.nc.u64 %rd1, [%r6];
127-
; PTX-NEXT: st.global.u64 [%r4], %rd1;
126+
; PTX-NEXT: ld.global.nc.b64 %rd1, [%r6];
127+
; PTX-NEXT: st.global.b64 [%r4], %rd1;
128128
; PTX-NEXT: ret;
129129
%g_global = addrspacecast ptr @G to ptr addrspace(1)
130130
%a_global = addrspacecast ptr %a to ptr addrspace(1)

0 commit comments

Comments
 (0)