19 complex<Float> temp[16];
20 complex<Float> *h_result = (complex<Float> *)(Float *)(h_result_);
21 complex<Float> I(0.0, 1.0);
23 for (
int site = 0; site <
V; site++) {
34 temp[G_idx] += h_result[idx + 4 * 0 + 0];
35 temp[G_idx] += h_result[idx + 4 * 1 + 1];
36 temp[G_idx] += h_result[idx + 4 * 2 + 2];
37 temp[G_idx] += h_result[idx + 4 * 3 + 3];
42 temp[G_idx] += I * h_result[idx + 4 * 0 + 3];
43 temp[G_idx] += I * h_result[idx + 4 * 1 + 2];
44 temp[G_idx] -= I * h_result[idx + 4 * 2 + 1];
45 temp[G_idx] -= I * h_result[idx + 4 * 3 + 0];
50 temp[G_idx] -= h_result[idx + 4 * 0 + 3];
51 temp[G_idx] += h_result[idx + 4 * 1 + 2];
52 temp[G_idx] += h_result[idx + 4 * 2 + 1];
53 temp[G_idx] -= h_result[idx + 4 * 3 + 0];
58 temp[G_idx] += I * h_result[idx + 4 * 0 + 2];
59 temp[G_idx] -= I * h_result[idx + 4 * 1 + 3];
60 temp[G_idx] -= I * h_result[idx + 4 * 2 + 0];
61 temp[G_idx] += I * h_result[idx + 4 * 3 + 1];
66 temp[G_idx] += h_result[idx + 4 * 0 + 2];
67 temp[G_idx] += h_result[idx + 4 * 1 + 3];
68 temp[G_idx] += h_result[idx + 4 * 2 + 0];
69 temp[G_idx] += h_result[idx + 4 * 3 + 1];
74 temp[G_idx] += h_result[idx + 4 * 0 + 0];
75 temp[G_idx] += h_result[idx + 4 * 1 + 1];
76 temp[G_idx] -= h_result[idx + 4 * 2 + 2];
77 temp[G_idx] -= h_result[idx + 4 * 3 + 3];
82 temp[G_idx] += I * h_result[idx + 4 * 0 + 3];
83 temp[G_idx] += I * h_result[idx + 4 * 1 + 2];
84 temp[G_idx] += I * h_result[idx + 4 * 2 + 1];
85 temp[G_idx] += I * h_result[idx + 4 * 3 + 0];
90 temp[G_idx] -= h_result[idx + 4 * 0 + 3];
91 temp[G_idx] += h_result[idx + 4 * 1 + 2];
92 temp[G_idx] -= h_result[idx + 4 * 2 + 1];
93 temp[G_idx] += h_result[idx + 4 * 3 + 0];
98 temp[G_idx] += I * h_result[idx + 4 * 0 + 2];
99 temp[G_idx] -= I * h_result[idx + 4 * 1 + 3];
100 temp[G_idx] += I * h_result[idx + 4 * 2 + 0];
101 temp[G_idx] -= I * h_result[idx + 4 * 3 + 1];
106 temp[G_idx] += h_result[idx + 4 * 0 + 2];
107 temp[G_idx] += h_result[idx + 4 * 1 + 3];
108 temp[G_idx] -= h_result[idx + 4 * 2 + 0];
109 temp[G_idx] -= h_result[idx + 4 * 3 + 1];
114 temp[G_idx] += h_result[idx + 4 * 0 + 0];
115 temp[G_idx] -= h_result[idx + 4 * 1 + 1];
116 temp[G_idx] += h_result[idx + 4 * 2 + 2];
117 temp[G_idx] -= h_result[idx + 4 * 3 + 3];
122 temp[G_idx] -= I * h_result[idx + 4 * 0 + 2];
123 temp[G_idx] -= I * h_result[idx + 4 * 1 + 3];
124 temp[G_idx] += I * h_result[idx + 4 * 2 + 0];
125 temp[G_idx] += I * h_result[idx + 4 * 3 + 1];
130 temp[G_idx] -= h_result[idx + 4 * 0 + 1];
131 temp[G_idx] -= h_result[idx + 4 * 1 + 0];
132 temp[G_idx] += h_result[idx + 4 * 2 + 3];
133 temp[G_idx] += h_result[idx + 4 * 3 + 2];
138 temp[G_idx] += h_result[idx + 4 * 0 + 1];
139 temp[G_idx] += h_result[idx + 4 * 1 + 0];
140 temp[G_idx] += h_result[idx + 4 * 2 + 3];
141 temp[G_idx] += h_result[idx + 4 * 3 + 2];
146 temp[G_idx] -= I * h_result[idx + 4 * 0 + 1];
147 temp[G_idx] += I * h_result[idx + 4 * 1 + 0];
148 temp[G_idx] += I * h_result[idx + 4 * 2 + 3];
149 temp[G_idx] -= I * h_result[idx + 4 * 3 + 2];
154 temp[G_idx] -= h_result[idx + 4 * 0 + 0];
155 temp[G_idx] -= h_result[idx + 4 * 1 + 1];
156 temp[G_idx] += h_result[idx + 4 * 2 + 2];
157 temp[G_idx] += h_result[idx + 4 * 3 + 3];
161 for (
int i = 0; i < 16; i++) h_result[idx + i] = temp[i];
165 template <
typename Float>
void contractColor(Float *spinorX, Float *spinorY, Float *h_result)
168 Float re = 0.0, im = 0.0;
171 for (
int i = 0; i <
V; i++) {
172 for (
int s1 = 0; s1 < 4; s1++) {
173 for (
int s2 = 0; s2 < 4; s2++) {
176 for (
int c = 0; c < 3; c++) {
177 re += (((Float *)spinorX)[24 * i + 6 * s1 + 2 * c + 0] * ((Float *)spinorY)[24 * i + 6 * s2 + 2 * c + 0]
178 + ((Float *)spinorX)[24 * i + 6 * s1 + 2 * c + 1] * ((Float *)spinorY)[24 * i + 6 * s2 + 2 * c + 1]);
180 im += (((Float *)spinorX)[24 * i + 6 * s1 + 2 * c + 0] * ((Float *)spinorY)[24 * i + 6 * s2 + 2 * c + 1]
181 - ((Float *)spinorX)[24 * i + 6 * s1 + 2 * c + 1] * ((Float *)spinorY)[24 * i + 6 * s2 + 2 * c + 0]);
184 ((Float *)h_result)[2 * (i * 16 + 4 * s1 + s2) + 0] = re;
185 ((Float *)h_result)[2 * (i * 16 + 4 * s1 + s2) + 1] = im;
191 template <
typename Float>
196 Float
tol = (
sizeof(Float) ==
sizeof(
double) ? 1e-9 : 2e-5);
197 void *h_result = malloc(
V * 2 * 16 *
sizeof(Float));
206 for (
int j = 0; j < 16; j++) {
208 for (
int i = 0; i <
V; i++) {
209 if (
abs(((Float *)h_result)[32 * i + 2 * j] - ((Float *)d_result)[32 * i + 2 * j]) > tol) {
217 if (
abs(((Float *)h_result)[32 * i + 2 * j + 1] - ((Float *)d_result)[32 * i + 2 * j + 1]) > tol) {
int contraction_reference(Float *spinorX, Float *spinorY, Float *d_result, QudaContractType cType, int X[])
void contractDegrandRossi(Float *h_result_)
enum QudaContractType_s QudaContractType
__host__ __device__ ValueType abs(ValueType x)
void contractColor(Float *spinorX, Float *spinorY, Float *h_result)