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