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

Question

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

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 过滤器时，它的计算时间是 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

这将很好地合并。经线中的相邻线程将读取内存中的相邻位置。这是“好案例”。

现在如果 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 活动。这是“坏情况”。

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

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