解释 3D 数组在内存中的间距、宽度、高度、深度

Question

我正在 python 中使用 CUDA 和 3D 纹理（使用 pycuda）。有一个名为 Memcpy3D which has the same members as Memcpy2D 的函数加上一些额外的功能。它在其中调用您描述诸如 width_in_bytes、src_pitch、src_height、height 和 copy_depth 之类的内容。这就是我正在努力解决的问题（在 3D 中）及其与 C 或 F 样式索引的相关性。例如，如果我在下面的工作示例中简单地将顺序从 F 更改为 C，它将停止工作 - 我不知道为什么。

首先，我将间距理解为在 threadIdx.x（或 x 方向，或一列）中移动一个索引所需的内存字节数。因此，对于 C 形 (3,2,4) 的 float32 数组，要在 x 中移动一个值，我希望在内存中移动 4 个值 (as the indexing goes down the z axis first?)。因此我的间距是 4*32 位。
我理解 height 是行数。（在这个例子中，3）
我理解 width 是cols的数量。（在这个例子中，2）
我理解 depth 是 z 切片的数量。（在这个例子中，4）
我理解 width_in_bytes 是行的宽度 x 包括它后面的 z 元素的，即行切片，(0,:,: ).这将是在 y 方向上横穿一个元素所需的内存地址数。

因此，当我在下面的代码中将顺序从 F 更改为 C，并相应地调整代码以更改 height/width 值时，它仍然不起作用。它只是呈现逻辑错误，让我觉得我没有正确理解间距、宽度、高度、深度的概念。

请赐教

下面是一个完整的工作脚本，它将一个数组作为纹理复制到 GPU 并将内容复制回来。

import pycuda.driver as drv
import pycuda.gpuarray as gpuarray
import pycuda.autoinit
from pycuda.compiler import SourceModule
import numpy as np

w = 2
h = 3
d = 4
shape = (w, h, d)

a = np.arange(24).reshape(*shape,order='F').astype('float32')
print(a.shape,a.strides)
print(a)


descr = drv.ArrayDescriptor3D()
descr.width = w
descr.height = h
descr.depth = d
descr.format = drv.dtype_to_array_format(a.dtype)
descr.num_channels = 1
descr.flags = 0

ary = drv.Array(descr)

copy = drv.Memcpy3D()
copy.set_src_host(a)
copy.set_dst_array(ary)
copy.width_in_bytes = copy.src_pitch = a.strides[1]
copy.src_height = copy.height = h
copy.depth = d

copy()

mod = SourceModule("""
    texture<float, 3, cudaReadModeElementType> mtx_tex;

    __global__ void copy_texture(float *dest)
    {
      int x = threadIdx.x;
      int y = threadIdx.y;
      int z = threadIdx.z;
      int dx = blockDim.x;
      int dy = blockDim.y;
      int i = (z*dy + y)*dx + x;
      dest[i] = tex3D(mtx_tex, x, y, z);
    }
""")

copy_texture = mod.get_function("copy_texture")
mtx_tex = mod.get_texref("mtx_tex")

mtx_tex.set_array(ary)

dest = np.zeros(shape, dtype=np.float32, order="F")
copy_texture(drv.Out(dest), block=shape, texrefs=[mtx_tex])

print(dest)

Answer 1

不确定我是否完全理解您的代码中的问题，但我会尝试澄清。

在CUDA中，width（x）指的是fastest-changing维度，height（y）是中间维度，depth (z) 是 slowest-changing 维度。 pitch 指的是沿 y 维度在值之间步进所需的步幅（以字节为单位）。

在Numpy中，定义为np.empty(shape=(3,2,4), dtype=np.float32, order="C")的数组有strides=(32, 16, 4)，对应width=4、height=2、depth=3、pitch=16 .

在 Numpy 中使用 "F" 排序意味着维度的顺序在内存中是相反的。

如果我进行以下更改，您的代码似乎可以正常工作：

#shape = (w, h, d)
shape = (d, h, w)

#a = np.arange(24).reshape(*shape,order='F').astype('float32')
a = np.arange(24).reshape(*shape,order='C').astype('float32')
...
#dest = np.zeros(shape, dtype=np.float32, order="F")
dest = np.zeros(shape, dtype=np.float32, order="C")
#copy_texture(drv.Out(dest), block=shape, texrefs=[mtx_tex])
copy_texture(drv.Out(dest), block=(w,h,d), texrefs=[mtx_tex])

解释 3D 数组在内存中的间距、宽度、高度、深度

Explain pitch, width, height, depth in memory for 3D arrays

python

cuda

pycuda