现在的位置: 首页 > 综合 > 正文

opencl学习(六)——local memory使用

2013年09月03日 ⁄ 综合 ⁄ 共 861字 ⁄ 字号 评论关闭

由于统一个workgroup中的所有work-item可以共用本地内存(local),也可以通过它进行同组work-item之间的通信,因此我们整理一下local memory的用法。
本地内存可以在kernel内部定义,也可以通过参数传递。下面就来说明一下这两种方式。

下面是一段来自《opencl异构计算》的代码:

__kernel void localAccess(
	__global float* A,
	__global float* B,
	__local float* C)
{
	__local float aLocalArray[1];
	if(get_local_id(0) == 0)
	{
		aLocalArray[0] = A[0];
	}
	C[get_local_id(0)] = A[get_global_id(0)];
	barrier( CLK_LOCAL_MEM_FENCE );
	float neighborSum = C[get_local_id(0)] + aLocalArray[0];
	if( get_local_id(0)>0 )
		neighborSum = neighborSum  + C[get_local_id(0)-1];
	B[get_local_id(0)] = neighborSum;
}

1.第一种创建方式:以内核函数内部变量的形式创建,如上代码中aLocalArray。

规律一:如果一个工作组中的一个工作项或多个工作项读取全局内存中的同一个位置,那么局部内存能大大提高性能。

本例中,aLocalArray是每一个工作项都需要的数据,这样使用使用本地内存,对与每一个workgroup只使用了一个32位本地内存,本组内所有工作项都可访问它。

2.第二种创建方式:以内核参数的形式传递,如上代码中的C。

这种情况下的使用跟内核函数的其他参数使用方式类似,要通过clSetKernelArg函数进行参数配置,但最后一个参数,地址空间设为NULL。

clErrNum = clSetKernelArg(kernel object, parameter index, size in bytes, NULL);

抱歉!评论已关闭.