第7章 OpenCL设备端内存模型 - 7.6 统一地址空间

之前的OpenCL标准中，编程者常常需要为了不同地址空间的数据写很多版本的OpenCL C函数(其实就是参数中的地址描述符不同)。考虑下，下面两种数据数组，第一个在全局变量中，另一个缓存在局部内存中。两种方式都很简单，对于全局内存指针来说，函数可以直接使用(就像在有自动缓存系统的CPU上使用一样)，另一种就是使用局部内存(GPU会将数组存储在快速便签式内存中)。OpenCL 2.0之前的标准中，要对这两种数组进行同样的操作，就需要将一个函数写两遍，如代码清单7.1所示。

{%ace edit=false, lang=’c_cpp’%}
void doDoubleGlobal(
__global float *data,
int index){

data[index] *= 2;
}

void doDoubleGlobal(
__local float *data,
int index){

data[index] *= 2;
}

__kernel
void doubleData(
global float globalData, // the data
local float localData, // local storage
int useLocal){ // whether or not to use local memory

int globalId = get_global_id(0);
int localId = get_local_id(0);

if (useLocal){
// copy data to local memroy
localData[localId] = globalData[globalId];

doDoubleLocal(localData, localId);
globalData[globalId] = localData[localId];

} else {
doDoubleGlobal(globalData, globalId);
}
}
{%endace%}

代码清单7.1 在OpenCl 1.x中，需要对不同的寻址空间定义不同版本的函数

OpenCL 2.0开始，就不需要在这样定义函数了，同样的一个函数，可以通过统一内存地址覆盖所有内存空间。代码参见代码清单7.2。

{%ace edit=false, lang=’c_cpp’%}
void doDoubleGlobal(
float *data,
int index){

data[index] *= 2;
}

__kernel
void doubleData(
global float globalData, // the data
local float localData, // local storage
int useLocal){ // whether or not to use local memory

int globalId = get_global_id(0);
int localId = get_local_id(0);

generic float *data; // generic keyword not required
int myIndex;

if (useLocal){
// copy data to local memroy
localData[localId] = globalData[globalId];

// set data to local address space
data = localData;
myIndex = localId;

} else {
// set data to global address space
data = globalData;
myIndex = globalId;
}

doDouble(data, myIndex);

if (useLocal){
globalData[globalId] = localData[localId];
}
}
{%endace%}

代码清单7.2 OpenCL 2.0中使用统一地址空间将7.1中代码重写

统一地址空间包括全局、局部和私有地址空间。2.0标准中，并未将常量地址划分到同一地址空间中。虽然，常量空间在逻辑上是全局便令的一部分，不过在某些处理器上(特别是图像处理器)，常量数据会映射到特定的硬件单元上，并不能使用指令对其进行动态指定。