@@ -584,44 +584,25 @@ define ptx_kernel void @test_select(ptr byval(i32) align 4 %input1, ptr byval(i3
584584; COPY-NEXT: store i32 [[VALLOADED]], ptr [[OUT]], align 4
585585; COPY-NEXT: ret void
586586;
587- ; PTX_60-LABEL: test_select(
588- ; PTX_60: {
589- ; PTX_60-NEXT: .reg .pred %p<2>;
590- ; PTX_60-NEXT: .reg .b16 %rs<3>;
591- ; PTX_60-NEXT: .reg .b32 %r<4>;
592- ; PTX_60-NEXT: .reg .b64 %rd<3>;
593- ; PTX_60-EMPTY:
594- ; PTX_60-NEXT: // %bb.0: // %bb
595- ; PTX_60-NEXT: ld.param.b8 %rs1, [test_select_param_3];
596- ; PTX_60-NEXT: and.b16 %rs2, %rs1, 1;
597- ; PTX_60-NEXT: setp.ne.b16 %p1, %rs2, 0;
598- ; PTX_60-NEXT: ld.param.b64 %rd1, [test_select_param_2];
599- ; PTX_60-NEXT: cvta.to.global.u64 %rd2, %rd1;
600- ; PTX_60-NEXT: ld.param.b32 %r1, [test_select_param_1];
601- ; PTX_60-NEXT: ld.param.b32 %r2, [test_select_param_0];
602- ; PTX_60-NEXT: selp.b32 %r3, %r2, %r1, %p1;
603- ; PTX_60-NEXT: st.global.b32 [%rd2], %r3;
604- ; PTX_60-NEXT: ret;
605- ;
606- ; PTX_70-LABEL: test_select(
607- ; PTX_70: {
608- ; PTX_70-NEXT: .reg .pred %p<2>;
609- ; PTX_70-NEXT: .reg .b16 %rs<3>;
610- ; PTX_70-NEXT: .reg .b32 %r<2>;
611- ; PTX_70-NEXT: .reg .b64 %rd<6>;
612- ; PTX_70-EMPTY:
613- ; PTX_70-NEXT: // %bb.0: // %bb
614- ; PTX_70-NEXT: ld.param.b8 %rs1, [test_select_param_3];
615- ; PTX_70-NEXT: and.b16 %rs2, %rs1, 1;
616- ; PTX_70-NEXT: setp.ne.b16 %p1, %rs2, 0;
617- ; PTX_70-NEXT: mov.b64 %rd1, test_select_param_0;
618- ; PTX_70-NEXT: ld.param.b64 %rd2, [test_select_param_2];
619- ; PTX_70-NEXT: cvta.to.global.u64 %rd3, %rd2;
620- ; PTX_70-NEXT: mov.b64 %rd4, test_select_param_1;
621- ; PTX_70-NEXT: selp.b64 %rd5, %rd1, %rd4, %p1;
622- ; PTX_70-NEXT: ld.param.b32 %r1, [%rd5];
623- ; PTX_70-NEXT: st.global.b32 [%rd3], %r1;
624- ; PTX_70-NEXT: ret;
587+ ; PTX-LABEL: test_select(
588+ ; PTX: {
589+ ; PTX-NEXT: .reg .pred %p<2>;
590+ ; PTX-NEXT: .reg .b16 %rs<3>;
591+ ; PTX-NEXT: .reg .b32 %r<2>;
592+ ; PTX-NEXT: .reg .b64 %rd<6>;
593+ ; PTX-EMPTY:
594+ ; PTX-NEXT: // %bb.0: // %bb
595+ ; PTX-NEXT: ld.param.b8 %rs1, [test_select_param_3];
596+ ; PTX-NEXT: and.b16 %rs2, %rs1, 1;
597+ ; PTX-NEXT: setp.ne.b16 %p1, %rs2, 0;
598+ ; PTX-NEXT: mov.b64 %rd1, test_select_param_0;
599+ ; PTX-NEXT: ld.param.b64 %rd2, [test_select_param_2];
600+ ; PTX-NEXT: cvta.to.global.u64 %rd3, %rd2;
601+ ; PTX-NEXT: mov.b64 %rd4, test_select_param_1;
602+ ; PTX-NEXT: selp.b64 %rd5, %rd1, %rd4, %p1;
603+ ; PTX-NEXT: ld.param.b32 %r1, [%rd5];
604+ ; PTX-NEXT: st.global.b32 [%rd3], %r1;
605+ ; PTX-NEXT: ret;
625606bb:
626607 %ptrnew = select i1 %cond , ptr %input1 , ptr %input2
627608 %valloaded = load i32 , ptr %ptrnew , align 4
0 commit comments