QUDA  0.9.0
tmc_dslash_def.h
Go to the documentation of this file.
1 // tmc_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 = twistedCloverDslash
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 twistedCloverDslash
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_PREC == 0)
59 #define DD_PREC_F D
60 #elif (DD_PREC == 1)
61 #define DD_PREC_F S
62 #else
63 #define DD_PREC_F H
64 #endif
65 
66 #if (DD_TWIST==0) // twisted input
67 #define DD_NAME_F twistedCloverInvDslash
68 #define CLOVER_TWIST_INV_DSLASH
69 #else
70 #define DD_NAME_F twistedCloverDslash
71 #endif
72 
74 #if (DD_PREC == 0)
75  #if (defined DIRECT_ACCESS_CLOVER) || (defined FERMI_NO_DBLE_TEX)
76 #define TMCLOVERTEX (param.clover)
77  #ifndef DYNAMIC_CLOVER
78 #define TM_INV_CLOVERTEX (param.cloverInv)
79  #endif
80  #define READ_CLOVER READ_CLOVER_DOUBLE_STR
81  #define ASSN_CLOVER ASSN_CLOVER_DOUBLE_STR
82  #else
83  #ifdef USE_TEXTURE_OBJECTS
84  #define TMCLOVERTEX (param.cloverTex)
85  #ifndef DYNAMIC_CLOVER
86  #define TM_INV_CLOVERTEX (param.cloverInvTex)
87  #endif
88  #else
89  #define TMCLOVERTEX cloverTexDouble
90  #ifndef DYNAMIC_CLOVER
91  #define TM_INV_CLOVERTEX cloverInvTexDouble
92  #endif
93  #endif
94  #define READ_CLOVER READ_CLOVER_DOUBLE_TEX
95  #define ASSN_CLOVER ASSN_CLOVER_DOUBLE_TEX
96  #endif
97  #define CLOVER_DOUBLE
98 #elif (DD_PREC == 1)
99  #ifdef DIRECT_ACCESS_CLOVER
100 #define TMCLOVERTEX (param.clover)
101  #ifndef DYNAMIC_CLOVER
102 #define TM_INV_CLOVERTEX (param.cloverInv)
103  #endif
104  #define READ_CLOVER READ_CLOVER_SINGLE
105  #define ASSN_CLOVER ASSN_CLOVER_SINGLE
106  #else
107  #ifdef USE_TEXTURE_OBJECTS
108  #define TMCLOVERTEX (param.cloverTex)
109  #ifndef DYNAMIC_CLOVER
110  #define TM_INV_CLOVERTEX (param.cloverInvTex)
111  #endif
112  #else
113  #define TMCLOVERTEX cloverTexSingle
114  #ifndef DYNAMIC_CLOVER
115  #define TM_INV_CLOVERTEX cloverInvTexSingle
116  #endif
117  #endif
118  #define READ_CLOVER READ_CLOVER_SINGLE_TEX
119  #define ASSN_CLOVER ASSN_CLOVER_SINGLE_TEX
120  #endif
121 #else // half-precision fields
122  #ifdef DIRECT_ACCESS_CLOVER
123  #define TMCLOVERTEX param.clover
124  #ifndef DYNAMIC_CLOVER
125  #define TM_INV_CLOVERTEX param.cloverInv
126  #endif
127  #define READ_CLOVER READ_CLOVER_HALF
128  #define ASSN_CLOVER ASSN_CLOVER_HALF
129  #else
130  #ifdef USE_TEXTURE_OBJECTS
131  #define TMCLOVERTEX (param.cloverTex)
132  #define TMCLOVERTEXNORM (param.cloverNormTex)
133  #ifndef DYNAMIC_CLOVER
134  #define TM_INV_CLOVERTEX (param.cloverInvTex)
135  #define TM_INV_CLOVERTEXNORM (param.cloverInvNormTex)
136  #endif
137  #else
138  #define TMCLOVERTEX cloverTexHalf
139  #define TMCLOVERTEXNORM cloverTexNorm
140  #ifndef DYNAMIC_CLOVER
141  #define TM_INV_CLOVERTEX cloverInvTexHalf
142  #define TM_INV_CLOVERTEXNORM cloverInvTexNorm
143  #endif
144  #endif
145  #define READ_CLOVER READ_CLOVER_HALF_TEX
146  #define ASSN_CLOVER ASSN_CLOVER_HALF_TEX
147  #endif
148 
149 #endif //End clover defs
150 
151 
152 #if (DD_RECON==0) // reconstruct from 8 reals
153 #define DD_RECON_F 8
154 
155 #if (DD_PREC==0)
156 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_8_DOUBLE
157 #ifdef DIRECT_ACCESS_LINK
158 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_DOUBLE2
159 #else
160 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_DOUBLE2_TEX
161 #endif // DIRECT_ACCESS_LINK
162 
163 #elif (DD_PREC==1)
164 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_8_SINGLE
165 #ifdef DIRECT_ACCESS_LINK
166 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_FLOAT4
167 #else
168 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_FLOAT4_TEX
169 #endif // DIRECT_ACCESS_LINK
170 
171 #else
172 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_8_SINGLE
173 #ifdef DIRECT_ACCESS_LINK
174 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_SHORT4
175 #else
176 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_8_SHORT4_TEX
177 #endif // DIRECT_ACCESS_LINK
178 #endif // DD_PREC
179 #elif (DD_RECON==1) // reconstruct from 12 reals
180 #define DD_RECON_F 12
181 
182 #if (DD_PREC==0)
183 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_12_DOUBLE
184 #ifdef DIRECT_ACCESS_LINK
185 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_DOUBLE2
186 #else
187 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_DOUBLE2_TEX
188 #endif // DIRECT_ACCESS_LINK
189 
190 #elif (DD_PREC==1)
191 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_12_SINGLE
192 #ifdef DIRECT_ACCESS_LINK
193 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_FLOAT4
194 #else
195 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_FLOAT4_TEX
196 #endif // DIRECT_ACCESS_LINK
197 
198 #else
199 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_12_SINGLE
200 #ifdef DIRECT_ACCESS_LINK
201 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_SHORT4
202 #else
203 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_12_SHORT4_TEX
204 #endif // DIRECT_ACCESS_LINK
205 #endif // DD_PREC
206 #else // no reconstruct, load all components
207 #define DD_RECON_F 18
208 #define GAUGE_FLOAT2
209 #if (DD_PREC==0)
210 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_18_DOUBLE
211 #ifdef DIRECT_ACCESS_LINK
212 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_DOUBLE2
213 #else
214 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_DOUBLE2_TEX
215 #endif // DIRECT_ACCESS_LINK
216 
217 #elif (DD_PREC==1)
218 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_18_SINGLE
219 #ifdef DIRECT_ACCESS_LINK
220 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_FLOAT2
221 #else
222 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_FLOAT2_TEX
223 #endif // DIRECT_ACCESS_LINK
224 
225 #else
226 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_MATRIX_18_SINGLE
227 #ifdef DIRECT_ACCESS_LINK
228 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_SHORT2
229 #else
230 #define READ_GAUGE_MATRIX READ_GAUGE_MATRIX_18_SHORT2_TEX
231 #endif //DIRECT_ACCESS_LINK
232 #endif
233 #endif
234 
235 #if (DD_PREC==0) // double-precision fields
236 
237 #define TPROJSCALE param.tProjScale
238 
239 // double-precision gauge field
240 #if (defined DIRECT_ACCESS_LINK) || (defined FERMI_NO_DBLE_TEX)
241 #define GAUGE0TEX param.gauge0
242 #define GAUGE1TEX param.gauge1
243 #else
244 #ifdef USE_TEXTURE_OBJECTS
245 #define GAUGE0TEX param.gauge0Tex
246 #define GAUGE1TEX param.gauge1Tex
247 #else
248 #define GAUGE0TEX gauge0TexDouble2
249 #define GAUGE1TEX gauge1TexDouble2
250 #endif // USE_TEXTURE_OBJECTS
251 #endif
252 
253 #define GAUGE_FLOAT2
254 
255 // double-precision spinor fields
256 #if (defined DIRECT_ACCESS_WILSON_SPINOR) || (defined FERMI_NO_DBLE_TEX)
257 #define READ_SPINOR READ_SPINOR_DOUBLE
258 #define READ_SPINOR_GHOST READ_SPINOR_GHOST_DOUBLE
259 #define READ_SPINOR_UP READ_SPINOR_DOUBLE_UP
260 #define READ_SPINOR_DOWN READ_SPINOR_DOUBLE_DOWN
261 #define SPINORTEX param.in
262 #else
263 #define READ_SPINOR READ_SPINOR_DOUBLE_TEX
264 #define READ_SPINOR_GHOST READ_SPINOR_GHOST_DOUBLE_TEX
265 #define READ_SPINOR_UP READ_SPINOR_DOUBLE_UP_TEX
266 #define READ_SPINOR_DOWN READ_SPINOR_DOUBLE_DOWN_TEX
267 #ifdef USE_TEXTURE_OBJECTS
268 #define SPINORTEX param.inTex
269 #define GHOSTSPINORTEX param.ghostTex
270 #else
271 #define SPINORTEX spinorTexDouble
272 #define GHOSTSPINORTEX ghostSpinorTexDouble
273 #endif // USE_TEXTURE_OBJECTS
274 #endif
275 #if (defined DIRECT_ACCESS_WILSON_INTER) || (defined FERMI_NO_DBLE_TEX)
276 #define READ_INTERMEDIATE_SPINOR READ_SPINOR_DOUBLE
277 #define INTERTEX param.out
278 #else
279 #define READ_INTERMEDIATE_SPINOR READ_SPINOR_DOUBLE_TEX
280 #ifdef USE_TEXTURE_OBJECTS
281 #define INTERTEX param.outTex
282 #else
283 #define INTERTEX interTexDouble
284 #endif
285 #endif
286 #define WRITE_SPINOR WRITE_SPINOR_DOUBLE2
287 #define SPINOR_DOUBLE
288 #if (DD_XPAY!=0)
289 #if (defined DIRECT_ACCESS_WILSON_ACCUM) || (defined FERMI_NO_DBLE_TEX)
290 #define ACCUMTEX param.x
291 #define READ_ACCUM READ_ACCUM_DOUBLE
292 #else
293 #ifdef USE_TEXTURE_OBJECTS
294 #define ACCUMTEX param.xTex
295 #else
296 #define ACCUMTEX accumTexDouble
297 #endif // USE_TEXTURE_OBJECTS
298 #define READ_ACCUM READ_ACCUM_DOUBLE_TEX
299 #endif
300 
301 #endif
302 
303 #define SPINOR_HOP 12
304 
305 #elif (DD_PREC==1) // single-precision fields
306 
307 #define TPROJSCALE param.tProjScale_f
308 
309 // single-precision gauge field
310 #ifdef DIRECT_ACCESS_LINK
311 #define GAUGE0TEX param.gauge0
312 #define GAUGE1TEX param.gauge1
313 #else
314 #ifdef USE_TEXTURE_OBJECTS
315 #define GAUGE0TEX param.gauge0Tex
316 #define GAUGE1TEX param.gauge1Tex
317 #else
318 #if (DD_RECON_F == 18)
319 #define GAUGE0TEX gauge0TexSingle2
320 #define GAUGE1TEX gauge1TexSingle2
321 #else
322 #define GAUGE0TEX gauge0TexSingle4
323 #define GAUGE1TEX gauge1TexSingle4
324 #endif
325 #endif // USE_TEXTURE_OBJECTS
326 #endif
327 
328 
329 // single-precision spinor fields
330 #ifdef DIRECT_ACCESS_WILSON_SPINOR
331 #define READ_SPINOR READ_SPINOR_SINGLE
332 #define READ_SPINOR_GHOST READ_SPINOR_GHOST_SINGLE
333 #define READ_SPINOR_UP READ_SPINOR_SINGLE_UP
334 #define READ_SPINOR_DOWN READ_SPINOR_SINGLE_DOWN
335 #define SPINORTEX param.in
336 #else
337 #define READ_SPINOR READ_SPINOR_SINGLE_TEX
338 #define READ_SPINOR_GHOST READ_SPINOR_GHOST_SINGLE_TEX
339 #define READ_SPINOR_UP READ_SPINOR_SINGLE_UP_TEX
340 #define READ_SPINOR_DOWN READ_SPINOR_SINGLE_DOWN_TEX
341 #ifdef USE_TEXTURE_OBJECTS
342 #define SPINORTEX param.inTex
343 #define GHOSTSPINORTEX param.ghostTex
344 #else
345 #define SPINORTEX spinorTexSingle
346 #define GHOSTSPINORTEX ghostSpinorTexSingle
347 #endif // USE_TEXTURE_OBJECTS
348 #endif
349 #ifdef DIRECT_ACCESS_WILSON_INTER
350 #define READ_INTERMEDIATE_SPINOR READ_SPINOR_SINGLE
351 #define INTERTEX param.out
352 #else
353 #define READ_INTERMEDIATE_SPINOR READ_SPINOR_SINGLE_TEX
354 #ifdef USE_TEXTURE_OBJECTS
355 #define INTERTEX param.outTex
356 #else
357 #define INTERTEX interTexSingle
358 #endif // USE_TEXTURE_OBJECTS
359 #endif
360 #define WRITE_SPINOR WRITE_SPINOR_FLOAT4
361 #if (DD_XPAY!=0)
362 #ifdef DIRECT_ACCESS_WILSON_ACCUM
363 #define ACCUMTEX param.x
364 #define READ_ACCUM READ_ACCUM_SINGLE
365 #else
366 #ifdef USE_TEXTURE_OBJECTS
367 #define ACCUMTEX param.xTex
368 #else
369 #define ACCUMTEX accumTexSingle
370 #endif // USE_TEXTURE_OBJECTS
371 #define READ_ACCUM READ_ACCUM_SINGLE_TEX
372 #endif
373 #endif
374 
375 #define SPINOR_HOP 6
376 
377 #else // half-precision fields
378 
379 #define TPROJSCALE param.tProjScale_f
380 
381 // half-precision gauge field
382 #ifdef DIRECT_ACCESS_LINK
383 #define GAUGE0TEX param.gauge0
384 #define GAUGE1TEX param.gauge1
385 #else
386 #ifdef USE_TEXTURE_OBJECTS
387 #define GAUGE0TEX param.gauge0Tex
388 #define GAUGE1TEX param.gauge1Tex
389 #else
390 #if (DD_RECON_F == 18)
391 #define GAUGE0TEX gauge0TexHalf2
392 #define GAUGE1TEX gauge1TexHalf2
393 #else
394 #define GAUGE0TEX gauge0TexHalf4
395 #define GAUGE1TEX gauge1TexHalf4
396 #endif
397 #endif // USE_TEXTURE_OBJECTS
398 #endif
399 
400 
401 // half-precision spinor fields
402 #ifdef DIRECT_ACCESS_WILSON_SPINOR
403 #define READ_SPINOR READ_SPINOR_HALF
404 #define READ_SPINOR_GHOST READ_SPINOR_GHOST_HALF
405 #define READ_SPINOR_UP READ_SPINOR_HALF_UP
406 #define READ_SPINOR_DOWN READ_SPINOR_HALF_DOWN
407 #define SPINORTEX param.in
408 #else
409 #define READ_SPINOR READ_SPINOR_HALF_TEX
410 #define READ_SPINOR_GHOST READ_SPINOR_GHOST_HALF_TEX
411 #define READ_SPINOR_UP READ_SPINOR_HALF_UP_TEX
412 #define READ_SPINOR_DOWN READ_SPINOR_HALF_DOWN_TEX
413 #ifdef USE_TEXTURE_OBJECTS
414 #define SPINORTEX param.inTex
415 #define GHOSTSPINORTEX param.ghostTex
416 #else
417 #define SPINORTEX spinorTexHalf
418 #define GHOSTSPINORTEX ghostSpinorTexHalf
419 #endif // USE_TEXTURE_OBJECTS
420 #endif
421 #ifdef DIRECT_ACCESS_WILSON_INTER
422 #define READ_INTERMEDIATE_SPINOR READ_SPINOR_HALF
423 #define INTERTEX param.out
424 #else
425 #define READ_INTERMEDIATE_SPINOR READ_SPINOR_HALF_TEX
426 #ifdef USE_TEXTURE_OBJECTS
427 #define INTERTEX param.outTex
428 #else
429 #define INTERTEX interTexHalf
430 #endif // USE_TEXTURE_OBJECTS
431 #endif
432 #define WRITE_SPINOR WRITE_SPINOR_SHORT4
433 #if (DD_XPAY!=0)
435 #ifdef DIRECT_ACCESS_WILSON_ACCUM
436 #define ACCUMTEX param.x
437 #define READ_ACCUM READ_ACCUM_HALF
438 #else
439 #ifdef USE_TEXTURE_OBJECTS
440 #define ACCUMTEX param.xTex
441 #else
442 #define ACCUMTEX accumTexHalf
443 #endif // USE_TEXTURE_OBJECTS
444 #define READ_ACCUM READ_ACCUM_HALF_TEX
445 #endif
446 
447 #endif
448 
449 #define SPINOR_HOP 6
450 
451 #endif
452 
453 #define DD_CONCAT(n,p,r,d,x) n ## p ## r ## d ## x ## Kernel
454 #define DD_FUNC(n,p,r,d,x) DD_CONCAT(n,p,r,d,x)
455 
456 // define the kernel
457 template <KernelType kernel_type>
458 __global__ void DD_FUNC(DD_NAME_F, DD_PREC_F, DD_RECON_F, DD_DAG_F, DD_XPAY_F)(const DslashParam param) {
459 
460 #ifdef GPU_TWISTED_CLOVER_DIRAC
461 
462 #ifdef SHARED_WILSON_DSLASH // Fermi optimal code
463 
464 #if DD_DAG
465 #include "tmc_dslash_dagger_fermi_core.h"
466 #else
467 #include "tmc_dslash_fermi_core.h"
468 #endif
469 
470 #else // no shared memory blocking
471 
472 #if DD_DAG
473 #include "tmc_dslash_dagger_gt200_core.h"
474 #else
475 #include "tmc_dslash_gt200_core.h"
476 #endif
477 
478 #endif
479 
480 }
481 
482 #ifdef MULTI_GPU
483 template <>
484 __global__ void DD_FUNC(DD_NAME_F, DD_PREC_F, DD_RECON_F, DD_DAG_F, DD_XPAY_F)<EXTERIOR_KERNEL_ALL>
485 (const DslashParam param) {
486 
487 #ifdef GPU_TWISTED_CLOVER_DIRAC
488 
489 #ifdef SHARED_WILSON_DSLASH // Fermi optimal code
490 
491 #if DD_DAG
492 #include "tmc_fused_exterior_dslash_dagger_fermi_core.h"
493 #else
494 #include "tmc_fused_exterior_dslash_fermi_core.h"
495 #endif
496 
497 #else // no shared memory blocking
498 
499 #if DD_DAG
500 #include "tmc_fused_exterior_dslash_dagger_gt200_core.h"
501 #else
502 #include "tmc_fused_exterior_dslash_gt200_core.h"
503 #endif
504 
505 #endif
506 
507 #endif
508 
509 }
510 #endif // MULTI_GPU
511 
512 
513 //NEW
514 #if (DD_XPAY==1) && (DD_TWIST==1)
515 #define CLOVER_TWIST_XPAY
516 
517 //redefine kernel name:
518 #undef DD_NAME_F
519 #define DD_NAME_F twistedCloverDslashTwist
520 
521 template <KernelType kernel_type>
522 __global__ void DD_FUNC(DD_NAME_F, DD_PREC_F, DD_RECON_F, DD_DAG_F, DD_XPAY_F)(const DslashParam param) {
523 
524 #ifdef GPU_TWISTED_CLOVER_DIRAC
525 
526 #ifdef SHARED_WILSON_DSLASH // Fermi optimal code
527 
528 #if DD_DAG
529 #include "tmc_dslash_dagger_fermi_core.h"
530 #else
531 #include "tmc_dslash_fermi_core.h"
532 #endif
533 
534 #else // no shared memory blocking
535 
536 #if DD_DAG
537 #include "tmc_dslash_dagger_gt200_core.h"
538 #else
539 #include "tmc_dslash_gt200_core.h"
540 #endif
541 
542 #endif
543 
544 #endif
545 
546 }
547 
548 #ifdef MULTI_GPU
549 template <>
550 __global__ void DD_FUNC(DD_NAME_F, DD_PREC_F, DD_RECON_F, DD_DAG_F, DD_XPAY_F)<EXTERIOR_KERNEL_ALL>
551 (const DslashParam param) {
552 
553 #ifdef GPU_TWISTED_CLOVER_DIRAC
554 
555 #ifdef SHARED_WILSON_DSLASH // Fermi optimal code
556 
557 #if DD_DAG
558 #include "tmc_fused_exterior_dslash_dagger_fermi_core.h"
559 #else
560 #include "tmc_fused_exterior_dslash_fermi_core.h"
561 #endif
562 
563 #else // no shared memory blocking
564 
565 #if DD_DAG
566 #include "tmc_fused_exterior_dslash_dagger_gt200_core.h"
567 #else
568 #include "tmc_fused_exterior_dslash_gt200_core.h"
569 #endif
570 
571 #endif
572 
573 #endif
574 
575 }
576 #endif // MULTI_GPU
577 
578 #undef CLOVER_TWIST_XPAY
579 #endif //(DD_XPAY==0) && (DD_TWIST==1)
580 
581 
582 //BEGIN DUMMY KERNEL (remove it later)
583 #if (DD_XPAY==0) && (DD_TWIST==1)
584 #define CLOVER_TWIST_XPAY
585 
586 //redefine kernel name:
587 #undef DD_NAME_F
588 #define DD_NAME_F twistedCloverDslashTwist
589 
590 template <KernelType kernel_type>
591 __global__ void DD_FUNC(DD_NAME_F, DD_PREC_F, DD_RECON_F, DD_DAG_F, DD_XPAY_F)
592 (DslashParam param) {
593 
594 #ifdef GPU_TWISTED_CLOVER_DIRAC
595 
596 #endif
597 
598 }
599 #undef CLOVER_TWIST_XPAY
600 #endif //(DD_XPAY==0) && (DD_TWIST==1)
601 //END DUMMY KERNEL
602 
603 #endif
604 
605 // clean up
606 
607 #undef DD_PREC_F
608 #undef DD_NAME_F
609 #undef DD_RECON_F
610 #undef DD_DAG_F
611 #undef DD_XPAY_F
612 #undef DD_CONCAT
613 #undef DD_FUNC
614 
615 #undef DSLASH_XPAY
616 
618 #undef CLOVER_TWIST_INV_DSLASH
619 #undef READ_GAUGE_MATRIX
621 #undef RECONSTRUCT_GAUGE_MATRIX
622 #undef GAUGE0TEX
623 #undef GAUGE1TEX
624 #undef READ_SPINOR
625 #undef READ_SPINOR_GHOST
626 #undef READ_SPINOR_UP
627 #undef READ_SPINOR_DOWN
628 #undef SPINORTEX
629 #undef GHOSTSPINORTEX
630 #undef READ_INTERMEDIATE_SPINOR
631 #undef INTERTEX
632 #undef READ_ACCUM
633 #undef ACCUMTEX
634 #undef WRITE_SPINOR
635 #undef GAUGE_FLOAT2
636 #undef SPINOR_DOUBLE
637 
638 #undef READ_CLOVER
639 #undef ASSN_CLOVER
640 #undef TMCLOVERTEX
641 #undef TMCLOVERTEXNORM
642 #ifndef DYNAMIC_CLOVER
643  #undef TM_INV_CLOVERTEX
644  #undef TM_INV_CLOVERTEXNORM
645 #endif
646 #undef CLOVER_DOUBLE
647 
648 #undef SPINOR_HOP
649 
650 #undef TPROJSCALE
651 
652 // prepare next set of options, or clean up after final iteration
653 
654 #if (DD_DAG==0)
655 #undef DD_DAG
656 #define DD_DAG 1
657 #else
658 #undef DD_DAG
659 #define DD_DAG 0
660 
661 #if (DD_TWIST==0)
662 #undef DD_TWIST
663 #define DD_TWIST 1
664 #else
665 #undef DD_TWIST
666 #define DD_TWIST 0
667 
668 #if (DD_XPAY==0)
669 #undef DD_XPAY
670 #define DD_XPAY 1
671 #else
672 #undef DD_XPAY
673 #define DD_XPAY 0
674 
675 #if (DD_RECON==0)
676 #undef DD_RECON
677 #define DD_RECON 1
678 #elif (DD_RECON==1)
679 #undef DD_RECON
680 #define DD_RECON 2
681 #else
682 #undef DD_RECON
683 #define DD_RECON 0
684 
685 #if (DD_PREC==0)
686 #undef DD_PREC
687 #define DD_PREC 1
688 #elif (DD_PREC==1)
689 #undef DD_PREC
690 #define DD_PREC 2
691 
692 #else
693 
694 #undef DD_LOOP
695 #undef DD_DAG
696 #undef DD_TWIST
697 #undef DD_XPAY
698 #undef DD_RECON
699 #undef DD_PREC
700 
701 #endif // DD_PREC
702 #endif // DD_RECON
703 #endif // DD_XPAY
704 #endif // DD_TWIST
705 #endif // DD_DAG
706 
707 #ifdef DD_LOOP
708 #include "tmc_dslash_def.h"
709 #endif
DD_DAG_F
#define DD_DAG_F
Definition: tmc_dslash_def.h:45
tmc_dslash_def.h
tmc_dslash_dagger_fermi_core.h
DD_RECON_F
#define DD_RECON_F
Definition: tmc_dslash_def.h:153
tmc_dslash_fermi_core.h
param
QudaGaugeParam param
Definition: pack_test.cpp:17
tmc_dslash_dagger_gt200_core.h
DD_PREC_F
#define DD_PREC_F
Definition: tmc_dslash_def.h:59
tmc_fused_exterior_dslash_fermi_core.h
EXTERIOR_KERNEL_ALL
Definition: dslash_constants.h:7
tmc_dslash_gt200_core.h
DD_FUNC
#define DD_FUNC(n, p, r, d, x)
Definition: tmc_dslash_def.h:454
tmc_fused_exterior_dslash_dagger_fermi_core.h
tmc_fused_exterior_dslash_gt200_core.h
DD_XPAY_F
#define DD_XPAY_F
Definition: tmc_dslash_def.h:52
DD_NAME_F
#define DD_NAME_F
Definition: tmc_dslash_def.h:67
tmc_fused_exterior_dslash_dagger_gt200_core.h
Generated by   doxygen 1.8.14