QUDA
v0.5.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
staggered_dslash_def.h
Go to the documentation of this file.
1
// staggered_dslash_def.h - staggered Dslash kernel definitions
2
//
3
// See comments in wilson_dslash_def.h
4
5
// initialize on first iteration
6
7
#ifndef DD_LOOP
8
#define DD_LOOP
9
10
#define DD_AXPY 0
11
#define DD_RECON 0
12
#define DD_PREC 0
13
#endif
14
15
// set options for current iteration
16
17
#define DD_FNAME staggeredDslash
18
19
#if (DD_AXPY==0) // no axpy
20
#define DD_AXPY_F
21
#else // axpy
22
#define DD_AXPY_F Axpy
23
#define DSLASH_AXPY
24
#endif
25
26
#if (DD_PREC == 0)
27
#define DD_PARAM_AXPY const double2 *x, const float *xNorm, const double a, const DslashParam param
28
#elif (DD_PREC == 1)
29
#define DD_PARAM_AXPY const float2 *x, const float *xNorm, const float a, const DslashParam param
30
#else
31
#define DD_PARAM_AXPY const short2 *x, const float *xNorm, const float a, const DslashParam param
32
#endif
33
34
35
#if (DD_RECON==0) // reconstruct from 8 reals
36
#define DD_RECON_F 8
37
38
#if (DD_PREC==0) // DOUBLE PRECISION
39
#define DD_PARAM_GAUGE const double2 *fatGauge0, const double2 *fatGauge1, const double2* longGauge0, const double2* longGauge1
40
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_GAUGE_MATRIX_8_DOUBLE
41
42
#ifdef DIRECT_ACCESS_FAT_LINK
43
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2(FAT, gauge, dir, idx, fat_ga_stride)
44
#else
45
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2_TEX(FAT, gauge, dir, idx, fat_ga_stride)
46
#endif // DIRECT_ACCESS_FAT_LINK
47
#ifdef DIRECT_ACCESS_LONG_LINK
48
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_8_DOUBLE2(LONG, gauge, dir, idx, long_ga_stride)
49
#else
50
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_8_DOUBLE2_TEX(LONG, gauge, dir, idx, long_ga_stride)
51
#endif // DIRECT_ACCESS_LONG_LINK
52
53
#elif (DD_PREC==1) // SINGLE PRECISION
54
#define DD_PARAM_GAUGE const float2 *fatGauge0, const float2 *fatGauge1, const float4* longGauge0, const float4* longGauge1
55
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_GAUGE_MATRIX_8_SINGLE
56
57
#ifdef DIRECT_ACCESS_FAT_LINK
58
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2(FAT, gauge, dir, idx, fat_ga_stride)
59
#else
60
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2_TEX(FAT, gauge, dir, idx, fat_ga_stride)
61
#endif // DIRECT_ACCESS_FAT_LINK
62
#ifdef DIRECT_ACCESS_LONG_LINK
63
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_8_FLOAT4(LONG, gauge, dir, idx, long_ga_stride)
64
#else
65
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_8_FLOAT4_TEX(LONG, gauge, dir, idx, long_ga_stride)
66
#endif // DIRECT_ACCESS_LONG_LINK
67
68
#else // HALF PRECISION
69
#define DD_PARAM_GAUGE const short2 *fatGauge0, const short2* fatGauge1, const short4* longGauge0, const short4* longGauge1
70
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_GAUGE_MATRIX_8_SINGLE
71
72
/*#ifdef DIRECT_ACCESS_FAT_LINK
73
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2(FAT, gauge, dir, idx, fat_ga_stride); RESCALE2(FAT, fat_ga_max);
74
#else*/
75
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2_TEX(FAT, gauge, dir, idx, fat_ga_stride); RESCALE2(FAT, fat_ga_max);
76
/*#endif // DIRECT_ACCESS_FAT_LINK
77
#ifdef DIRECT_ACCESS_LONG_LINK
78
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_8_SHORT4(LONG, gauge, dir, idx, long_ga_stride)
79
#else*/
80
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_8_SHORT4_TEX(LONG, gauge, dir, idx, long_ga_stride)
81
//#endif // DIRECT_ACCESS_LONG_LINK
82
83
#endif // DD_PREC
84
85
#elif (DD_RECON ==1)// reconstruct from 12 reals
86
87
#define DD_RECON_F 12
88
89
#if (DD_PREC==0) // DOUBLE PRECISION
90
#define DD_PARAM_GAUGE const double2 *fatGauge0, const double2 *fatGauge1, const double2* longGauge0, const double2* longGauge1
91
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_GAUGE_MATRIX_12_DOUBLE
92
93
#ifdef DIRECT_ACCESS_FAT_LINK
94
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2(FAT, gauge, dir, idx, fat_ga_stride)
95
#else
96
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2_TEX(FAT, gauge, dir, idx, fat_ga_stride)
97
#endif // DIRECT_ACCESS_FAT_LINK
98
#ifdef DIRECT_ACCESS_LONG_LINK
99
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_12_DOUBLE2(LONG, gauge, dir, idx, long_ga_stride)
100
#else
101
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_12_DOUBLE2_TEX(LONG, gauge, dir, idx, long_ga_stride)
102
#endif // DIRECT_ACCESS_LONG_LINK
103
104
#elif (DD_PREC==1) // SINGLE PRECISION
105
#define DD_PARAM_GAUGE const float2 *fatGauge0, const float2 *fatGauge1, const float4* longGauge0, const float4* longGauge1
106
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_GAUGE_MATRIX_12_SINGLE
107
108
#ifdef DIRECT_ACCESS_FAT_LINK
109
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2(FAT, gauge, dir, idx, fat_ga_stride)
110
#else
111
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2_TEX(FAT, gauge, dir, idx, fat_ga_stride)
112
#endif // DIRECT_ACCESS_FAT_LINK
113
#ifdef DIRECT_ACCESS_LONG_LINK
114
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_12_FLOAT4(LONG, gauge, dir, idx, long_ga_stride)
115
#else
116
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_12_FLOAT4_TEX(LONG, gauge, dir, idx, long_ga_stride)
117
#endif // DIRECT_ACCESS_LONG_LINK
118
119
#else // HALF PRECISION
120
#define DD_PARAM_GAUGE const short2 *fatGauge0, const short2 *fatGauge1, const short4* longGauge0, const short4* longGauge1
121
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_GAUGE_MATRIX_12_SINGLE
122
123
/*#ifdef DIRECT_ACCESS_FAT_LINK
124
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2(FAT, gauge, dir, idx, fat_ga_stride); RESCALE2(FAT, fat_ga_max);
125
#else*/
126
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2_TEX(FAT, gauge, dir, idx, fat_ga_stride); RESCALE2(FAT, fat_ga_max);
127
/*#endif // DIRECT_ACCCESS_FAT_LINK
128
#ifdef DIRECT_ACCESS_LONG_LINK
129
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_12_SHORT4(LONG, gauge, dir, idx, long_ga_stride)
130
#else*/
131
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_12_SHORT4_TEX(LONG, gauge, dir, idx, long_ga_stride)
132
//#endif // DIRECT_ACCCESS_LONG_LINK
133
134
#endif // DD_PREC
135
136
#else //18 reconstruct
137
#define DD_RECON_F 18
138
#define RECONSTRUCT_GAUGE_MATRIX(dir, gauge, idx, sign)
139
140
#if (DD_PREC==0) // DOUBLE PRECISION
141
#define DD_PARAM_GAUGE const double2 *fatGauge0, const double2 *fatGauge1, const double2* longGauge0, const double2* longGauge1
142
143
#ifdef DIRECT_ACCESS_FAT_LINK
144
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2(FAT, gauge, dir, idx, fat_ga_stride)
145
#else
146
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2_TEX(FAT, gauge, dir, idx, fat_ga_stride)
147
#endif // DIRECT_ACCCESS_FAT_LINK
148
#ifdef DIRECT_ACCESS_LONG_LINK
149
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2(LONG, gauge, dir, idx, long_ga_stride)
150
#else
151
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2_TEX(LONG, gauge, dir, idx, long_ga_stride)
152
#endif // DIRECT_ACCCESS_LONG_LINK
153
154
#elif (DD_PREC==1) // SINGLE PRECISION
155
156
#define DD_PARAM_GAUGE const float2 *fatGauge0, const float2 *fatGauge1, const float4* longGauge0, const float4* longGauge1
157
158
#ifdef DIRECT_ACCESS_FAT_LINK
159
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2(FAT, gauge, dir, idx, fat_ga_stride)
160
#else
161
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2_TEX(FAT, gauge, dir, idx, fat_ga_stride)
162
#endif // DIRECT_ACCCESS_FAT_LINK
163
#ifdef DIRECT_ACCESS_LONG_LINK
164
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2(LONG, gauge, dir, idx, long_ga_stride)
165
#else
166
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2_TEX(LONG, gauge, dir, idx, long_ga_stride)
167
#endif // DIRECT_ACCCESS_LONG_LINK
168
169
#else // HALF PRECISION
170
171
#define DD_PARAM_GAUGE const short2 *fatGauge0, const short2 *fatGauge1, const short4* longGauge0, const short4* longGauge1
172
173
/*#ifdef DIRECT_ACCESS_FAT_LINK
174
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2(FAT, gauge, dir, idx, fat_ga_stride); RESCALE2(FAT, fat_ga_max);
175
#else*/
176
#define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2_TEX(FAT, gauge, dir, idx, fat_ga_stride); RESCALE2(FAT, fat_ga_max);
177
/*#endif // DIRECT_ACCESS_FAT_LINK
178
#ifdef DIRECT_ACCESS_LONG_LINK
179
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2(LONG, gauge, dir, idx, long_ga_stride)
180
#else*/
181
#define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2_TEX(LONG, gauge, dir, idx, long_ga_stride)
182
//#endif // DIRECT_ACCCESS_LONG_LINK
183
184
#endif // DD_PREC
185
186
#endif // DD_RECON
187
188
#if (DD_PREC==0) // double-precision fields
189
190
// gauge field
191
#define DD_PREC_F D
192
#if (defined DIRECT_ACCESS_FAT_LINK) || (defined FERMI_NO_DBLE_TEX)
193
#define FATLINK0TEX fatGauge0
194
#define FATLINK1TEX fatGauge1
195
#else
196
#ifdef USE_TEXTURE_OBJECTS
197
#define FATLINK0TEX param.gauge0Tex
198
#define FATLINK1TEX param.gauge1Tex
199
#else
200
#define FATLINK0TEX fatGauge0TexDouble
201
#define FATLINK1TEX fatGauge1TexDouble
202
#endif // USE_TEXTURE_OBJECTS
203
#endif
204
205
#if (defined DIRECT_ACCESS_LONG_LINK) || (defined FERMI_NO_DBLE_TEX)
206
#define LONGLINK0TEX longGauge0
207
#define LONGLINK1TEX longGauge1
208
#else
209
#ifdef USE_TEXTURE_OBJECTS
210
#define LONGLINK0TEX param.longGauge0Tex
211
#define LONGLINK1TEX param.longGauge1Tex
212
#else
213
#define LONGLINK0TEX longGauge0TexDouble
214
#define LONGLINK1TEX longGauge1TexDouble
215
#endif // USE_TEXTURE_OBJECTS
216
#endif
217
218
#define GAUGE_DOUBLE
219
220
// spinor fields
221
#define DD_PARAM_OUT double2* out, float *null1
222
#define DD_PARAM_IN const double2* in, const float *null4
223
#if (defined DIRECT_ACCESS_SPINOR) || (defined FERMI_NO_DBLE_TEX)
224
#define SPINORTEX in
225
#define READ_1ST_NBR_SPINOR READ_1ST_NBR_SPINOR_DOUBLE
226
#define READ_3RD_NBR_SPINOR READ_3RD_NBR_SPINOR_DOUBLE
227
#else
228
#ifdef USE_TEXTURE_OBJECTS
229
#define SPINORTEX param.inTex
230
#else
231
#define SPINORTEX spinorTexDouble
232
#endif // USE_TEXTURE_OBJECTS
233
#define READ_1ST_NBR_SPINOR READ_1ST_NBR_SPINOR_DOUBLE_TEX
234
#define READ_3RD_NBR_SPINOR READ_3RD_NBR_SPINOR_DOUBLE_TEX
235
#endif
236
#if (defined DIRECT_ACCESS_INTER) || (defined FERMI_NO_DBLE_TEX)
237
#define READ_AND_SUM_SPINOR READ_AND_SUM_ST_SPINOR
238
#define INTERTEX out
239
#else
240
#define READ_AND_SUM_SPINOR READ_AND_SUM_ST_SPINOR_DOUBLE_TEX
241
#ifdef USE_TEXTURE_OBJECTS
242
#define INTERTEX param.outTex
243
#else
244
#define INTERTEX interTexDouble
245
#endif
246
#endif
247
#define WRITE_SPINOR WRITE_ST_SPINOR_DOUBLE2
248
#define SPINOR_DOUBLE
249
#if (DD_AXPY==1)
250
#if (defined DIRECT_ACCESS_ACCUM) || (defined FERMI_NO_DBLE_TEX)
251
#define ACCUMTEX x
252
#define READ_ACCUM READ_ST_ACCUM_DOUBLE
253
#else
254
#ifdef USE_TEXTURE_OBJECTS
255
#define ACCUMTEX param.xTex
256
#else
257
#define ACCUMTEX accumTexDouble
258
#endif // USE_TEXTURE_OBJECTS
259
#define READ_ACCUM READ_ST_ACCUM_DOUBLE_TEX
260
#endif
261
#endif // DD_AXPY
262
263
264
#elif (DD_PREC==1) // single-precision fields
265
266
// gauge fields
267
#define DD_PREC_F S
268
269
#ifndef DIRECT_ACCESS_FAT_LINK
270
#ifdef USE_TEXTURE_OBJECTS
271
#define FATLINK0TEX param.gauge0Tex
272
#define FATLINK1TEX param.gauge1Tex
273
#else
274
#define FATLINK0TEX fatGauge0TexSingle
275
#define FATLINK1TEX fatGauge1TexSingle
276
#endif
277
#else
278
#define FATLINK0TEX fatGauge0
279
#define FATLINK1TEX fatGauge1
280
#endif
281
282
#ifndef DIRECT_ACCESS_LONG_LINK //longlink access
283
#ifdef USE_TEXTURE_OBJECTS
284
#define LONGLINK0TEX param.longGauge0Tex
285
#define LONGLINK1TEX param.longGauge1Tex
286
#else
287
#if (DD_RECON ==2)
288
#define LONGLINK0TEX longGauge0TexSingle_norecon
289
#define LONGLINK1TEX longGauge1TexSingle_norecon
290
#else
291
#define LONGLINK0TEX longGauge0TexSingle
292
#define LONGLINK1TEX longGauge1TexSingle
293
#endif
294
#endif // USE_TEXTURE_OBJECTS
295
#else
296
#define LONGLINK0TEX longGauge0
297
#define LONGLINK1TEX longGauge1
298
#endif
299
300
// spinor fields
301
#define DD_PARAM_OUT float2* out, float *null1
302
#define DD_PARAM_IN const float2* in, const float *null4
303
#ifndef DIRECT_ACCESS_SPINOR
304
#ifdef USE_TEXTURE_OBJECTS
305
#define SPINORTEX param.inTex
306
#else
307
#define SPINORTEX spinorTexSingle2
308
#endif // USE_TEXTURE_OBJECTS
309
#define READ_1ST_NBR_SPINOR READ_1ST_NBR_SPINOR_SINGLE_TEX
310
#define READ_3RD_NBR_SPINOR READ_3RD_NBR_SPINOR_SINGLE_TEX
311
#else
312
#define SPINORTEX in
313
#define READ_1ST_NBR_SPINOR READ_1ST_NBR_SPINOR_SINGLE
314
#define READ_3RD_NBR_SPINOR READ_3RD_NBR_SPINOR_SINGLE
315
#endif
316
#if (defined DIRECT_ACCESS_INTER)
317
#define READ_AND_SUM_SPINOR READ_AND_SUM_ST_SPINOR
318
#define INTERTEX out
319
#else
320
#define READ_AND_SUM_SPINOR READ_AND_SUM_ST_SPINOR_SINGLE_TEX
321
#ifdef USE_TEXTURE_OBJECTS
322
#define INTERTEX param.outTex
323
#else
324
#define INTERTEX interTexSingle2
325
#endif // USE_TEXTURE_OBJECTS
326
#endif
327
#define WRITE_SPINOR WRITE_ST_SPINOR_FLOAT2
328
#if (DD_AXPY==1)
329
#if (defined DIRECT_ACCESS_ACCUM)
330
#define ACCUMTEX x
331
#define READ_ACCUM READ_ST_ACCUM_SINGLE
332
#else
333
#ifdef USE_TEXTURE_OBJECTS
334
#define ACCUMTEX param.xTex
335
#else
336
#define ACCUMTEX accumTexSingle2
337
#endif // USE_TEXTURE_OBJECTS
338
#define READ_ACCUM READ_ST_ACCUM_SINGLE_TEX
339
#endif
340
#endif // DD_AXPY
341
342
343
#else // half-precision fields
344
345
// all reads done through texture cache regardless
346
347
// gauge fields
348
#define DD_PREC_F H
349
#ifdef USE_TEXTURE_OBJECTS
350
#define FATLINK0TEX param.gauge0Tex
351
#define FATLINK1TEX param.gauge1Tex
352
#define LONGLINK0TEX param.longGauge0Tex
353
#define LONGLINK1TEX param.longGauge1Tex
354
#else
355
#define FATLINK0TEX fatGauge0TexHalf
356
#define FATLINK1TEX fatGauge1TexHalf
357
#if (DD_RECON ==2)
358
#define LONGLINK0TEX longGauge0TexHalf_norecon
359
#define LONGLINK1TEX longGauge1TexHalf_norecon
360
#else
361
#define LONGLINK0TEX longGauge0TexHalf
362
#define LONGLINK1TEX longGauge1TexHalf
363
#endif
364
#endif // USE_TEXTURE_OBJECTS
365
366
#define READ_1ST_NBR_SPINOR READ_1ST_NBR_SPINOR_HALF_TEX
367
#define READ_3RD_NBR_SPINOR READ_3RD_NBR_SPINOR_HALF_TEX
368
#ifdef USE_TEXTURE_OBJECTS
369
#define SPINORTEX param.inTex
370
#else
371
#define SPINORTEX spinorTexHalf2
372
#endif // USE_TEXTURE_OBJECTS
373
#define DD_PARAM_OUT short2* out, float *outNorm
374
#define DD_PARAM_IN const short2* in, const float *inNorm
375
#if (defined DIRECT_ACCESS_INTER)
376
#define READ_AND_SUM_SPINOR READ_AND_SUM_ST_SPINOR_HALF
377
#define INTERTEX out
378
#else
379
#define READ_AND_SUM_SPINOR READ_AND_SUM_ST_SPINOR_HALF_TEX
380
#ifdef USE_TEXTURE_OBJECTS
381
#define INTERTEX param.outTex
382
#else
383
#define INTERTEX interTexHalf2
384
#endif // USE_TEXTURE_OBJECTS
385
#endif
386
#define WRITE_SPINOR WRITE_ST_SPINOR_SHORT2
387
#if (DD_AXPY==1)
388
#ifdef USE_TEXTURE_OBJECTS
389
#define ACCUMTEX param.xTex
390
#else
391
#define ACCUMTEX accumTexHalf2
392
#endif // USE_TEXTURE_OBJECTS
393
#define READ_ACCUM READ_ST_ACCUM_HALF_TEX
394
#endif // DD_AXPY
395
396
#endif
397
398
// only build double precision if supported
399
#if !(__COMPUTE_CAPABILITY__ < 130 && DD_PREC == 0)
400
401
#define DD_CONCAT(n,r,x) n ## r ## x ## Kernel
402
#define DD_FUNC(n,r,x) DD_CONCAT(n,r,x)
403
404
// define the kernel
405
406
template
<KernelType kernel_type>
407
__global__
void
DD_FUNC
(
DD_FNAME
,
DD_RECON_F
,
DD_AXPY_F
)
408
(
DD_PARAM_OUT
,
DD_PARAM_GAUGE
,
DD_PARAM_IN
,
DD_PARAM_AXPY
) {
409
#ifdef GPU_STAGGERED_DIRAC
410
#include "
staggered_dslash_core.h
"
411
#endif
412
}
413
414
#endif // !(__COMPUTE_CAPABILITY__ < 130 && DD_PREC == 0)
415
416
417
// clean up
418
419
#undef DD_PREC_F
420
#undef DD_RECON_F
421
#undef DD_AXPY_F
422
#undef DD_PARAM_OUT
423
#undef DD_PARAM_GAUGE
424
#undef DD_PARAM_IN
425
#undef DD_PARAM_AXPY
426
#undef DD_FNAME
427
#undef DD_CONCAT
428
#undef DD_FUNC
429
430
#undef DSLASH_AXPY
431
#undef READ_GAUGE_MATRIX
432
#undef RECONSTRUCT_GAUGE_MATRIX
433
#undef FATLINK0TEX
434
#undef FATLINK1TEX
435
#undef LONGLINK0TEX
436
#undef LONGLINK1TEX
437
#undef SPINORTEX
438
#undef WRITE_SPINOR
439
#undef READ_AND_SUM_SPINOR
440
#undef INTERTEX
441
#undef ACCUMTEX
442
#undef READ_ACCUM
443
#undef CLOVERTEX
444
#undef READ_CLOVER
445
#undef DSLASH_CLOVER
446
#undef GAUGE_DOUBLE
447
#undef SPINOR_DOUBLE
448
#undef CLOVER_DOUBLE
449
#undef READ_FAT_MATRIX
450
#undef READ_LONG_MATRIX
451
#undef READ_1ST_NBR_SPINOR
452
#undef READ_3RD_NBR_SPINOR
453
454
455
// prepare next set of options, or clean up after final iteration
456
457
#if (DD_AXPY==0)
458
#undef DD_AXPY
459
#define DD_AXPY 1
460
#else
461
#undef DD_AXPY
462
#define DD_AXPY 0
463
464
#if (DD_RECON==0)
465
#undef DD_RECON
466
#define DD_RECON 1
467
#elif (DD_RECON ==1)
468
#undef DD_RECON
469
#define DD_RECON 2
470
#else
471
#undef DD_RECON
472
#define DD_RECON 0
473
474
#if (DD_PREC==0)
475
#undef DD_PREC
476
#define DD_PREC 1
477
#elif (DD_PREC==1)
478
#undef DD_PREC
479
#define DD_PREC 2
480
#else
481
#undef DD_PREC
482
#define DD_PREC 0
483
484
#undef DD_LOOP
485
#undef DD_AXPY
486
#undef DD_RECON
487
#undef DD_PREC
488
489
#endif // DD_PREC
490
#endif // DD_RECON
491
#endif // DD_AXPY
492
493
#ifdef DD_LOOP
494
#include "
staggered_dslash_def.h
"
495
#endif
Generated on Wed Mar 20 2013 12:52:17 for QUDA by
1.8.2
1.8.2