cudamemcpy3d error

Hi,

I have a 3d array in host and i’m trying to copy a part of it to device and viceversa.
I didn’t find a good documentation for cudamemcpy3d in CUDA FORTRAN.
This is the pseudocode

REAL VEL1(M, N, W)
REAL, DEVICE :: VEL1D(M, N, W)
TYPE(CUDAMEMCPY3DPARMS) :: CUDA3DPARMS

CUDA3DPARMS%srcPtr%ptr = UXVEL1(0,0,0);
CUDA3DPARMS%srcPtr%pitch = ACTUAL_N !actual_n < n
CUDA3DPARMS%srcPtr%xsize = ACTUAL_N !actual_m < m
CUDA3DPARMS%srcPtr%ysize = ACTUAL_M
CUDA3DPARMS%dstPtr%ptr = UXVEL1D(0,0,0);
CUDA3DPARMS%dstPtr%pitch = ACTUAL_N
CUDA3DPARMS%dstPtr%xsize = ACTUAL_N
CUDA3DPARMS%dstPtr%ysize = ACTUAL_M
CUDA3DPARMS%extent%width = ACTUAL_N
CUDA3DPARMS%extent%height = ACTUAL_M
CUDA3DPARMS%extent%depth = KMMAI !actual_k < k

CALL KERNEL

COPY BACK THE 3D ARRAY

The output of the compiler is:
PGF90-S-0148-Reference to derived type required

for the instructions
CUDA3DPARMS%srcPtr%ptr = UXVEL1(0,0,0)
CUDA3DPARMS%dstPtr%ptr = UXVEL1D(0,0,0)

any idea?

-Alberto

Well, we don’t get many users of cudaMemcpy3D() anymore. But, I think you want to do this:

CUDA3DPARMS%srcPtr%ptr = c_loc(UXVEL1(0,0,0))
CUDA3DPARMS%dstPtr%ptr = c_devloc(UTXVELD1(0,0,0))

Or something like that.

It is probably much more readable to use slice notation in your array assignment. Is this just a sub-cube (in every dimension) of a larger cube?

Is this just a sub-cube (in every dimension) of a larger cube?

Yes, it is.

So, cudamemcpy3D is very particular about the leading dimension (the pitch) and we usually recommend you not use it for general Fortran arrays. Here’s two ways to do it. In recent PGI compilers, 2D slicing makes efficient use of cudaMemcpy2D (which does not have the same restrictions on the pitch) and performs better than a “sub-cube” Fortran slice, which I show in the first case:

program do3d
use cudafor
integer, parameter :: n1 = 128, n2 = 128, n3 = 128
integer, parameter :: k1 = 64, k2 = 80, k3 = 96
integer, parameter :: ntimes = 100
integer, dimension(n1,n2,n3) :: x, xd
attributes(device) :: xd

xd = 0
x = 1
print ,“easy, slow way”
do i = 1, ntimes
xd(1:k1,1:k2,1:k3) = x(1:k1,1:k2,1:k3)
end do
print ,sum(xd).eq.k1k2
k3
print ,sum(xd(1:k1,1:k2,1:k3)).eq.k1k2*k3

xd = 0
print ,“2d slicing way”
do i = 1, ntimes
if (k2.le.k3) then
do k = 1, k2
xd(1:k1,k,1:k3) = x(1:k1,k,1:k3)
end do
else
do k = 1, k3
xd(1:k1,1:k2,k) = x(1:k1,1:k2,k)
end do
end if
end do
print ,sum(xd).eq.k1k2
k3
print ,sum(xd(1:k1,1:k2,1:k3)).eq.k1k2*k3
end