OpenCL编程(2)- 程序基本流程,以向量相加为例

下面以一个OpenCL的HelloWorld——向量相加,来说明OpenCL程序的基本流程

1.查询平台信息

OpenCL标准中定义平台模型的概念,一个平台会指定一个主机(host)和对应的一个或多个设备(device)处理器。比如如果是一个CPU+GPU的环境,则CPU是host,GPU是device
查询平台信息分为两步,均调用clGetPlatformIDs接口
第一步,查询平台个数
第二部,查询平台对象
示例代码如下:

	cl_int			iStatus = 0;				// 函数返回状态
	cl_uint			uiNumPlatforms = 0;			// 平台个数
 	// 查询可用的平台个数,并返回状态
	iStatus = clGetPlatformIDs(0, NULL, &uiNumPlatforms);
	if (CL_SUCCESS != iStatus)
	{
		cout << "Error: Getting platforms error" << endl;
		return 0;
	}
	// 获得平台地址
	if (uiNumPlatforms > 0)  // 如果有可用平台
	{
		// 根据平台个数为平台对象分配内存空间
		cl_platform_id *pPlatforms = (cl_platform_id *)malloc(uiNumPlatforms * sizeof(cl_platform_id));
 
		// 获得可用的平台
		iStatus = clGetPlatformIDs(uiNumPlatforms, pPlatforms, NULL);
		Platform = pPlatforms[0];	// 获得第一个平台的地址
		free(pPlatforms);			// 释放平台占用的内存空间
	}

另外,可通过 clGetPlatformInfo接口获取平台具体信息。

2.查询设备信息

查询设备也分为两步,均调用clGetDeviceIDs接口:
第一步,获取设备数量
第二部,获取设备对象
示例代码如下:

	cl_uint			uiNumDevices = 0;				// 设备数量
	cl_device_id	*pDevices = NULL;				// 设备
	// 获得GPU设备数量
	iStatus = clGetDeviceIDs(Platform, CL_DEVICE_TYPE_GPU, 0, NULL, &uiNumDevices);
	if (uiNumDevices > 0)	// 如果有GPU设备
	{
		// 根据设备个数为设备对象分配内存空间
		pDevices = (cl_device_id *)malloc(uiNumDevices * sizeof(cl_device_id));
 		// 获得可用的设备
		iStatus = clGetDeviceIDs(Platform, CL_DEVICE_TYPE_GPU, uiNumDevices, pDevices, NULL);
	} 
3.创建上下文

上下文(Context)是OpenCL执行模型的一部分,上下文用于协调主机与设备的交互,管理设备端可用的内存对象,并持续跟踪在设备上创建出来的程序对象和内核对象。
示例代码如下:

	cl_context		Context = NULL;				// 设备环境
	// 创建上下文
	Context = clCreateContext(NULL, 1, pDevices, NULL, NULL, NULL);
	if (NULL == Context)
	{
		cout << "Error: Can not create context" << endl;
		return 0;
	}
4.创建命令队列

命令队列(command queue)也是OpenCL执行模型的一部分,命令队列是一种通信机制,可以让host发请求到对应的device;命令队列通过上下文来创建,每个设备上有自己独立的命令队列;如果平台上有多个设备时,就需要每个设备上创建
一个命令队列。
示例代码如下:

	cl_context		Context = NULL;				// 设备环境
	// 创建第1个设备的命令队列
	CommandQueue = clCreateCommandQueue(Context, pDevices[0], 0, NULL);
	if (NULL == CommandQueue)
	{
		cout << "Error: Can not create CommandQueue" << endl;
		return 0;
	}
5.创建程序对象

创建程序对象首先要以字符串形式读入内核程序文件(以.cl为后缀),向量相加的内核程序文件内容如下:
VecAdd.cl

__kernel void adder(__global const float* a, __global const float* b, __global float* result)
{
	int idx = get_global_id(0);
	result[idx] = a[idx] +b[idx];
}

然后通过代码字符串创建程序对象:

	cl_program	Program = NULL;				// 程序对象
	string		strSource = "";				// 用于存储cl文件中的代码
	const char	*pSource;							// 代码字符串指针
	size_t		uiArrSourceSize[] = { 0 };			// 代码字符串长度
	// 将cl文件中的代码以字符串形式读入
	iStatus = ConvertToString("VecAdd.cl", strSource);
	pSource = strSource.c_str();			// 获得strSource指针
	uiArrSourceSize[0] = strlen(pSource);	// 字符串大小

	// 创建程序对象
	Program = clCreateProgramWithSource(Context, 1, &pSource, uiArrSourceSize, NULL);
	if (NULL == Program)
	{
		cout << "Error: Can not create program" << endl;
		return 0;
	}
5.编译内核程序

