Opencl - 减少工作项数量的速度
Opencl - reduction in the rate of the number of work-item
我正在尝试为多线程计算 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,时间会略有增加,但时间增加了两倍。为什么?
(我很抱歉我的英语很糟糕)。
while( j != i )
len += lenghtIn[ j++ ];
部分运行 get_global_id( 0 ) 次。
当它是 128 时,最近要完成的工作项正在执行 128 次循环迭代。
当它是 256 时,它正在进行 256 次迭代,因此从内存的角度来看应该增加 %100,但仅限于最后一个工作项。当我们整合所有工人的总内存访问次数时,
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²)) 的影响。
我正在尝试为多线程计算 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,时间会略有增加,但时间增加了两倍。为什么? (我很抱歉我的英语很糟糕)。
while( j != i )
len += lenghtIn[ j++ ];
部分运行 get_global_id( 0 ) 次。
当它是 128 时,最近要完成的工作项正在执行 128 次循环迭代。
当它是 256 时,它正在进行 256 次迭代,因此从内存的角度来看应该增加 %100,但仅限于最后一个工作项。当我们整合所有工人的总内存访问次数时,
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²)) 的影响。