Opencl - 减少工作项数量的比率

Question

我正在尝试为多线程计算 crc32。我正在尝试使用 OpenCL。 GPU代码是：

__kernel void crc32_Sarwate( __global int* lenghtIn, 
                         __global unsigned char *In, 
                         __global int *OutCrc32,
                            int size ) {
int i, j, len;

i = get_global_id( 0 ); 
if( i >= size )
    return;
len = j = 0;
while( j != i )
    len += lenghtIn[ j++ ];
OutCrc32[ i ] = crc32( In + len, lenghtIn[ i ] ); }

我收到了这个结果（时间），重复了一千次：
4 人使用工作项目：29.82
8 人使用工作项目：29.9
16 人使用工作项目：35.16
对于 32 使用工作项：35.93
对于 64 使用工作项：38.69
对于 128 使用工作项：52.83
对于 256 使用工作项：152.08
对于 512 使用工作项：333.63

我有 350 MHz 的英特尔高清显卡和 3 个工作组和 256 个工作项 每个工作组。 我假设通过将工作项 128 增加到 256 的数量会稍微增加时间，但时间增加了两倍。为什么？ （对不起，我的英语很糟糕）。

Answer 1

while( j != i )
    len += lenghtIn[ j++ ];

部分运行 get_global_id( 0 ) 次。

当它是 128 时，最新完成的工作项正在做 128 次循环迭代。

当它是 256 时，它正在执行 256 次迭代，因此从内存的角度来看，它应该增加 %100，但仅限于最后一个工作项。当我们整合所有worker的总内存访问次数时，

 1 item from 0 to 0 ---> 1 access
 2 item from 0 to 0 and 0 to 1  ---> 3 access
 4 item from 0 to 0 and 0 to 1 and 0 to 2 and 0 to 3---> 10 access
 8 items: SUM(1 to 8) => 36 accesses  
 16 items: SUM(1 to 16) => 136 accesses (even more than + %200)
 32 items: => 528 (~ %400)
 64 items: => 2080 ( ~%400)
 128 items: => 8256 (~%400) (cache of your igpu starts failing here)
 256 items: => 32896 (~400%) (now caching is saturated and you start )
                              ( seeing %400 per doubling of work items)

 512 => uses second compute unit too! But %400 work is done 
  so it is not only %200 time consuming.

因此，每次将工作项增加 %100 时，都会增加总内存 访问 %400 。但是缓存在某种程度上有所帮助。当你越过它时，内存访问会严重增加。此外，执行开销（驱动程序，..）变得不重要。

您正在非并行访问内存。您需要先缓存它，但在该硬件中可能无法实现，因此您应该在工作项之间平均分配作业，并使内核之间的内存访问连续（矢量化）。这应该会提供更高的性能。

目前，每个向量单元都可以：

unit        :   v0 v1 v2 v3 v4 ... v7
read address:   0  0  0  0  0      0
                -  1  1  1  1      1
                -  -  2  2  2      2
                -  -  -  3  3      3
                -  -  -  -  4      4
                     ....  
                -  -  -  -  - ...  7

在 8 个流式内核上分 8 个步骤完成。

在最后一步，只有单个工作项实际上正在计算一些东西。这应该是这样的：

一些优化

unit        :   v0 v1 v2 v3  no need other work items
read address:   0  0  0  0  \
                1  1  1  1   \
                2  2  2  2    \
                3  3  3  3    / this is 5th work item's work
                4  4  4  4   /
                5  5  5  0   \
                6  6  0  1    \ this is 0 to 3 as 4th work
                7  0  1  2    /
 first item<--  0  1  2  3   /

仅在 4 个流核心中分 8 个步骤完成，并且第一个执行相同的工作 一半（可能更快）。

进一步优化建议

我认为在进入 crc32() 部分之前，在另一个内核上使用前缀扫描（总和）算法会更好。 （这个例子可能只需要 3 步，而不是 8 步，而且效率更高）

使用预先计算的值

while( j != i )
        len += lenghtIn[ j++ ];

应该使 crc32 不受当前算法复杂度 (O(n²)) 的影响。