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
wilson_dslash_def.h
Go to the documentation of this file.
1
// wilson_dslash_def.h - Dslash kernel definitions
2
3
// There are currently 72 different variants of the Wilson Dslash
4
// kernel, each one characterized by a set of 5 options, where each
5
// option can take one of several values (2*3*2*2*3 = 72). This file
6
// is structured so that the C preprocessor loops through all 72
7
// variants (in a manner resembling a counter), sets the appropriate
8
// macros, and defines the corresponding functions.
9
//
10
// As an example of the function naming conventions, consider
11
//
12
// cloverDslash12DaggerXpayKernel(float4* out, ...).
13
//
14
// This is a clover Dslash^dagger kernel where the result is
15
// multiplied by "a" and summed with an input vector (Xpay), and the
16
// gauge matrix is reconstructed from 12 real numbers. More
17
// generally, each function name is given by the concatenation of the
18
// following 4 fields, with "Kernel" at the end:
19
//
20
// DD_NAME_F = dslash, cloverDslash
21
// DD_RECON_F = 8, 12, 18
22
// DD_DAG_F = Dagger, [blank]
23
// DD_XPAY_F = Xpay, [blank]
24
//
25
// In addition, the kernels are templated on the precision of the
26
// fields (double, single, or half).
27
28
// initialize on first iteration
29
30
#ifndef DD_LOOP
31
#define DD_LOOP
32
#define DD_DAG 0
33
#define DD_XPAY 0
34
#define DD_RECON 0
35
#define DD_PREC 0
36
#define DD_CLOVER 0
37
#endif
38
39
// set options for current iteration
40
41
#if (DD_CLOVER==0) // no clover
42
#define DD_NAME_F dslash
43
#elif (DD_CLOVER==1) // clover
44
#define DSLASH_CLOVER
45
#define DD_NAME_F cloverDslash
46
#else
47
#define DSLASH_CLOVER
48
#define DSLASH_CLOVER_XPAY
49
#define DD_NAME_F asymCloverDslash
50
#endif
51
52
#if (DD_DAG==0) // no dagger
53
#define DD_DAG_F
54
#else // dagger
55
#define DD_DAG_F Dagger
56
#endif
57
58
// DSLASH_CLOVER_XPAY implies DD_XPAY=1
59
#if (DD_XPAY==0) && defined(DSLASH_CLOVER_XPAY)
60
#undef DD_XPAY
61
#define DD_XPAY 1
62
#endif
63
64
#if (DD_XPAY==0) // no xpay
65
#define DD_XPAY_F
66
#else // xpay
67
#define DD_XPAY_F Xpay
68
#define DSLASH_XPAY
69
#endif
70
71
#if (DD_PREC == 0)
72
#define DD_PARAM_XPAY const double2 *x, const float *xNorm, const double a,
73
#elif (DD_PREC == 1)
74
#define DD_PARAM_XPAY const float4 *x, const float *xNorm, const float a,
75
#else
76
#define DD_PARAM_XPAY const short4 *x, const float *xNorm, const float a,
77
#endif
78
79
#if (DD_RECON==0) // reconstruct from 8 reals
80
#define DD_RECON_F 8
81
82
#if (DD_PREC==0)
83
#define DD_PARAM_GAUGE const double2 *gauge0, const double2 *gauge1,
84
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_8_DOUBLE
85
#ifdef DIRECT_ACCESS_LINK
86
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_DOUBLE2
87
#else
88
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_DOUBLE2_TEX
89
#endif // DIRECT_ACCESS_LINK
90
91
#elif (DD_PREC==1)
92
#define DD_PARAM_GAUGE const float4 *gauge0, const float4 *gauge1,
93
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_8_SINGLE
94
#ifdef DIRECT_ACCESS_LINK
95
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_FLOAT4
96
#else
97
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_FLOAT4_TEX
98
#endif // DIRECT_ACCESS_LINK
99
100
#else
101
#define DD_PARAM_GAUGE const short4 *gauge0, const short4* gauge1,
102
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_8_SINGLE
103
#ifdef DIRECT_ACCESS_LINK
104
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_SHORT4
105
#else
106
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_SHORT4_TEX
107
#endif // DIRECT_ACCESS_LINK
108
#endif // DD_PREC
109
#elif (DD_RECON==1) // reconstruct from 12 reals
110
#define DD_RECON_F 12
111
112
#if (DD_PREC==0)
113
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_12_DOUBLE
114
#ifdef DIRECT_ACCESS_LINK
115
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_DOUBLE2
116
#else
117
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_DOUBLE2_TEX
118
#endif // DIRECT_ACCESS_LINK
119
#define DD_PARAM_GAUGE const double2 *gauge0, const double2 *gauge1,
120
121
#elif (DD_PREC==1)
122
#define DD_PARAM_GAUGE const float4 *gauge0, const float4 *gauge1,
123
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_12_SINGLE
124
#ifdef DIRECT_ACCESS_LINK
125
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_FLOAT4
126
#else
127
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_FLOAT4_TEX
128
#endif // DIRECT_ACCESS_LINK
129
130
#else
131
#define DD_PARAM_GAUGE const short4 *gauge0, const short4 *gauge1,
132
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_12_SINGLE
133
#ifdef DIRECT_ACCESS_LINK
134
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_SHORT4
135
#else
136
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_SHORT4_TEX
137
#endif // DIRECT_ACCESS_LINK
138
#endif // DD_PREC
139
#else // no reconstruct, load all components
140
#define DD_RECON_F 18
141
#define GAUGE_FLOAT2
142
#if (DD_PREC==0)
143
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_18_DOUBLE
144
#ifdef DIRECT_ACCESS_LINK
145
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_DOUBLE2
146
#else
147
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_DOUBLE2_TEX
148
#endif // DIRECT_ACCESS_LINK
149
#define DD_PARAM_GAUGE const double2 *gauge0, const double2 *gauge1,
150
151
#elif (DD_PREC==1)
152
#define DD_PARAM_GAUGE const float4 *gauge0, const float4 *gauge1, // FIXME for direct reading, really float2
153
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_18_SINGLE
154
#ifdef DIRECT_ACCESS_LINK
155
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_FLOAT2
156
#else
157
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_FLOAT2_TEX
158
#endif // DIRECT_ACCESS_LINK
159
160
#else
161
#define DD_PARAM_GAUGE const short4 *gauge0, const short4 *gauge1, // FIXME for direct reading, really short2
162
#define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_18_SINGLE
163
#ifdef DIRECT_ACCESS_LINK
164
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_SHORT2
165
#else
166
#define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_SHORT2_TEX
167
#endif //DIRECT_ACCESS_LINK
168
#endif
169
#endif
170
171
#if (DD_PREC==0) // double-precision fields
172
173
#define TPROJSCALE tProjScale
174
175
// double-precision gauge field
176
#if (defined DIRECT_ACCESS_LINK) || (defined FERMI_NO_DBLE_TEX)
177
#define GAUGE0TEX gauge0
178
#define GAUGE1TEX gauge1
179
#else
180
#ifdef USE_TEXTURE_OBJECTS
181
#define GAUGE0TEX param.gauge0Tex
182
#define GAUGE1TEX param.gauge1Tex
183
#else
184
#define GAUGE0TEX gauge0TexDouble2
185
#define GAUGE1TEX gauge1TexDouble2
186
#endif
187
#endif
188
189
#define GAUGE_FLOAT2
190
191
// double-precision spinor fields
192
#define DD_PARAM_OUT double2* out, float *null1,
193
#define DD_PARAM_IN const double2* in, const float *null4,
194
195
#if (defined DIRECT_ACCESS_WILSON_SPINOR) || (defined FERMI_NO_DBLE_TEX)
196
#define READ_SPINOR READ_SPINOR_DOUBLE
197
#define READ_SPINOR_UP READ_SPINOR_DOUBLE_UP
198
#define READ_SPINOR_DOWN READ_SPINOR_DOUBLE_DOWN
199
#define SPINORTEX in
200
#else
201
#define READ_SPINOR READ_SPINOR_DOUBLE_TEX
202
#define READ_SPINOR_UP READ_SPINOR_DOUBLE_UP_TEX
203
#define READ_SPINOR_DOWN READ_SPINOR_DOUBLE_DOWN_TEX
204
#ifdef USE_TEXTURE_OBJECTS
205
#define SPINORTEX param.inTex
206
#else
207
#define SPINORTEX spinorTexDouble
208
#endif // USE_TEXTURE_OBJECTS
209
#endif
210
#if (defined DIRECT_ACCESS_WILSON_INTER) || (defined FERMI_NO_DBLE_TEX)
211
#define READ_INTERMEDIATE_SPINOR READ_SPINOR_DOUBLE
212
#define INTERTEX out
213
#else
214
#define READ_INTERMEDIATE_SPINOR READ_SPINOR_DOUBLE_TEX
215
#ifdef USE_TEXTURE_OBJECTS
216
#define INTERTEX param.outTex
217
#else
218
#define INTERTEX interTexDouble
219
#endif
220
#endif
221
#define WRITE_SPINOR WRITE_SPINOR_DOUBLE2_STR
222
#define SPINOR_DOUBLE
223
#if (DD_XPAY==1)
224
#if (defined DIRECT_ACCESS_WILSON_ACCUM) || (defined FERMI_NO_DBLE_TEX)
225
#define ACCUMTEX x
226
#define READ_ACCUM READ_ACCUM_DOUBLE
227
#else
228
#ifdef USE_TEXTURE_OBJECTS
229
#define ACCUMTEX param.xTex
230
#else
231
#define ACCUMTEX accumTexDouble
232
#endif // USE_TEXTURE_OBJECTS
233
#define READ_ACCUM READ_ACCUM_DOUBLE_TEX
234
#endif
235
236
#endif
237
238
#define SPINOR_HOP 12
239
240
// double-precision clover field
241
#if (DD_CLOVER==0)
242
#define DD_PARAM_CLOVER
243
#else
244
#define DD_PARAM_CLOVER const double2 *clover, const float *null3,
245
#endif
246
#if (defined DIRECT_ACCESS_CLOVER) || (defined FERMI_NO_DBLE_TEX)
247
#define CLOVERTEX clover
248
#define READ_CLOVER READ_CLOVER_DOUBLE_STR
249
#else
250
#ifdef USE_TEXTURE_OBJECTS
251
#define CLOVERTEX (param.cloverTex)
252
#else
253
#define CLOVERTEX cloverTexDouble
254
#endif
255
#define READ_CLOVER READ_CLOVER_DOUBLE_TEX
256
#endif
257
#define CLOVER_DOUBLE
258
259
#elif (DD_PREC==1) // single-precision fields
260
261
#define TPROJSCALE tProjScale_f
262
263
// single-precision gauge field
264
#ifdef DIRECT_ACCESS_LINK
265
#define GAUGE0TEX gauge0
266
#define GAUGE1TEX gauge1
267
#else
268
#ifdef USE_TEXTURE_OBJECTS
269
#define GAUGE0TEX param.gauge0Tex
270
#define GAUGE1TEX param.gauge1Tex
271
#else
272
#if (DD_RECON_F == 18)
273
#define GAUGE0TEX gauge0TexSingle2
274
#define GAUGE1TEX gauge1TexSingle2
275
#else
276
#define GAUGE0TEX gauge0TexSingle4
277
#define GAUGE1TEX gauge1TexSingle4
278
#endif
279
#endif // USE_TEXTURE_OBJECTS
280
#endif
281
282
283
// single-precision spinor fields
284
#define DD_PARAM_OUT float4* out, float *null1,
285
#define DD_PARAM_IN const float4* in, const float *null4,
286
#ifdef DIRECT_ACCESS_WILSON_SPINOR
287
#define READ_SPINOR READ_SPINOR_SINGLE
288
#define READ_SPINOR_UP READ_SPINOR_SINGLE_UP
289
#define READ_SPINOR_DOWN READ_SPINOR_SINGLE_DOWN
290
#define SPINORTEX in
291
#else
292
#define READ_SPINOR READ_SPINOR_SINGLE_TEX
293
#define READ_SPINOR_UP READ_SPINOR_SINGLE_UP_TEX
294
#define READ_SPINOR_DOWN READ_SPINOR_SINGLE_DOWN_TEX
295
#ifdef USE_TEXTURE_OBJECTS
296
#define SPINORTEX param.inTex
297
#else
298
#define SPINORTEX spinorTexSingle
299
#endif // USE_TEXTURE_OBJECTS
300
#endif
301
#ifdef DIRECT_ACCESS_WILSON_INTER
302
#define READ_INTERMEDIATE_SPINOR READ_SPINOR_SINGLE
303
#define INTERTEX out
304
#else
305
#define READ_INTERMEDIATE_SPINOR READ_SPINOR_SINGLE_TEX
306
#ifdef USE_TEXTURE_OBJECTS
307
#define INTERTEX param.outTex
308
#else
309
#define INTERTEX interTexSingle
310
#endif // USE_TEXTURE_OBJECTS
311
#endif
312
#define WRITE_SPINOR WRITE_SPINOR_FLOAT4_STR
313
#if (DD_XPAY==1)
314
#ifdef DIRECT_ACCESS_WILSON_ACCUM
315
#define ACCUMTEX x
316
#define READ_ACCUM READ_ACCUM_SINGLE
317
#else
318
#ifdef USE_TEXTURE_OBJECTS
319
#define ACCUMTEX param.xTex
320
#else
321
#define ACCUMTEX accumTexSingle
322
#endif // USE_TEXTURE_OBJECTS
323
#define READ_ACCUM READ_ACCUM_SINGLE_TEX
324
#endif
325
#endif
326
327
#define SPINOR_HOP 6
328
329
// single-precision clover field
330
#if (DD_CLOVER==0)
331
#define DD_PARAM_CLOVER
332
#else
333
#define DD_PARAM_CLOVER const float4 *clover, const float *null3,
334
#endif
335
#ifdef DIRECT_ACCESS_CLOVER
336
#define CLOVERTEX clover
337
#define READ_CLOVER READ_CLOVER_SINGLE
338
#else
339
#ifdef USE_TEXTURE_OBJECTS
340
#define CLOVERTEX (param.cloverTex)
341
#else
342
#define CLOVERTEX cloverTexSingle
343
#endif
344
#define READ_CLOVER READ_CLOVER_SINGLE_TEX
345
#endif
346
347
#else // half-precision fields
348
349
#define TPROJSCALE tProjScale_f
350
351
// half-precision gauge field
352
#ifdef DIRECT_ACCESS_LINK
353
#define GAUGE0TEX gauge0
354
#define GAUGE1TEX gauge1
355
#else
356
#ifdef USE_TEXTURE_OBJECTS
357
#define GAUGE0TEX param.gauge0Tex
358
#define GAUGE1TEX param.gauge1Tex
359
#else
360
#if (DD_RECON_F == 18)
361
#define GAUGE0TEX gauge0TexHalf2
362
#define GAUGE1TEX gauge1TexHalf2
363
#else
364
#define GAUGE0TEX gauge0TexHalf4
365
#define GAUGE1TEX gauge1TexHalf4
366
#endif
367
#endif // USE_TEXTURE_OBJECTS
368
#endif
369
370
371
// half-precision spinor fields
372
#ifdef DIRECT_ACCESS_WILSON_SPINOR
373
#define READ_SPINOR READ_SPINOR_HALF
374
#define READ_SPINOR_UP READ_SPINOR_HALF_UP
375
#define READ_SPINOR_DOWN READ_SPINOR_HALF_DOWN
376
#define SPINORTEX in
377
#else
378
#define READ_SPINOR READ_SPINOR_HALF_TEX
379
#define READ_SPINOR_UP READ_SPINOR_HALF_UP_TEX
380
#define READ_SPINOR_DOWN READ_SPINOR_HALF_DOWN_TEX
381
#ifdef USE_TEXTURE_OBJECTS
382
#define SPINORTEX param.inTex
383
#else
384
#define SPINORTEX spinorTexHalf
385
#endif // USE_TEXTURE_OBJECTS
386
#endif
387
#ifdef DIRECT_ACCESS_WILSON_INTER
388
#define READ_INTERMEDIATE_SPINOR READ_SPINOR_HALF
389
#define INTERTEX out
390
#else
391
#define READ_INTERMEDIATE_SPINOR READ_SPINOR_HALF_TEX
392
#ifdef USE_TEXTURE_OBJECTS
393
#define INTERTEX param.outTex
394
#else
395
#define INTERTEX interTexHalf
396
#endif // USE_TEXTURE_OBJECTS
397
#endif
398
#define DD_PARAM_OUT short4* out, float *outNorm,
399
#define DD_PARAM_IN const short4* in, const float *inNorm,
400
#define WRITE_SPINOR WRITE_SPINOR_SHORT4_STR
401
#if (DD_XPAY==1)
402
#ifdef DIRECT_ACCESS_WILSON_ACCUM
403
#define ACCUMTEX x
404
#define READ_ACCUM READ_ACCUM_HALF
405
#else
406
#ifdef USE_TEXTURE_OBJECTS
407
#define ACCUMTEX param.xTex
408
#else
409
#define ACCUMTEX accumTexHalf
410
#endif // USE_TEXTURE_OBJECTS
411
#define READ_ACCUM READ_ACCUM_HALF_TEX
412
#endif
413
#endif
414
415
#define SPINOR_HOP 6
416
417
// half-precision clover field
418
#if (DD_CLOVER==0)
419
#define DD_PARAM_CLOVER
420
#else
421
#define DD_PARAM_CLOVER const short4 *clover, const float *cloverNorm,
422
#endif
423
#ifdef DIRECT_ACCESS_CLOVER
424
#define CLOVERTEX clover
425
#define READ_CLOVER READ_CLOVER_HALF
426
#else
427
#ifdef USE_TEXTURE_OBJECTS
428
#define CLOVERTEX (param.cloverTex)
429
#define CLOVERTEXNORM (param.cloverNormTex)
430
#else
431
#define CLOVERTEX cloverTexHalf
432
#define CLOVERTEXNORM cloverTexNorm
433
#endif
434
#define READ_CLOVER READ_CLOVER_HALF_TEX
435
#endif
436
437
#endif
438
439
// only build double precision if supported
440
#if !(__COMPUTE_CAPABILITY__ < 130 && DD_PREC == 0)
441
442
#define DD_CONCAT(n,r,d,x) n ## r ## d ## x ## Kernel
443
#define DD_FUNC(n,r,d,x) DD_CONCAT(n,r,d,x)
444
445
// define the kernel
446
447
template
<KernelType kernel_type>
448
__global__
void
DD_FUNC
(
DD_NAME_F
,
DD_RECON_F
,
DD_DAG_F
,
DD_XPAY_F
)
449
(
DD_PARAM_OUT
DD_PARAM_GAUGE
DD_PARAM_CLOVER
DD_PARAM_IN
DD_PARAM_XPAY
const
DslashParam
param
) {
450
451
// build Wilson or clover as appropriate
452
#if ((DD_CLOVER==0 && defined(GPU_WILSON_DIRAC)) || ((DD_CLOVER==1 || DD_CLOVER==2) && defined(GPU_CLOVER_DIRAC)))
453
454
#if (__COMPUTE_CAPABILITY__ >= 200 && defined(SHARED_WILSON_DSLASH)) // Fermi optimal code
455
456
#ifdef DSLASH_CLOVER_XPAY
457
458
#if DD_DAG
459
#include "
asym_wilson_clover_dslash_dagger_fermi_core.h
"
460
#else
461
#include "
asym_wilson_clover_dslash_fermi_core.h
"
462
#endif
463
464
#else
465
466
#if DD_DAG
467
#include "
wilson_dslash_dagger_fermi_core.h
"
468
#else
469
#include "
wilson_dslash_fermi_core.h
"
470
#endif
471
472
#endif
473
474
#elif (__COMPUTE_CAPABILITY__ >= 120) // GT200 optimal code
475
476
#ifdef DSLASH_CLOVER_XPAY
477
478
#if DD_DAG
479
#include "
asym_wilson_clover_dslash_dagger_gt200_core.h
"
480
#else
481
#include "
asym_wilson_clover_dslash_gt200_core.h
"
482
#endif
483
484
#else
485
486
#if DD_DAG
487
#include "
wilson_dslash_dagger_gt200_core.h
"
488
#else
489
#include "
wilson_dslash_gt200_core.h
"
490
#endif
491
492
#endif
493
494
#else // fall-back is original G80
495
496
#ifdef DSLASH_CLOVER_XPAY
497
498
#if DD_DAG
499
#include "
asym_wilson_clover_dslash_dagger_g80_core.h
"
500
#else
501
#include "
asym_wilson_clover_dslash_g80_core.h
"
502
#endif
503
504
#else
505
506
#if DD_DAG
507
#include "
wilson_dslash_dagger_g80_core.h
"
508
#else
509
#include "
wilson_dslash_g80_core.h
"
510
#endif
511
512
#endif // DSLASH_CLOVER_XPAY
513
514
515
#endif // __COMPUTE_CAPABILITY
516
517
518
#endif // DD_CLOVER
519
520
}
521
522
#endif
523
524
// clean up
525
526
#undef DD_NAME_F
527
#undef DD_RECON_F
528
#undef DD_DAG_F
529
#undef DD_XPAY_F
530
#undef DD_PARAM_OUT
531
#undef DD_PARAM_GAUGE
532
#undef DD_PARAM_CLOVER
533
#undef DD_PARAM_IN
534
#undef DD_PARAM_XPAY
535
#undef DD_CONCAT
536
#undef DD_FUNC
537
538
#undef DSLASH_XPAY
539
#undef READ_GAUGE_MATRIX
540
#undef RECONSTRUCT_GAUGE_MATRIX
541
#undef GAUGE0TEX
542
#undef GAUGE1TEX
543
#undef READ_SPINOR
544
#undef READ_SPINOR_UP
545
#undef READ_SPINOR_DOWN
546
#undef SPINORTEX
547
#undef READ_INTERMEDIATE_SPINOR
548
#undef INTERTEX
549
#undef WRITE_SPINOR
550
#undef READ_ACCUM
551
#undef ACCUMTEX
552
#undef READ_CLOVER
553
#undef CLOVERTEX
554
#undef DSLASH_CLOVER
555
#undef DSLASH_CLOVER_XPAY
556
#undef GAUGE_FLOAT2
557
#undef SPINOR_DOUBLE
558
#undef CLOVER_DOUBLE
559
#undef SPINOR_HOP
560
561
#undef TPROJSCALE
562
563
// prepare next set of options, or clean up after final iteration
564
565
#if (DD_DAG==0)
566
#undef DD_DAG
567
#define DD_DAG 1
568
#else
569
#undef DD_DAG
570
#define DD_DAG 0
571
572
#if (DD_XPAY==0)
573
#undef DD_XPAY
574
#define DD_XPAY 1
575
#else
576
#undef DD_XPAY
577
#define DD_XPAY 0
578
579
#if (DD_RECON==0)
580
#undef DD_RECON
581
#define DD_RECON 1
582
#elif (DD_RECON==1)
583
#undef DD_RECON
584
#define DD_RECON 2
585
#else
586
#undef DD_RECON
587
#define DD_RECON 0
588
589
#if (DD_PREC==0)
590
#undef DD_PREC
591
#define DD_PREC 1
592
#elif (DD_PREC==1)
593
#undef DD_PREC
594
#define DD_PREC 2
595
#else
596
#undef DD_PREC
597
#define DD_PREC 0
598
599
#if (DD_CLOVER==0)
600
#undef DD_CLOVER
601
#define DD_CLOVER 1
602
#elif (DD_CLOVER==1)
603
#undef DD_CLOVER
604
#define DD_CLOVER 2
605
606
#else
607
608
#undef DD_LOOP
609
#undef DD_DAG
610
#undef DD_XPAY
611
#undef DD_RECON
612
#undef DD_PREC
613
#undef DD_CLOVER
614
615
#endif // DD_CLOVER
616
#endif // DD_PREC
617
#endif // DD_RECON
618
#endif // DD_XPAY
619
#endif // DD_DAG
620
621
#ifdef DD_LOOP
622
#include "
wilson_dslash_def.h
"
623
#endif
Generated on Wed Mar 20 2013 12:52:17 for QUDA by
1.8.2
1.8.2