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
|
/*
* Copyright 1993-2015 NVIDIA Corporation. All rights reserved.
*
* Please refer to the NVIDIA end user license agreement (EULA) associated
* with this source code for terms and conditions that govern your use of
* this software. Any use, reproduction, disclosure, or distribution of
* this software and related documentation outside the terms of the EULA
* is strictly prohibited.
*
*/
#ifndef _PTXJIT_H_
#define _PTXJIT_H_
/*
* PTX is equivalent to the following kernel:
*
* __global__ void myKernel(int *data)
* {
* int tid = blockIdx.x * blockDim.x + threadIdx.x;
* data[tid] = tid;
* }
*
*/
char myPtx64[] = "\n\
.version 3.2\n\
.target sm_20\n\
.address_size 64\n\
.visible .entry _Z8myKernelPi(\n\
.param .u64 _Z8myKernelPi_param_0\n\
)\n\
{\n\
.reg .s32 %r<5>;\n\
.reg .s64 %rd<5>;\n\
ld.param.u64 %rd1, [_Z8myKernelPi_param_0];\n\
cvta.to.global.u64 %rd2, %rd1;\n\
.loc 1 3 1\n\
mov.u32 %r1, %ntid.x;\n\
mov.u32 %r2, %ctaid.x;\n\
mov.u32 %r3, %tid.x;\n\
mad.lo.s32 %r4, %r1, %r2, %r3;\n\
mul.wide.s32 %rd3, %r4, 4;\n\
add.s64 %rd4, %rd2, %rd3;\n\
.loc 1 4 1\n\
st.global.u32 [%rd4], %r4;\n\
.loc 1 5 2\n\
ret;\n\
}\n\
";
char myPtx32[] = "\n\
.version 3.2\n\
.target sm_20\n\
.address_size 32\n\
.visible .entry _Z8myKernelPi(\n\
.param .u32 _Z8myKernelPi_param_0\n\
)\n\
{\n\
.reg .s32 %r<9>;\n\
ld.param.u32 %r1, [_Z8myKernelPi_param_0];\n\
cvta.to.global.u32 %r2, %r1;\n\
.loc 1 3 1\n\
mov.u32 %r3, %ntid.x;\n\
mov.u32 %r4, %ctaid.x;\n\
mov.u32 %r5, %tid.x;\n\
mad.lo.s32 %r6, %r3, %r4, %r5;\n\
.loc 1 4 1\n\
shl.b32 %r7, %r6, 2;\n\
add.s32 %r8, %r2, %r7;\n\
st.global.u32 [%r8], %r6;\n\
.loc 1 5 2\n\
ret;\n\
}\n\
";
#endif
|
