-
Notifications
You must be signed in to change notification settings - Fork 0
/
Copy pathnonphysical_core.ptx
119 lines (114 loc) · 3.73 KB
/
nonphysical_core.ptx
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
110
111
112
113
114
115
116
117
118
119
//
// Generated by LLVM NVPTX Back-End
//
.version 7.0
.target sm_80
.address_size 64
// .globl _ZN16nonphysical_core6random3pcg29PermutedCongruentialGenerator13shuffle_usize17h4098009cdef9b8bcE
.visible .func _ZN4core9panicking18panic_bounds_check17h4e5c286e8d48ca34E
(
.param .b64 _ZN4core9panicking18panic_bounds_check17h4e5c286e8d48ca34E_param_0,
.param .b64 _ZN4core9panicking18panic_bounds_check17h4e5c286e8d48ca34E_param_1,
.param .b64 _ZN4core9panicking18panic_bounds_check17h4e5c286e8d48ca34E_param_2
)
.noreturn{
trap;
exit;
}
.global .align 1 .b8 anon_$_2cbddb30228dc0cec36ed13501cf80e7_$_0[34] = {110, 111, 110, 112, 104, 121, 115, 105, 99, 97, 108, 95, 99, 111, 114, 101, 92, 115, 114, 99, 92, 114, 97, 110, 100, 111, 109, 92, 112, 99, 103, 46, 114, 115};
.global .align 8 .u64 anon_$_2cbddb30228dc0cec36ed13501cf80e7_$_1[3] = {generic(anon_$_2cbddb30228dc0cec36ed13501cf80e7_$_0), 34, 64424509497};
.visible .func _ZN16nonphysical_core6random3pcg29PermutedCongruentialGenerator13shuffle_usize17h4098009cdef9b8bcE(
.param .b64 _ZN16nonphysical_core6random3pcg29PermutedCongruentialGenerator13shuffle_usize17h4098009cdef9b8bcE_param_0,
.param .b64 _ZN16nonphysical_core6random3pcg29PermutedCongruentialGenerator13shuffle_usize17h4098009cdef9b8bcE_param_1,
.param .b64 _ZN16nonphysical_core6random3pcg29PermutedCongruentialGenerator13shuffle_usize17h4098009cdef9b8bcE_param_2
)
{
.reg .pred %p<6>;
.reg .b32 %r<16>;
.reg .b64 %rd<25>;
ld.param.u64 %rd11, [_ZN16nonphysical_core6random3pcg29PermutedCongruentialGenerator13shuffle_usize17h4098009cdef9b8bcE_param_2];
setp.lt.u64 %p1, %rd11, 2;
@%p1 bra $L__BB0_8;
ld.param.u64 %rd10, [_ZN16nonphysical_core6random3pcg29PermutedCongruentialGenerator13shuffle_usize17h4098009cdef9b8bcE_param_1];
ld.param.u64 %rd9, [_ZN16nonphysical_core6random3pcg29PermutedCongruentialGenerator13shuffle_usize17h4098009cdef9b8bcE_param_0];
ld.u32 %r4, [%rd9];
ld.u32 %r2, [%rd9+4];
shl.b64 %rd12, %rd11, 3;
add.s64 %rd13, %rd12, %rd10;
add.s64 %rd22, %rd13, -8;
mov.u64 %rd23, %rd11;
$L__BB0_2:
add.s64 %rd23, %rd23, -1;
shr.u32 %r5, %r4, 28;
add.s32 %r6, %r5, 4;
shr.u32 %r7, %r4, %r6;
xor.b32 %r8, %r7, %r4;
mul.lo.s32 %r9, %r8, 277803737;
shr.u32 %r10, %r9, 22;
xor.b32 %r11, %r10, %r9;
cvt.u64.u32 %rd7, %r11;
setp.lt.u64 %p2, %rd7, %rd23;
@%p2 bra $L__BB0_4;
cvt.u32.u64 %r12, %rd23;
cvt.u32.u64 %r13, %rd7;
rem.u32 %r14, %r13, %r12;
cvt.u64.u32 %rd7, %r14;
$L__BB0_4:
mad.lo.s32 %r4, %r4, 747796405, %r2;
setp.ge.u64 %p3, %rd23, %rd11;
@%p3 bra $L__BB0_10;
setp.lt.u64 %p4, %rd7, %rd11;
@%p4 bra $L__BB0_6;
bra.uni $L__BB0_9;
$L__BB0_6:
shl.b64 %rd18, %rd7, 3;
add.s64 %rd19, %rd10, %rd18;
ld.u64 %rd20, [%rd22];
ld.u64 %rd21, [%rd19];
st.u64 [%rd22], %rd21;
st.u64 [%rd19], %rd20;
add.s64 %rd22, %rd22, -8;
setp.gt.u64 %p5, %rd23, 1;
@%p5 bra $L__BB0_2;
st.u32 [%rd9], %r4;
$L__BB0_8:
ret;
$L__BB0_9:
st.u32 [%rd9], %r4;
mov.u64 %rd16, anon_$_2cbddb30228dc0cec36ed13501cf80e7_$_1;
cvta.global.u64 %rd17, %rd16;
{ // callseq 1, 0
.param .b64 param0;
st.param.b64 [param0+0], %rd7;
.param .b64 param1;
st.param.b64 [param1+0], %rd11;
.param .b64 param2;
st.param.b64 [param2+0], %rd17;
call.uni
_ZN4core9panicking18panic_bounds_check17h4e5c286e8d48ca34E,
(
param0,
param1,
param2
);
} // callseq 1
$L__BB0_10:
st.u32 [%rd9], %r4;
mov.u64 %rd14, anon_$_2cbddb30228dc0cec36ed13501cf80e7_$_1;
cvta.global.u64 %rd15, %rd14;
{ // callseq 0, 0
.param .b64 param0;
st.param.b64 [param0+0], %rd23;
.param .b64 param1;
st.param.b64 [param1+0], %rd11;
.param .b64 param2;
st.param.b64 [param2+0], %rd15;
call.uni
_ZN4core9panicking18panic_bounds_check17h4e5c286e8d48ca34E,
(
param0,
param1,
param2
);
} // callseq 0
}