这两个Cuda核函数及其调用代码，执行效率可否优化提高？

Question

自己写的核函数，功能是实现YUV格式的图像转RGB格式的图像，基本上是纯数学计算，感觉执行速度太慢了，不知如何优化？
第一个核函数代码如下：

__global__ void WholeUyvyToRgbKernel(uint8_t *yuv_buffer, uint8_t *rgb_buffer) {
    int bSize = blockDim.z * blockDim.y * blockDim.x;
    int bIndex = blockIdx.z * gridDim.x * gridDim.y + \
               blockIdx.y * gridDim.x + \
               blockIdx.x;
    int tIndex = threadIdx.z * blockDim.x * blockDim.y + \
               threadIdx.y * blockDim.x + \
               threadIdx.x;
    int index = bIndex * bSize + tIndex;
    //printf("block idx: %3d, thread idx in block: %3d, thread idx, %3d\n",
    //     bIndex, tIndex, index);
    uint8_t *yuv_start_ptr = yuv_buffer + index * 4;
    uint8_t *rgb_start_ptr = rgb_buffer + index * 6;
    int U = *(yuv_start_ptr) - 128;
    int Y1 = *(yuv_start_ptr + 1);
    int V = *(yuv_start_ptr + 2) - 128;
    int Y2 = *(yuv_start_ptr + 3);
    int R1 = Y1 + 1.4075 * V;
    int G1 = Y1 - 0.3455 * U - 0.7169 * V;
    int B1 = Y1 + 1.779 * U;
    int R2 = Y2 + 1.4075 * V;
    int G2 = Y2 - 0.3455 * U - 0.7169 * V;
    int B2 = Y2 + 1.779 * U;
    *(rgb_start_ptr + 0) = min(max(R1, 0), 255);
    *(rgb_start_ptr + 1) = min(max(G1, 0), 255);
    *(rgb_start_ptr + 2) = min(max(B1, 0), 255);
    *(rgb_start_ptr + 3) = min(max(R2, 0), 255);
    *(rgb_start_ptr + 4) = min(max(G2, 0), 255);
    *(rgb_start_ptr + 5) = min(max(B2, 0), 255);
}

第一个核函数调用代码：

dim3 block_size_dim(32, 32, 1);
dim3 grid_size_dim = dim3(2616, 1, 1);  // 实际执行中都是固定宽高的图像
WholeUyvyToRgbKernel <<< grid_size_dim, block_size_dim, 0, stream >>>(WholeUyvyBufferDevice, WholeRgbBufferDevice);

第二个核函数代码如下：

__global__ void SubYuvToRgbKernel(uint8_t *sub_buffer_device, int sub_width) {
    int position = blockDim.x * blockIdx.x + threadIdx.x;
    int col_index = position % sub_width;
    int row_index = position / sub_width;
    int sub_rgb_index = row_index * sub_width + col_index;
    
    auto y = sub_buffer_device[sub_rgb_index * 3 + 0];
    auto u = sub_buffer_device[sub_rgb_index * 3 + 1] - 128;
    auto v = sub_buffer_device[sub_rgb_index * 3 + 2] - 128;
    auto r = static_cast<int>(y + 1.4075 * v);
    auto g = static_cast<int>(y - 0.3455 * u - 0.7169 * v);
    auto b = static_cast<int>(y + 1.779 * u);
    sub_buffer_device[sub_rgb_index * 3 + 0] = static_cast<unsigned char>(min(max(r, 0), 255));
    sub_buffer_device[sub_rgb_index * 3 + 1] = static_cast<unsigned char>(min(max(g, 0), 255));
    sub_buffer_device[sub_rgb_index * 3 + 2] = static_cast<unsigned char>(min(max(b, 0), 255));
}

第二个核函数调用代码：

int threads = (std::min)(src_  width, src_height);   // 两个变量 src_  width, src_height 的取值范围均大于64小于512，一般在128左右
int blocks = (std::max)(src_width, src_height);
SubYuvToRgbKernel <<<blocks, threads, 0, stream >>>(SubBufferDevice, src_width);                

输入的uint8_t *类型的参数，都是从内存使用cudaMemCpy到显存的，好像是在Global Memory上，没做其他处理。可否指点一下，如何加速？ 先行谢过。

Answer 1

引自免费微信小程序：皆我百晓生

针对您的两个CUDA核函数及其调用代码，这里有一些建议可以帮助您优化执行效率：

1. 数据布局优化:

   • 考虑使用 pitched memory layout 或者 cudaMallocPitch 来分配YUV和RGB缓冲区。这样可以更高效地访问相邻像素，减少全局内存访问开销。

2. 线程块大小调整:

   • 调整线程块大小以适应GPU的纹理单元（TPC）和 warp 结构。对于现代GPU，通常推荐使用32x32或16x16的线程块大小。

