OpenCL 第6课:矩阵转置

上一节我们写了个一维向量相加的程序。这节我们来看一个4×4矩阵转置程序。

4X4矩阵我们采用二维数组进行存储，在程序设计上，我们让转置过程分4次转置完成，就是一次转一行。注意这里的OpenCL的工作维数是二维。（当然用一维的方式也可以，只是在CL代码中要用到循环，效率不高）

程序分两部份：

（1）transposition.cl代码

?
__kernel void transposition(__global int* A,
                    __global int* B)
{
    //获取索引号，这里是二维的，所以可以取两个
    //否则另一个永远是0
    int col = get_global_id(0);
    int row = get_global_id(1);
    B[col*4+row] = A[row*4+col];
}

（2）main.cpp代码

?
#include <iostream>
#include <stdio.h>
#include <string.h>
#include <string>
#include <CL/cl.h>//包含CL的头文件
using namespace std;
//4x4数组
#define dim_x 4
#define dim_y 4
//从外部文件获取cl内核代码
bool GetFileData(const char* fname,string& str)
{
    FILE* fp = fopen(fname,"r");
    if(fp==NULL)
    {
        printf("no found file\n");
        return false;
    }
    int n=0;
    while(feof(fp)==0)
    {
        str += fgetc(fp);
    }
    return true;
}
int  main()
{
    //先读外部CL核心代码，如果失败则退出。
    //代码存buf_code里面
    string code_file;
    if(false == GetFileData("transposition.cl",code_file))
        return 0;
    char* buf_code = new char[code_file.size()];
    strcpy(buf_code,code_file.c_str());
    buf_code[code_file.size()-1] = NULL;
    //声明CL所需变量。
    cl_device_id device;
    cl_platform_id platform_id = NULL;
    cl_context context;
    cl_command_queue cmdQueue;
    cl_mem bufferA,bufferB,bufferC;
    cl_program program;
    cl_kernel kernel = NULL;
    //我们使用的是二维向量
    //设定向量大小（维数）
    size_t globalWorkSize[2];
    globalWorkSize[0] = dim_x ;
    globalWorkSize[1] = dim_y;
    cl_int err;
    /*
        定义输入变量和输出变量，并设定初值
    */
    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;
        }
    }
    //step 1:初始化OpenCL
    err = clGetPlatformIDs(1,&platform_id,NULL);
    if(err!=CL_SUCCESS)
    {
        cout<<"clGetPlatformIDs error"<<endl;
        return 0;
    }
    //这次我们只用CPU来进行并行运算，当然你也可以该成GPU
    clGetDeviceIDs(platform_id,CL_DEVICE_TYPE_GPU,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**)&buf_code,
                                        NULL,NULL);
    clBuildProgram(program,1,&device,NULL,NULL,NULL);
    kernel = clCreateKernel(program,"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++)
        {
            cout<< buf_A[m][n] <<",";
        }
        cout<<endl;
    }
    cout<<endl<<"====transposition===="<<endl<<endl;
    for(n=0;n<dim_x;n++)
    {
        for(m=0;m<dim_y;m++)
        {
            cout<< buf_B[m][n] <<",";
        }
        cout<<endl;
    }
    //释放所有调用和内存
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseCommandQueue(cmdQueue);
    clReleaseMemObject(bufferA);
    clReleaseMemObject(bufferB);
    clReleaseContext(context);
    delete buf_code;
    return 0;
}