128 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s | FileCheck %s3 4target triple = "nvptx64-nvidia-cuda"5 6define i32 @test_ld_param_const(ptr byval(i32) %a) {7; CHECK-LABEL: test_ld_param_const(8; CHECK: {9; CHECK-NEXT: .reg .b32 %r<2>;10; CHECK-EMPTY:11; CHECK-NEXT: // %bb.0:12; CHECK-NEXT: ld.param.b32 %r1, [test_ld_param_const_param_0+4];13; CHECK-NEXT: st.param.b32 [func_retval0], %r1;14; CHECK-NEXT: ret;15 %p2 = getelementptr i32, ptr %a, i32 116 %ld = load i32, ptr %p217 ret i32 %ld18}19 20define i32 @test_ld_param_non_const(ptr byval([10 x i32]) %a, i32 %b) {21; CHECK-LABEL: test_ld_param_non_const(22; CHECK: {23; CHECK-NEXT: .reg .b32 %r<2>;24; CHECK-NEXT: .reg .b64 %rd<4>;25; CHECK-EMPTY:26; CHECK-NEXT: // %bb.0:27; CHECK-NEXT: mov.b64 %rd1, test_ld_param_non_const_param_0;28; CHECK-NEXT: ld.param.s32 %rd2, [test_ld_param_non_const_param_1];29; CHECK-NEXT: add.s64 %rd3, %rd1, %rd2;30; CHECK-NEXT: ld.local.b32 %r1, [%rd3];31; CHECK-NEXT: st.param.b32 [func_retval0], %r1;32; CHECK-NEXT: ret;33 %p2 = getelementptr i8, ptr %a, i32 %b34 %ld = load i32, ptr %p235 ret i32 %ld36}37 38declare void @escape(ptr)39declare void @byval_user(ptr byval(i32))40 41define void @test_ld_param_escaping(ptr byval(i32) %a) {42; CHECK-LABEL: test_ld_param_escaping(43; CHECK: {44; CHECK-NEXT: .reg .b64 %rd<3>;45; CHECK-EMPTY:46; CHECK-NEXT: // %bb.0:47; CHECK-NEXT: mov.b64 %rd1, test_ld_param_escaping_param_0;48; CHECK-NEXT: cvta.local.u64 %rd2, %rd1;49; CHECK-NEXT: { // callseq 0, 050; CHECK-NEXT: .param .b64 param0;51; CHECK-NEXT: st.param.b64 [param0], %rd2;52; CHECK-NEXT: call.uni escape, (param0);53; CHECK-NEXT: } // callseq 054; CHECK-NEXT: ret;55 call void @escape(ptr %a)56 ret void57}58 59define void @test_ld_param_byval(ptr byval(i32) %a) {60; CHECK-LABEL: test_ld_param_byval(61; CHECK: {62; CHECK-NEXT: .reg .b32 %r<2>;63; CHECK-EMPTY:64; CHECK-NEXT: // %bb.0:65; CHECK-NEXT: { // callseq 1, 066; CHECK-NEXT: .param .align 4 .b8 param0[4];67; CHECK-NEXT: ld.param.b32 %r1, [test_ld_param_byval_param_0];68; CHECK-NEXT: st.param.b32 [param0], %r1;69; CHECK-NEXT: call.uni byval_user, (param0);70; CHECK-NEXT: } // callseq 171; CHECK-NEXT: ret;72 call void @byval_user(ptr %a)73 ret void74}75 76define i32 @test_modify_param(ptr byval([10 x i32]) %a, i32 %b, i32 %c ) {77; CHECK-LABEL: test_modify_param(78; CHECK: {79; CHECK-NEXT: .reg .b32 %r<3>;80; CHECK-NEXT: .reg .b64 %rd<2>;81; CHECK-EMPTY:82; CHECK-NEXT: // %bb.0:83; CHECK-NEXT: mov.b64 %rd1, test_modify_param_param_0;84; CHECK-NEXT: ld.param.b32 %r1, [test_modify_param_param_1];85; CHECK-NEXT: ld.param.b32 %r2, [test_modify_param_param_2];86; CHECK-NEXT: st.local.b32 [%rd1+2], %r1;87; CHECK-NEXT: st.param.b32 [func_retval0], %r2;88; CHECK-NEXT: ret;89 %p2 = getelementptr i8, ptr %a, i32 290 store volatile i32 %b, ptr %p291 ret i32 %c92}93 94define i32 @test_multi_block(ptr byval([10 x i32]) %a, i1 %p) {95; CHECK-LABEL: test_multi_block(96; CHECK: {97; CHECK-NEXT: .reg .pred %p<3>;98; CHECK-NEXT: .reg .b16 %rs<3>;99; CHECK-NEXT: .reg .b32 %r<2>;100; CHECK-EMPTY:101; CHECK-NEXT: // %bb.0:102; CHECK-NEXT: ld.param.b8 %rs1, [test_multi_block_param_1];103; CHECK-NEXT: and.b16 %rs2, %rs1, 1;104; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;105; CHECK-NEXT: not.pred %p2, %p1;106; CHECK-NEXT: @%p2 bra $L__BB5_2;107; CHECK-NEXT: // %bb.1: // %if108; CHECK-NEXT: ld.param.b32 %r1, [test_multi_block_param_0+4];109; CHECK-NEXT: bra.uni $L__BB5_3;110; CHECK-NEXT: $L__BB5_2: // %else111; CHECK-NEXT: ld.param.b32 %r1, [test_multi_block_param_0+8];112; CHECK-NEXT: $L__BB5_3: // %end113; CHECK-NEXT: st.param.b32 [func_retval0], %r1;114; CHECK-NEXT: ret;115 br i1 %p, label %if, label %else116if:117 %p2 = getelementptr i8, ptr %a, i32 4118 %v2 = load i32, ptr %p2119 br label %end120else:121 %p3 = getelementptr i8, ptr %a, i32 8122 %v3 = load i32, ptr %p3123 br label %end124end:125 %v = phi i32 [ %v2, %if ], [ %v3, %else ]126 ret i32 %v127}128