QUDA  0.9.0
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 #endif
37 
38 // set options for current iteration
39 
40 #if (DD_CLOVER==0) // no clover
41 #define DD_NAME_F dslash
42 #elif (DD_CLOVER==1) // clover
43 #define DSLASH_CLOVER
44 #define DD_NAME_F cloverDslash
45 #else
46 #define DSLASH_CLOVER
47 #define DSLASH_CLOVER_XPAY
48 #define DD_NAME_F asymCloverDslash
49 #endif
50 
51 #if (DD_DAG==0) // no dagger
52 #define DD_DAG_F
53 #else // dagger
54 #define DD_DAG_F Dagger
55 #endif
56 
57 // DSLASH_CLOVER_XPAY implies DD_XPAY=1
58 #if (DD_XPAY==0) && defined(DSLASH_CLOVER_XPAY)
59 #undef DD_XPAY
60 #define DD_XPAY 1
61 #endif
62 
63 #if (DD_XPAY==0) // no xpay
64 #define DD_XPAY_F
65 #else // xpay
66 #define DD_XPAY_F Xpay
67 #define DSLASH_XPAY
68 #endif
69 
70 #if (DD_PREC == 0)
71 #define DD_PREC_F D
72 #elif (DD_PREC == 1)
73 #define DD_PREC_F S
74 #else
75 #define DD_PREC_F H
76 #endif
77 
78 #if (DD_RECON==0) // reconstruct from 8 reals
79 #define DD_RECON_F 8
80 
81 #if (DD_PREC==0)
82 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_8_DOUBLE
83 #ifdef DIRECT_ACCESS_LINK
84 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_DOUBLE2
85 #else
86 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_DOUBLE2_TEX
87 #endif // DIRECT_ACCESS_LINK
88 
89 #elif (DD_PREC==1)
90 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_8_SINGLE
91 #ifdef DIRECT_ACCESS_LINK
92 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_FLOAT4
93 #else
94 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_FLOAT4_TEX
95 #endif // DIRECT_ACCESS_LINK
96 
97 #else
98 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_8_SINGLE
99 #ifdef DIRECT_ACCESS_LINK
100 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_SHORT4
101 #else
102 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_SHORT4_TEX
103 #endif // DIRECT_ACCESS_LINK
104 #endif // DD_PREC
105 #elif (DD_RECON==1) // reconstruct from 12 reals
106 #define DD_RECON_F 12
107 
108 #if (DD_PREC==0)
109 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_12_DOUBLE
110 #ifdef DIRECT_ACCESS_LINK
111 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_DOUBLE2
112 #else
113 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_DOUBLE2_TEX
114 #endif // DIRECT_ACCESS_LINK
115 
116 #elif (DD_PREC==1)
117 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_12_SINGLE
118 #ifdef DIRECT_ACCESS_LINK
119 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_FLOAT4
120 #else
121 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_FLOAT4_TEX
122 #endif // DIRECT_ACCESS_LINK
123 
124 #else
125 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_12_SINGLE
126 #ifdef DIRECT_ACCESS_LINK
127 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_SHORT4
128 #else
129 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_SHORT4_TEX
130 #endif // DIRECT_ACCESS_LINK
131 #endif // DD_PREC
132 #else // no reconstruct, load all components
133 #define DD_RECON_F 18
134 #define GAUGE_FLOAT2
135 #if (DD_PREC==0)
136 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_18_DOUBLE
137 #ifdef DIRECT_ACCESS_LINK
138 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_DOUBLE2
139 #else
140 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_DOUBLE2_TEX
141 #endif // DIRECT_ACCESS_LINK
142 
143 #elif (DD_PREC==1)
144 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_18_SINGLE
145 #ifdef DIRECT_ACCESS_LINK
146 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_FLOAT2
147 #else
148 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_FLOAT2_TEX
149 #endif // DIRECT_ACCESS_LINK
150 
151 #else
152 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_18_SINGLE
153 #ifdef DIRECT_ACCESS_LINK
154 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_SHORT2
155 #else
156 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_SHORT2_TEX
157 #endif //DIRECT_ACCESS_LINK
158 #endif
159 #endif
160 
161 #if (DD_PREC==0) // double-precision fields
162 
163 #define TPROJSCALE param.tProjScale
164 
165 // double-precision gauge field
166 #if (defined DIRECT_ACCESS_LINK) || (defined FERMI_NO_DBLE_TEX)
167 #define GAUGE0TEX param.gauge0
168 #define GAUGE1TEX param.gauge1
169 #else
170 #ifdef USE_TEXTURE_OBJECTS
171 #define GAUGE0TEX param.gauge0Tex
172 #define GAUGE1TEX param.gauge1Tex
173 #else
174 #define GAUGE0TEX gauge0TexDouble2
175 #define GAUGE1TEX gauge1TexDouble2
176 #endif
177 #endif
178 
179 #define GAUGE_FLOAT2
180 
181 // double-precision spinor fields
182 #if (defined DIRECT_ACCESS_WILSON_SPINOR) || (defined FERMI_NO_DBLE_TEX)
183 #define READ_SPINOR READ_SPINOR_DOUBLE
184 #define READ_SPINOR_GHOST READ_SPINOR_GHOST_DOUBLE
185 #define READ_SPINOR_UP READ_SPINOR_DOUBLE_UP
186 #define READ_SPINOR_DOWN READ_SPINOR_DOUBLE_DOWN
187 #define SPINORTEX param.in
188 #define GHOSTSPINORTEX param.ghost
189 #else
190 #define READ_SPINOR READ_SPINOR_DOUBLE_TEX
191 #define READ_SPINOR_GHOST READ_SPINOR_GHOST_DOUBLE_TEX
192 #define READ_SPINOR_UP READ_SPINOR_DOUBLE_UP_TEX
193 #define READ_SPINOR_DOWN READ_SPINOR_DOUBLE_DOWN_TEX
194 #ifdef USE_TEXTURE_OBJECTS
195 #define SPINORTEX param.inTex
196 #define GHOSTSPINORTEX param.ghostTex
197 #else
198 #define SPINORTEX spinorTexDouble
199 #define GHOSTSPINORTEX ghostSpinorTexDouble
200 #endif // USE_TEXTURE_OBJECTS
201 #endif
202 #if (defined DIRECT_ACCESS_WILSON_INTER) || (defined FERMI_NO_DBLE_TEX)
203 #define READ_INTERMEDIATE_SPINOR READ_SPINOR_DOUBLE
204 #define INTERTEX param.out
205 #else
206 #define READ_INTERMEDIATE_SPINOR READ_SPINOR_DOUBLE_TEX
207 #ifdef USE_TEXTURE_OBJECTS
208 #define INTERTEX param.outTex
209 #else
210 #define INTERTEX interTexDouble
211 #endif
212 #endif
213 #define WRITE_SPINOR WRITE_SPINOR_DOUBLE2_STR
214 #define SPINOR_DOUBLE
215 #if (DD_XPAY==1)
216 #if (defined DIRECT_ACCESS_WILSON_ACCUM) || (defined FERMI_NO_DBLE_TEX)
217 #define ACCUMTEX param.x
218 #define READ_ACCUM READ_ACCUM_DOUBLE
219 #else
220 #ifdef USE_TEXTURE_OBJECTS
221 #define ACCUMTEX param.xTex
222 #else
223 #define ACCUMTEX accumTexDouble
224 #endif // USE_TEXTURE_OBJECTS
225 #define READ_ACCUM READ_ACCUM_DOUBLE_TEX
226 #endif
227 
228 #endif
229 
230 #define SPINOR_HOP 12
231 
232 // double-precision clover field
233 #if (defined DIRECT_ACCESS_CLOVER) || (defined FERMI_NO_DBLE_TEX)
234 #define CLOVERTEX param.clover
235 #if (DD_CLOVER==2)
236 #define READ_CLOVER READ_CLOVER2_DOUBLE_STR
237 #else
238 #define READ_CLOVER READ_CLOVER_DOUBLE_STR
239 #endif
240 #else
241 #ifdef USE_TEXTURE_OBJECTS
242 #define CLOVERTEX (param.cloverTex)
243 #else
244 #define CLOVERTEX cloverTexDouble
245 #endif
246 
247 #if (DD_CLOVER==2)
248 #define READ_CLOVER READ_CLOVER2_DOUBLE_TEX
249 #else
250 #define READ_CLOVER READ_CLOVER_DOUBLE_TEX
251 #endif
252 
253 #endif
254 #define CLOVER_DOUBLE
255 
256 #elif (DD_PREC==1) // single-precision fields
257 
258 #define TPROJSCALE param.tProjScale_f
259 
260 // single-precision gauge field
261 #ifdef DIRECT_ACCESS_LINK
262 #define GAUGE0TEX param.gauge0
263 #define GAUGE1TEX param.gauge1
264 #else
265 #ifdef USE_TEXTURE_OBJECTS
266 #define GAUGE0TEX param.gauge0Tex
267 #define GAUGE1TEX param.gauge1Tex
268 #else
269 #if (DD_RECON_F == 18)
270 #define GAUGE0TEX gauge0TexSingle2
271 #define GAUGE1TEX gauge1TexSingle2
272 #else
273 #define GAUGE0TEX gauge0TexSingle4
274 #define GAUGE1TEX gauge1TexSingle4
275 #endif
276 #endif // USE_TEXTURE_OBJECTS
277 #endif
278 
279 
280 // single-precision spinor fields
281 #ifdef DIRECT_ACCESS_WILSON_SPINOR
282 #define READ_SPINOR READ_SPINOR_SINGLE
283 #define READ_SPINOR_GHOST READ_SPINOR_GHOST_SINGLE
284 #define READ_SPINOR_UP READ_SPINOR_SINGLE_UP
285 #define READ_SPINOR_DOWN READ_SPINOR_SINGLE_DOWN
286 #define SPINORTEX param.in
287 #define GHOSTSPINORTEX param.ghost
288 #else
289 #define READ_SPINOR READ_SPINOR_SINGLE_TEX
290 #define READ_SPINOR_GHOST READ_SPINOR_GHOST_SINGLE_TEX
291 #define READ_SPINOR_UP READ_SPINOR_SINGLE_UP_TEX
292 #define READ_SPINOR_DOWN READ_SPINOR_SINGLE_DOWN_TEX
293 #ifdef USE_TEXTURE_OBJECTS
294 #define SPINORTEX param.inTex
295 #define GHOSTSPINORTEX param.ghostTex
296 #else
297 #define SPINORTEX spinorTexSingle
298 #define GHOSTSPINORTEX ghostSpinorTexSingle
299 #endif // USE_TEXTURE_OBJECTS
300 #endif
301 #ifdef DIRECT_ACCESS_WILSON_INTER
302 #define READ_INTERMEDIATE_SPINOR READ_SPINOR_SINGLE
303 #define INTERTEX param.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 param.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 #ifdef DIRECT_ACCESS_CLOVER
331 #define CLOVERTEX param.clover
332 #if (DD_CLOVER==2)
333 #define READ_CLOVER READ_CLOVER2_SINGLE
334 #else
335 #define READ_CLOVER READ_CLOVER_SINGLE
336 #endif
337 #else
338 #ifdef USE_TEXTURE_OBJECTS
339 #define CLOVERTEX (param.cloverTex)
340 #else
341 #define CLOVERTEX cloverTexSingle
342 #endif
343 #if (DD_CLOVER==2)
344 #define READ_CLOVER READ_CLOVER2_SINGLE_TEX
345 #else
346 #define READ_CLOVER READ_CLOVER_SINGLE_TEX
347 #endif
348 
349 #endif
350 
351 #else // half-precision fields
352 
353 #define TPROJSCALE param.tProjScale_f
354 
355 // half-precision gauge field
356 #ifdef DIRECT_ACCESS_LINK
357 #define GAUGE0TEX param.gauge0
358 #define GAUGE1TEX param.gauge1
359 #else
360 #ifdef USE_TEXTURE_OBJECTS
361 #define GAUGE0TEX param.gauge0Tex
362 #define GAUGE1TEX param.gauge1Tex
363 #else
364 #if (DD_RECON_F == 18)
365 #define GAUGE0TEX gauge0TexHalf2
366 #define GAUGE1TEX gauge1TexHalf2
367 #else
368 #define GAUGE0TEX gauge0TexHalf4
369 #define GAUGE1TEX gauge1TexHalf4
370 #endif
371 #endif // USE_TEXTURE_OBJECTS
372 #endif
373 
374 
375 // half-precision spinor fields
376 #ifdef DIRECT_ACCESS_WILSON_SPINOR
377 #define READ_SPINOR READ_SPINOR_HALF
378 #define READ_SPINOR_GHOST READ_SPINOR_GHOST_HALF
379 #define READ_SPINOR_UP READ_SPINOR_HALF_UP
380 #define READ_SPINOR_DOWN READ_SPINOR_HALF_DOWN
381 #define SPINORTEX param.in
382 #define GHOSTSPINORTEX param.ghost
383 #else
384 #define READ_SPINOR READ_SPINOR_HALF_TEX
385 #define READ_SPINOR_GHOST READ_SPINOR_GHOST_HALF_TEX
386 #define READ_SPINOR_UP READ_SPINOR_HALF_UP_TEX
387 #define READ_SPINOR_DOWN READ_SPINOR_HALF_DOWN_TEX
388 #ifdef USE_TEXTURE_OBJECTS
389 #define SPINORTEX param.inTex
390 #define GHOSTSPINORTEX param.ghostTex
391 #else
392 #define SPINORTEX spinorTexHalf
393 #define GHOSTSPINORTEX ghostSpinorTexHalf
394 #endif // USE_TEXTURE_OBJECTS
395 #endif
396 #ifdef DIRECT_ACCESS_WILSON_INTER
397 #define READ_INTERMEDIATE_SPINOR READ_SPINOR_HALF
398 #define INTERTEX param.out
399 #else
400 #define READ_INTERMEDIATE_SPINOR READ_SPINOR_HALF_TEX
401 #ifdef USE_TEXTURE_OBJECTS
402 #define INTERTEX param.outTex
403 #else
404 #define INTERTEX interTexHalf
405 #endif // USE_TEXTURE_OBJECTS
406 #endif
407 #define WRITE_SPINOR WRITE_SPINOR_SHORT4_STR
408 #if (DD_XPAY==1)
409 #ifdef DIRECT_ACCESS_WILSON_ACCUM
410 #define ACCUMTEX param.x
411 #define READ_ACCUM READ_ACCUM_HALF
412 #else
413 #ifdef USE_TEXTURE_OBJECTS
414 #define ACCUMTEX param.xTex
415 #else
416 #define ACCUMTEX accumTexHalf
417 #endif // USE_TEXTURE_OBJECTS
418 #define READ_ACCUM READ_ACCUM_HALF_TEX
419 #endif
420 #endif
421 
422 #define SPINOR_HOP 6
423 
424 // half-precision clover field
425 #ifdef DIRECT_ACCESS_CLOVER
426 #define CLOVERTEX param.clover
427 #ifdef (DD_CLOVER==2)
428 #define READ_CLOVER READ_CLOVER2_HALF
429 #else
430 #define READ_CLOVER READ_CLOVER_HALF
431 #endif
432 #define CLOVERTEXNORM (param.cloverNorm)
433 #else
434 #ifdef USE_TEXTURE_OBJECTS
435 #define CLOVERTEX (param.cloverTex)
436 #define CLOVERTEXNORM (param.cloverNormTex)
437 #else
438 #define CLOVERTEX cloverTexHalf
439 #define CLOVERTEXNORM cloverTexNorm
440 #endif
441 #if (DD_CLOVER==2)
442 #define READ_CLOVER READ_CLOVER2_HALF_TEX
443 #else
444 #define READ_CLOVER READ_CLOVER_HALF_TEX
445 #endif
446 #endif
447 
448 #endif
449 
450 #define DD_CONCAT(n,p,r,d,x) n ## p ## r ## d ## x ## Kernel
451 #define DD_FUNC(n,p,r,d,x) DD_CONCAT(n,p,r,d,x)
452 
453 // define the kernel
454 
455 template <KernelType kernel_type>
456  __global__ void DD_FUNC(DD_NAME_F, DD_PREC_F, DD_RECON_F, DD_DAG_F, DD_XPAY_F)(const DslashParam param) {
457 
458  // build Wilson or clover as appropriate
459 #if ((DD_CLOVER==0 && defined(GPU_WILSON_DIRAC)) || ((DD_CLOVER==1 || DD_CLOVER==2) && defined(GPU_CLOVER_DIRAC)))
460 
461 #ifdef SHARED_WILSON_DSLASH // Fermi optimal code
462 
463 #ifdef DSLASH_CLOVER_XPAY
464 
465 #if DD_DAG
466 #include "asym_wilson_clover_dslash_dagger_fermi_core.h"
467 #else
468 #include "asym_wilson_clover_dslash_fermi_core.h"
469 #endif
470 
471 #else
472 
473 #if DD_DAG
474 #include "wilson_dslash_dagger_fermi_core.h"
475 #else
476 #include "wilson_dslash_fermi_core.h"
477 #endif
478 
479 #endif
480 
481 #else // no shared-memory blocking
482 
483 #ifdef DSLASH_CLOVER_XPAY
484 
485 #if DD_DAG
486 #include "asym_wilson_clover_dslash_dagger_gt200_core.h"
487 #else
488 #include "asym_wilson_clover_dslash_gt200_core.h"
489 #endif
490 
491 #else
492 
493 #if DD_DAG
494 #include "wilson_dslash_dagger_gt200_core.h"
495 #else
496 #include "wilson_dslash_gt200_core.h"
497 #endif
498 
499 #endif
500 
501 #endif // SHARED_WILSON_DSLASH
502 
503 
504 #endif // DD_CLOVER
505 
506 }
507 
508 #ifdef MULTI_GPU
509 template <>
510 __global__ void DD_FUNC(DD_NAME_F, DD_PREC_F, DD_RECON_F, DD_DAG_F, DD_XPAY_F)<EXTERIOR_KERNEL_ALL>
511  (const DslashParam param) {
512 
513  // build Wilson or clover as appropriate
514 #if ((DD_CLOVER==0 && defined(GPU_WILSON_DIRAC)) || ((DD_CLOVER==1 || DD_CLOVER==2) && defined(GPU_CLOVER_DIRAC)))
515 
516 #ifdef SHARED_WILSON_DSLASH // Fermi optimal code
517 
518 #ifdef DSLASH_CLOVER_XPAY
519 
520 #if DD_DAG
521 #include "asym_wilson_clover_fused_exterior_dslash_dagger_fermi_core.h"
522 #else
523 #include "asym_wilson_clover_fused_exterior_dslash_fermi_core.h"
524 #endif
525 
526 #else
527 
528 #if DD_DAG
529 #include "wilson_fused_exterior_dslash_dagger_fermi_core.h"
530 #else
531 #include "wilson_fused_exterior_dslash_fermi_core.h"
532 #endif
533 
534 #endif
535 
536 #else // no shared-memory blocking
537 
538 #ifdef DSLASH_CLOVER_XPAY
539 
540 #if DD_DAG
541 #include "asym_wilson_clover_fused_exterior_dslash_dagger_gt200_core.h"
542 #else
543 #include "asym_wilson_clover_fused_exterior_dslash_gt200_core.h"
544 #endif
545 
546 #else
547 
548 #if DD_DAG
549 #include "wilson_fused_exterior_dslash_dagger_gt200_core.h"
550 #else
551 #include "wilson_fused_exterior_dslash_gt200_core.h"
552 #endif
553 
554 #endif
555 
556 #endif // SHARED_WILSON_DSLASH
557 
558 
559 #endif // DD_CLOVER
560 
561 }
562 #endif // MULTI_GPU
563 
564 // clean up
565 
566 #undef DD_PREC_F
567 #undef DD_NAME_F
568 #undef DD_RECON_F
569 #undef DD_DAG_F
570 #undef DD_XPAY_F
571 #undef DD_CONCAT
572 #undef DD_FUNC
573 
574 #undef DSLASH_XPAY
575 #undef READ_GAUGE_MATRIX
576 #undef RECONSTRUCT_GAUGE_MATRIX
577 #undef GAUGE0TEX
578 #undef GAUGE1TEX
579 #undef READ_SPINOR
580 #undef READ_SPINOR_GHOST
581 #undef READ_SPINOR_UP
582 #undef READ_SPINOR_DOWN
583 #undef SPINORTEX
584 #undef GHOSTSPINORTEX
585 #undef READ_INTERMEDIATE_SPINOR
586 #undef INTERTEX
587 #undef WRITE_SPINOR
588 #undef READ_ACCUM
589 #undef ACCUMTEX
590 #undef READ_CLOVER
591 #undef CLOVERTEX
592 #undef DSLASH_CLOVER
593 #undef DSLASH_CLOVER_XPAY
594 #undef GAUGE_FLOAT2
595 #undef SPINOR_DOUBLE
596 #undef CLOVER_DOUBLE
597 #undef SPINOR_HOP
598 
599 #undef TPROJSCALE
600 
601 // prepare next set of options, or clean up after final iteration
602 
603 #if (DD_DAG==0)
604 #undef DD_DAG
605 #define DD_DAG 1
606 #else
607 #undef DD_DAG
608 #define DD_DAG 0
609 
610 #if (DD_XPAY==0)
611 #undef DD_XPAY
612 #define DD_XPAY 1
613 #else
614 #undef DD_XPAY
615 #define DD_XPAY 0
616 
617 #if (DD_RECON==0)
618 #undef DD_RECON
619 #define DD_RECON 1
620 #elif (DD_RECON==1)
621 #undef DD_RECON
622 #define DD_RECON 2
623 #else
624 #undef DD_RECON
625 #define DD_RECON 0
626 
627 #if (DD_PREC==0)
628 #undef DD_PREC
629 #define DD_PREC 1
630 #elif (DD_PREC==1)
631 #undef DD_PREC
632 #define DD_PREC 2
633 #else
634 
635 #undef DD_LOOP
636 #undef DD_DAG
637 #undef DD_XPAY
638 #undef DD_RECON
639 #undef DD_PREC
640 
641 #endif // DD_PREC
642 #endif // DD_RECON
643 #endif // DD_XPAY
644 #endif // DD_DAG
645 
646 #ifdef DD_LOOP
647 #include "wilson_dslash_def.h"
648 #endif
DD_RECON_F
#define DD_RECON_F
Definition: wilson_dslash_def.h:79
asym_wilson_clover_fused_exterior_dslash_dagger_gt200_core.h
DD_XPAY_F
#define DD_XPAY_F
Definition: wilson_dslash_def.h:64
wilson_fused_exterior_dslash_gt200_core.h
wilson_fused_exterior_dslash_dagger_gt200_core.h
DD_DAG_F
#define DD_DAG_F
Definition: wilson_dslash_def.h:52
wilson_dslash_gt200_core.h
wilson_dslash_dagger_fermi_core.h
param
QudaGaugeParam param
Definition: pack_test.cpp:17
DD_FUNC
#define DD_FUNC(n, p, r, d, x)
Definition: wilson_dslash_def.h:451
wilson_dslash_dagger_gt200_core.h
asym_wilson_clover_fused_exterior_dslash_dagger_fermi_core.h
wilson_fused_exterior_dslash_dagger_fermi_core.h
DD_PREC_F
#define DD_PREC_F
Definition: wilson_dslash_def.h:71
wilson_dslash_fermi_core.h
asym_wilson_clover_dslash_gt200_core.h
asym_wilson_clover_dslash_dagger_gt200_core.h
wilson_dslash_def.h
EXTERIOR_KERNEL_ALL
Definition: dslash_constants.h:7
wilson_fused_exterior_dslash_fermi_core.h
asym_wilson_clover_dslash_fermi_core.h
DD_NAME_F
#define DD_NAME_F
Definition: wilson_dslash_def.h:41
asym_wilson_clover_dslash_dagger_fermi_core.h
asym_wilson_clover_fused_exterior_dslash_fermi_core.h
asym_wilson_clover_fused_exterior_dslash_gt200_core.h
Generated by   doxygen 1.8.14