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

Question

我正在 python 中处理 CUDA 和 3D 纹理（使用 pycuda）。有一个名为Memcpy3D 的函数与Memcpy2D 具有相同的成员以及一些额外的成员。在其中它要求您描述诸如width_in_bytes、src_pitch、src_height、height 和copy_depth 之类的东西。这就是我正在努力解决的问题（在 3D 中）以及它与 C 或 F 样式索引的相关性。例如，如果我在下面的工作示例中简单地将排序从 F 更改为 C，它就会停止工作 - 我不知道为什么。

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

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

请教我。

下面是一个完整的工作脚本，它将数组作为纹理复制到 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) 是指变化最快的维度，height (y) 是中间维度，depth (z) 是变化最慢的维度. 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])