3. 共享内存利用:

   • 使用共享内存来存储最近使用的像素值，降低全局内存访问次数。

4. 消除不必要的计算:

   • 由于每个线程处理两组像素，您可以考虑合并R1、G1、B1和R2、G2、B2的计算，减少冗余操作。

5. 预加载数据:

   • 避免使用指针间接访问数组元素，通过索引一次性加载数据到寄存器中，减少访存延迟。

以下是经过优化后的核函数代码示例：

首先，为YUV和RGB创建pitched内存布局：

// 分配 pitched 内存
size_t pitch_yuv, pitch_rgb;
cudaMallocPitch(&WholeYuvBufferDevice, &pitch_yuv, sizeof(uint8_t)*4, img_width*img_height/2);
cudaMallocPitch(&WholeRgbBufferDevice, &pitch_rgb, sizeof(uint8_t)*6, img_width*img_height);

// 调用转换函数
WholeUyvyToRgbKernel<<<grid_size_dim, block_size_dim, 0, stream>>>(pitch_yuv, pitch_rgb, img_width, img_height);

接下来是优化过的核函数 WholeUyvyToRgbKernel：

__global__ void WholeUyvyToRgbKernel(size_t pitch_yuv, size_t pitch_rgb, int img_width, int img_height) {
    __shared__ uint8_t shared_data[32][32][4];  // 基于线程块大小调整共享内存大小

    int bX = blockIdx.x;
    int bY = blockIdx.y;
    int bZ = blockIdx.z;
    int tX = threadIdx.x;
    int tY = threadIdx.y;

    int x = bX * 32 + tX;
    int y = bY * 32 + tY;
    if (x < img_width && y < img_height) {
        int yuv_offset = (y / 2) * pitch_yuv + (x * 4) + (y % 2) * 2; // 计算YUV偏移量
        int rgb_offset = y * pitch_rgb + x * 6; // 计算RGB偏移量

        uint8_t y1, u, y2, v;
        if (y % 2 == 0) {
            y1 = WholeYuvBufferDevice[yuv_offset];
            u = WholeYuvBufferDevice[yuv_offset + 1];
            y2 = WholeYuvBufferDevice[yuv_offset + 2];
            v = WholeYuvBufferDevice[yuv_offset + 3];
        } else {
            y2 = WholeYuvBufferDevice[yuv_offset];
            v = WholeYuvBufferDevice[yuv_offset + 1];
            y1 = WholeYuvBufferDevice[yuv_offset + 2];
            u = WholeYuvBufferDevice[yuv_offset + 3];
        }

        shared_data[tY][tX][0] = y1;
        shared_data[tY][tX][1] = u;
        shared_data[tY][tX][2] = y2;
        shared_data[tY][tX][3] = v;

        __syncthreads();

        int R1 = y1 + 1.4075 * shared_data[tY][tX][1];
        int G1 = y1 - 0.3455 * shared_data[tY][tX][0] - 0.7169 * shared_data[tY][tX][1];
        int B1 = y1 + 1.779 * shared_data[tY][tX][0];

        int R2 = y2 + 1.4075 * shared_data[tY][tX][3];
        int G2 = y2 - 0.3455 * shared_data[tY][tX][2] - 0.7169 * shared_data[tY][tX][3];
        int B2 = y2 + 1.779 * shared_data[tY][tX][2];

        R1 = min(max(R1, 0), 255);
        G1 = min(max(G1, 0), 255);
        B1 = min(max(B1, 0), 255);
        R2 = min(max(R2, 0), 255);
        G2 = min(max(G2, 0), 255);
        B2 = min(max(B2, 0), 255);

        WholeRgbBufferDevice[rgb_offset + 0] = static_cast<uint8_t>(R1);
        WholeRgbBufferDevice[rgb_offset + 1] = static_cast<uint8_t>(G1);
        WholeRgbBufferDevice[rgb_offset + 2] = static_cast<uint8_t>(B1);
        WholeRgbBufferDevice[rgb_offset + 3] = static_cast<uint8_t>(R2);
        WholeRgbBufferDevice[rgb_offset + 4] = static_cast<uint8_t>(G2);
        WholeRgbBufferDevice[rgb_offset + 5] = static_cast<uint8_t>(B2);
    }
}

对于SubYuvToRgbKernel，由于其处理的是小区域，可能不需要进行大规模的数据布局优化。不过，仍可以考虑使用线程块大小调整和预加载数据来提高效率。这里不再给出优化后的代码，因为优化方法与上面类似。

最后，确保您在调用核函数之前正确设置stream，并在适当的时候同步stream以确保所有工作完成。

注意：这些优化建议可能需要根据您的具体GPU硬件性能进行微调，以便达到最佳效果。此外，使用nvprof或其他工具对程序进行性能分析，以确定是否存在其他瓶颈并进行针对性优化。