我Mean Filter
为CPU serial
版本和NVIDIA
GPU parallel
版本实现了图像代码。我得到的运行时间(请参见结果的测试用例和规格的设备。为什么case 2
有highest
加速并case 3
有lowest
加速?
GPU执行配置
int block_size = 32;
int grid_size = width/block_size; //width of the image in pixels
dim3 dimBlock(block_size, block_size, 1);
dim3 dimGrid(grid_size, grid_size, 1);
GPU 代码的时间测量
clock_t start_d=clock();
meanFilter_d <<< dimGrid, dimBlock >>> (image_data_d, result_image_data_d, width, height, half_window);
cudaThreadSynchronize();
clock_d end_d=clock();
CPU 代码的时间测量(单线程)
clock_t start_h = clock();
meanFilter_h(data, result_image_data_h1, width, height, window_size);
clock_t end_h = clock();
主机代码
void meanFilter_h(unsigned char* raw_image_matrix,unsigned char* filtered_image_data,int image_width, int image_height, int window_size)
{
// int size = 3 * image_width * image_height;
int half_window = (window_size-window_size % 2)/2;
for(int i = 0; i < image_height; i += 1){
for(int j = 0; j < image_width; j += 1){
int k = 3*(i*image_height+j);
int top, bottom, left, right;
if(i-half_window >= 0){top = i-half_window;}else{top = 0;}// top limit
if(i+half_window <= image_height-1){bottom = i+half_window;}else{bottom = image_height-1;}// bottom limit
if(j-half_window >= 0){left = j-half_window;}else{left = 0;}// left limit
if(j+half_window <= image_width-1){right = j+half_window;}else{right = image_width-1;}// right limit
double first_byte = 0;
double second_byte = 0;
double third_byte = 0;
// move inside the window
for(int x = top; x <= bottom; x++){
for(int y = left; y <= right; y++){
int pos = 3*(x*image_height + y); // three bytes
first_byte += raw_image_matrix[pos];
second_byte += raw_image_matrix[pos+1];
third_byte += raw_image_matrix[pos+2];
}
}
int effective_window_size = (bottom-top+1)*(right-left+1);
filtered_image_data[k] = first_byte/effective_window_size;
filtered_image_data[k+1] = second_byte/effective_window_size;
filtered_image_data[k+2] =third_byte/effective_window_size;
}
}
}
设备代码
__global__ void meanFilter_d(unsigned char* raw_image_matrix, unsigned char* filtered_image_data, int image_width, int image_height, int half_window)
{
int j = blockIdx.x * blockDim.x + threadIdx.x;
int i = blockIdx.y * blockDim.y + threadIdx.y;
if (i < image_height && j < image_width){
int k = 3*(i*image_height+j);
int top, bottom, left, right;
if(i-half_window >= 0){top = i-half_window;}else{top = 0;}// top limit
if(i+half_window <= image_height-1){bottom = i+half_window;}else{bottom = image_height-1;}// bottom limit
if(j-half_window >= 0){left = j-half_window;}else{left = 0;}// left limit
if(j+half_window <= image_width-1){right = j+half_window;}else{right = image_width-1;}// right limit
double first_byte = 0;
double second_byte = 0;
double third_byte = 0;
// move inside the window
for(int x = top; x <= bottom; x++){
for(int y = left; y <= right; y++){
int pos = 3*(x*image_height + y); // three bytes
first_byte += raw_image_matrix[pos];
second_byte += raw_image_matrix[pos+1];
third_byte += raw_image_matrix[pos+2];
}
}
int effective_window_size = (bottom-top+1)*(right-left+1);
filtered_image_data[k] = first_byte/effective_window_size;
filtered_image_data[k+1] = second_byte/effective_window_size;
filtered_image_data[k+2] =third_byte/effective_window_size;
}
}
可以看出,有3×3
内核的两种图像大小都比5*5
内核慢。由于图像尺寸较大,情况 1 比情况 3 具有更多的并行度。因此,案例1的设备利用率高于案例3。但我不知道进一步解释。请给我一些见解。
首先要指出的是:你在测量什么,最重要的是如何测量?从您的问题中无法特别推断出如何。
无论如何,我强烈建议您看一看,这是 Mark Harris 的一篇非常简单且有用的文章,它解释了一些对设备端代码(即 CUDA 内存传输、内核等)的执行时间进行采样的良好实践。
顺便说一下,试图获得CPU/GPU 加速是一个非常棘手的话题,这是由于两种架构的本质不同。即使您的 CPU 和 GPU 代码显然在做同样的事情,您也可能需要考虑很多因素(例如 CPU 内核、GPU 流多处理器和每个 SM 的内核)。在这里, Robert Crovella 对类似的问题给出了一个很好的答案,正如他所说:
如果您声称“GPU 比 CPU 快 XX”,那么 IMO 建议您只比较执行相同工作并高效地使用底层架构(对于 CPU 和 GPU)的代码。例如,在 CPU 情况下,您当然应该使用多线程代码,以便利用大多数现代 CPU 提供的多个 CPU 内核。无论如何,这些类型的声明可能会受到怀疑,因此最好避免它们,除非它是您意图的关键。
我建议你也看看这个讨论。
在经过一些前提之后,我认为您不能认为这些加速是可靠的(实际上这些对我来说似乎有点奇怪)。
试图解释你想说的话:
可以看出,使用 3×3 内核的两种图像大小都较慢
也许你想说在 3x3 中你获得了一个较小的加速比 5x5 窗口大小。尝试更准确。
为什么情况 2 的加速比最高,而情况 3 的加速比最低?
好吧,通过您提供的糟糕信息很难推断出某些内容。
请添加:一些代码以查看您在做什么以及您如何在设备和主机情况下实现问题,描述您正在测量的方式和内容。
编辑 :
嗯,我认为你应该采取更准确的措施。
clock()
. 看看这里的答案和 C++ 参考,我建议你考虑使用std::chrono::system_clock::now()
std::chrono::high_resolution_clock::now();
使用主机-设备同步点(例如 )的一个问题
cudaDeviceSynchronize()
是它们会导致 GPU 流水线停滞。出于这个原因,CUDA 通过 CUDA 事件 API 提供了一个相对轻量级的 CPU 计时器替代方案。CUDA 事件 API 包括创建和销毁事件、记录事件以及计算两个记录事件之间经过的时间(以毫秒为单位)的调用。
这意味着您提供的措施的实际结果可能会因使用cudaDeviceSynchronize()
. 此外,如果您使用 simple cudaMemcpy
,则没有必要使用同步机制,因为它是一个同步调用。
我认为您应该按照上述建议对新措施进行抽样,并对获得的新措施进行考虑。
顺便说一下
由于图像尺寸较大,情况 1 比情况 3 具有更多的并行度。因此,情况 1 的设备利用率高于情况 3。
我不同意,因为你已经 int grid_size = width/block_size;
案例 1:grid_size = 640/32 = 20
情况 2:grid_size = 1280/32 = 40
因此,在第 2 种情况下,您的并行度更高。但由于您只有 2 个 SM,这可能是时间可能比您预期的要长的原因。换句话说,您有更多块 (40*40) 等待计算两个 SM。
本文收集自互联网,转载请注明来源。
如有侵权,请联系 [email protected] 删除。
我来说两句