Объяснение высоты тона, ширины, высоты, глубины в памяти для 3D-массивов

Question

Я работаю с CUDA и 3D-текстурами в python (используя pycuda). Существует функция, называемая Memcpy3D, которая имеет те же элементы, что и Memcpy2D, плюс несколько дополнительных функций. В нем он призывает вас описать такие вещи, как width_in_bytes, src_pitch, src_height, height и copy_depth. Это то, с чем я борюсь (в 3D) и его актуальность с индексацией стиля C или F. Например, если я просто изменяю порядок от F до C в рабочем примере ниже, он перестает работать - и я не знаю почему.

1. Прежде всего, я понимаю шаг, чтобы быть сколько байтов в памяти, необходимое для перемещения одного индекса по в threadIdx.x (или направление х, или столбец). Итак, для массива float32 формы C (3,2,4), чтобы переместить одно значение в x, я ожидаю перемещение 4 значений в памяти (as the indexing goes down the z axis first?). Поэтому мой танец будет 4 * 32 бит.
2. Я понимаю height как количество строк. (В этом примере 3)
3. Я понимаю width, чтобы быть числом cols. (В этом примере 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.

Использование "F" Заказ в Numpy означает, что порядок измерений изменяется в памяти.

Ваш код, кажется, работает, если я делаю следующие изменения:

#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])