如何并行化OpenCL程序：一个数据处理的例子
如何并行化OpenCL程序：一个数据处理的例子
==============================================================
目录结构
1、滑动平均滤波器C程序
2、滑动平均滤波器OpenCL程序
3、滑动平均滤波器OpenCL程序向量操作（Vector Operations）
4、滑动平均滤波器OpenCL程序——基于向量操作
5、滑动平均滤波器OpenCL程序——数据并行模式
6、滑动平均滤波器OpenCL程序——任务并行模式
7、参考
==============================================================
关键词：并行化 OpenCL
本节将通过从标准的C语言程序到OpenCLC程序例子分析股票价格数据，功能是通过对不同的股票数据进行滑动平均处理，这将对如何并行化代码有所帮助。
滑动平均滤波器是一个低通滤波器，主要用于信号处理和图像处理，简单讲就是求和取平均。
1、滑动平均滤波器C程序
主函数：moving_average.c
void moving_average(int *values,
float *average,
int length,
int width)
{
int i;
int add_value;
/* Compute sum for the first"width" elements */
add_value = 0;
for (i=0; i < width; i++) {
add_value += values[i];
}
average[width-1] = (float)add_value;
/* Compute sum for the (width)th ～(length-1)th elements */
for (i=width; i < length; i++) {
add_value = add_value -values[i-width] + values[i];
average[i] =(float)(add_value);
}
/* Insert zeros to 0th ～(width-2)th element */
for (i=0; i < width-1; i++) {
average[i] = 0.0f;
}
/* Compute average from the sum */
for (i=width-1; i < length; i++){
average[i] /=(float)width;
}
}此处不关注语法的情况下，主要看下该函数的功能：
（1）将前width个数求和；
（2）将width到数据长度length的数求和；
（3）初始化前width个数的平均值为0（满width个数才求均值）；
（4）计算width到length数据的数的平均值。
测试数据：stock_array1.txt
100,
109,
98,
104,
107,
...
50
主函数：main.cpp
#include<stdio.h>
#include<stdlib.h>
/*Read Stock data */
intstock_array1[] = {
#include"stock_array1.txt"
};
/*Define width for the moving average */
#define WINDOW_SIZE (13)
int main(int argc, char *argv[])
{
float*result;
int data_num = sizeof(stock_array1) / sizeof(stock_array1[0]);
int window_num = (int)WINDOW_SIZE;
int i;
/*Allocate space for the result */
result= (float *)malloc(data_num*sizeof(float));
/*Call the moving average function */
moving_average(stock_array1,
result,
data_num,
window_num);
/*Print result */
for(i=0; i < data_num; i++) {
printf("result[%d]= %f￥n", i, result[i]);
}
/*Deallocate memory */
free(result);
}2、滑动平均滤波器OpenCL程序
将1中的C语言的核与主函数分别翻译为OpenCL程序，如下。
核函数：moving_average.cl
__kernel void moving_average(__global int *values,
__global float *average,
int length,
int width)
{
int i;
int add_value;
/* Compute sum for the first"width" elements */
add_value = 0;
for (i=0; i < width; i++) {
add_value += values[i];
}
average[width-1] = (float)add_value;
/* Compute sum for the (width)th ～(length-1)th elements */
for (i=width; i < length; i++) {
add_value = add_value -values[i-width] + values[i];
average[i] =(float)(add_value);
}
/* Insert zeros to 0th ~ (width-2)thelements */
for (i=0; i < width-1; i++) {
average[i] = 0.0f;
}
/* Compute average of (width-1) ~(length-1) elements */
for (i=width-1; i < length; i++){
average[i] /=(float)width;
}
}测试数据：stock_array1.txt
主函数：main.cpp
https://github.com/yywyz/OpenCL-Programming-Examples/blob/master/moving_average_filter1/main.cpp
3、滑动平均滤波器OpenCL程序向量操作（VectorOperations）
为了加速计算，可以将多个股票数据导入SIMD单元，该计算单元可以并行操作。假设该SIMD单元有128比特的SIMD单元，并行操作4个32比特的数据，在OpenCL中可以使用int4, float4数据类型。
多个股票数据（stock_array_many.txt）
100, 212, 315, 1098, 763, 995, ...,12
109, 210, 313,1100, 783, 983, ..., 15
98, 209, 310, 1089,790, 990, ..., 18
104, 200, 319,1098, 792, 985, ..., 21
107, 100, 321,1105, 788, 971, ..., 18
…
50, 33, 259, 980, 687, 950, ..., 9
4个股票数据（stock_array_4.txt）
100, 212, 315, 1098,
109, 210, 313,1100,
98, 209, 310, 1089,
104, 200, 319,1098,
107, 100, 321,1105,
…
50, 33, 259, 980
1次处理4值数据，可以将原核函数int和float类型分别换成int4和float4，如下。
核函数：moving_average_vec4.cl
__kernel void moving_average_vec4(__global int4 *values,
__global float4 *average,
int length,
int width)
{
int i;
int4 add_value; /* A vector to hold4 components */
/* Compute sum for the first"width" elements for 4 stocks */
add_value = (int4)0;
for (i=0; i < width; i++) {
add_value += values[i];
}
average[width-1] =convert_float4(add_value);
/* Compute sum for the (width)th ～(length-1)th elements for 4 stocks */
for (i=width; i < length; i++) {
add_value = add_value -values[i-width] + values[i];
average[i] =convert_float4(add_value);
}
/* Insert zeros to 0th ～(width-2)th element for 4 stocks*/
for (i=0; i < width-1; i++) {
average[i] =(float4)(0.0f);
}
/* Compute average of (width-1) ~(length-1) elements for 4 stocks */
for (i=width-1; i < length; i++){
average[i] /= (float4)width;
}
}主函数: main.cpp
https://github.com/yywyz/OpenCL-Programming-Examples/blob/master/moving_average_filter2/main.cpp
4、滑动平均滤波器OpenCL程序——基于向量操作
假设处理的数据是4的倍数，即可以基于章节3中向量操作改进并行度，本次程序与3中程序不同主要在于核函数以及主函数修改调用核函数和数据存储大小。核函数如下：
核函数：moving_average_many.cl
__kernel void moving_average_many(__global int4 *values,
__global float4 *average,
int length,
int name_num,
int width)
{
int i, j;
int loop_num = name_num / 4; /* computethe number of times to loop */
int4 add_value;
for (j=0; j < loop_num; j++) {
add_value = (int4)0;
for (i=0; i < width;i++) {
add_value +=values[i*loop_num+j];
}
average[(width-1)*loop_num+j]= convert_float4(add_value);
for (i=width; i <length; i++) {
add_value =add_value - values[(i-width)*loop_num+j] + values[i*loop_num+j];
average[i*loop_num+j]= convert_float4(add_value);
}
for (i=0; i < width-1;i++) {
average[i*loop_num+j]= (float4)(0.0f);
}
for (i=width-1; i <length; i++) {
average[i*loop_num+j]/= (float4)width;
}
}
}主函数: main.cpp
https://github.com/yywyz/OpenCL-Programming-Examples/blob/master/moving_average_filter3/main.cpp
测试数据：stock_array_many.txt
5、滑动平均滤波器OpenCL程序——数据并行模式
这一步将使用SIMD的方法来扩展计算单元容限。只使用一个核函数，处理所有的并行数据。为了同时计算多个处理单元，必须实例化多个核函数。这些核函数可以是同一个核函数（数据并行），也可以是并行的不同核（任务并行）。这个例子将使用数据并行模型，因为这个方法对于并行处理更加合适。
例：基于上节处理8路股票数据。
由于处理4路股票数据一次操作即可完成，所以，可以使用2个计算单元处理8路数据，这将通过设置工作组为2来实现。
为了使用数据并行模式，每一个例化的核函数必须知道自己的地址索引。
get_global_id( )可用于获取核实例的全局ID，与上节核函数中的j类似。
核函数：moving_average_vec4_para.cl
__kernel void moving_average_vec4_para(__global int4 *values,
__global float4 *average,
int length,
int name_num,
int width)
{
int i, j;
int loop_num = name_num / 4; /*compute the number of times to loop */
int4 add_value;
j = get_global_id(0);
add_value = (int4)0;
for (i=0; i < width; i++) {
add_value +=values[i*loop_num+j];
}
average[(width-1)*loop_num+j] =convert_float4(add_value);
for (i=width; i < length; i++) {
add_value = add_value -values[(i-width)*loop_num+j] + values[i*loop_num+j];
average[i*loop_num+j] =convert_float4(add_value);
}
for (i=0; i < width-1; i++) {
average[i*loop_num+j] =(float4)(0.0f);
}
for (i=width-1; i < length; i++){
average[i*loop_num+j] /=(float4)width;
}
}主函数：main.cpp
https://github.com/yywyz/OpenCL-Programming-Examples/blob/master/moving_average_filter4/main.cpp
测试文件：stock_array_vec4_par.txt
6、滑动平均滤波器OpenCL程序——任务并行模式
在处理不同参数时，可通过本例实现并行计算，在主程序中配置不同的参数，如配置不同的滤波器的阶数。
核函数：moving_average_vec4.cl，该核函数与4路数据处理的核函数一致
__kernel void moving_average_vec4(__global int4 *values,
__global float4 *average,
int length,
int width)
{
int i;
int4 add_value; /* A vector to hold4 components */
/* Compute sum for the first"width" elements for 4 stocks */
add_value = (int4)0;
for (i=0; i < width; i++) {
add_value += values[i];
}
average[width-1] =convert_float4(add_value);
/* Compute sum for the (width)th ～(length-1)th elements for 4 stocks */
for (i=width; i < length; i++) {
add_value = add_value -values[i-width] + values[i];
average[i] =convert_float4(add_value);
}
/* Insert zeros to 0th ～(width-2)th element for 4 stocks*/
for (i=0; i < width-1; i++) {
average[i] =(float4)(0.0f);
}
/* Compute average of (width-1) ~(length-1) elements for 4 stocks */
for (i=width-1; i < length; i++){
average[i] /= (float4)width;
}
}主函数：main.cpp
https://github.com/yywyz/OpenCL-Programming-Examples/blob/master/moving_average_filter5/main.cpp
测试数据：stock_array_4.txt
main函数中特别需要注意的是：
Command_queue= clCreateCommandQueue(context, device_id,
CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &ret);
该函数中的CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE可以将当前需排队的任务发送给可用的计算单元，这可以避免任务排队进行。
6、总结
本节的例子逐渐展示了如何并行化程序的过程。
（1）滑动平均滤波器C程序
最原始的C语言算法程序，滤波器的乘加运算。
（2）滑动平均滤波器OpenCL程序
将原始C程序翻译为OpenCL的程序，此时只能处理1路数据。
（3）滑动平均滤波器OpenCL程序向量操作（Vector Operations）
在2的基础上，通过将核函数的数据类型改为int4、float4等数据向量的方式将程序并行化。此时1次可以同时处理4路数据。
（4）滑动平均滤波器OpenCL程序——基于向量操作
基于数据向量的类型，处理4的倍数的数据，如可处理8路数据。
（5）滑动平均滤波器OpenCL程序——数据并行模式
使用设置工作组的方式，利用数据并行模式处理数据。
（6）滑动平均滤波器OpenCL程序——任务并行模式
使用相同的核函数，在主程序中对核函数设置不同的参数，对相同的数据进行不同的操作。
7、参考
《The OpenCL Programming Boook》
如有错误，请指出，谢谢。
途次客
2018年5月20日