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