Hope

如果爱,请深爱。

学习OpenCL的一些心得

这篇文章记录我对于前一周学习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;
}


当然,这里是直接使用第0个平台第0个设备,如果第0个平台中没有GPU,程序会出错。对于你想选择第几个GPU来执行,直接在上面那个ocl_Info导出的列表中看就好了,执行结果就是:



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编程指南         --------     对于详细的东西可以在这个上面找到

阅读更多
文章标签: OpenCL心得
个人分类: OpenCL
想对作者说点什么? 我来说一句

没有更多推荐了,返回首页

加入CSDN,享受更精准的内容推荐,与500万程序员共同成长!
关闭
关闭