xref: /llvm-project/llvm/test/CodeGen/NVPTX/stacksaverestore.ll (revision b279f6b098d3849f7f1c1f539b108307d5f8ae2d) 
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2; RUN: llc < %s -mtriple=nvptx -mcpu=sm_60 -mattr=+ptx73 | FileCheck %s --check-prefix=CHECK-32
3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_60 -mattr=+ptx73 | FileCheck %s --check-prefix=CHECK-64
4; RUN: llc < %s -mtriple=nvptx64 -nvptx-short-ptr -mcpu=sm_60 -mattr=+ptx73 | FileCheck %s --check-prefix=CHECK-MIXED
5; RUN: %if ptxas && ptxas-12.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_60 -mattr=+ptx73 | %ptxas-verify %}
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 %1
43}
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.u32 %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.u64 %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.u64 %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 void
80}
81
82declare ptr @llvm.stacksave()
83declare void @llvm.stackrestore(ptr)
84