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