优化 CUDA 二维数组访问

Posted 2023-02-16

【中文标题】优化 CUDA 二维数组访问

【英文标题】：optimizing CUDA 2d array access

【发布时间】：2013-12-06 09:18:41

【问题描述】：

我正在 gpu 上进行一些非常基本的图像处理。我传递了一个填充了像素 RGB 值的数组，它们一个接一个地存储。

我像

一样启动我的内核

kernel<<<numBlocks,numThreadsPerBlock>>>(unsigned char * imageData, int val)

网格大小取决于我的图像有多大，但通常会像数千块 512-1024 线程

在内核中，我进行一些基本计算并直接比较全局内存中的值。

int blueIdx = (blockIdx.x * blockDim.x + threadIdx.x) * 3;
    int greenIdx = blueIdx + 1;
    int redIdx = greenIdx + 1;
float ypx = 0.299 * imageData[redIdx] + 0.587 * imageData[greenIdx] + 0.114 * imageData[blueIdx];

if( ypx > val) 
    imageData[blueIdx] = 255;

有没有办法优化这种访问方式？本质上是从二维数组中读取值，执行简单的静态乘法，执行比较，然后将新值保存回全局内存。

我已经尝试过共享内存，但我的实现速度较慢。我假设由于线程是独立的共享内存将无济于事。

【问题讨论】：

如需更好的答案，请说明您的问题的内核签名（类型）和大小（宽度/高度）。还有完整的内核启动参数。

根据您所展示的内容，我可以想象的主要优化好处是确保您的读取和写入合并。根据您到目前为止所展示的内容，我不会假设“每个块都已经以合并的方式从全局内存中加载图像数据”。无论这是字节数组还是其他大小，通过偏移处理数据都需要特别注意以实现最佳合并。如果没有数据重用（或线程间通信），共享内存将无济于事。

8 位和 16 位访问永远不会表现出良好的性能，因此您应该考虑使用 32 位 ARGB 像素。合并的好处将超过未使用内存带宽的 30% 损失。另外，使用 cudaMallocPitch() 和 family 来分配 arra。已将倾斜的内存分配添加到 CUDA 中，以便在您访问相邻行时合并全局内存操作。

【参考方案1】：

我还没有尝试过，但取决于使用 uchar3 甚至 uchar4（用于 32 位对齐）的编译器的聪明程度（即缺乏它）可能会更快。我必须查看它为您的内核生成的 PTX 代码才能确定。

int idx = blockIdx.x * blockDim.x + threadIdx.x;
uchar3 rgb = imageData[idx];
float ypx = 0.299 * rgb.x + 0.587 * rgb.y + 0.114 * rgb.z;

if( ypx > val) 
    rgb.z = 255;
    imageData[idx] = rgb;

请注意，只有一个读取操作。它只是没有比这更简单的了。好吧，由于更容易合并，将写操作移出条件分支可能会更快，但由于更多的写操作，它也可能会更慢。实验。

但实际上，事情应该已经非常快了。我假设您可以在现代游戏 GPU 上大约 5-10 毫秒内在 1GB 数据上运行此内核。这对你来说太慢了吗？确定是内核耗时过长？

【讨论】：

我得到以下 uchar3 rgb = imageData[idx]; 的编译器错误没有合适的构造函数可以将“unsigned char”转换为“uchar3”

@smashbourne 您需要更改 imageData 的类型、分配、初始化并将结果返回到 uchar3。