下面对4x4矩阵进行转置,调用openCL,4X4矩阵采用二维数组进行存储,在程序设计上,让转置过程分10次转置完成,就是一次转一行,因此OpenCL的工作维数是二维。
openCL程序分为两个部份,一部份是内核代码,负责具体算法。另一部份是主程序负责初始化OpenCL和准备数据。主程序加载内核代码,并按照既定方法进行运算。
核函数代码如下
__kernel void matrix_transposition(__global int* a, __global int* b)
{
int col = get_global_id(0);
int row = get_global_id(1);
b[col*4+row] = a[row*4+col];
}
主机端代码如下
int main()
{
cl_device_id device;
cl_platform_id platform_id = NULL;
cl_context context;
cl_command_queue cmdQueue;
cl_mem bufferA,bufferB;
cl_program program;
cl_kernel kernel = NULL;
size_t ocl_string_size;
char *ocl_string;
ocl_string = (char *)malloc(1024*1024);
//我们使用的是二维向量
//设定向量大小(维数)
size_t globalWorkSize[2];
globalWorkSize[0] = dim_x ;
globalWorkSize[1] = dim_y;
/*
定义输入变量和输出变量,并设定初值
*/
int buf_A[dim_x][dim_y];
int buf_B[dim_x][dim_y];
size_t datasize = sizeof(int) * dim_x * dim_y;
int n=0;
int m=0;
for(n=0;n<dim_x;n++)
{
for(m=0;m<dim_y;m++)
{
buf_A[m][n] = m + n*dim_x;
}
}
ocl_string_size = get_ocl_string("matrix_transposition.cl", ocl_string);
//step 1:初始化OpenCL
clGetPlatformIDs(1, &platform_id, NULL);
clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_ACCELERATOR, 1, &device, NULL);
//step 2:创建上下文
context = clCreateContext(NULL,1,&device,NULL,NULL,NULL);
//step 3:创建命令队列
cmdQueue = clCreateCommandQueue(context,device,0,NULL);
//step 4:创建数据缓冲区
bufferA = clCreateBuffer(context,
CL_MEM_READ_ONLY,
datasize,NULL,NULL);
bufferB = clCreateBuffer(context,
CL_MEM_WRITE_ONLY,
datasize,NULL,NULL);
//step 5:将数据上传到缓冲区
clEnqueueWriteBuffer(cmdQueue,
bufferA,CL_FALSE,
0,datasize,
buf_A,0,
NULL,NULL);
//step 6:加载编译代码,创建内核调用函数
program = clCreateProgramWithSource(context,1,
(const char**)&ocl_string,
&ocl_string_size,NULL);
clBuildProgram(program,1,&device,NULL,NULL,NULL);
kernel = clCreateKernel(program,"matrix_transposition",NULL);
//step 7:设置参数,执行内核
clSetKernelArg(kernel,0,sizeof(cl_mem),&bufferA);
clSetKernelArg(kernel,1,sizeof(cl_mem),&bufferB);
//<span style="color: #ff0000;"><strong>注意这里第三个参数已经改成2,表示二维数据。</strong></span>
clEnqueueNDRangeKernel(cmdQueue,kernel,
2,NULL,
globalWorkSize,
NULL,0,NULL,NULL);
//step 8:取回计算结果
clEnqueueReadBuffer(cmdQueue,bufferB,CL_TRUE,0,
datasize,buf_B,0,NULL,NULL);
//输出计算结果
for(n=0;n<dim_x;n++)
{
for(m=0;m<dim_y;m++)
{
printf("%d ", buf_A[m][n]);
}
printf("\n");
}
for(n=0;n<dim_x;n++)
{
for(m=0;m<dim_y;m++)
{
printf("%d ", buf_B[m][n]);
}
printf("\n");
}
//释放所有调用和内存
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseCommandQueue(cmdQueue);
clReleaseMemObject(bufferA);
clReleaseMemObject(bufferB);
clReleaseContext(context);
return 0;
}
将上述程序编译,结果如下
0 1 2 3
4 5 6 7
8 9 10 11
12 13 14 25
矩阵转置后
0 4 8 12
1 5 9 13
2 6 10 14
3 7 11 15
该矩阵的行和列转换成功,即第 n 行变成第 n 列