7.6 统一地址空间

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

  1. void doDoubleGlobal(
  2.   __global float *data,
  3.   int index){
  5.   data[index] *= 2;
  6. }
  8. void doDoubleGlobal(
  9.   __local float *data,
  10.   int index){
  12.   data[index] *= 2;
  13. }
  15. __kernel
  16. void doubleData(
  17.   global float *globalData, // the data
  18.   local float *localData, // local storage
  19.   int useLocal){ // whether or not to use local memory
  21.   int globalId = get_global_id(0);
  22.   int localId = get_local_id(0);
  24.   if (useLocal){
  25.     // copy data to local memroy
  26.     localData[localId] = globalData[globalId];
  28.     doDoubleLocal(localData, localId);
  30.     globalData[globalId] = localData[localId];
  31.   } else {
  32.     doDoubleGlobal(globalData, globalId);
  33.   }
  34. }

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

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

  1. void doDoubleGlobal(
  2.   float *data,
  3.   int index){
  5.   data[index] *= 2;
  6. }
  8. __kernel
  9. void doubleData(
  10.   global float *globalData, // the data
  11.   local float *localData, // local storage
  12.   int useLocal){ // whether or not to use local memory
  14.   int globalId = get_global_id(0);
  15.   int localId = get_local_id(0);
  17.   generic float *data; // generic keyword not required
  18.   int myIndex;
  20.   if (useLocal){
  21.     // copy data to local memroy
  22.     localData[localId] = globalData[globalId];
  24.     // set data to local address space
  25.     data = localData;
  26.     myIndex = localId;
  27.   } else {
  28.     // set data to global address space
  29.     data = globalData;
  30.     myIndex = globalId;
  31.   }
  33.   doDouble(data, myIndex);
  35.   if (useLocal){
  36.     globalData[globalId] = localData[localId];
  37.   }
  38. }

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

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