GPU并行计算OpenCL(2)——矩阵卷积
我们完成了helloworld之后,会感慨一个helloworld就要2-3百行代码。但由于OpenCL需要完成一系列启用GPU并行计算的过程,所以我们之后许多编码会和helloworld十分类似。比如这次,我们将实现非常常见的功能,矩阵的卷积计算,该功能便可以体会到GPU并行计算的强大之处了。
我们这次我们将输入信号作为一个矩阵,完成对信号的卷积,我们的例子就是将3*3的卷积核应用到8*8的输入信号矩阵中。然后获取一个6*6的矩阵。完成过程如下图:
该算法非常简单,步骤为:
1)将模版置于输入信号上。
2)将输入值与模版中的相应元素相乘。
3)将第二步的结果累加为一个和,写至相应的输出位置。
因此我们的.cl内核代码为:
__kernel void convolve(const __global uint *const input,
__constant uint* const mask,
__global float *const output,
const int inputWidth,
const int maskWidth){
const int x=get_global_id(0);
const int y=get_global_id(1);
uint sum=0;
for(int r=0;r<maskWidth;r++){
const int idxIntpm=(y+r)*inputWidth+x;
for(int c=0;c<maskWidth;c++){
sum+=mask[(r*maskWidth)+c]*input[idxIntpm+c];
}
}
output[y*get_global_size(0)+x]=sum;
}
这里的x,y为我们在并行计算的时候,获取的全局索引。其中get_global_id(i),表示获取第i维中当前的索引。下面的sum则为卷积的每个项的结果。(这里注意的是:虽然这里用y去获取第二维的索引,但我们的代码最后还是通过一维实现的,所以所有的y获得的都是0,x为0-35)。
接下来就是我们的宿主机代码(.cpp文件)了。这里的代码流程是与helloworld基本一致的。
1)创建上下文,创建上下文代码:
cl_context CreatContext(){
cl_int errNum;
cl_uint numPlatforms;
cl_platform_id firstPlatformID;
cl_context context=nullptr;
/*首先,选择要运行的OpenCL平台。 对于这个例子,我们
只需选择第一个可用平台。通常,我们也可以
查询所有可用平台,然后选择最合适的平台。
*/
errNum=clGetPlatformIDs(1, &firstPlatformID, &numPlatforms);
if (errNum != CL_SUCCESS || numPlatforms <= 0)
{
std::cerr << "Failed to find any OpenCL platforms." << std::endl;
return NULL;
}
/*接下来,在平台上创建一个OpenCL上下文。 尝试
创建基于GPU的上下文,如果失败,请尝试创建
基于CPU的上下文。
*/
//创建上下文需要的资源 属性 属性值 0结束
cl_context_properties contextProperties[]={
CL_CONTEXT_PLATFORM,
(cl_context_properties)firstPlatformID,
0
};
context=clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, nullptr, nullptr, &errNum);
if (errNum != CL_SUCCESS)
{
std::cout << "Could not create GPU context, trying CPU..." << std::endl;
context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU,
NULL, NULL, &errNum);
if (errNum != CL_SUCCESS)
{
std::cerr << "Failed to create an OpenCL GPU or CPU context." << std::endl;
return NULL;
}
}
return context;
}
2)创建命令队列,创建命令队列的函数:
cl_command_queue CreateCommandQueue(cl_context context, cl_device_id *device)
{
cl_int errNum;
cl_device_id *devices;
cl_command_queue commandQueue = NULL;
size_t deviceBufferSize = -1;
// 获取设备缓存的尺寸
errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &deviceBufferSize);
if (errNum != CL_SUCCESS)
{
std::cerr << "Failed call to clGetContextInfo(...,GL_CONTEXT_DEVICES,...)";
return NULL;
}
if (deviceBufferSize <= 0)
{
std::cerr << "No devices available.";
return NULL;
}
// 给设备缓存分配空间
devices = new cl_device_id[deviceBufferSize / sizeof(cl_device_id)];
errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, deviceBufferSize, devices, NULL);
if (errNum != CL_SUCCESS)
{
delete [] devices;
std::cerr << "Failed to get device IDs";
return NULL;
}
// 选择第一个设备和上下文,创建出一个命令队列
commandQueue = clCreateCommandQueue(context, devices[0], 0, NULL);
if (commandQueue == NULL)
{
delete [] devices;
std::cerr << "Failed to create commandQueue for device 0";
return NULL;
}
*device = devices[0];
delete [] devices;
return commandQueue;
}
3)创建程序对象,函数代码:
cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName)
{
cl_int errNum;
cl_program program;
std::ifstream kernelFile(fileName, std::ios::in);
if (!kernelFile.is_open())
{
std::cerr << "Failed to open file for reading: " << fileName << std::endl;
return NULL;
}
std::ostringstream oss;
oss << kernelFile.rdbuf();//oss输出kernelFile指向的流缓冲
std::string srcStdStr = oss.str();
const char *srcStr = srcStdStr.c_str();
//在context上下文上创建程序对象(字符串个数为1)
program = clCreateProgramWithSource(context, 1,
(const char**)&srcStr,
NULL, NULL);
if (program == NULL)
{
std::cerr << "Failed to create CL program from source." << std::endl;
return NULL;
}
//编译内核源码
errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
if (errNum != CL_SUCCESS)
{
// 输出编译错误信息
char buildLog[16384];
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
sizeof(buildLog), buildLog, NULL);
std::cerr << "Error in kernel: " << std::endl;
std::cerr << buildLog;
clReleaseProgram(program);//释放程序对象空间
return NULL;
}
return program;
}
到这里,我们的代码和helloworld是一致的。我们来看.cpp文件:
#include <iostream>
#include <OpenCL/OpenCL.h>
#include <fstream>
#include <sstream>
//输入矩阵长宽
const unsigned int inputSignalWidth = 8;
const unsigned int inputSignalHeight = 8;
//矩阵值
cl_int inputSignal[inputSignalWidth][inputSignalHeight] =
{
{3, 1, 1, 4, 8, 2, 1, 3},
{4, 2, 1, 1, 2, 1, 2, 3},
{4, 4, 4, 4, 3, 2, 2, 2},
{9, 8, 3, 8, 9, 0, 0, 0},
{9, 3, 3, 9, 0, 0, 0, 0},
{0, 9, 0, 8, 0, 0, 0, 0},
{3, 0, 8, 8, 9, 4, 4, 4},
{5, 9, 8, 1, 8, 1, 1, 1}
};
//输出长宽
const unsigned int outputSignalWidth = 6;
const unsigned int outputSignalHeight = 6;
//输出矩阵值
cl_float outputSignal[outputSignalWidth][outputSignalHeight];
const unsigned int maskWidth = 3;
const unsigned int maskHeight = 3;
cl_int mask[maskWidth][maskHeight] =
{
{1, 1, 1}, {1, 0, 1}, {1, 1, 1},
};
/*
1 选择平台创建OpenCL上下文
2 选择设备创建命令队列
3 加载内核文件(hello_world.cl)并构建到程序对象中
4 为内核函数创建hello_kernel()创建一个内核对象
5 为内核参数创建内存对象(result,a,b)
6 将待执行的内核排队
7 将内核结果读回结果缓冲区
*/
cl_context CreatContext();
cl_command_queue CreateCommandQueue(cl_context context, cl_device_id *device);
cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName);
bool CreateMemObjects(cl_context context, cl_mem memObjects[3], cl_int (*a)[inputSignalWidth], cl_int (*b)[maskWidth]);
void Cleanup(cl_context context, cl_command_queue commandQueue,cl_program program, cl_kernel kernel, cl_mem memObjects[3]);
int main(int argc, const char * argv[]) {
cl_context context=0;
cl_command_queue commandQueue=0;
cl_program program=0;
cl_device_id device=0;
cl_kernel kernel=0;
cl_mem memObjects[3]={0,0,0};
cl_int errNum;
//创建上下文
context=CreatContext();
if (context == NULL)
{
std::cerr << "Failed to create OpenCL context." << std::endl;
return 1;
}
//创建命令队列
commandQueue=CreateCommandQueue(context, &device);
if (commandQueue == NULL)
{
Cleanup(context, commandQueue, program, kernel, memObjects);
return 1;
}
//创建程序对象
program=CreateProgram(context, device, "convolve.cl");
if (program == NULL)
{
Cleanup(context, commandQueue, program, kernel, memObjects);
return 1;
}
//创建OpenCL核
kernel=clCreateKernel(program, "convolve", nullptr);
if (kernel == NULL)
{
std::cerr << "Failed to create kernel" << std::endl;
Cleanup(context, commandQueue, program, kernel, memObjects);
return 1;
}
//不同
//创建内存对象
if (!CreateMemObjects(context, memObjects, inputSignal, mask))
{
Cleanup(context, commandQueue, program, kernel, memObjects);
return 1;
}
//设置内核参数
errNum=clSetKernelArg(kernel, 0, sizeof(cl_mem), &memObjects[0]);
errNum|=clSetKernelArg(kernel, 1, sizeof(cl_mem), &memObjects[1]);
errNum|=clSetKernelArg(kernel, 2, sizeof(cl_mem), &memObjects[2]);
errNum|=clSetKernelArg(kernel, 3, sizeof(cl_uint), &inputSignalWidth);
errNum|=clSetKernelArg(kernel, 4, sizeof(cl_uint), &maskWidth);
if (errNum != CL_SUCCESS)
{
std::cerr << "Error setting kernel arguments." << std::endl;
Cleanup(context, commandQueue, program, kernel, memObjects);
return 1;
}
size_t globalWorkSize[1] = { outputSignalWidth * outputSignalHeight };
size_t localWorkSize[1] = { 1 };
//为将在设备上执行的内核排队
errNum=clEnqueueNDRangeKernel(commandQueue, kernel, 1, nullptr, globalWorkSize, localWorkSize, 0, nullptr, nullptr);
if (errNum != CL_SUCCESS)
{
std::cerr << "Error queuing kernel for execution." << std::endl;
Cleanup(context, commandQueue, program, kernel, memObjects);
return 1;
}
//执行内核并读出数据
errNum = clEnqueueReadBuffer(commandQueue, memObjects[2], CL_TRUE,
0, outputSignalHeight * outputSignalWidth * sizeof(cl_int), outputSignal,
0, NULL, NULL);
if (errNum != CL_SUCCESS)
{
std::cerr << "Error reading result buffer." << std::endl;
Cleanup(context, commandQueue, program, kernel, memObjects);
return 1;
}
// Output the result buffer
for (int i = 0; i < outputSignalHeight; i++)
{
for (int j=0; j<outputSignalWidth; j++) {
std::cout << outputSignal[i][j] << " ";
}
std::cout<<std::endl;
}
std::cout << std::endl;
std::cout << "Executed program succesfully." << std::endl;
Cleanup(context, commandQueue, program, kernel, memObjects);
return 0;
}
4)我们创建OpenCL核,然后创建内存对象。我们创建内存对象不同了,因为内存对象是需要我们根据算法合理创建的。创建内存对象以及清空缓存代码为:
bool CreateMemObjects(cl_context context, cl_mem memObjects[3], cl_int (*a)[inputSignalWidth], cl_int (*b)[maskWidth])
{
memObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(cl_int) * inputSignalHeight * inputSignalWidth,
a, NULL);
memObjects[1] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(cl_int) * maskHeight * maskWidth,
b, NULL);
memObjects[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(cl_float) * outputSignalHeight * outputSignalWidth, NULL, NULL);
if (memObjects[0] == NULL || memObjects[1] == NULL || memObjects[2] == NULL)
{
std::cerr << "Error creating memory objects." << std::endl;
return false;
}
return true;
}
//释放OpenCL资源
void Cleanup(cl_context context, cl_command_queue commandQueue,
cl_program program, cl_kernel kernel, cl_mem memObjects[3])
{
for (int i = 0; i < 3; i++)
{
if (memObjects[i] != 0)
clReleaseMemObject(memObjects[i]);
}
if (commandQueue != 0)
clReleaseCommandQueue(commandQueue);
if (kernel != 0)
clReleaseKernel(kernel);
if (program != 0)
clReleaseProgram(program);
if (context != 0)
clReleaseContext(context);
}
这里我们内存对象依旧为3个,分别为输入信号矩阵,卷积核,以及输出矩阵。
5)为内核设置参数以及在设备上为内核排队。该代码:
errNum=clEnqueueNDRangeKernel(commandQueue, kernel, 1, nullptr, globalWorkSize, localWorkSize, 0, nullptr, nullptr);
第三个参数正是我们全局工作项的维度,指明为1,则表明我们实际仍是以1维的方式进行计算的。
由于我们的内核有5个参数,因此我们设置参数应该为它们分别设置。为内核排完队,我们就可以成功读出输出矩阵,代码输出如果为:
22 21 27 25 22 16
27 35 35 31 31 27
19 10 29 33 39 43
35 26 16 6 22 31
41 48 31 34 9 0
12 24 26 48 37 38
那我们的GPU计算矩阵卷积就成功了。