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