如何解释 GPU 和 CPU 串行版本的均值滤波器的这些结果?

阿米拉·塞纳德希拉

Mean FilterCPU serial版本和NVIDIA GPU parallel版本实现了图像代码我得到的运行时间(请参见结果的测试用例和规格的设备。为什么case 2highest加速并case 3lowest加速?

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则没有必要使用同步机制,因为它是一个同步调用。

  • 还要考虑包括 H2D/D2H 传输,据我说,在 CPU/GPU 比较中考虑这种开销很重要(但这个选择取决于你);
  • 关于您在图中给出的度量,它们是直接结果还是重复不同执行的平均值(可能会丢弃外层值)?

我认为您应该按照上述建议对新措施进行抽样,并对获得的新措施进行考虑。

顺便说一下

由于图像尺寸较大,情况 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] 删除。

编辑于
0

我来说两句

0 条评论
登录 后参与评论

相关文章