shared_memory_cache_helper.cuh

Go to the documentation of this file.

#pragma once

namespace quda
{

  template <typename real, typename Vector> class VectorCache
  {

    __device__ inline real *cache()
    {
      extern __shared__ int cache_[];
      return reinterpret_cast<real *>(cache_);
    }

public:
    __device__ inline void save(const Vector &a)
    {
      int j = (threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + threadIdx.x;
#pragma unroll
      for (int i = 0; i < 2 * a.size; i++) {
        cache()[j] = *(reinterpret_cast<const real *>(a.data) + i);
        j += blockDim.z * blockDim.y * blockDim.x;
      }
    }

    __device__ inline Vector load(int x, int y, int z)
    {
      Vector a;
      int j = (z * blockDim.y + y) * blockDim.x + x;
#pragma unroll
      for (int i = 0; i < 2 * a.size; i++) {
        *(reinterpret_cast<real *>(a.data) + i) = cache()[j];
        j += blockDim.z * blockDim.y * blockDim.x;
      }
      return a;
    }

    __device__ inline void sync() { __syncthreads(); }
  };

} // namespace quda