QUDA
v0.7.0
A library for QCD on GPUs
Main Page
Namespaces
Classes
Files
File List
File Members
All
Classes
Namespaces
Files
Functions
Variables
Typedefs
Enumerations
Enumerator
Friends
Macros
Pages
quda
lib
covDev.h
Go to the documentation of this file.
1
// DD_NAME_F = covDevM#
2
// DD_RECON_F = 8, 12, 18
3
// DD_DAG_F = Dagger, [blank]
4
// covDevM012() (no dagger)
5
//
6
// In addition, the kernels are templated on the precision of the
7
// fields (double, single, or half).
8
9
// initialize on first iteration
10
11
#ifndef DD_LOOP
12
#define DD_LOOP
13
#define DD_DAG 0
14
#define DD_RECON 0
15
#define DD_PREC 0
16
#endif
17
18
// set options for current iteration
19
20
#if (DD_PREC == 0)
21
#define DD_PARAM2 const double2 *gauge0, const double2 *gauge1
22
#else
23
#define DD_PARAM2 const float4 *gauge0, const float4 *gauge1
24
#endif
25
26
#if (DD_DAG==0) // no dagger
27
#define DD_DAG_F
28
#else // dagger
29
#define DD_DAG_F Dagger
30
#endif
31
32
#define DD_PARAM4 const DslashParam param
33
34
35
#if (DD_RECON==0) // reconstruct from 8 reals
36
#define DD_RECON_F 8
37
38
#if (DD_PREC == 0)
39
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_8_DOUBLE
40
41
#ifdef DIRECT_ACCESS_LINK
42
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_DOUBLE2
43
#else
44
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_DOUBLE2_TEX
45
#endif // DIRECT_ACCESS_LINK
46
#elif (DD_PREC == 1)
47
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_8_SINGLE
48
49
#ifdef DIRECT_ACCESS_LINK
50
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_FLOAT4
51
#else
52
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_FLOAT4_TEX
53
#endif
54
#endif
55
#elif (DD_RECON==1) // reconstruct from 12 reals
56
#define DD_RECON_F 12
57
58
#if (DD_PREC == 0)
59
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_12_DOUBLE
60
61
#ifdef DIRECT_ACCESS_LINK
62
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_DOUBLE2
63
#else
64
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_DOUBLE2_TEX
65
#endif // DIRECT_ACCESS_LINK
66
#elif (DD_PREC == 1)
67
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_12_SINGLE
68
69
#ifdef DIRECT_ACCESS_LINK
70
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_FLOAT4
71
#else
72
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_FLOAT4_TEX
73
#endif
74
#endif
75
#else // no reconstruct, load all components
76
#define DD_RECON_F 18
77
#define GAUGE_FLOAT2
78
79
#if (DD_PREC == 0)
80
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_18_DOUBLE
81
82
#ifdef DIRECT_ACCESS_LINK
83
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_DOUBLE2
84
#else
85
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_DOUBLE2_TEX
86
#endif // DIRECT_ACCESS_LINK
87
#elif (DD_PREC == 1)
88
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_18_SINGLE
89
90
#ifdef DIRECT_ACCESS_LINK
91
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_FLOAT2
92
#else
93
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_FLOAT2_TEX
94
#endif
95
#endif
96
#endif
97
98
99
#if (DD_PREC==0)
100
101
#define TPROJSCALE tProjScale
102
103
// double-precision gauge field
104
#if (defined DIRECT_ACCESS_LINK) || (defined FERMI_NO_DBLE_TEX)
105
#define GAUGE0TEX gauge0
106
#define GAUGE1TEX gauge1
107
#else
108
#ifdef USE_TEXTURE_OBJECTS
109
#define GAUGE0TEX param.gauge0Tex
110
#define GAUGE1TEX param.gauge1Tex
111
#else
112
#define GAUGE0TEX gauge0TexDouble2
113
#define GAUGE1TEX gauge1TexDouble2
114
#endif
115
#endif
116
117
#define GAUGE_FLOAT2
118
119
// double-precision spinor fields
120
#define DD_PARAM1 double2* out
121
#define DD_PARAM3 const double2* in
122
123
#if (defined DIRECT_ACCESS_WILSON_SPINOR) || (defined FERMI_NO_DBLE_TEX)
124
#define READ_SPINOR READ_SPINOR_DOUBLE
125
#define READ_SPINOR_UP READ_SPINOR_DOUBLE_UP
126
#define READ_SPINOR_DOWN READ_SPINOR_DOUBLE_DOWN
127
#define SPINORTEX in
128
#else
129
#define READ_SPINOR READ_SPINOR_DOUBLE_TEX
130
#define READ_SPINOR_UP READ_SPINOR_DOUBLE_UP_TEX
131
#define READ_SPINOR_DOWN READ_SPINOR_DOUBLE_DOWN_TEX
132
133
#ifdef USE_TEXTURE_OBJECTS
134
#define SPINORTEX param.inTex
135
#else
136
#define SPINORTEX spinorTexDouble
137
#endif // USE_TEXTURE_OBJECTS
138
#endif
139
140
#if (defined DIRECT_ACCESS_WILSON_INTER) || (defined FERMI_NO_DBLE_TEX)
141
#define READ_INTERMEDIATE_SPINOR READ_SPINOR_DOUBLE
142
#define INTERTEX out
143
#else
144
#define READ_INTERMEDIATE_SPINOR READ_SPINOR_DOUBLE_TEX
145
146
#ifdef USE_TEXTURE_OBJECTS
147
#define INTERTEX param.outTex
148
#else
149
#define INTERTEX interTexDouble
150
#endif
151
#endif
152
153
154
#define WRITE_SPINOR WRITE_SPINOR_DOUBLE2_STR
155
//#define WRITE_SPINOR WRITE_SPINOR_DOUBLE2
156
#define SPINOR_DOUBLE
157
158
#define SPINOR_HOP 12
159
160
#else
161
#define TPROJSCALE tProjScale_f
162
163
// single-precision gauge field
164
#ifdef DIRECT_ACCESS_LINK
165
#define GAUGE0TEX gauge0
166
#define GAUGE1TEX gauge1
167
#else
168
#ifdef USE_TEXTURE_OBJECTS
169
#define GAUGE0TEX param.gauge0Tex
170
#define GAUGE1TEX param.gauge1Tex
171
#else
172
#if (DD_RECON_F == 18)
173
#define GAUGE0TEX gauge0TexSingle2
174
#define GAUGE1TEX gauge1TexSingle2
175
#else
176
#define GAUGE0TEX gauge0TexSingle4
177
#define GAUGE1TEX gauge1TexSingle4
178
#endif
179
#endif // USE_TEXTURE_OBJECTS
180
#endif
181
182
183
// single-precision spinor fields
184
#define DD_PARAM1 float4* out
185
#define DD_PARAM3 const float4* in
186
187
#ifdef DIRECT_ACCESS_WILSON_SPINOR
188
#define READ_SPINOR READ_SPINOR_SINGLE
189
#define READ_SPINOR_UP READ_SPINOR_SINGLE_UP
190
#define READ_SPINOR_DOWN READ_SPINOR_SINGLE_DOWN
191
#define SPINORTEX in
192
#else
193
#define READ_SPINOR READ_SPINOR_SINGLE_TEX
194
#define READ_SPINOR_UP READ_SPINOR_SINGLE_UP_TEX
195
#define READ_SPINOR_DOWN READ_SPINOR_SINGLE_DOWN_TEX
196
197
#ifdef USE_TEXTURE_OBJECTS
198
#define SPINORTEX param.inTex
199
#else
200
#define SPINORTEX spinorTexSingle
201
#endif // USE_TEXTURE_OBJECTS
202
#endif
203
204
#ifdef DIRECT_ACCESS_WILSON_INTER
205
#define READ_INTERMEDIATE_SPINOR READ_SPINOR_SINGLE
206
#define INTERTEX out
207
#else
208
#define READ_INTERMEDIATE_SPINOR READ_SPINOR_SINGLE_TEX
209
210
#ifdef USE_TEXTURE_OBJECTS
211
#define INTERTEX param.outTex
212
#else
213
#define INTERTEX interTexSingle
214
#endif // USE_TEXTURE_OBJECTS
215
#endif
216
217
#define WRITE_SPINOR WRITE_SPINOR_FLOAT4_STR
218
219
#define SPINOR_HOP 6
220
#endif
221
222
// only build double precision if supported
223
#if !(__COMPUTE_CAPABILITY__ < 130 && DD_PREC == 0)
224
225
#define DD_CONCAT(n,r,d) n ## r ## d ## Kernel
226
#define DD_FUNC(n,r,d) DD_CONCAT(n,r,d)
227
228
// define the kernels
229
#define DD_NAME_F covDevM0
230
template
<KernelType kernel_type>
231
__global__
void
DD_FUNC
(
DD_NAME_F
,
DD_RECON_F
,
DD_DAG_F
) (
DD_PARAM1
,
DD_PARAM2
,
DD_PARAM3
,
DD_PARAM4
)
232
{
233
#if DD_DAG
234
#include "
covDev_mu0_dagger_core.h
"
235
#else
236
#include "
covDev_mu0_core.h
"
237
#endif
238
}
239
240
#undef DD_NAME_F
241
#define DD_NAME_F covDevM1
242
template
<KernelType kernel_type>
243
__global__
void
DD_FUNC
(
DD_NAME_F
,
DD_RECON_F
,
DD_DAG_F
) (
DD_PARAM1
,
DD_PARAM2
,
DD_PARAM3
,
DD_PARAM4
)
244
{
245
#if DD_DAG
246
#include "
covDev_mu1_dagger_core.h
"
247
#else
248
#include "
covDev_mu1_core.h
"
249
#endif
250
}
251
252
#undef DD_NAME_F
253
#define DD_NAME_F covDevM2
254
template
<KernelType kernel_type>
255
__global__
void
DD_FUNC
(
DD_NAME_F
,
DD_RECON_F
,
DD_DAG_F
) (
DD_PARAM1
,
DD_PARAM2
,
DD_PARAM3
,
DD_PARAM4
)
256
{
257
#if DD_DAG
258
#include "
covDev_mu2_dagger_core.h
"
259
#else
260
#include "
covDev_mu2_core.h
"
261
#endif
262
}
263
264
#undef DD_NAME_F
265
#define DD_NAME_F covDevM3
266
template
<KernelType kernel_type>
267
__global__
void
DD_FUNC
(
DD_NAME_F
,
DD_RECON_F
,
DD_DAG_F
) (
DD_PARAM1
,
DD_PARAM2
,
DD_PARAM3
,
DD_PARAM4
)
268
{
269
#if DD_DAG
270
#include "
covDev_mu3_dagger_core.h
"
271
#else
272
#include "
covDev_mu3_core.h
"
273
#endif
274
}
275
276
#undef DD_NAME_F
277
#endif
278
279
// clean up
280
281
#undef DD_NAME_F
282
#undef DD_RECON_F
283
#undef DD_DAG_F
284
#undef DD_PARAM1
285
#undef DD_PARAM2
286
#undef DD_PARAM3
287
#undef DD_PARAM4
288
#undef DD_CONCAT
289
#undef DD_FUNC
290
291
#undef DSLASH_XPAY
292
#undef READ_GAUGE_MATRIX
293
#undef RECONSTRUCT_GAUGE_MATRIX
294
#undef GAUGE0TEX
295
#undef GAUGE1TEX
296
#undef READ_SPINOR
297
#undef READ_SPINOR_UP
298
#undef READ_SPINOR_DOWN
299
#undef SPINORTEX
300
#undef READ_INTERMEDIATE_SPINOR
301
#undef INTERTEX
302
#undef ACCUMTEX
303
#undef WRITE_SPINOR
304
#undef GAUGE_FLOAT2
305
#undef SPINOR_DOUBLE
306
307
#undef SPINOR_HOP
308
309
#undef TPROJSCALE
310
311
// prepare next set of options, or clean up after final iteration
312
313
#if (DD_PREC==0)
314
#undef DD_PREC
315
#define DD_PREC 1
316
#else
317
#undef DD_PREC
318
#define DD_PREC 0
319
320
#if (DD_DAG==0)
321
#undef DD_DAG
322
#define DD_DAG 1
323
#else
324
#undef DD_DAG
325
#define DD_DAG 0
326
327
#if (DD_RECON==0)
328
#undef DD_RECON
329
#define DD_RECON 1
330
#elif (DD_RECON==1)
331
#undef DD_RECON
332
#define DD_RECON 2
333
#else
334
#undef DD_LOOP
335
#undef DD_DAG
336
#undef DD_RECON
337
#undef DD_PREC
338
#endif // DD_RECON
339
#endif // DD_DAG
340
#endif // DD_PREC
341
342
#ifdef DD_LOOP
343
#include "
covDev.h
"
344
#endif
345
covDev_mu0_core.h
DD_PARAM4
#define DD_PARAM4
Definition:
covDev.h:32
DD_DAG_F
#define DD_DAG_F
Definition:
covDev.h:27
DD_FUNC
#define DD_FUNC(x)
Definition:
clover_def.h:141
covDev_mu3_core.h
covDev_mu1_dagger_core.h
covDev_mu1_core.h
covDev_mu0_dagger_core.h
covDev.h
DD_RECON_F
#define DD_RECON_F
Definition:
covDev.h:36
covDev_mu2_dagger_core.h
covDev_mu2_core.h
DD_PARAM2
#define DD_PARAM2
Definition:
covDev.h:21
covDev_mu3_dagger_core.h
DD_NAME_F
#define DD_NAME_F
Definition:
dw_dslash4_def.h:40
DD_PARAM1
#define DD_PARAM1
Definition:
covDev.h:120
DD_PARAM3
#define DD_PARAM3
Definition:
covDev.h:121
Generated on Wed Feb 4 2015 17:00:05 for QUDA by
1.8.6
1.8.6