OpenCL 工作组函数

在OpenCL中，提供了一些工作组级别的内建函数，下表列出了这些函数。执行内核的一个工作组内所有的内核都要执行这些内建函数。在开始这些函数之前，我们对表格中的数据类型做一些约定：用通用类型名gentype指示这些数学函数参数可以取half、int、uint、long、ulong、f loat或double类型。表格中的〈op〉指的是add、min或max。

假设a[4]={0,1,2,3,4},只有一个工作组，工作组内有4个工作项，如上几个工作组函数，我们通过一个例子来讲解它们的用法：

kernel void DoubleTest(global int ＊a, global int ＊b)
{
    int id = get_global_id(0);
    int predicate = a[id];
    b[0] = work_group_reduce_add(predicate);
    b[1] = work_group_reduce_min(predicate);
    b[2] = work_group_reduce_max(predicate);
    b[3] = work_group_any(predicate);
    b[4] = work_group_all(predicate);
    int value = work_group_scan_exclusive_add(predicate);
    int value1 = work_group_scan_inclusive_add(predicate);
    if(id == 3)
    {
        b[5] = value;
        //b[5]= work_group_scan_exclusive_add(predicate); 错误
        b[6] = value1;
    }
    int global_id;
    if(0 == id)
    {
        global_id = 100;
    }
    global_id = work_group_broadcast(global_id, 0);
    if(id == 3)
b[7] = global_id;
}

上述代码中：

• 对于work_group_reduce_add（predicate），返回的是b[0]=0+1+2+3=6。
• 对于work_group_reduce_min（predicate），返回的是最小值b[1]=0。
• 对于work_group_reduce_max（predicate），返回的是最大值b[2]=3。
• 对于work_group_any（predicate），因为工作项ID为1、2、3的predicate值非零，故返回值b[3]=1。
• 对于work_group_all（predicate），因为工作项ID为0的predicate值为零，故返回值b[4]=0。
• 对于work_group_scan_exclusive_add（predicate），如图4-6所示，b[5]=3。
• 对于work_group_scan_inclusive_add（predicate），如图4-7所示，b[6]=6。
• 对于work_group_broadcast（global_id,0），广播global_id值，所以b[7]=100。