opencl最佳组大小

Question

我在OpenCL上运行mandelbrot生成器(来自静态参数的2D图像).该计划很简单:

__kernel
void mandelbrot(__global uchar * output, 
                const float xstep,
                const float xoffset,
                const float ystep,
                const float yoffset,
                const int maxiter)
{
    int gid_y = get_global_id(1);
    int gid_x = get_global_id(0);

    //calculate x and y on the fly for every pixel. 
    //This is just as fast as reading precalculated rulers from global memory.
    float x = gid_x * xstep + xoffset;
    float y = gid_y * ystep + yoffset;

    float real = 0;
    float imag = 0;

    int out = 0;

    for(int curiter = 0; curiter < maxiter; curiter++) {
        float nreal = real*real - imag*imag + x;
        imag = 2* real*imag + y;
        real = nreal;

        if (real*real + imag*imag > 4.0f) {
            out = curiter;
            break;
        }
    }

    //normalize output
    out *= 256.0 / (float)maxiter;
    output[gid_y * get_global_size(0) + gid_x] = out;

}

[编辑] [发布完整内核,并根据建议交换行和列.这样我在AMD上获得了18%的表现,但在NVidia上获得了0%的表现.原始代码是

output[get_global_id(0) * get_global_size(1) + get_global_id(1)] = out;

[/编辑]

我在我的Nvidia Quadro 1000M上运行它,它有2个计算单元和96个CUDA核心(每个计算单元48个核心).

在排队内核时,我正在改变本地组大小.这些是我在生成400万像素图像时获得的不同尺寸的性能结果.所有数字都来自OpenCL分析器,并将最终内存副本排除回操作系统.图像为40992x10272 - 高度和宽度均可被48整除.

rows x columns
8x8: 397 MPixel/s
8x12: 505 MPixel/s
8x16: 523 MPixel/s
8x24: 521 MPixel/s
8x32: 520 MPixel/s
8x48: 520 MPixel/s

1x48: 321 MPixel/s
2x32: 424 MPixel/s
2x48: 523 MPixel/s
4x24: 519 MPixel/s
3x32: 525 MPixel/s
4x32: 525 MPixel/s
4x48: 525 MPixel/s

12x8: 490 MPixel/s
12x12:464 MPixel/s
12x24:505 MPixel/s
12x32:508 MPixel/s
12x48:433 MPixel/s

16x8: 499 MPixel/s
16x12:499 MPixel/s
16x16:472 MPixel/s
16x24:450 MPixel/s
16x32:440 MPixel/s
16x48:418 MPixel/s

其中一些数字让我感到困惑.虽然很明显为什么我用48列获得最佳结果(感谢SIMD操作如何工作),但我不明白:

1. 当我每组使用16行时,为什么性能会急剧下降？
2. 为什么我的1x48表现不佳？
3. 为什么在天堂我能获得3x32,4x32和8x32的顶级性能？!？我原以为33%的SIMD处理器处于空闲状态,而工作组看起来就像是在两个计算单元之间？
4. 为什么PREFERRED_WORK_GROUP_SIZE_MULTIPLE返回32而不是48？
5. 是否有一种非实证的方法可以在任何GPU(ATI/Nvidia/Intel HD)上找出最佳性能的几何结构,只考虑我从OpenCL信息结构中获得的内容？

提前致谢

Answer 1

我在这里回答了一个类似的问题,在阅读以下文章之前你可能会感兴趣.

当我每组使用16行时,为什么性能会急剧下降？

实际上,当你使用12行时,它已经降级了.内存访问按事务处理.事务将一次性获取一定数量的字节.现在,如果多个工作项尝试访问数组中的几个连续元素,则意味着一个事务可能足以为它们提供服务.

因为您以这种方式访问​​内存:

output[get_global_id(0) * get_global_size(1) + get_global_id(1)] = out;

这意味着本地大小在维度0中越大,事务的数量就越大,因为您必须访问非连续元素(由get_global_size(1)元素分隔).全局内存访问很昂贵.

因此,对于12/16行,您至少需要12/16个事务.这导致了你的第二个问题:

为什么我的1x48表现不佳？

基于我之前刚才所说的,似乎性能应该很好,因为事务的数量很少.

但是这里出现了空转线程的问题.关于每个SM的48个核心所获得的信息是错误的,正如其他人已经指出的那样.在NVIDIA硬件上,线程在32组(在NVIDIA中称为warp)中执行.请注意,这些组称为wavefront,AMD最多可以有64个线程.由于在这种情况下您有一个由48个线程(1乘48)组成的工作组,这意味着计划了64个线程.由于您无法执行一小部分扭曲,因此它始终是多个32的多个线程.

因此,在这种情况下,您有四分之一的线程什么都不做.实际上,当您与2x32(仍为64个线程 - 2个经线,但已充分利用)获得的结果进行比较时,321 MPixel/s几乎是424 MPixel/s的3/4.

值得注意的是这个结果:2x48:523 MPixel/s.在这种情况下,您的工作组大小为96,是32的倍数.因此没有空闲线程.

为什么在天堂我能获得3x32,4x32和8x32的顶级性能？!？

好吧,答案来自之前的两个:你使用32的倍数,并保持维度0中的线程数相对较小.但是让我们仔细看看你的结果:

2x32:  424 MPixel/s
3x32:  525 MPixel/s
4x32:  525 MPixel/s
8x32:  520 MPixel/s
16x32: 440 MPixel/s

最后两行的性能下降很容易用上述内容解释.但是,第一行和第二行之间的性能提高不是.

在这种情况下,性能的提高来自其他地方.事实上,在第二种情况下,足够的warp在同一SM上运行以隐藏访问内存延迟.您会看到REFERRED_WORK_GROUP_SIZE_MULTIPLE值仅表示您应尝试使用此值的MULTIPLE以获得最佳性能.可以在同一个SM上同时安排几个warp.

那么它是怎样工作的？我们来看3x32案例吧.您有一个由3个warp组成的工作组.由于它们属于同一工作组,因此它们按照OpenCL标准的要求安排在相同的SM上(如果不是这种情况,则无法在工作组内的线程之间进行同步).

第一个warp开始运行,直到它停止,因为需要内存访问.同时warp 1等待内存事务完成,warp 2可以开始运行.由于SM上有很多寄存器,因此SM可以轻松快速地切换上下文以运行其他warp.warp 1的所有变量都保留在分配给warp 1的寄存器上.然后warp 2命中了需要内存访问的行并停止.那一刻,下一个准备好运行的warp就可以开始运行了.如果内存访问完成,它可能是warp 3,也可能是warp 1.在你的情况下,它似乎是warp 3运行,因为你有2x32和3x32之间的差异.在第一种情况下,没有足够的warp被安排隐藏内存访问,尽管在第二种情况下有.

事实上,这个影响以及问题2中1x48尺寸的不良表现.

为什么PREFERRED_WORK_GROUP_SIZE_MULTIPLE返回32而不是48？

已经回答了.

是否有一种非实证的方法可以在任何GPU(ATI/Nvidia/Intel HD)上找出最佳性能的几何结构,只考虑我从OpenCL信息结构中获得的内容？

它就像任何其他语言一样.如果您知道它是如何工作的,它可以帮助您生成良好的第一个代码.但是你仍然需要对它进行基准测试,并经历一个试验和错误的过程来调整它.记住我刚刚写的东西只是对性能至关重要的一小部分.从OpenCL查询一些信息并结合对CPU/GPU的良好理解显然会有所帮助......但就是这样.

因为很多影响性能的参数都是对手,所以你在另一方获得的东西会在另一方面失去.

因此,保持基准;).