这篇文章记录我对于前一周学习OpenCL的心得和对于一些东西的理解。
1.GPU
对于最开始,完全不清楚GPU是个啥,感觉好像很神秘似得,只知道需要话好几千块RMB买一个的。所以在一段时间里面都没有去碰OpenCL,认为是没有GPU,无法学习并测试。直到后来才发现,所谓GPU就是平常所谓的显卡,带计算功能的显卡(以前有的显卡不支持计算,现在看nVIDIA官网说,他们的所有卡都实现了对cuda的支持)。对于查看显卡是否支持OpenCL运算,可以下载一个GPU-Z软件,打开后可以在下面看到,如果支持OpenCL,那么下面会打个勾勾。
2.OpenCL
让OpenCL跑起来很简单,也就那么几句话,只是优化它则需要知道更多的细节。
SDK:我尝试过Cuda Toolkit里面的OpenCL和Amd App里面的OpenCL,发现写的代码可以通用(当然,OpenCL的目的也就是跨平台的),只是一些小的细节上有点区别,比如有的错误代码在Amd App里面有而在Cuda Toolkit里面没有,之类的不太影响使用。所以当然也就可以下载Cuda Toolkit或者Amd App了。
让OpenCL运行起来
a.这里贴一个后面需要用到的工具 ocl_Info,它是查询机器上OpenCL设备信息的东西,将它下载下来并编译出来,在控制台中:
c:\User\Administrator>ocl_info >1.txt
然后可以在C:/User/Administrator/目录下找到1.txt,里面记录了设备信息,如下图,如果没有找到,软件会提示的:
Number of platforms: 1 表示现在有1个平台
Number of devices: 2 表示有2个设备
在OpenCL中,它规定有4种模型(平台模型,执行模型,内存模型,编程模型),当然只需要了解平台模型就好了,其他暂时放一边。其中平台模型是指:
一台电脑有多个平台,一个平台有多个设备(比如:1个GPU + 1个CPU),一个设备又有多个计算单元,一个计算单元又有多个处理单元(就是所谓的SIMD,可以理解成一个线程)和一块属于计算单元自己的局部内存。 知道这些就够了,然后下面开始是代码。
b.代码:
Simple test.txt
__kernel void SimpleTest(__global uint* indata, __global uint* outdata)
{
uint idx = get_global_id(0);
outdata[idx] = indata[idx] * 2;
}
main.cpp
#include <stdlib.h>
#include <stdio.h>
#include <cl.h>
#pragma comment(lib, "OpenCL.lib")
char* readFile(const char* file, size_t* len)
{
FILE* f = fopen(file, "rb");
if(f == 0) return 0;
fseek(f, 0, SEEK_END);
*len = ftell(f);
fseek(f, 0, SEEK_SET);
char* data = new char[*len];
fread(data, 1, *len, f);
fclose(f);
return data;
}
int main(int argc, char** argv)
{
cl_uint uarr[] = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10};
cl_uint uout[sizeof(uarr) / sizeof(uarr[0])] = {0};
cl_uint status;
cl_platform_id platform;
status = clGetPlatformIDs(1, &platform, 0);
cl_device_id device;
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, 0);
cl_context context = clCreateContext(0, 1, &device, 0, 0, 0);
cl_command_queue queue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, 0);
cl_mem clArrBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(uarr), uarr, 0);
cl_mem clArrOutBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(uout), uout, 0);
size_t len = 0;
char* code = readFile("Simple test.txt", &len);
if(code == 0)
{
printf("Can not open \"Simple test.txt\"\n");
return 0;
}
cl_program program = clCreateProgramWithSource(context, 1, (const char**)&code, &len, 0);
delete code, code = 0;
status = clBuildProgram(program, 1, &device, 0, 0, 0);
if (status != CL_SUCCESS)
{
printf("Build failed: %d\n", status);
char tbuf[1024];
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 1024, tbuf, NULL);
printf("\n%s\n", tbuf);
return -1;
}
cl_uint sizeOfArr = sizeof(uarr) / sizeof(uarr[0]);
cl_kernel kernel = clCreateKernel(program, "SimpleTest", 0);
clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&clArrBuffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&clArrOutBuffer);
clSetKernelArg(kernel, 2, sizeof(cl_uint), (void*)&sizeOfArr);
size_t global_work_size = sizeOfArr;
clEnqueueNDRangeKernel(queue, kernel, 1, 0, &global_work_size, 0, 0, 0, 0);
clEnqueueReadBuffer(queue, clArrOutBuffer, CL_TRUE, 0, sizeof(uout), uout, 0, 0, 0);
clReleaseMemObject(clArrOutBuffer);
clReleaseMemObject(clArrBuffer);
clReleaseKernel(kernel);
clReleaseCommandQueue(queue);
clReleaseContext(context);
for (int i = 0; i < sizeof(uout) / sizeof(uout[0]); ++i)
printf("%d * 2 = %d\n", i + 1, uout[i]);
system("pause");
return 0;
}
c.使用的函数,列举下,这个程序中无非就是使用了那么几个函数而已:
clGetPlatformIDs---------------------------获取平台ID
clGetDeviceIDs-----------------------------获取设备ID
clCreateContext----------------------------创建上下文
clCreateCommandQueue--------------创建命令队列
clCreateBuffer------------------------------创建设备内存
clCreateProgramWithSource----------创建程序
clBuildProgram-----------------------------编译程序
clGetProgramBuildInfo-------------------获取编译信息
clCreateKernel-----------------------------创建核
clSetKernelArg-----------------------------设置核参数
clEnqueueNDRangeKernel------------执行核
clEnqueueReadBuffer-------------------读取设备内存
clReleaseMemObject--------------------释放内存对象
clReleaseKernel---------------------------释放核
clReleaseCommandQueue------------释放命令队列
clReleaseContext--------------------------释放上下文
对于各个函数的参数是什么意思,建议看这个网站(这个是clGetPlatformIDs函数的,查其他函数信息,直接修改末尾为xxx.html就好了):
http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clGetPlatformIDs.html
这里,着重说明几个函数:
clCreateKernel:所谓Kernel(核),可以看成是函数,一个Kernel对应一个函数,不同的设备,需要分别创建Kernel。Kernel编译后可以重复使用。
clEnqueueNDRangeKernel:执行Kernel。它可以同步,也可以异步执行。对于异步,看其最后的3个参数说明就好了。当程序在GPU中执行完成后需要Copy结果到内存,就使用clEnqueueReadBuffer函数,当然也有对应的clEnqueueWriteBuffer函数写内存。
clCreateBuffer:创建设备内存,也可以说是分配设备内存,这里分配的是全局内存。在这里,每次对设备分配内存的最大尺寸是这个设备信息中CL_DEVICE_MAX_MEM_ALLOC_SIZE指定的值,对一个设备允许分配内存的总最大尺寸是CL_DEVICE_GLOBAL_MEM_SIZE指定,单位都是字节。这些值可以在ocl_Info导出的设备信息中找到。分配失败后其最后一个参数为错误代码的返回值。这里需要注意一点是,我测试发现,即使分配尺寸超过最大值,有时它的返回值依然是CL_SUCCESS,但是到clEnqueueReadBuffer函数的时候却出错(CL_OUT_OF_RESOURCES),所以分配的时候不要超过最大值了。
对于编写OpenCL中用的程序,其结构是类C语言,如下:
__kernel void SimpleTest(__global uint* indata, __global uint* outdata)
{
uint idx = get_global_id(0);
outdata[idx] = indata[idx] * 2;
}
1.每个程序返回值必定是void,函数前面必定有前缀__kernel,__global描述符说明指针是来自于公用内存(OpenCL中还有私有内存和局部内存的概念)。
2.get_global_id获取线程索引,其取值范围是0 到 global_work_size - 1,如下面,通常global_work_size的取值为并行任务数。
size_t global_work_size = sizeOfArr;
clEnqueueNDRangeKernel(queue, kernel, 1, 0, &global_work_size, 0, 0, 0, 0);
最后是参考资料:
百度文库里面的OpenCL资料:http://wenku.baidu.com/view/53e216a4b0717fd5360cdc24.html
书籍:OpenCL编程指南 -------- 对于详细的东西可以在这个上面找到