我正在学习在python中使用opencl,我想优化其中一个函数。我了解到,这可以通过在本地内存中存储全局内存来完成。但是,它不能正常工作,它的持续时间是原来的两倍。干得好吗?我可以对这段代码进行更多的优化吗?
__kernel void sumOP( __global float *input,
__global float *weights,
int layer_size,
__global float *partialSums,__local float* cache)
{
private const int i = get_global_id(0);
private const int in_layer_s = layer_size;
private const int item_id = get_local_id(0);
private const int group_id = get_group_id(0);
private const int group_count = get_num_groups(0);
const int localsize = get_local_size(0);
for ( int x = 0; x < in_layer_s; x++ )
{
cache[x] = weights[i*in_layer_s + x];
}
float total1 = 0;
for ( int x = 0; x < in_layer_s; x++ )
{
total1 += cache[x] *input[x];
}
partialSums[i] = sigmoid(total1);
}
Python调用
l = opencl.LocalMemory(len(inputs))
event = program.sumOP(queue, output.shape, np.random.randn(6,).shape, inputs.data, weights.data,np.int32(len(inputs)),output.data,l)
谢谢你的建议
发布于 2017-09-29 13:48:34
除了通过组的所有工作项(如Dithermaster所说)写入相同共享内存地址cache[x]
的数据写入争用条件和缺少屏障()函数之外,还可以在修复后添加一些优化:
内核中的第一个循环
for ( int x = 0; x < in_layer_s; x++ )
{
cache[x] = weights[i*in_layer_s + x];
}
扫描每个工作项的不同内存区域,每次只扫描一个元素。这在全局内存性能方面可能是错误的,因为每个工作项在自己的循环中都可能使用相同的内存通道,甚至是相同的内存库,因此,所有工作项都依次访问该通道或银行。如果in_layer_s得到更大的值,特别是当它是2的幂时,情况就更糟了。为了解决这个问题,所有的工作项都应该访问相邻的连续地址。当全局内存与工作项一致访问时,GPU工作得更好。在本地内存中,随机访问或在工作项之间存在空白的问题较少。这就是为什么它建议在全局上使用统一的保存/加载,而在本地进行随机/分散/聚集。
内核中的第二循环
for ( int x = 0; x < in_layer_s; x++ )
{
total1 += cache[x] *input[x];
}
只使用单个累加器。这是一个依赖链,需要在进入下一个循环之前完成每个循环周期。至少使用2个临时“总计”变量并展开循环。在这里,如果in_layer_s足够小,则可以将input
数组移动到本地或常量内存中,以更快地访问它(所有工作项都重复访问它,因为所有工作项都访问相同的输入数组)(可能一半输入到常量内存,另一半输入本地内存以增加总带宽)。
weights[i*in_layer_s + x];
是一个结构数组吗?如果是的话,您可以通过使其成为数组的结构来实现加速,并且完全摆脱第一个循环的优化,同时增加主机端的代码膨胀,但是如果优先级是速度,那么在gpu端数组的结构是更快和更易读的。这也使得只从主机端向gpu上传必要的权重数据(SOA数组)成为可能,从而进一步减少了总延迟(上载+计算+下载)。
您还可以尝试异步本地<->全局传输函数,使每个工作项组的加载和计算重叠,作为最后手段隐藏更多的延迟。copy.html
https://stackoverflow.com/questions/46489375
复制相似问题