OpenCL 第6课:矩阵转置

来源：互联网发布：萌化软件编辑：程序博客网时间：2024/05/17 01:50

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

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

程序分两部份：

（1）transposition.cl代码

?
1
2
3
4
5
6
7
8
9
10
11
__kernel voidtransposition(__global int* A,
                    __globalint* B)
{
    //获取索引号，这里是二维的，所以可以取两个
    //否则另一个永远是0
    intcol = get_global_id(0);
    introw = get_global_id(1);
 
    B[col*4+row] = A[row*4+col];
 
}

（2）main.cpp代码

?
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
#include <iostream>
#include <stdio.h>
#include <string.h>
#include <string>
#include <CL/cl.h>//包含CL的头文件
 
usingnamespace std;
 
//4x4数组
#define dim_x 4
#define dim_y 4
 
//从外部文件获取cl内核代码
boolGetFileData(constchar* fname,string& str)
{
    FILE* fp =fopen(fname,"r");
    if(fp==NULL)
    {
        printf("no found file\n");
        returnfalse;
    }
 
    intn=0;
    while(feof(fp)==0)
    {
        str +=fgetc(fp);
    }
 
    returntrue;
}
 
int main()
{
    //先读外部CL核心代码，如果失败则退出。
    //代码存buf_code里面
    string code_file;
 
    if(false== GetFileData("transposition.cl",code_file))
        return0;
 
    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_tglobalWorkSize[2];
    globalWorkSize[0] = dim_x ;
    globalWorkSize[1] = dim_y;
 
    cl_int err;
 
    /*
        定义输入变量和输出变量，并设定初值
    */
    intbuf_A[dim_x][dim_y];
    intbuf_B[dim_x][dim_y];
 
    size_tdatasize = sizeof(int) * dim_x * dim_y;
 
    intn=0;
    intm=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;
        return0;
    }
 
    //这次我们只用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,
                                        (constchar**)&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);
 
    deletebuf_code;
 
    return0;
}

运算结果：

123

0 0