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
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
|
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-27506705
// Cuda compilation tools, release 10.2, V10.2.89
// Based on LLVM 3.4svn
//
.version 6.5
.target sm_30
.address_size 64
// .globl _Z5tracev
.visible .global .align 1 .b8 var_buffer[1];
.visible .global .align 4 .b8 bvh[4];
.visible .global .align 8 .b8 launch_index[8];
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo3bvhE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo12launch_indexE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 8 .u64 _ZN21rti_internal_register20reg_bitness_detectorE;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail0E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail1E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail2E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail3E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail4E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail5E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail6E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail7E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail8E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail9E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail0E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail1E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail2E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail3E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail4E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail5E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail6E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail7E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail8E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail9E;
.visible .global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_xE;
.visible .global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_yE;
.visible .global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_zE;
.visible .global .align 1 .b8 _ZN21rti_internal_typename3bvhE[9] = {114, 116, 79, 98, 106, 101, 99, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename12launch_indexE[6] = {117, 105, 110, 116, 50, 0};
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum3bvhE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum12launch_indexE = 4919;
.visible .global .align 1 .b8 _ZN21rti_internal_semantic3bvhE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic12launch_indexE[14] = {114, 116, 76, 97, 117, 110, 99, 104, 73, 110, 100, 101, 120, 0};
.visible .global .align 1 .b8 _ZN23rti_internal_annotation3bvhE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation12launch_indexE[1];
.visible .func _Z5tracev(
)
{
.local .align 1 .b8 __local_depot0[1];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .b16 %rs<2>;
.reg .f32 %f<9>;
.reg .b32 %r<7>;
.reg .b64 %rd<3>;
mov.u64 %SPL, __local_depot0;
cvta.local.u64 %SP, %SPL;
add.u64 %rd1, %SP, 0;
add.u64 %rd2, %SPL, 0;
ld.global.u32 %r6, [launch_index];
cvt.rn.f32.u32 %f1, %r6;
mov.u16 %rs1, 0;
st.local.u8 [%rd2], %rs1;
ld.global.u32 %r1, [bvh];
mov.u32 %r3, 255;
mov.u32 %r4, 0;
mov.u32 %r5, 1;
mov.f32 %f3, 0fBF800000;
mov.f32 %f6, 0f3F800000;
mov.f32 %f7, 0f00000000;
mov.f32 %f8, 0f6C4ECB8F;
// inline asm
call _rt_trace_mask_flags_64, (%r1, %f1, %f7, %f3, %f7, %f7, %f6, %r4, %f7, %f8, %r3, %r4, %rd1, %r5);
// inline asm
ret;
}
// .globl _Z5startv
.visible .entry _Z5startv(
)
{
// Callseq Start 0
{
.reg .b32 temp_param_reg;
// <end>}
call.uni
_Z5tracev,
(
);
//{
}// Callseq End 0
ret;
}
// .globl _Z6throw_v
.visible .entry _Z6throw_v(
)
{
.reg .b32 %r<2>;
mov.u32 %r1, 1024;
// inline asm
call _rt_throw, (%r1);
// inline asm
ret;
}
// .globl _Z9exceptionv
.visible .entry _Z9exceptionv(
)
{
.reg .b32 %r<4>;
.reg .b64 %rd<8>;
// inline asm
call (%r1), _rt_get_exception_code, ();
// inline asm
mov.u64 %rd7, var_buffer;
cvta.global.u64 %rd2, %rd7;
mov.u32 %r2, 1;
mov.u32 %r3, 4;
mov.u64 %rd6, 0;
// inline asm
call (%rd1), _rt_buffer_get_64, (%rd2, %r2, %r3, %rd6, %rd6, %rd6, %rd6);
// inline asm
st.u32 [%rd1], %r1;
ret;
}
|
