PyOpenCL示例－05.数值归并

>>数值归并例子2

在OpenCL里，有一个float4向量数据类型，它一次存4个浮点数。我们可以把8192个元素分成2048个float4来进行归并计算。最后再把向量中4个值相加。这种方法大大提高运算效率。速度是之前的4倍多。

主程序（main.py）

import pyopencl as cl
import numpy as np

if __name__ == '__main__':
    
    #要计算的元素个数
    Array_size = 8192    
    
    #step 0:获取GPU计算设备
    platform  = cl.get_platforms()[0]
    devices = platform.get_devices()
    
    #step 1:选择设备并创建上下文
    context = cl.create_some_context() 
    
    #step 2:创建命令队列，选择计算设备
    queue = cl.CommandQueue(context,devices[0],cl.command_queue_properties.PROFILING_ENABLE)
    
    #获取系统信息,我的K600中max_work_group_size=1024
    max_work_group_size = devices[0].get_info(cl.device_info.MAX_WORK_GROUP_SIZE)
    
    #分组
    num_groups = Array_size // max_work_group_size
    
    #创建一个包含8192个1的浮点数组
    buf_num = np.ones(Array_size,dtype=np.float32)
    #存储每个向量的结果
    buf_out = np.zeros(num_groups//4,dtype=np.float32)
    
    num_buf = cl.Buffer(context, cl.mem_flags.READ_ONLY | cl.mem_flags.COPY_HOST_PTR, hostbuf=buf_num)
    out_buf = cl.Buffer(context, cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, hostbuf=buf_out)
    
    #为每个组准备一个存放中间计算值的空间,由于是4个浮点数的向量，所以空间4倍
    tmp_buf = cl.LocalMemory(Array_size*4)
    
    #加载并创建CL程序
    f_cl = open("reduction_vector.cl","r")
    f_buf = f_cl.read()
    prg = cl.Program(context,f_buf).build()
     
    #运行
    prg.reduction_vector(queue,(Array_size//4,),(max_work_group_size,),num_buf,tmp_buf,out_buf)
    
    #读取计算值
    cl.enqueue_copy(queue,buf_out,out_buf)
    
    sum=0
    for i in range(0,num_groups//4):
        sum += buf_out[i]
        
    print(sum)
    
    num_buf.release()
    out_buf.release()
    
    queue.finish()
    

内核程序（reduction_vector.py）

__kernel void reduction_vector(__global float4* data, __local float4* partial_sums, __global float* output){
         int cid = get_local_id(0);
         int group_size = get_local_size(0);
                 
         partial_sums[cid] = data[get_global_id(0)];
         barrier(CLK_LOCAL_MEM_FENCE);
                 
         for(int i=group_size/2;i>0;i/=2){
             if(cid<i){
                 partial_sums[cid] += partial_sums[cid+i];
             } 
             barrier(CLK_LOCAL_MEM_FENCE);
         }
                 
         if(cid==0){
             output[get_group_id(0)] = dot(partial_sums[0], (float4)(1.0f));
         }                        
    }