python - 解释 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. 我将高度理解为行数。 (在这个例子中，3)
3. 我理解 width 是列数。 (在这个例子中，2)
4. 我理解 depth 是 z 切片的数量。 (在这个例子中，4)
5. 我将 width_in_bytes 理解为行的宽度 x 包括 后面的 z 元素，即行切片 (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])