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); }