CUDA 6.0 统一寻址

1. 介绍

CUDA6.0发布已经有一段时间了，最引人注目的依然是最新引入的统一寻址（unified memory managed），网上关于统一寻址的介绍也不少了，但是我还是想补充一下，因为我看到的都是从官方文档直接拷贝的代码，甚至未考虑是否能执行。当然，官方的肯定没问题，但是本着严谨的态度，至少需要把情况说清楚猜对，不然让我这样的小白走了多少的弯路。

好吧，我也先说下所谓的统一寻址，统一寻址就是说：

使用函数cudaMallocManaged()开辟一块存储空间，无论是在Kernel函数中还是main函数中，都可以使用这块内存，达到了统一寻址的目的。

通过这种方式大大的简化了代码的复杂度，因为CUDA6之前没有统一寻址，进行GPU计算的步骤稍许麻烦： 1. 在显存上开辟空间 2. 将内存上的数据拷贝到显存 3. 调用CUDA核进行计算 4. 将显存上处理过的数据拷贝到内存上

而统一寻址的最大优势就是避免了人为的数据拷贝，为什么说人为呢，是因为即使是统一寻址也是要进行数据拷贝的，只不过现在这一部分有程序自动完成，而不用程序员操心了。因此，统一寻址后程序的执行效率并不会显著改善，仅仅是为了方便而已。

现在再说其他人没有谈到过的，官方文档有这样一段话：

Unified Memory has three basic requirements:

‣ a GPU with SM architecture 3.0 or higher (Kepler class or newer)

‣ a 64-bit host application and operating system, except on Android

‣ Linux or Windows

简单的翻译下来如下所示（英文水平有限，还请见谅）：

使用统一寻址需要满足三个要求:

‣ Kepler 架构或者是最新架构的GPU，并且计算能力至少是3.0

‣ 64位的系统，64位的应用

‣ Linux 或 Windows

尼玛，统一寻址的使用是有诸多限制的！！！！第三条可以忽略，但是第一、二条至关重要。怪不得我拿着实验室的GT520试了许久不行，第一条都不满足，果断换了笔记本的GT750M跑，结果试了半天依旧不行，才发现官方文档中的这段话，原来是要64位编译！！

废话不多说了，下面直接上代码。

2. 未使用统一寻址

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <iostream>
#include <Windows.h>

using namespace std;

__global__ void AplusB(int *ret, int a, int b)
{
ret[threadIdx.x] = a + b + threadIdx.x;
}

// 非统一寻址
int main()
{
int *ret;

cudaMalloc(&ret, 10 * sizeof(int));
AplusB<<< 1, 10 >>>(ret, 10, 100);

int *host_ret = (int *)malloc(10 * sizeof(int));
cudaMemcpy(host_ret, ret, 10 * sizeof(int), cudaMemcpyDeviceToHost);

for(int i=0; i<10; i++)
{
cout<< "A+B = " << host_ret[i] << endl;
}

cudaFree(ret);
free(host_ret);
Sleep(20000);
return 0;
}

3. 使用统一寻址

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <iostream>
#include <Windows.h>

using namespace std;

__global__ void AplusB(int *ret, int a, int b)
{
ret[threadIdx.x] = a + b + threadIdx.x;
}

// 统一寻址
int main()
{
int *ret;

cudaMallocManaged(&ret, 10 * sizeof(int));
AplusB<<< 1, 10 >>>(ret, 10, 100);
cudaDeviceSynchronize();

for(int i=0; i<10; i++)
{
cout<< "A+B = " << ret[i] << endl;
}

cudaFree(ret);
Sleep(20000);
return 0;
}

4. 使用 __managed__

__managed__引入了一种全局变量，在内存和显存中都可以使用，这是非常方便的，但是也有诸多限制，我测试发现只有把sm_10改成sm_30才能使用。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <iostream>
#include <Windows.h>

using namespace std;

// 使用 __managed__
// 要求sm_30
__device__ __managed__ int ret[10];
__global__ void AplusB(int a, int b)
{
ret[threadIdx.x] = a + b + 2 * ret[threadIdx.x];
}
int main()
{
for(int i=0; i<10; i++)
{
ret[i] = i;
}

AplusB<<< 1, 10 >>>(10, 100);
cudaDeviceSynchronize();
for(int i=0; i<10; i++)
{
cout<< "A+B = " << ret[i] << endl;
}

Sleep(20000);
return 0;
}