forked from JuliaGPU/CUDA.jl
-
Notifications
You must be signed in to change notification settings - Fork 0
/
vadd.ptx
47 lines (40 loc) · 990 Bytes
/
vadd.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
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-19856038
// Cuda compilation tools, release 7.5, V7.5.17
// Based on LLVM 3.4svn
//
.version 4.3
.target sm_20
.address_size 64
// .globl kernel_vadd
.visible .entry kernel_vadd(
.param .u64 kernel_vadd_param_0,
.param .u64 kernel_vadd_param_1,
.param .u64 kernel_vadd_param_2
)
{
.reg .f32 %f<4>;
.reg .b32 %r<5>;
.reg .b64 %rd<11>;
ld.param.u64 %rd1, [kernel_vadd_param_0];
ld.param.u64 %rd2, [kernel_vadd_param_1];
ld.param.u64 %rd3, [kernel_vadd_param_2];
cvta.to.global.u64 %rd4, %rd3;
cvta.to.global.u64 %rd5, %rd2;
cvta.to.global.u64 %rd6, %rd1;
mov.u32 %r1, %ctaid.x;
mov.u32 %r2, %ntid.x;
mov.u32 %r3, %tid.x;
mad.lo.s32 %r4, %r2, %r1, %r3;
mul.wide.s32 %rd7, %r4, 4;
add.s64 %rd8, %rd6, %rd7;
ld.global.f32 %f1, [%rd8];
add.s64 %rd9, %rd5, %rd7;
ld.global.f32 %f2, [%rd9];
add.f32 %f3, %f1, %f2;
add.s64 %rd10, %rd4, %rd7;
st.global.f32 [%rd10], %f3;
ret;
}