197 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=PTX323; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=PTX644; RUN: %if ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}5; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}6 7; Ensure we access the local stack properly8 9define void @foo(i32 %a) {10; PTX32-LABEL: foo(11; PTX32: {12; PTX32-NEXT: .local .align 4 .b8 __local_depot0[4];13; PTX32-NEXT: .reg .b32 %SP;14; PTX32-NEXT: .reg .b32 %SPL;15; PTX32-NEXT: .reg .b32 %r<3>;16; PTX32-EMPTY:17; PTX32-NEXT: // %bb.0:18; PTX32-NEXT: mov.b32 %SPL, __local_depot0;19; PTX32-NEXT: ld.param.b32 %r1, [foo_param_0];20; PTX32-NEXT: add.u32 %r2, %SPL, 0;21; PTX32-NEXT: st.local.b32 [%r2], %r1;22; PTX32-NEXT: ret;23;24; PTX64-LABEL: foo(25; PTX64: {26; PTX64-NEXT: .local .align 4 .b8 __local_depot0[4];27; PTX64-NEXT: .reg .b64 %SP;28; PTX64-NEXT: .reg .b64 %SPL;29; PTX64-NEXT: .reg .b32 %r<2>;30; PTX64-NEXT: .reg .b64 %rd<2>;31; PTX64-EMPTY:32; PTX64-NEXT: // %bb.0:33; PTX64-NEXT: mov.b64 %SPL, __local_depot0;34; PTX64-NEXT: ld.param.b32 %r1, [foo_param_0];35; PTX64-NEXT: add.u64 %rd1, %SPL, 0;36; PTX64-NEXT: st.local.b32 [%rd1], %r1;37; PTX64-NEXT: ret;38 %local = alloca i32, align 439 store volatile i32 %a, ptr %local40 ret void41}42 43define ptx_kernel void @foo2(i32 %a) {44; PTX32-LABEL: foo2(45; PTX32: {46; PTX32-NEXT: .local .align 4 .b8 __local_depot1[4];47; PTX32-NEXT: .reg .b32 %SP;48; PTX32-NEXT: .reg .b32 %SPL;49; PTX32-NEXT: .reg .b32 %r<4>;50; PTX32-EMPTY:51; PTX32-NEXT: // %bb.0:52; PTX32-NEXT: mov.b32 %SPL, __local_depot1;53; PTX32-NEXT: cvta.local.u32 %SP, %SPL;54; PTX32-NEXT: ld.param.b32 %r1, [foo2_param_0];55; PTX32-NEXT: add.u32 %r2, %SP, 0;56; PTX32-NEXT: add.u32 %r3, %SPL, 0;57; PTX32-NEXT: st.local.b32 [%r3], %r1;58; PTX32-NEXT: { // callseq 0, 059; PTX32-NEXT: .param .b32 param0;60; PTX32-NEXT: st.param.b32 [param0], %r2;61; PTX32-NEXT: call.uni bar, (param0);62; PTX32-NEXT: } // callseq 063; PTX32-NEXT: ret;64;65; PTX64-LABEL: foo2(66; PTX64: {67; PTX64-NEXT: .local .align 4 .b8 __local_depot1[4];68; PTX64-NEXT: .reg .b64 %SP;69; PTX64-NEXT: .reg .b64 %SPL;70; PTX64-NEXT: .reg .b32 %r<2>;71; PTX64-NEXT: .reg .b64 %rd<3>;72; PTX64-EMPTY:73; PTX64-NEXT: // %bb.0:74; PTX64-NEXT: mov.b64 %SPL, __local_depot1;75; PTX64-NEXT: cvta.local.u64 %SP, %SPL;76; PTX64-NEXT: ld.param.b32 %r1, [foo2_param_0];77; PTX64-NEXT: add.u64 %rd1, %SP, 0;78; PTX64-NEXT: add.u64 %rd2, %SPL, 0;79; PTX64-NEXT: st.local.b32 [%rd2], %r1;80; PTX64-NEXT: { // callseq 0, 081; PTX64-NEXT: .param .b64 param0;82; PTX64-NEXT: st.param.b64 [param0], %rd1;83; PTX64-NEXT: call.uni bar, (param0);84; PTX64-NEXT: } // callseq 085; PTX64-NEXT: ret;86 %local = alloca i32, align 487 store i32 %a, ptr %local88 call void @bar(ptr %local)89 ret void90}91 92declare void @bar(ptr %a)93 94define void @foo3(i32 %a) {95; PTX32-LABEL: foo3(96; PTX32: {97; PTX32-NEXT: .local .align 4 .b8 __local_depot2[12];98; PTX32-NEXT: .reg .b32 %SP;99; PTX32-NEXT: .reg .b32 %SPL;100; PTX32-NEXT: .reg .b32 %r<5>;101; PTX32-EMPTY:102; PTX32-NEXT: // %bb.0:103; PTX32-NEXT: mov.b32 %SPL, __local_depot2;104; PTX32-NEXT: ld.param.b32 %r1, [foo3_param_0];105; PTX32-NEXT: add.u32 %r2, %SPL, 0;106; PTX32-NEXT: shl.b32 %r3, %r1, 2;107; PTX32-NEXT: add.s32 %r4, %r2, %r3;108; PTX32-NEXT: st.local.b32 [%r4], %r1;109; PTX32-NEXT: ret;110;111; PTX64-LABEL: foo3(112; PTX64: {113; PTX64-NEXT: .local .align 4 .b8 __local_depot2[12];114; PTX64-NEXT: .reg .b64 %SP;115; PTX64-NEXT: .reg .b64 %SPL;116; PTX64-NEXT: .reg .b32 %r<2>;117; PTX64-NEXT: .reg .b64 %rd<4>;118; PTX64-EMPTY:119; PTX64-NEXT: // %bb.0:120; PTX64-NEXT: mov.b64 %SPL, __local_depot2;121; PTX64-NEXT: ld.param.b32 %r1, [foo3_param_0];122; PTX64-NEXT: add.u64 %rd1, %SPL, 0;123; PTX64-NEXT: mul.wide.s32 %rd2, %r1, 4;124; PTX64-NEXT: add.s64 %rd3, %rd1, %rd2;125; PTX64-NEXT: st.local.b32 [%rd3], %r1;126; PTX64-NEXT: ret;127 %local = alloca [3 x i32], align 4128 %1 = getelementptr inbounds i32, ptr %local, i32 %a129 store i32 %a, ptr %1130 ret void131}132 133define void @foo4() {134; PTX32-LABEL: foo4(135; PTX32: {136; PTX32-NEXT: .local .align 4 .b8 __local_depot3[8];137; PTX32-NEXT: .reg .b32 %SP;138; PTX32-NEXT: .reg .b32 %SPL;139; PTX32-NEXT: .reg .b32 %r<5>;140; PTX32-EMPTY:141; PTX32-NEXT: // %bb.0:142; PTX32-NEXT: mov.b32 %SPL, __local_depot3;143; PTX32-NEXT: cvta.local.u32 %SP, %SPL;144; PTX32-NEXT: add.u32 %r1, %SP, 0;145; PTX32-NEXT: add.u32 %r2, %SPL, 0;146; PTX32-NEXT: add.u32 %r3, %SP, 4;147; PTX32-NEXT: add.u32 %r4, %SPL, 4;148; PTX32-NEXT: st.local.b32 [%r2], 0;149; PTX32-NEXT: st.local.b32 [%r4], 0;150; PTX32-NEXT: { // callseq 1, 0151; PTX32-NEXT: .param .b32 param0;152; PTX32-NEXT: st.param.b32 [param0], %r1;153; PTX32-NEXT: call.uni bar, (param0);154; PTX32-NEXT: } // callseq 1155; PTX32-NEXT: { // callseq 2, 0156; PTX32-NEXT: .param .b32 param0;157; PTX32-NEXT: st.param.b32 [param0], %r3;158; PTX32-NEXT: call.uni bar, (param0);159; PTX32-NEXT: } // callseq 2160; PTX32-NEXT: ret;161;162; PTX64-LABEL: foo4(163; PTX64: {164; PTX64-NEXT: .local .align 4 .b8 __local_depot3[8];165; PTX64-NEXT: .reg .b64 %SP;166; PTX64-NEXT: .reg .b64 %SPL;167; PTX64-NEXT: .reg .b64 %rd<5>;168; PTX64-EMPTY:169; PTX64-NEXT: // %bb.0:170; PTX64-NEXT: mov.b64 %SPL, __local_depot3;171; PTX64-NEXT: cvta.local.u64 %SP, %SPL;172; PTX64-NEXT: add.u64 %rd1, %SP, 0;173; PTX64-NEXT: add.u64 %rd2, %SPL, 0;174; PTX64-NEXT: add.u64 %rd3, %SP, 4;175; PTX64-NEXT: add.u64 %rd4, %SPL, 4;176; PTX64-NEXT: st.local.b32 [%rd2], 0;177; PTX64-NEXT: st.local.b32 [%rd4], 0;178; PTX64-NEXT: { // callseq 1, 0179; PTX64-NEXT: .param .b64 param0;180; PTX64-NEXT: st.param.b64 [param0], %rd1;181; PTX64-NEXT: call.uni bar, (param0);182; PTX64-NEXT: } // callseq 1183; PTX64-NEXT: { // callseq 2, 0184; PTX64-NEXT: .param .b64 param0;185; PTX64-NEXT: st.param.b64 [param0], %rd3;186; PTX64-NEXT: call.uni bar, (param0);187; PTX64-NEXT: } // callseq 2188; PTX64-NEXT: ret;189 %A = alloca i32190 %B = alloca i32191 store i32 0, ptr %A192 store i32 0, ptr %B193 call void @bar(ptr %A)194 call void @bar(ptr %B)195 ret void196}197