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
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
|
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}
target triple = "nvptx64-nvidia-cuda"
@u128_max = internal addrspace(1) global i128 0, align 16
@u128_zero = internal addrspace(1) global i128 0, align 16
@i128_max = internal addrspace(1) global i128 0, align 16
@i128_min = internal addrspace(1) global i128 0, align 16
@v_u128_max = internal addrspace(1) global i128 0, align 16
@v_u128_zero = internal addrspace(1) global i128 0, align 16
@v_i128_max = internal addrspace(1) global i128 0, align 16
@v_i128_min = internal addrspace(1) global i128 0, align 16
@v64 = internal addrspace(1) global ptr null, align 8
define void @test_corner_values() {
; CHECK-LABEL: test_corner_values(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<24>;
; CHECK-NEXT: .reg .b128 %rq<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.global.u64 %rd1, [v64];
; CHECK-NEXT: add.s64 %rd2, %rd1, 8;
; CHECK-NEXT: mov.u64 %rd13, -1;
; CHECK-NEXT: mov.b128 %rq1, {%rd13, %rd13};
; CHECK-NEXT: mov.u64 %rd14, v_u128_max;
; CHECK-NEXT: cvta.global.u64 %rd3, %rd14;
; CHECK-NEXT: // begin inline asm
; CHECK-NEXT: {
; CHECK-NEXT: .reg .b64 hi;
; CHECK-NEXT: .reg .b64 lo;
; CHECK-NEXT: mov.b128 {lo, hi}, %rq1;
; CHECK-NEXT: st.b64 [%rd1], lo;
; CHECK-NEXT: st.b64 [%rd2], hi;
; CHECK-NEXT: st.b128 [%rd3], %rq1;
; CHECK-NEXT: }
; CHECK-NEXT: // end inline asm
; CHECK-NEXT: ld.global.u64 %rd15, [v64];
; CHECK-NEXT: add.s64 %rd4, %rd15, 16;
; CHECK-NEXT: add.s64 %rd5, %rd15, 24;
; CHECK-NEXT: mov.u64 %rd16, 9223372036854775807;
; CHECK-NEXT: mov.b128 %rq2, {%rd13, %rd16};
; CHECK-NEXT: mov.u64 %rd17, v_i128_max;
; CHECK-NEXT: cvta.global.u64 %rd6, %rd17;
; CHECK-NEXT: // begin inline asm
; CHECK-NEXT: {
; CHECK-NEXT: .reg .b64 hi;
; CHECK-NEXT: .reg .b64 lo;
; CHECK-NEXT: mov.b128 {lo, hi}, %rq2;
; CHECK-NEXT: st.b64 [%rd4], lo;
; CHECK-NEXT: st.b64 [%rd5], hi;
; CHECK-NEXT: st.b128 [%rd6], %rq2;
; CHECK-NEXT: }
; CHECK-NEXT: // end inline asm
; CHECK-NEXT: ld.global.u64 %rd18, [v64];
; CHECK-NEXT: add.s64 %rd7, %rd18, 32;
; CHECK-NEXT: add.s64 %rd8, %rd18, 40;
; CHECK-NEXT: mov.u64 %rd19, -9223372036854775808;
; CHECK-NEXT: mov.u64 %rd20, 0;
; CHECK-NEXT: mov.b128 %rq3, {%rd20, %rd19};
; CHECK-NEXT: mov.u64 %rd21, v_i128_min;
; CHECK-NEXT: cvta.global.u64 %rd9, %rd21;
; CHECK-NEXT: // begin inline asm
; CHECK-NEXT: {
; CHECK-NEXT: .reg .b64 hi;
; CHECK-NEXT: .reg .b64 lo;
; CHECK-NEXT: mov.b128 {lo, hi}, %rq3;
; CHECK-NEXT: st.b64 [%rd7], lo;
; CHECK-NEXT: st.b64 [%rd8], hi;
; CHECK-NEXT: st.b128 [%rd9], %rq3;
; CHECK-NEXT: }
; CHECK-NEXT: // end inline asm
; CHECK-NEXT: ld.global.u64 %rd22, [v64];
; CHECK-NEXT: add.s64 %rd10, %rd22, 48;
; CHECK-NEXT: add.s64 %rd11, %rd22, 56;
; CHECK-NEXT: mov.b128 %rq4, {%rd20, %rd20};
; CHECK-NEXT: mov.u64 %rd23, v_u128_zero;
; CHECK-NEXT: cvta.global.u64 %rd12, %rd23;
; CHECK-NEXT: // begin inline asm
; CHECK-NEXT: {
; CHECK-NEXT: .reg .b64 hi;
; CHECK-NEXT: .reg .b64 lo;
; CHECK-NEXT: mov.b128 {lo, hi}, %rq4;
; CHECK-NEXT: st.b64 [%rd10], lo;
; CHECK-NEXT: st.b64 [%rd11], hi;
; CHECK-NEXT: st.b128 [%rd12], %rq4;
; CHECK-NEXT: }
; CHECK-NEXT: // end inline asm
; CHECK-NEXT: ret;
%1 = load ptr, ptr addrspace(1) @v64, align 8
%2 = getelementptr inbounds i64, ptr %1, i64 1
tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 -1, ptr %1, ptr nonnull %2, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_max to ptr))
%3 = load ptr, ptr addrspace(1) @v64, align 8
%4 = getelementptr inbounds i64, ptr %3, i64 2
%5 = getelementptr inbounds i64, ptr %3, i64 3
tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 170141183460469231731687303715884105727, ptr nonnull %4, ptr nonnull %5, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_max to ptr))
%6 = load ptr, ptr addrspace(1) @v64, align 8
%7 = getelementptr inbounds i64, ptr %6, i64 4
%8 = getelementptr inbounds i64, ptr %6, i64 5
tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 -170141183460469231731687303715884105728, ptr nonnull %7, ptr nonnull %8, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_min to ptr))
%9 = load ptr, ptr addrspace(1) @v64, align 8
%10 = getelementptr inbounds i64, ptr %9, i64 6
%11 = getelementptr inbounds i64, ptr %9, i64 7
tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 0, ptr nonnull %10, ptr nonnull %11, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_zero to ptr))
ret void
}
|
