MVT*_*VTC 2 c++ cuda reduction opencl
我已经经历了一些例子,将一个元素数组减少到一个元素,但没有成功.有人在NVIDIA论坛上发布了此消息.我已经从浮点变量更改为整数.
__kernel void sum(__global const short *A,__global unsigned long *C,uint size, __local unsigned long *L) {
unsigned long sum=0;
for(int i=get_local_id(0);i<size;i+=get_local_size(0))
sum+=A[i];
L[get_local_id(0)]=sum;
for(uint c=get_local_size(0)/2;c>0;c/=2)
{
barrier(CLK_LOCAL_MEM_FENCE);
if(c>get_local_id(0))
L[get_local_id(0)]+=L[get_local_id(0)+c];
}
if(get_local_id(0)==0)
C[0]=L[0];
barrier(CLK_LOCAL_MEM_FENCE);
}
这看起来不错吗？第三个参数"大小",应该是本地工作规模,还是全球工作规模？
我设置了这样的论点,
clSetKernelArg(ocReduce, 0, sizeof(cl_mem), (void*) &DevA);
clSetKernelArg(ocReduce, 1, sizeof(cl_mem), (void*) &DevC);
clSetKernelArg(ocReduce, 2, sizeof(uint), (void*) &size);
clSetKernelArg(ocReduce, 3, LocalWorkSize * sizeof(unsigned long), NULL);
第一个参数是输入,我试图保留从之前启动的内核的输出.
clRetainMemObject(DevA);
clEnqueueNDRangeKernel(hCmdQueue[Plat-1][Dev-1], ocKernel, 1, NULL, &GlobalWorkSize, &LocalWorkSize, 0, NULL, NULL);
//the device memory object DevA now has the data to be reduced
clEnqueueNDRangeKernel(hCmdQueue[Plat-1][Dev-1], ocReduce, 1, NULL, &GlobalWorkSize, &LocalWorkSize, 0, NULL, NULL);
clEnqueueReadBuffer(hCmdQueue[Plat-1][Dev-1],DevRE, CL_TRUE, 0, sizeof(unsigned long)*512,(void*) RE , 0, NULL, NULL);
今天我打算尝试将以下cuda缩减示例转换为openCL.
__global__ voidreduce1(int*g_idata, int*g_odata){
extern __shared__ intsdata[];
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;
sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];
__syncthreads();
for(unsigned int s=blockDim.x/2; s>0; s>>=1) {
if (tid < s) {
sdata[tid] += sdata[tid + s];
}
__syncthreads();
}
// write result for this block to global mem
if(tid == 0) g_odata[blockIdx.x] = sdata[0];
}
有一个更优化的(每个线程完全展开+多个元素).
http://developer.download.nvidia.com/compute/cuda/1_1/Website/projects/reduction/doc/reduction.pdf
这可能使用openCL吗？
前几天Grizzly给了我这个建议,
"......使用一个对n元素进行操作的简化内核,并将它们缩减为n/16(或任何其他数字).然后你迭代地调用那个内核直到你得到一个元素,这就是你的结果"
我也想尝试这个,但我不知道从哪里开始,我想先得到一些工作.
只要只有一个工作组正在进行减少(因此get_global_size(0) == get_local_size(0)),您提供的第一个减少代码应该有效.在这种情况下,size内核的参数将是元素的数量A(与全局或本地工作量没有真正的相关性).虽然这是一个可行的解决方案,但在进行简化时让大部分gpu空闲时间看起来仍然是浪费,这正是我提出迭代调用还原内核的原因.只需对代码稍作修改即可实现:
__kernel void sum(__global const short *A, __global unsigned long *C, uint size, __local unsigned long *L) {
unsigned long sum=0;
for(int i=get_global_id(0); i < size; i += get_global_size(0))
sum += A[i];
L[get_local_id(0)]=sum;
for(uint c=get_local_size(0)/2;c>0;c/=2)
{
barrier(CLK_LOCAL_MEM_FENCE);
if(c>get_local_id(0))
L[get_local_id(0)]+=L[get_local_id(0)+c];
}
if(get_local_id(0)==0)
C[get_group_id(0)]=L[0];
barrier(CLK_LOCAL_MEM_FENCE);
}
用GlobalWorkSize较小的size(例如4)调用它会将输入减少A一个因子4*LocalWorkSize,可以迭代(通过使用输出缓冲区作为下一次调用sum具有不同输出缓冲区的输入.实际上这不是真的,因为第二个(以及所有后续)迭代需要A是类型global const unsigned long*,所以你实际上需要内核,但你明白了.
关于cuda减少样本:你为什么要费心转换它,它的工作原理与我上面发布的opencl版本完全一样,除了每次迭代只减少一个硬编码的大小(2*LocalWorkSizeinsted of size/GlobalWorkSize*LocalWorkSize).
我个人使用几乎相同的方法进行缩减,虽然我将内核拆分为两部分,并且仅使用本地内存的路径进行最后一次迭代:
__kernel void reduction_step(__global const unsigned long* A, __global unsigned long * C, uint size) {
unsigned long sum=0;
for(int i=start; i < size; i += stride)
sum += A[i];
C[get_global_id(0)]= sum;
}
对于最后一步,使用了在工作组内进行缩减的完整版本.当然你需要第二个版本的reduction step服用global const short*,这段代码是对你的代码的未经测试的调整(我不能发布我自己的版本,令人遗憾).这种方法的优点是内核完成大部分工作的复杂性要小得多,而且wasted work由于分支分支的数量较少.这使它比其他变体快一点.但是我对最新的编译器版本和最新的硬件都没有结果,因此该点可能会或可能不再正确(尽管我怀疑它可能因为减少了不同的分支数量).
现在为您链接的文章:除了使用opencl不支持的模板之外,当然可以在opencl中使用该文章中建议的优化,因此块大小必须是硬编码的.当然opencl版本已经为每个内核添加了多个,如果你按照我上面提到的方法,通过本地内存展开减少并不会真正受益,因为这只能在最后一步完成,不应该采用足够大的输入的整个计算时间的重要部分.此外,我发现在展开的实现中缺乏同步有点麻烦.这只能起作用,因为进入该部分的所有线程都属于同一个warp.然而这不是'