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
149
150
|
// macros and whatnot for translator
#ifdef CUDA
#if 0
Types
1. Array stored in device memory
1.a fully dimensional
Example 3D
name: qc
how allocated: argument
dimensionality: ims:ime,kms:kme,jms:jme
index as: P3(ti,k,tj)
Example 2D
name: qc
how allocated: argument
dimensionality: ims:ime,jms:jme
index as: P2(ti,k,tj)
1.b 1 dimensional (vertical only), local storage
Example
name: w3
how allocated: local
dimensionality: constant (MKX)
index as: w3[k]
2. Array stored in shared memory
2.a fully dimensional
3D
Example
name: qc_s
how allocated: as offsets into sm[SM_SIZE]
dimensionality: bx*by*kx
index as: S3(ti,k,tj)
2D
Example
name: <none>_s
how allocated: as offsets into sm[SM_SIZE]
dimensionality: bx*by
index as: S2(ti,tj)
#endif
#ifndef MKX
-- intentional syntax error -- need a defined constant MKX that is static number of levels --
#endif
#ifdef PROMOTE
# define float double
#endif
#define SM_SIZE (0x1000-0xd4)
#define MAX_THREADS_PER_BLOCK 512
#define bi blockIdx.x
#define bj blockIdx.y
#define bx blockDim.x
#define by blockDim.y
#define ti threadIdx.x
#define tj threadIdx.y
# define ix (ime-ims+1)
# define jx (jme-jms+1)
# define kx (kme-kms+1)
// basic indexing macros. indices are always given as global indices
// in undecompsed Domain(ids:ide,jds:jde)
//
// That is, given IJ (global index), the global Index mapped to
// a local index on a Patch(0:nx-1,0:ny-1) in Device Memory as:
//
// I - (ips-ims) + nx * ( J - (jps-jms) )
//
// where ips is the global index of the start of the patch (the -1 is
// for translating from WRF fortran indices).
//
// The global index I is mapped to a local index on a GPU Block's
// shared memory (0:bx-1, 0:by-1) as:
//
// I - (ips-ims) - bi * bx + by * ( J - (jps-jms) - bj * by )
//
// Where bi is the index into the GPU Block, and bx is the
// GPU Block Width.
// global to patch index converter
#define GtoP(i,p,P) ((i)-(p)+(P))
#define GtoB(i,n,N,p,P) ((i)-(p)+(P)-(n)*(N))
// thread index to local memory index = i + bi * bx + ips - ims
#define TtoP(i,a,b,c,d) ((i)+(a)*(b)+(c)-(d))
//#define MAX(x,y) ((x)>(y)?(x):(y))
//#define MIN(x,y) ((x)<(y)?(x):(y))
#define MAX(x,y) max(x,y)
#define MIN(x,y) min(x,y)
// basic indexing macros
//#define I2(i,j,m) ((i)+((j)*(m)))
//#define I3(i,j,m,k,n) (I2(i,j,m)+((k)*(m)*(n)))
#define I2(i,j,m) ((i)+(__mul24((j),(m))))
#define I3(i,j,m,k,n) (I2(i,j,m)+(__mul24((k),__mul24((m),(n)))))
#if (FLOAT_4 != 4)
// index into a patch stored on device memory - 1
# define P2(i,j) I2(TtoP(i,bi,bx,ips,ims),TtoP(j,bj,by,jps,jms),ime-ims+1)
# define P3(i,k,j) I3(TtoP(i,bi,bx,ips,ims),k,ime-ims+1,TtoP(j,bj,by,jps,jms),kme-kms+1)
// index into a block stored on shared memory
# define S2(i,j) I2(i,j,bx)
//# define S3(i,k,j) I3(i,k,bx,j,kme-kms+1)
//# define S3(i,k,j) I3(i,j,bx,k,by)
# define S3(i,k,j) I3(k,i,kx,j,bx)
// Local arrays in device mem
# define LOCDM(a,s) float * a ; cudaMalloc( (void**) & a , (s)*sizeof(float)) ;
// Local scratch arrays in shared memory
# define LOCSM(a,s) __shared__ float * a ; a = &(sm[isize]) ; isize += (s) ;
#define ig (TtoP(ti,bi,bx,ips,ims))
#define jg (TtoP(tj,bj,by,jps,jms))
#else
// index into a patch stored on device memory - 4
# define P2(i,j) (I2(TtoP((i)*4,bi,(bx)*4,ips,ims),TtoP((j),bj,by,jps,jms),ime-ims+1)/4)
# define P3(i,k,j) (I3(TtoP((i)*4,bi,(bx)*4,ips,ims),k,ime-ims+1,TtoP((j),bj,by,jps,jms),kme-kms+1)/4)
// index into a block stored on shared memory
# define S2(i,j) (I2(i*4,j,(bx)*4)/4)
# define S3(i,k,j) (I3(i*4,k,(bx)*4,j,kme-kms+1)/4)
// Local arrays in device mem
# define LOCDM(a,s) float4 * a ; cudaMalloc( (void**) & a , (s)*sizeof(float4)) ;
// Local scratch arrays in shared memory
# define LOCSM(a,s) __shared__ Float4 * a ; a = &(sm[isize]) ; isize += (s)*4 ;
#define ig (TtoP((ti)*4,bi,(bx)*4,ips,ims))
#define jg (TtoP(tj,bj,by,jps,jms))
#endif
#endif
|
