QUDA: quda/lib/staggered_dslash

QUDA v0.4.0
A library for QCD on GPUs
00001 // staggered_dslash_def.h - staggered Dslash kernel definitions
00002 //
00003 // See comments in wilson_dslash_def.h
00004 
00005 // initialize on first iteration
00006 
00007 #ifndef DD_LOOP
00008 #define DD_LOOP
00009 #define DD_AXPY 0
00010 #define DD_RECON 0
00011 #define DD_PREC 0
00012 #endif
00013 
00014 // set options for current iteration
00015 
00016 #define DD_FNAME staggeredDslash
00017 
00018 #if (DD_AXPY==0) // no axpy
00019 #define DD_AXPY_F 
00020 #else            // axpy
00021 #define DD_AXPY_F Axpy
00022 #define DSLASH_AXPY
00023 #endif
00024 
00025 #if (DD_PREC == 0)
00026 #define DD_PARAM_AXPY const double2 *x, const float *xNorm, const double a, const DslashParam param
00027 #elif (DD_PREC == 1) 
00028 #define DD_PARAM_AXPY const float2 *x, const float *xNorm, const float a, const DslashParam param
00029 #else
00030 #define DD_PARAM_AXPY const short2 *x, const float *xNorm, const float a, const DslashParam param
00031 #endif
00032 
00033 
00034 #if (DD_RECON==0) // reconstruct from 8 reals
00035 #define DD_RECON_F 8
00036 
00037 #if (DD_PREC==0) // DOUBLE PRECISION
00038 #define DD_PARAM_GAUGE const double2 *fatGauge0, const double2 *fatGauge1, const double2* longGauge0, const double2* longGauge1
00039 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_GAUGE_MATRIX_8_DOUBLE
00040 
00041 #ifdef DIRECT_ACCESS_FAT_LINK
00042 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2(FAT, gauge, dir, idx, fat_ga_stride)
00043 #else
00044 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2_TEX(FAT, gauge, dir, idx, fat_ga_stride)
00045 #endif // DIRECT_ACCESS_FAT_LINK
00046 #ifdef DIRECT_ACCESS_LONG_LINK
00047 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_8_DOUBLE2(LONG, gauge, dir, idx, long_ga_stride)
00048 #else
00049 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_8_DOUBLE2_TEX(LONG, gauge, dir, idx, long_ga_stride)
00050 #endif // DIRECT_ACCESS_LONG_LINK
00051 
00052 #elif (DD_PREC==1) // SINGLE PRECISION
00053 #define DD_PARAM_GAUGE const float2 *fatGauge0, const float2 *fatGauge1, const float4* longGauge0, const float4* longGauge1
00054 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_GAUGE_MATRIX_8_SINGLE
00055 
00056 #ifdef DIRECT_ACCESS_FAT_LINK
00057 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2(FAT, gauge, dir, idx, fat_ga_stride)
00058 #else
00059 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2_TEX(FAT, gauge, dir, idx, fat_ga_stride)
00060 #endif // DIRECT_ACCESS_FAT_LINK
00061 #ifdef DIRECT_ACCESS_LONG_LINK
00062 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_8_FLOAT4(LONG, gauge, dir, idx, long_ga_stride)
00063 #else
00064 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_8_FLOAT4_TEX(LONG, gauge, dir, idx, long_ga_stride)
00065 #endif // DIRECT_ACCESS_LONG_LINK
00066 
00067 #else // HALF PRECISION
00068 #define DD_PARAM_GAUGE const short2 *fatGauge0, const short2* fatGauge1, const short4* longGauge0, const short4* longGauge1
00069 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_GAUGE_MATRIX_8_SINGLE
00070 
00071 /*#ifdef DIRECT_ACCESS_FAT_LINK
00072 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2(FAT, gauge, dir, idx, fat_ga_stride); RESCALE2(FAT, fat_ga_max);
00073 #else*/
00074 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2_TEX(FAT, gauge, dir, idx, fat_ga_stride); RESCALE2(FAT, fat_ga_max);
00075 /*#endif // DIRECT_ACCESS_FAT_LINK
00076 #ifdef DIRECT_ACCESS_LONG_LINK
00077 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_8_SHORT4(LONG, gauge, dir, idx, long_ga_stride)
00078 #else*/
00079 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_8_SHORT4_TEX(LONG, gauge, dir, idx, long_ga_stride)
00080 //#endif // DIRECT_ACCESS_LONG_LINK
00081 
00082 #endif // DD_PREC
00083 
00084 #elif (DD_RECON ==1)// reconstruct from 12 reals
00085 
00086 #define DD_RECON_F 12
00087 
00088 #if (DD_PREC==0) // DOUBLE PRECISION
00089 #define DD_PARAM_GAUGE const double2 *fatGauge0, const double2 *fatGauge1,  const double2* longGauge0, const double2* longGauge1
00090 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_GAUGE_MATRIX_12_DOUBLE
00091 
00092 #ifdef DIRECT_ACCESS_FAT_LINK
00093 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2(FAT, gauge, dir, idx, fat_ga_stride)
00094 #else
00095 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2_TEX(FAT, gauge, dir, idx, fat_ga_stride)
00096 #endif // DIRECT_ACCESS_FAT_LINK
00097 #ifdef DIRECT_ACCESS_LONG_LINK
00098 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_12_DOUBLE2(LONG, gauge, dir, idx, long_ga_stride)
00099 #else
00100 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_12_DOUBLE2_TEX(LONG, gauge, dir, idx, long_ga_stride)
00101 #endif // DIRECT_ACCESS_LONG_LINK
00102 
00103 #elif (DD_PREC==1) // SINGLE PRECISION
00104 #define DD_PARAM_GAUGE const float2 *fatGauge0, const float2 *fatGauge1, const float4* longGauge0, const float4* longGauge1
00105 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_GAUGE_MATRIX_12_SINGLE
00106 
00107 #ifdef DIRECT_ACCESS_FAT_LINK
00108 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2(FAT, gauge, dir, idx, fat_ga_stride)
00109 #else
00110 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2_TEX(FAT, gauge, dir, idx, fat_ga_stride)
00111 #endif // DIRECT_ACCESS_FAT_LINK
00112 #ifdef DIRECT_ACCESS_LONG_LINK
00113 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_12_FLOAT4(LONG, gauge, dir, idx, long_ga_stride)
00114 #else
00115 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_12_FLOAT4_TEX(LONG, gauge, dir, idx, long_ga_stride)
00116 #endif // DIRECT_ACCESS_LONG_LINK
00117 
00118 #else // HALF PRECISION
00119 #define DD_PARAM_GAUGE const short2 *fatGauge0, const short2 *fatGauge1, const short4* longGauge0, const short4* longGauge1
00120 #define RECONSTRUCT_GAUGE_MATRIX RECONSTRUCT_GAUGE_MATRIX_12_SINGLE
00121 
00122 /*#ifdef DIRECT_ACCESS_FAT_LINK
00123 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2(FAT, gauge, dir, idx, fat_ga_stride); RESCALE2(FAT, fat_ga_max);
00124 #else*/
00125 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2_TEX(FAT, gauge, dir, idx, fat_ga_stride); RESCALE2(FAT, fat_ga_max);
00126 /*#endif // DIRECT_ACCCESS_FAT_LINK
00127 #ifdef DIRECT_ACCESS_LONG_LINK
00128 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_12_SHORT4(LONG, gauge, dir, idx, long_ga_stride)
00129 #else*/
00130 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_12_SHORT4_TEX(LONG, gauge, dir, idx, long_ga_stride)
00131                                                                         //#endif // DIRECT_ACCCESS_LONG_LINK
00132 
00133 #endif // DD_PREC
00134 
00135 #else //18 reconstruct
00136 #define DD_RECON_F 18
00137 #define RECONSTRUCT_GAUGE_MATRIX(dir, gauge, idx, sign)
00138 
00139 #if (DD_PREC==0) // DOUBLE PRECISION
00140 #define DD_PARAM_GAUGE const double2 *fatGauge0, const double2 *fatGauge1,  const double2* longGauge0, const double2* longGauge1
00141 
00142 #ifdef DIRECT_ACCESS_FAT_LINK
00143 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2(FAT, gauge, dir, idx, fat_ga_stride)
00144 #else
00145 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2_TEX(FAT, gauge, dir, idx, fat_ga_stride)
00146 #endif // DIRECT_ACCCESS_FAT_LINK
00147 #ifdef DIRECT_ACCESS_LONG_LINK
00148 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2(LONG, gauge, dir, idx, long_ga_stride)
00149 #else
00150 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_DOUBLE2_TEX(LONG, gauge, dir, idx, long_ga_stride)
00151 #endif // DIRECT_ACCCESS_LONG_LINK
00152 
00153 #elif (DD_PREC==1) // SINGLE PRECISION
00154 
00155 #define DD_PARAM_GAUGE const float2 *fatGauge0, const float2 *fatGauge1, const float4* longGauge0, const float4* longGauge1
00156 
00157 #ifdef DIRECT_ACCESS_FAT_LINK
00158 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2(FAT, gauge, dir, idx, fat_ga_stride)
00159 #else
00160 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2_TEX(FAT, gauge, dir, idx, fat_ga_stride)
00161 #endif // DIRECT_ACCCESS_FAT_LINK
00162 #ifdef DIRECT_ACCESS_LONG_LINK
00163 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2(LONG, gauge, dir, idx, long_ga_stride)
00164 #else
00165 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_FLOAT2_TEX(LONG, gauge, dir, idx, long_ga_stride)
00166 #endif // DIRECT_ACCCESS_LONG_LINK
00167 
00168 #else  // HALF PRECISION
00169 
00170 #define DD_PARAM_GAUGE const short2 *fatGauge0, const short2 *fatGauge1, const short4* longGauge0, const short4* longGauge1
00171 
00172 /*#ifdef DIRECT_ACCESS_FAT_LINK
00173 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2(FAT, gauge, dir, idx, fat_ga_stride); RESCALE2(FAT, fat_ga_max);
00174 #else*/
00175 #define READ_FAT_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2_TEX(FAT, gauge, dir, idx, fat_ga_stride); RESCALE2(FAT, fat_ga_max);
00176                                                                          /*#endif // DIRECT_ACCESS_FAT_LINK
00177 #ifdef DIRECT_ACCESS_LONG_LINK
00178 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2(LONG, gauge, dir, idx, long_ga_stride)
00179 #else*/
00180 #define READ_LONG_MATRIX(gauge, dir, idx) READ_GAUGE_MATRIX_18_SHORT2_TEX(LONG, gauge, dir, idx, long_ga_stride)
00181                                                                          //#endif // DIRECT_ACCCESS_LONG_LINK
00182 
00183 #endif // DD_PREC
00184 
00185 #endif // DD_RECON
00186 
00187 #if (DD_PREC==0) // double-precision fields
00188 
00189 // gauge field
00190 #define DD_PREC_F D
00191 #if (defined DIRECT_ACCESS_FAT_LINK) || (defined FERMI_NO_DBLE_TEX)
00192 #define FATLINK0TEX fatGauge0
00193 #define FATLINK1TEX fatGauge1
00194 #else
00195 #define FATLINK0TEX fatGauge0TexDouble
00196 #define FATLINK1TEX fatGauge1TexDouble
00197 #endif
00198 
00199 #if (defined DIRECT_ACCESS_LONG_LINK) || (defined FERMI_NO_DBLE_TEX)
00200 #define LONGLINK0TEX longGauge0
00201 #define LONGLINK1TEX longGauge1
00202 #else
00203 #define LONGLINK0TEX longGauge0TexDouble
00204 #define LONGLINK1TEX longGauge1TexDouble
00205 #endif
00206 
00207 #define GAUGE_DOUBLE
00208 
00209 // spinor fields
00210 #define DD_PARAM_OUT double2* out, float *null1
00211 #define DD_PARAM_IN const double2* in, const float *null4
00212 #if (defined DIRECT_ACCESS_SPINOR) || (defined FERMI_NO_DBLE_TEX)
00213 #define SPINORTEX in
00214 #define READ_1ST_NBR_SPINOR READ_1ST_NBR_SPINOR_DOUBLE
00215 #define READ_3RD_NBR_SPINOR READ_3RD_NBR_SPINOR_DOUBLE
00216 #else
00217 #define SPINORTEX spinorTexDouble
00218 #define READ_1ST_NBR_SPINOR READ_1ST_NBR_SPINOR_DOUBLE_TEX
00219 #define READ_3RD_NBR_SPINOR READ_3RD_NBR_SPINOR_DOUBLE_TEX
00220 #endif
00221 #if (defined DIRECT_ACCESS_INTER) || (defined FERMI_NO_DBLE_TEX)
00222 #define READ_AND_SUM_SPINOR READ_AND_SUM_ST_SPINOR
00223 #define INTERTEX out
00224 #else
00225 #define READ_AND_SUM_SPINOR READ_AND_SUM_ST_SPINOR_DOUBLE_TEX
00226 #define INTERTEX interTexDouble
00227 #endif
00228 #define WRITE_SPINOR WRITE_ST_SPINOR_DOUBLE2
00229 #define SPINOR_DOUBLE
00230 #if (DD_AXPY==1)
00231 #if (defined DIRECT_ACCESS_ACCUM) || (defined FERMI_NO_DBLE_TEX)
00232 #define ACCUMTEX x
00233 #define READ_ACCUM READ_ST_ACCUM_DOUBLE
00234 #else
00235 #define ACCUMTEX accumTexDouble
00236 #define READ_ACCUM READ_ST_ACCUM_DOUBLE_TEX
00237 #endif
00238 #endif // DD_AXPY
00239 
00240 
00241 #elif (DD_PREC==1) // single-precision fields
00242 
00243 // gauge fields
00244 #define DD_PREC_F S
00245 
00246 #ifndef DIRECT_ACCESS_FAT_LINK
00247 #define FATLINK0TEX fatGauge0TexSingle
00248 #define FATLINK1TEX fatGauge1TexSingle
00249 #else
00250 #define FATLINK0TEX fatGauge0
00251 #define FATLINK1TEX fatGauge1
00252 #endif
00253 
00254 #ifndef DIRECT_ACCESS_LONG_LINK //longlink access
00255 #if (DD_RECON ==2)
00256 #define LONGLINK0TEX longGauge0TexSingle_norecon
00257 #define LONGLINK1TEX longGauge1TexSingle_norecon
00258 #else
00259 #define LONGLINK0TEX longGauge0TexSingle
00260 #define LONGLINK1TEX longGauge1TexSingle
00261 #endif
00262 #else
00263 #define LONGLINK0TEX longGauge0
00264 #define LONGLINK1TEX longGauge1
00265 #endif
00266 
00267 // spinor fields
00268 #define DD_PARAM_OUT float2* out, float *null1
00269 #define DD_PARAM_IN const float2* in, const float *null4
00270 #ifndef DIRECT_ACCESS_SPINOR
00271 #define SPINORTEX spinorTexSingle2
00272 #define READ_1ST_NBR_SPINOR READ_1ST_NBR_SPINOR_SINGLE_TEX
00273 #define READ_3RD_NBR_SPINOR READ_3RD_NBR_SPINOR_SINGLE_TEX
00274 #else
00275 #define SPINORTEX in
00276 #define READ_1ST_NBR_SPINOR READ_1ST_NBR_SPINOR_SINGLE
00277 #define READ_3RD_NBR_SPINOR READ_3RD_NBR_SPINOR_SINGLE
00278 #endif
00279 #if (defined DIRECT_ACCESS_INTER)
00280 #define READ_AND_SUM_SPINOR READ_AND_SUM_ST_SPINOR
00281 #define INTERTEX out
00282 #else
00283 #define READ_AND_SUM_SPINOR READ_AND_SUM_ST_SPINOR_SINGLE_TEX
00284 #define INTERTEX interTexSingle2
00285 #endif
00286 #define WRITE_SPINOR WRITE_ST_SPINOR_FLOAT2
00287 #if (DD_AXPY==1)
00288 #if (defined DIRECT_ACCESS_ACCUM)
00289 #define ACCUMTEX x
00290 #define READ_ACCUM READ_ST_ACCUM_SINGLE
00291 #else
00292 #define ACCUMTEX accumTexSingle2
00293 #define READ_ACCUM READ_ST_ACCUM_SINGLE_TEX
00294 #endif
00295 #endif // DD_AXPY
00296 
00297 
00298 #else             // half-precision fields
00299 
00300 // all reads done through texture cache regardless
00301 
00302 // gauge fields
00303 #define DD_PREC_F H
00304 #define FATLINK0TEX fatGauge0TexHalf
00305 #define FATLINK1TEX fatGauge1TexHalf
00306 #if (DD_RECON ==2)
00307 #define LONGLINK0TEX longGauge0TexHalf_norecon
00308 #define LONGLINK1TEX longGauge1TexHalf_norecon
00309 #else
00310 #define LONGLINK0TEX longGauge0TexHalf
00311 #define LONGLINK1TEX longGauge1TexHalf
00312 #endif
00313 
00314 #define READ_1ST_NBR_SPINOR READ_1ST_NBR_SPINOR_HALF_TEX
00315 #define READ_3RD_NBR_SPINOR READ_3RD_NBR_SPINOR_HALF_TEX
00316 #define SPINORTEX spinorTexHalf2
00317 #define DD_PARAM_OUT short2* out, float *outNorm
00318 #define DD_PARAM_IN const short2* in, const float *inNorm
00319 #if (defined DIRECT_ACCESS_INTER)
00320 #define READ_AND_SUM_SPINOR READ_AND_SUM_ST_SPINOR_HALF
00321 #define INTERTEX out
00322 #else
00323 #define READ_AND_SUM_SPINOR READ_AND_SUM_ST_SPINOR_HALF_TEX
00324 #define INTERTEX interTexHalf2
00325 #endif
00326 #define WRITE_SPINOR WRITE_ST_SPINOR_SHORT2
00327 #if (DD_AXPY==1)
00328 #define ACCUMTEX accumTexHalf2
00329 #define READ_ACCUM READ_ST_ACCUM_HALF_TEX
00330 #endif // DD_AXPY
00331 
00332 #endif
00333 
00334 // only build double precision if supported
00335 #if !(__COMPUTE_CAPABILITY__ < 130 && DD_PREC == 0) 
00336 
00337 #define DD_CONCAT(n,r,x) n ## r ## x ## Kernel
00338 #define DD_FUNC(n,r,x) DD_CONCAT(n,r,x)
00339 
00340 // define the kernel
00341 
00342 template <KernelType kernel_type>
00343 __global__ void DD_FUNC(DD_FNAME, DD_RECON_F, DD_AXPY_F)
00344   (DD_PARAM_OUT, DD_PARAM_GAUGE, DD_PARAM_IN, DD_PARAM_AXPY) {
00345 #ifdef GPU_STAGGERED_DIRAC
00346   #include "staggered_dslash_core.h"
00347 #endif
00348 }
00349 
00350 #endif // !(__COMPUTE_CAPABILITY__ < 130 && DD_PREC == 0)
00351 
00352 
00353 // clean up
00354 
00355 #undef DD_PREC_F
00356 #undef DD_RECON_F
00357 #undef DD_AXPY_F
00358 #undef DD_PARAM_OUT
00359 #undef DD_PARAM_GAUGE
00360 #undef DD_PARAM_IN
00361 #undef DD_PARAM_AXPY
00362 #undef DD_FNAME
00363 #undef DD_CONCAT
00364 #undef DD_FUNC
00365 
00366 #undef DSLASH_AXPY
00367 #undef READ_GAUGE_MATRIX
00368 #undef RECONSTRUCT_GAUGE_MATRIX
00369 #undef FATLINK0TEX
00370 #undef FATLINK1TEX
00371 #undef LONGLINK0TEX
00372 #undef LONGLINK1TEX
00373 #undef SPINORTEX
00374 #undef WRITE_SPINOR
00375 #undef READ_AND_SUM_SPINOR
00376 #undef INTERTEX
00377 #undef ACCUMTEX
00378 #undef READ_ACCUM
00379 #undef CLOVERTEX
00380 #undef READ_CLOVER
00381 #undef DSLASH_CLOVER
00382 #undef GAUGE_DOUBLE
00383 #undef SPINOR_DOUBLE
00384 #undef CLOVER_DOUBLE
00385 #undef READ_FAT_MATRIX
00386 #undef READ_LONG_MATRIX
00387 #undef READ_1ST_NBR_SPINOR
00388 #undef READ_3RD_NBR_SPINOR
00389 
00390 
00391 // prepare next set of options, or clean up after final iteration
00392 
00393 #if (DD_AXPY==0)
00394 #undef DD_AXPY
00395 #define DD_AXPY 1
00396 #else
00397 #undef DD_AXPY
00398 #define DD_AXPY 0
00399 
00400 #if (DD_RECON==0)
00401 #undef DD_RECON
00402 #define DD_RECON 1
00403 #elif (DD_RECON ==1)
00404 #undef DD_RECON
00405 #define DD_RECON 2
00406 #else
00407 #undef DD_RECON
00408 #define DD_RECON 0
00409 
00410 #if (DD_PREC==0)
00411 #undef DD_PREC
00412 #define DD_PREC 1
00413 #elif (DD_PREC==1)
00414 #undef DD_PREC
00415 #define DD_PREC 2
00416 #else
00417 #undef DD_PREC
00418 #define DD_PREC 0
00419 
00420 #undef DD_LOOP
00421 #undef DD_AXPY
00422 #undef DD_RECON
00423 #undef DD_PREC
00424 
00425 #endif // DD_PREC
00426 #endif // DD_RECON
00427 #endif // DD_AXPY
00428 
00429 #ifdef DD_LOOP
00430 #include "staggered_dslash_def.h"
00431 #endif