void operator()(Type* dest, const math::Size_t<2u> pitchDest,
Type* source, const math::Size_t<2u> pitchSource, const math::Size_t<3>& size,
flags::Memcopy::Direction direction)
{
const cudaMemcpyKind kind[] = {cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost,
cudaMemcpyHostToHost, cudaMemcpyDeviceToDevice};
cudaPitchedPtr pitchedPtrDest;
pitchedPtrDest.pitch = pitchDest.x(); pitchedPtrDest.ptr = dest;
pitchedPtrDest.xsize = size.x(); pitchedPtrDest.ysize = size.y();
cudaPitchedPtr pitchedPtrSource;
pitchedPtrSource.pitch = pitchSource.x(); pitchedPtrSource.ptr = source;
pitchedPtrSource.xsize = size.x(); pitchedPtrSource.ysize = size.y();
cudaMemcpy3DParms params;
params.srcArray = NULL;
params.srcPos = make_cudaPos(0,0,0);
params.srcPtr = pitchedPtrSource;
params.dstArray = NULL;
params.dstPos = make_cudaPos(0,0,0);
params.dstPtr = pitchedPtrDest;
params.extent = make_cudaExtent(size.x() * sizeof(Type), size.y(), size.z());
params.kind = kind[direction];
CUDA_CHECK_NO_EXCEP(cudaMemcpy3D(¶ms));
}
dim3 cudaGridDim(const math::Size_t<3>& size) const
{
return dim3((size.x() / BlockDim::x::value) * (size.z() / BlockDim::z::value),
size.y() / BlockDim::y::value, 1);
}