84 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_60 -mattr=+ptx73 | FileCheck %s --check-prefix=CHECK-323; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_60 -mattr=+ptx73 | FileCheck %s --check-prefix=CHECK-644; RUN: llc < %s -mtriple=nvptx64 -nvptx-short-ptr -mcpu=sm_60 -mattr=+ptx73 | FileCheck %s --check-prefix=CHECK-MIXED5; RUN: %if ptxas-sm_60 && ptxas-isa-7.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_60 -mattr=+ptx73 | %ptxas-verify -arch=sm_60 %}6 7target triple = "nvptx64-nvidia-cuda"8 9define ptr @test_save() {10; CHECK-32-LABEL: test_save(11; CHECK-32: {12; CHECK-32-NEXT: .reg .b32 %r<3>;13; CHECK-32-EMPTY:14; CHECK-32-NEXT: // %bb.0:15; CHECK-32-NEXT: stacksave.u32 %r1;16; CHECK-32-NEXT: cvta.local.u32 %r2, %r1;17; CHECK-32-NEXT: st.param.b32 [func_retval0], %r2;18; CHECK-32-NEXT: ret;19;20; CHECK-64-LABEL: test_save(21; CHECK-64: {22; CHECK-64-NEXT: .reg .b64 %rd<3>;23; CHECK-64-EMPTY:24; CHECK-64-NEXT: // %bb.0:25; CHECK-64-NEXT: stacksave.u64 %rd1;26; CHECK-64-NEXT: cvta.local.u64 %rd2, %rd1;27; CHECK-64-NEXT: st.param.b64 [func_retval0], %rd2;28; CHECK-64-NEXT: ret;29;30; CHECK-MIXED-LABEL: test_save(31; CHECK-MIXED: {32; CHECK-MIXED-NEXT: .reg .b32 %r<2>;33; CHECK-MIXED-NEXT: .reg .b64 %rd<3>;34; CHECK-MIXED-EMPTY:35; CHECK-MIXED-NEXT: // %bb.0:36; CHECK-MIXED-NEXT: stacksave.u32 %r1;37; CHECK-MIXED-NEXT: cvt.u64.u32 %rd1, %r1;38; CHECK-MIXED-NEXT: cvta.local.u64 %rd2, %rd1;39; CHECK-MIXED-NEXT: st.param.b64 [func_retval0], %rd2;40; CHECK-MIXED-NEXT: ret;41 %1 = call ptr @llvm.stacksave()42 ret ptr %143}44 45 46define void @test_restore(ptr %p) {47; CHECK-32-LABEL: test_restore(48; CHECK-32: {49; CHECK-32-NEXT: .reg .b32 %r<3>;50; CHECK-32-EMPTY:51; CHECK-32-NEXT: // %bb.0:52; CHECK-32-NEXT: ld.param.b32 %r1, [test_restore_param_0];53; CHECK-32-NEXT: cvta.to.local.u32 %r2, %r1;54; CHECK-32-NEXT: stackrestore.u32 %r2;55; CHECK-32-NEXT: ret;56;57; CHECK-64-LABEL: test_restore(58; CHECK-64: {59; CHECK-64-NEXT: .reg .b64 %rd<3>;60; CHECK-64-EMPTY:61; CHECK-64-NEXT: // %bb.0:62; CHECK-64-NEXT: ld.param.b64 %rd1, [test_restore_param_0];63; CHECK-64-NEXT: cvta.to.local.u64 %rd2, %rd1;64; CHECK-64-NEXT: stackrestore.u64 %rd2;65; CHECK-64-NEXT: ret;66;67; CHECK-MIXED-LABEL: test_restore(68; CHECK-MIXED: {69; CHECK-MIXED-NEXT: .reg .b32 %r<2>;70; CHECK-MIXED-NEXT: .reg .b64 %rd<3>;71; CHECK-MIXED-EMPTY:72; CHECK-MIXED-NEXT: // %bb.0:73; CHECK-MIXED-NEXT: ld.param.b64 %rd1, [test_restore_param_0];74; CHECK-MIXED-NEXT: cvta.to.local.u64 %rd2, %rd1;75; CHECK-MIXED-NEXT: cvt.u32.u64 %r1, %rd2;76; CHECK-MIXED-NEXT: stackrestore.u32 %r1;77; CHECK-MIXED-NEXT: ret;78 call void @llvm.stackrestore(ptr %p)79 ret void80}81 82declare ptr @llvm.stacksave()83declare void @llvm.stackrestore(ptr)84