图像处理 - 使用 CUDA 在 3D 空间中过滤，水平访问比垂直访问更快

Question

我正在尝试在 3D 结构（体积）中重复应用过滤器 1x3、3x1。
例如，如果有 20（列）x 10（行）x 10（深度）结构，

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

对2D图像（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 过滤器时，它的计算时间为 46mSec。

  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

这将很好地融合在一起。 扭曲中的相邻线程将读取内存中的相邻位置。 这是"好案例"。

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

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

它变成：

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

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

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

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