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
clover_def.h
Go to the documentation of this file.
1
// clover_def.h - clover 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
#define DD_XPAY 0
10
#define DD_PREC 0
11
#endif
12
13
// set options for current iteration
14
15
#if (DD_XPAY==0) // no xpay
16
#define DD_XPAY_F
17
#define DD_PARAM4 DslashParam param
18
#else // xpay
19
#define DD_XPAY_F Xpay
20
#if (DD_PREC == 0)
21
#define DD_PARAM4 DslashParam param, double a
22
#else
23
#define DD_PARAM4 DslashParam param, float a
24
#endif
25
#define DSLASH_XPAY
26
#endif
27
28
#if (DD_PREC==0) // double-precision spinor field
29
#define DD_PREC_F D
30
#define DD_PARAM1 double2* out, float *null1
31
#define DD_PARAM3 const double2* in, const float *null3
32
#if (defined DIRECT_ACCESS_WILSON_SPINOR) || (defined FERMI_NO_DBLE_TEX)
33
#define READ_SPINOR READ_SPINOR_DOUBLE
34
#define SPINORTEX in
35
#else
36
#define READ_SPINOR READ_SPINOR_DOUBLE_TEX
37
#ifdef USE_TEXTURE_OBJECTS
38
#define SPINORTEX param.inTex
39
#else
40
#define SPINORTEX spinorTexDouble
41
#endif // USE_TEXTURE_OBJECTS
42
#endif
43
#if (DD_XPAY==1) // never used
44
#define ACCUMTEX accumTexDouble
45
#define READ_ACCUM READ_ACCUM_DOUBLE
46
#endif
47
#define SPINOR_DOUBLE
48
#define WRITE_SPINOR WRITE_SPINOR_DOUBLE2
49
#elif (DD_PREC==1) // single-precision spinor field
50
#define DD_PREC_F S
51
#define DD_PARAM1 float4* out, float *null1
52
#define DD_PARAM3 const float4* in, const float *null3
53
#ifdef DIRECT_ACCESS_WILSON_SPINOR
54
#define READ_SPINOR READ_SPINOR_SINGLE
55
#define SPINORTEX in
56
#else
57
#define READ_SPINOR READ_SPINOR_SINGLE_TEX
58
#ifdef USE_TEXTURE_OBJECTS
59
#define SPINORTEX param.inTex
60
#else
61
#define SPINORTEX spinorTexSingle
62
#endif // USE_TEXTURE_OBJECTS
63
#endif
64
#define WRITE_SPINOR WRITE_SPINOR_FLOAT4
65
#if (DD_XPAY==1)
66
#define ACCUMTEX accumTexSingle
67
#define READ_ACCUM READ_ACCUM_SINGLE
68
#endif
69
#else // half-precision spinor field
70
#define DD_PREC_F H
71
#ifdef DIRECT_ACCESS_WILSON_SPINOR
72
#define READ_SPINOR READ_SPINOR_HALF
73
#define SPINORTEX in
74
#else
75
#define READ_SPINOR READ_SPINOR_HALF_TEX
76
#ifdef USE_TEXTURE_OBJECTS
77
#define SPINORTEX param.inTex
78
#else
79
#define SPINORTEX spinorTexHalf
80
#endif // USE_TEXTURE_OBJECTS
81
#endif
82
#define DD_PARAM1 short4* out, float *outNorm
83
#define DD_PARAM3 const short4* in, const float *inNorm
84
#define WRITE_SPINOR WRITE_SPINOR_SHORT4
85
#if (DD_XPAY==1)
86
#define ACCUMTEX accumTexHalf
87
#define READ_ACCUM READ_ACCUM_HALF
88
#endif
89
#endif
90
91
#if (DD_PREC==0) // double-precision clover term
92
#define DD_PREC_F D
93
#define DD_PARAM2 const double2* clover, const float *null
94
#if (defined DIRECT_ACCESS_CLOVER) || (defined FERMI_NO_DBLE_TEX)
95
#define CLOVERTEX clover
96
#define READ_CLOVER READ_CLOVER_DOUBLE
97
#else
98
#ifdef USE_TEXTURE_OBJECTS
99
#define CLOVERTEX (param.cloverTex)
100
#else
101
#define CLOVERTEX cloverTexDouble
102
#endif
103
#define READ_CLOVER READ_CLOVER_DOUBLE_TEX
104
#endif
105
106
#define CLOVER_DOUBLE
107
#elif (DD_PREC==1) // single-precision clover term
108
#define DD_PREC_F S
109
#define DD_PARAM2 const float4* clover, const float *null
110
#ifdef DIRECT_ACCESS_CLOVER
111
#define CLOVERTEX clover
112
#define READ_CLOVER READ_CLOVER_SINGLE
113
#else
114
#ifdef USE_TEXTURE_OBJECTS
115
#define CLOVERTEX (param.cloverTex)
116
#else
117
#define CLOVERTEX cloverTexSingle
118
#endif
119
#define READ_CLOVER READ_CLOVER_SINGLE_TEX
120
#endif
121
#else // half-precision clover term
122
#define DD_PREC_F H
123
#define DD_PARAM2 const short4* clover, const float *cloverNorm
124
#ifdef DIRECT_ACCESS_CLOVER
125
#define CLOVERTEX clover
126
#define READ_CLOVER READ_CLOVER_HALF
127
#else
128
#ifdef USE_TEXTURE_OBJECTS
129
#define CLOVERTEX (param.cloverTex)
130
#define CLOVERTEXNORM (param.cloverNormTex)
131
#else
132
#define CLOVERTEX cloverTexHalf
133
#define CLOVERTEXNORM cloverTexNorm
134
#endif
135
#define READ_CLOVER READ_CLOVER_HALF_TEX
136
#endif
137
#endif
138
139
//#define DD_CONCAT(s,c,x) clover ## s ## c ## x ## Kernel
140
#define DD_CONCAT(x) clover ## x ## Kernel
141
#define DD_FUNC(x) DD_CONCAT(x)
142
143
// define the kernel
144
#if !(__COMPUTE_CAPABILITY__ < 130 && DD_PREC == 0)
145
146
__global__
void
DD_FUNC
(
DD_XPAY_F
)(
DD_PARAM1
,
DD_PARAM2
,
DD_PARAM3
,
DD_PARAM4
) {
147
148
#ifdef GPU_CLOVER_DIRAC
149
#include "
clover_core.h
"
150
#endif
151
152
}
153
154
#endif
155
156
// clean up
157
158
#undef DD_PREC_F
159
#undef DD_XPAY_F
160
#undef DD_PARAM1
161
#undef DD_PARAM2
162
#undef DD_PARAM3
163
#undef DD_PARAM4
164
#undef DD_CONCAT
165
#undef DD_FUNC
166
167
#undef DSLASH_XPAY
168
#undef READ_SPINOR
169
#undef SPINORTEX
170
#undef WRITE_SPINOR
171
#undef ACCUMTEX
172
#undef READ_ACCUM
173
#undef CLOVERTEX
174
#undef READ_CLOVER
175
#undef GAUGE_DOUBLE
176
#undef SPINOR_DOUBLE
177
#undef CLOVER_DOUBLE
178
179
// prepare next set of options, or clean up after final iteration
180
181
//#if (DD_XPAY==0) // xpay variant is not needed
182
//#undef DD_XPAY
183
//#define DD_XPAY 1
184
//#else
185
//#undef DD_XPAY
186
//#define DD_XPAY 0
187
188
#if (DD_PREC==0)
189
#undef DD_PREC
190
#define DD_PREC 1
191
#elif (DD_PREC==1)
192
#undef DD_PREC
193
#define DD_PREC 2
194
#else
195
196
#undef DD_LOOP
197
#undef DD_XPAY
198
#undef DD_PREC
199
200
#endif // DD_PREC
201
//#endif // DD_XPAY
202
203
#ifdef DD_LOOP
204
#include "
clover_def.h
"
205
#endif
DD_PARAM1
#define DD_PARAM1
Definition:
clover_def.h:30
DD_XPAY_F
#define DD_XPAY_F
Definition:
clover_def.h:16
clover_def.h
DD_FUNC
#define DD_FUNC(x)
Definition:
clover_def.h:141
DD_PARAM4
#define DD_PARAM4
Definition:
clover_def.h:17
DD_PARAM3
#define DD_PARAM3
Definition:
clover_def.h:31
DD_PARAM2
#define DD_PARAM2
Definition:
clover_def.h:93
clover_core.h
Generated on Wed Feb 4 2015 17:00:05 for QUDA by
1.8.6
1.8.6