大家知道，Nvidia公司发布的CUDA6开发包拥有一个新特性，就是“统一内存寻址”，那究竟统一内存寻址有什么特殊的地方呢？我们编写CUDA代码跟以前有什么区别呢？现在拥有的GPU架构能够很好的支持吗？我们带着这些问题开始我们的话题。
从名字上看，统一内存寻址就是将CPU端的内存同GPU显存统一起来，使得程序猿在编写代码的时候不用明显的使用诸如cudaMalloc或者cudaMemcpy等操作显存的指令，而能够在Kernel函数中直接使用定义的变量。
例子如下：
原始代码（CUDA6.0文档）
__global__
void
AplusB(
int
*ret,
int
a,
int
b)
{
ret[threadIdx.x] = a + b + threadIdx.x;
}
int
main()
{
int
*ret;
//**************************************
cudaMalloc(&ret, 1000 *
sizeof
(
int
));
AplusB<<<1, 1000>>>(ret, 10, 100);
//**************************************
int
*host_ret = (
int
*)
malloc
(1000 *
sizeof
(
int
));
cudaMemcpy(host_ret, ret, 1000 *
sizeof
(
int
), cudaMemcpyDefault);
for
(
int
i = 0; i < 1000; i++)
printf
(
"%d: A + B = %d\n"
, i, host_ret[i]);
free
(host_ret);
cudaFree(ret);
return
0;
}
统一寻址的代码：
__global__
void
AplusB(
int
*ret,
int
a,
int
b)
{
ret[threadIdx.x] = a + b + threadIdx.x;
}
int
main()
{
int
*ret;
//***********************************************
cudaMallocManaged(&ret, 1000 *
sizeof
(
int
));
AplusB<<<1, 1000>>>(ret, 10, 100);
//***********************************************
cudaDeviceSynchronize();
for
(
int
i = 0; i < 1000; i++)
printf
(
"%d: A + B = %d\n"
, i, ret[i]);
cudaFree(ret);
return
0;
}
从上面不同的代码可以看出，统一寻址后的代码更简洁，使用了函数cudaMallocManaged()开辟一块存储空间，无论是在Kernel函数中还是main函数中，都可以使用这块内存，达到了统一寻址的目的。
这里有一个需要注意的地方就是main函数在调用Kernel函数之后，使用了一个同步函数。仔细思考后就会有所领悟——既然这块存储空间既可以被Kernel函数访问，也可以被main函数访问，为了解决访问冲突的问题，因此使用了同步函数，使得在Kernel改变变量的值后，main函数才能使用该变量。
上面只是简单的举了个例子，但已经把统一寻址的问题描述的很清楚：
1. 开辟空间所使用的新函数；
2. 变量访问同步。
在CUDA6.0的文档中更加仔细的描述了相关的问题，包括流和多GPU条件下使用统一寻址的问题等，其主要针对的还是变量访问同步，大家有兴趣可以参看文档。
同AMD的APU不同，即使是Kepler架构的GPU，也不是真正意义上从硬件手段实现“统一内存寻址”，只是从软件上解决的一种方式，Nvidia公司的新一代GPU产品——Maxwell架构，才真正是从硬件层面实现这一特性。而速度，还有待测试与研究~~
既然“统一内存寻址”是GPU产品所前进的一个方向，势必会改变我们的编程方式，因此，这一话题还是有重大的意义的，希望大家热烈讨论，在讨论中真正掌握这一技术！
