【发布时间】：2015-07-21 13:25:17
【问题描述】：
我有一个大小为 1920 x 1080 的图像。我正在使用三个 CUDA 流从 H2D 传输、处理并从 D2H 传输回来，每个流负责处理总数据的 1/3。通过了解 SM、SP、warp 的概念，我能够优化块的尺寸和每个块的线程数。如果必须在内核中进行简单的计算，则代码可以令人满意地运行（需要 2 毫秒）。下面的简单计算代码从源图像中找到 R、G 和 B 值，然后将这些值放在同一个源图像中。
ptr_source[numChannels* (iw*y + x) + 0] = ptr_source[numChannels* (iw*y + x) + 0];
ptr_source[numChannels* (iw*y + x) + 1] = ptr_source[numChannels* (iw*y + x) + 1];
ptr_source[numChannels* (iw*y + x) + 2] = ptr_source[numChannels* (iw*y + x) + 2];
但是我必须执行 更多计算，这些计算独立于所有其他线程，然后计算时间增加了 6 毫秒，这对我的应用程序来说太多了。我已经尝试在constant memory 中声明最常用的常量值。这些计算的代码如下所示。在该代码中，我再次找到了 R、G 和 B 值。然后，我通过将旧值乘以一些常数来计算 R、G 和 B 的新值，最后我将这些新的 R、G 和 B 值再次放入同一源图像的相应位置。
__constant__ int iw = 1080;
__constant__ int ih = 1920;
__constant__ int numChannels = 3;
__global__ void cudaKernel(unsigned char *ptr_source, int numCudaStreams)
{
// Calculate our pixel's location
int x = (blockIdx.x * blockDim.x) + threadIdx.x;
int y = (blockIdx.y * blockDim.y) + threadIdx.y;
// Operate only if we are in the correct boundaries
if (x >= 0 && x < iw && y >= 0 && y < ih / numCudaStreams)
{
const int index_b = numChannels* (iw*y + x) + 0;
const int index_g = numChannels* (iw*y + x) + 1;
const int index_r = numChannels* (iw*y + x) + 2;
//GET VALUES: get the R,G and B values from Source image
unsigned char b_val = ptr_source[index_b];
unsigned char g_val = ptr_source[index_g];
unsigned char r_val = ptr_source[index_r];
float float_r_val = ((1.574090) * (float)r_val + (0.088825) * (float)g_val + (-0.1909) * (float)b_val);
float float_g_val = ((-0.344198) * (float)r_val + (1.579802) * (float)g_val + (-1.677604) * (float)b_val);
float float_b_val = ((-1.012951) * (float)r_val + (-1.781485) * (float)g_val + (2.404436) * (float)b_val);
unsigned char dst_r_val = (float_r_val > 255.0f) ? 255 : static_cast<unsigned char>(float_r_val);
unsigned char dst_g_val = (float_g_val > 255.0f) ? 255 : static_cast<unsigned char>(float_g_val);
unsigned char dst_b_val = (float_b_val > 255.0f) ? 255 : static_cast<unsigned char>(float_b_val);
//PUT VALUES---put the new calculated values of R,G and B
ptr_source[index_b] = dst_b_val;
ptr_source[index_g] = dst_g_val;
ptr_source[index_r] = dst_r_val;
}
}
问题：我认为将图像段（即ptr_src）传输到共享内存会有所帮助，但我对如何做到这一点感到很困惑。我的意思是，共享内存的范围只有一个块，我如何管理图像段到共享内存的传输。
PS：我的 GPU 是 Quadro K2000，计算 3.0，2 SM，每个 SM 192 SP。
【问题讨论】：
-
我们遗漏了您的一些实施细节。您是否使用三个流，然后在流之间划分每个颜色通道？此外，在此示例中使用共享内存不会获得任何性能改进，因为所有操作都是在像素级别完成的。事实上，总是渴望优化 nvcc 编译器可能会将计算转移到寄存器中。最后要注意的是，您对全局内存的内存访问模式没有合并，因为连续线程不会访问连续的内存位置。
-
@pQB：我已经更新了我的问题的第二句话。希望，现在我能够澄清 3 个流的使用。
-
我不确定将图像拆分为三个流是否有任何好处。为什么不让
y维度中的块完成所有工作？尽管如此，关于共享内存的评论仍然是一样的，以及合并的问题。这些改进（合并）可能会缩短执行时间。 -
@pQB：由于并行数据传输和内核执行，三个流肯定使我受益。我没看懂你
memory access pattern to global memory is not coalesct because consecutive threads do not access consecutive memory positions.这句话，因为我对这东西不太熟悉。 -
你确实意识到你在这个内核中有双精度算术，并且可以在不改变你的代码行的情况下获得两倍的免费加速（除了修复双精度常量）......跨度>