创建好的程序对象需要进行编译,编译之后的内核才能在OpenCL设备上运行。

	iStatus = clBuildProgram(Program, 1, pDevices, NULL, NULL, NULL);
	if (CL_SUCCESS != iStatus)	// 编译错误
	{
		cout << "Error: Can not build program" << endl;
		char szBuildLog[16384];
		clGetProgramBuildInfo(Program, *pDevices, CL_PROGRAM_BUILD_LOG, sizeof(szBuildLog), szBuildLog, NULL);
 
		cout << "Error in Kernel: " << endl << szBuildLog;
		clReleaseProgram(Program);
 
		return 0;
	}
6.创建输入输出内核内存对象

创建内核内存对象,即在设备上分配内存,并关联上主机上内存,数据将在主机内存和设备内存间传输

 	//host上的两个输入参数内存
	std::vector<float> a(DATA_SIZE), b(DATA_SIZE);
	for(int i = 0; i < DATA_SIZE; i++) {
		a[i] = i;
		b[i] = i;
	}
	//创建设备上的内核内存,内核程序对该内存只读,输入数据将从相应的host内存拷贝到device的内存上
	cl_mem cl_a = clCreateBuffer(Context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &a[0], NULL);
	cl_mem cl_b = clCreateBuffer(Context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &b[0], NULL);
	//在设备上分配的保存结果的输出内存,内核程序对该内存只写
	cl_mem cl_res = clCreateBuffer(Context, CL_MEM_WRITE_ONLY, sizeof(cl_float) * DATA_SIZE, NULL, NULL);
7.创建内核对象

创建内核对象,即从编译好的程序中指定需要执行的内核函数,与获取动态链接库中函数指针有点类似。一个程序对象可以提取多个内核对象。

	cl_kernel		Kernel = NULL;				// 内核对象
	Kernel = clCreateKernel(Program,
		"adder",  // cl文件中的入口函数
		NULL);
	if (NULL == Kernel)
	{
		cout << "Error: Can not create kernel" << endl;
		return 0;
	}
8.设置内核参数

设置内核参数,相当于为内核函数指定入参和出参。

	//设置好kernal的参数
	iStatus = clSetKernelArg(Kernel, 0, sizeof(cl_mem), &cl_a);
	iStatus |= clSetKernelArg(Kernel, 1, sizeof(cl_mem), &cl_b);
	iStatus |= clSetKernelArg(Kernel, 2, sizeof(cl_mem), &cl_res);
 
	if (CL_SUCCESS != iStatus)
	{
		cout << "Error setting kernel arguments" << endl;
	}
9.运行内核

顾名思义,运行内核就是在设备上执行内核程序。

	uiGlobal_Work_Size[0] = sizeof(cl_float) * DATA_SIZE;  // 输入字符串大小
	// 将内核放入命令队列,命令队列会调度内核并执行
	iStatus = clEnqueueNDRangeKernel(
		CommandQueue,
		Kernel,
		1,
		NULL,
		uiGlobal_Work_Size,  // 确定内核在设备上的多个处理单元间的分布
		NULL,				 // 确定内核在设备上的多个处理单元间的分布
		0,
		NULL,
		NULL);
 
	if (CL_SUCCESS != iStatus)
	{
		cout << "Error: Can not run kernel" << endl;
		return 0;
	}
10.从内核内存中读出结果到主机内存
	std::vector<float> res(DATA_SIZE); //主机上用于保存结果的内存
	iStatus = clEnqueueReadBuffer(
		CommandQueue,		// 命令队列
		cl_res,				// 输出内存对象
		CL_TRUE,			// 内核读取结束之前该函数不会返回
		0,
		sizeof(float) * DATA_SIZE,
		&res[0],
		0,
		NULL,
		NULL);

	if (CL_SUCCESS != iStatus)
	{
		cout << "Error: Can not reading result buffer" << endl;
		return 0;
	}
 
	for(int i = 0; i < res.size(); ++i)
	{
		printf("%f ", res[i]);
	}
11.释放资源

内核执行完成后,并且输出已经传出到主机端,分配的资源需要进行释放。由于数组和命
令队列都绑定在上下文上,因此上下文需要最后释放

	iStatus = clReleaseKernel(Kernel);
	iStatus = clReleaseProgram(Program);
	iStatus = clReleaseMemObject(memInutBuffer);
	iStatus = clReleaseMemObject(memOutputBuffer);
	iStatus = clReleaseCommandQueue(CommandQueue);
	iStatus = clReleaseContext(Context); 
	if (NULL != pOutput)
	{
		free(pOutput);
		pOutput = NULL;
	} 
	if (NULL != pDevices)
	{
		free(pDevices);
		pDevices = NULL;
	}
相关推荐
©️2020 CSDN 皮肤主题: 深蓝海洋 设计师:CSDN官方博客 返回首页