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
|
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-32267302
// Cuda compilation tools, release 12.0, V12.0.140
// Based on NVVM 7.0.1
//
.version 8.0
.target sm_75
.address_size 64
// .globl mail_kernel
.global .align 1 .b8 kMail[40] = {116, 104, 101, 32, 108, 105, 116, 116, 108, 101, 32, 103, 112, 117, 32, 109, 97, 110, 32, 115, 97, 121, 115, 32, 121, 111, 117, 39, 118, 101, 32, 103, 111, 116, 32, 109, 97, 105, 108, 0};
.visible .entry mail_kernel(
.param .u64 mail_kernel_param_0,
.param .u32 mail_kernel_param_1
)
{
.reg .pred %p<2>;
.reg .b16 %rs<3>;
.reg .b32 %r<6>;
.reg .b64 %rd<7>;
ld.param.u64 %rd1, [mail_kernel_param_0];
ld.param.u32 %r2, [mail_kernel_param_1];
mov.u32 %r3, %ctaid.x;
mov.u32 %r4, %ntid.x;
mov.u32 %r5, %tid.x;
mad.lo.s32 %r1, %r3, %r4, %r5;
setp.ge.s32 %p1, %r1, %r2;
@%p1 bra $L__BB0_2;
cvta.to.global.u64 %rd2, %rd1;
cvt.s64.s32 %rd3, %r1;
mov.u64 %rd4, kMail;
add.s64 %rd5, %rd4, %rd3;
ld.global.nc.u8 %rs1, [%rd5];
cvt.u16.u8 %rs2, %rs1;
add.s64 %rd6, %rd2, %rd3;
st.global.u8 [%rd6], %rs2;
$L__BB0_2:
ret;
}
// .globl vadd
.visible .entry vadd(
.param .u64 vadd_param_0,
.param .u64 vadd_param_1,
.param .u64 vadd_param_2,
.param .u32 vadd_param_3
)
{
.reg .pred %p<2>;
.reg .f32 %f<4>;
.reg .b32 %r<6>;
.reg .b64 %rd<11>;
ld.param.u64 %rd1, [vadd_param_0];
ld.param.u64 %rd2, [vadd_param_1];
ld.param.u64 %rd3, [vadd_param_2];
ld.param.u32 %r2, [vadd_param_3];
mov.u32 %r3, %ctaid.x;
mov.u32 %r4, %ntid.x;
mov.u32 %r5, %tid.x;
mad.lo.s32 %r1, %r3, %r4, %r5;
setp.ge.s32 %p1, %r1, %r2;
@%p1 bra $L__BB1_2;
cvta.to.global.u64 %rd4, %rd1;
mul.wide.s32 %rd5, %r1, 4;
add.s64 %rd6, %rd4, %rd5;
cvta.to.global.u64 %rd7, %rd2;
add.s64 %rd8, %rd7, %rd5;
ld.global.f32 %f1, [%rd8];
ld.global.f32 %f2, [%rd6];
add.f32 %f3, %f2, %f1;
cvta.to.global.u64 %rd9, %rd3;
add.s64 %rd10, %rd9, %rd5;
st.global.f32 [%rd10], %f3;
$L__BB1_2:
ret;
}
// .globl vadd_mutated
.visible .entry vadd_mutated(
.param .u64 vadd_mutated_param_0,
.param .u64 vadd_mutated_param_1,
.param .u64 vadd_mutated_param_2,
.param .u32 vadd_mutated_param_3
)
{
.reg .pred %p<2>;
.reg .f32 %f<5>;
.reg .b32 %r<6>;
.reg .b64 %rd<11>;
ld.param.u64 %rd1, [vadd_mutated_param_0];
ld.param.u64 %rd2, [vadd_mutated_param_1];
ld.param.u64 %rd3, [vadd_mutated_param_2];
ld.param.u32 %r2, [vadd_mutated_param_3];
mov.u32 %r3, %ctaid.x;
mov.u32 %r4, %ntid.x;
mov.u32 %r5, %tid.x;
mad.lo.s32 %r1, %r3, %r4, %r5;
setp.ge.s32 %p1, %r1, %r2;
@%p1 bra $L__BB2_2;
cvta.to.global.u64 %rd4, %rd1;
mul.wide.s32 %rd5, %r1, 4;
add.s64 %rd6, %rd4, %rd5;
cvta.to.global.u64 %rd7, %rd2;
add.s64 %rd8, %rd7, %rd5;
ld.global.f32 %f1, [%rd8];
ld.global.f32 %f2, [%rd6];
add.f32 %f3, %f2, %f1;
add.f32 %f4, %f3, 0f3F800000;
cvta.to.global.u64 %rd9, %rd3;
add.s64 %rd10, %rd9, %rd5;
st.global.f32 [%rd10], %f4;
$L__BB2_2:
ret;
}
|
