使用 CUDA 在 3D space 中过滤，水平访问比垂直访问快？

Question

我正在尝试在 3D 结构（体积）中分别应用过滤器 1x3、3x1。
例如，如果有 20(cols) x 10(rows) x 10(depth) 结构，

for(int depth = 0; depth < 10; depth++)
    Apply image filter(depth);

对二维图像 (20x10) 应用滤镜 10 次。每一张图片都是不同的。

首先，我分配 3D 结构，如

// COLS = 450, ROWS = 375, MAX_DISPARITY = 60
cudaPitchedPtr volume;
cudaExtent volumeExtent = make_cudaExtent(COLS, ROWS, MAX_DISPARITY);
HANDLE_ERROR(cudaMalloc3D(&volume, volumeExtent ));

并将内存设置为零以实现稳定输出。到目前为止一切顺利，直到将图像复制到卷中。

像下面这样应用 3x1 过滤器时，计算时间为 6 毫秒。

Apply_3by1 << <ROWS, COLS, COLS>> > (volume, COLS, ROWS);

__global__ void Apply_3by1 (cudaPitchedPtr src, unsigned int COLS, unsigned int ROWS)
{
    const unsigned int x = threadIdx.x;
    const unsigned int y = blockIdx.x;

    extern __shared__ unsigned char SharedMemory[];
    for (int dispCnt = 0; dispCnt < MAX_DISPARITY; dispCnt++)
    {
        if (x < dispCnt) continue;//exception for my algorithm.

        unsigned char dst_val = *GET_UCHAR_PTR_3D(src, x, y, dispCnt);
        SharedMemory[x] = dst_val;
        __syncthreads();


        unsigned char left;
        int leftIdx = x - 3;
        if (leftIdx < 0)//index underflow
            left = 0;
        else
            left = SharedMemory[leftIdx];


        unsigned char right;//index overflow
        int rightIdx = x + 3;
        if (COLS < rightIdx)
            right = 0;
        else
            right = SharedMemory[rightIdx];


        *GET_UCHAR_PTR_3D(src, x, y, dispCnt) += left + right;
    }
}

但是当我应用垂直方向的 1x3 过滤器时，它的计算时间是 46 毫秒。

  Apply_1by3 << <COLS, ROWS, ROWS >> > (volume, COLS, ROWS);

__global__ void Apply_1by3 (cudaPitchedPtr src, unsigned int COLS, unsigned int ROWS)
{
    const unsigned int x = threadIdx.x;
    const unsigned int y = blockIdx.x;

    extern __shared__ unsigned char SharedMemory[];
    for (int dispCnt = 0; dispCnt < MAX_DISPARITY; dispCnt++)
    {
        unsigned char my_val = *GET_UCHAR_PTR_3D(src, y, x, dispCnt);
        SharedMemory[x] = my_val;
        __syncthreads();

        if (y < dispCnt) continue;

        int topIdx = x - 3;
        unsigned char top_value;
        if (topIdx < 0)
            top_value = 0;
        else
            top_value = SharedMemory[topIdx];

        int bottomIdx = x + 3;
        unsigned char bottom_value;
        if (ROWS <= bottomIdx)
            bottom_value = 0;
        else
            bottom_value = SharedMemory[bottomIdx];

        *GET_UCHAR_PTR_3D(src, y, x, dispCnt) += bottom_value + top_value;
    }
}

我不知道为什么垂直方向访问比水平访问慢，差不多8倍。如果你知道为什么访问时间不同，请赐教。

对不起，我忘了补充

#define GET_UCHAR_PTR_3D(pptr, x, y, d) \
(unsigned char*)((char*)(pptr).ptr + (sizeof(unsigned char)* x) + ((pptr).pitch * y) + ((pptr).pitch * (pptr).ysize * d))

Answer 1

考虑两种情况之间的全局内存访问和合并行为。是否考虑加载操作无所谓：

    unsigned char my_val = *GET_UCHAR_PTR_3D(src, y, x, dispCnt);

或商店操作：

    *GET_UCHAR_PTR_3D(src, y, x, dispCnt) += bottom_value + top_value;

让我们解压缩您的宏，并在每种情况下替换为 x 和 y 的实际值：

define GET_UCHAR_PTR_3D(pptr, x, y, d) \
(unsigned char*)((char*)(pptr).ptr + (sizeof(unsigned char)* x) + ((pptr).pitch * y) + ((pptr).pitch * (pptr).ysize * d))

我们有：

  (a pointer) + (1*x) + (pitch*y) + offset

现在，如果 x = threadIdx.x 和 y = blockIdx.x 我们有：

  (a pointer) + (1*threadIdx.x) + (pitch*blockIdx.x) + offset

变成：

  (a pointer) + (some offset) + threadIdx.x

这将很好地合并。 warp 中的相邻线程将读取内存中的相邻位置。这是"good case".

现在如果 x = blockIdx.x 和 y = threadIdx.x 会发生什么？我们有：

  (a pointer) + (1*blockIdx.x) + (pitch*threadIdx.x) + offset

变成：

  (a pointer) + (some offset) + (pitch*threadIdx.x)

这意味着 warp 中的相邻线程不是读取内存中的相邻位置，而是读取由 pitch 值分隔的位置。这将不会合并，并将转化为更多的全局请求来满足 warp activity。这是 "bad case".

GPU 就像 "horizontal" 扭曲中的内存访问。他们不喜欢 "vertical" 扭曲内的内存访问。这将导致两种情况之间的性能差异非常大。在这两种情况下看到 10 倍的性能差异并不少见，理论上它可能高达 32 倍的性能差异。

如果您想了解有关合并全局内存访问优化的更多背景知识，请尝试 this 演示文稿，尤其是幻灯片 30-48。

使用 CUDA 在 3D space 中过滤，水平访问比垂直访问快？

Filtering in 3D space with CUDA, horizontal access faster than vertical access?

image-processing

3d

cuda

stereo-3d

imagefilter