QUDA  v1.1.0
A library for QCD on GPUs
matrix_field.h
Go to the documentation of this file.
1 #pragma once
2 
9 // trove requires the warp shuffle instructions introduced with Kepler
10 #if __COMPUTE_CAPABILITY__ >= 300
11 #include <trove/ptr.h>
12 #else
13 #define DISABLE_TROVE
14 #endif
15 
16 #include <quda_matrix.h>
17 
18 namespace quda
19 {
20 
21  template <typename T, int n> struct matrix_field {
22  T *field;
23  int volume_cb;
24 
26 
27  __device__ __host__ inline void load(Matrix<T, n> &A, int x_cb, int parity) const
28  {
29  int idx = parity * volume_cb + x_cb;
30 #ifdef __CUDA_ARCH__
32  A = field_[idx];
33 #else
34 #pragma unroll
35  for (int i = 0; i < n; i++)
36 #pragma unroll
37  for (int j = 0; j < n; j++) A(i, j) = field[(n * idx + i) * n + j] = A(i, j);
38 #endif
39  }
40 
41  __device__ __host__ inline void save(const Matrix<T, n> &A, int x_cb, int parity)
42  {
43  int idx = parity * volume_cb + x_cb;
44 #ifdef __CUDA_ARCH__
46  field_[idx] = A;
47 #else
48 #pragma unroll
49  for (int i = 0; i < n; i++)
50 #pragma unroll
51  for (int j = 0; j < n; j++) field[(n * idx + i) * n + j] = A(i, j);
52 #endif
53  }
54  };
55 
56 } // namespace quda
QudaParity parity
Definition: covdev_test.cpp:40
__device__ __host__ void save(const Matrix< T, n > &A, int x_cb, int parity)
Definition: matrix_field.h:41
matrix_field(T *field, int volume_cb)
Definition: matrix_field.h:25
__device__ __host__ void load(Matrix< T, n > &A, int x_cb, int parity) const
Definition: matrix_field.h:27