CUDA Hello World 程序

CUDA（Compute Unified Device Architecture）是NVIDIA（英伟达）公司基于其生产的图形处理器GPU（Graphics Processing Unit）开发的一个并行计算平台和编程模型。

基于CPU编程，程序都是运行在CPU上的；基于GPU编写的程序，程序运行在GPU上。在原有CPU平台上添加GPU设备后，通过增加额外的计算单元，可以有效的提高程序的运行效率。
一般将CPU端叫做Host, 将GPU端叫做Device. 如何写一个运行在GPU端的Hello world程序呢？
简单的看了Manual后，可能会这么写：

#include <stdio.h>

__global__ void hello() {
    printf("Hello world from device");
}

int main() {
    hello<<<1, 1>>>();
    printf("Hello world from host");
    return 0;
}

将上面文件命名为helloWorld.cu，用nvcc编译：

nvcc -o helloWorld helloWorld.cu

运行./helloWorld，发现只会打印出"hello world from host"。
查看NVIDIA的开发文档，device端的printf其实和设备端的printf函数是不同的实现，那为什么include一个标准库的头文件就可以使用device端的printf函数？其实这是源于cuda实现的一些技术细节，标准库的头文件只是给出了一个函数原型的提示，当在不同架构的计算机上编译时，通过函数重载，会调用不同的函数，这里有人给出了详细解释。而device端的输出其实是输出到device端的一个环形buffer(超出size后会覆写之前的内容），要flush到host端，需要采用下面任一操作：

1. Kernel launch via <<<>>> or cuLaunchKernel()
2. Synchronization via cudaDeviceSynchronize(), cuCtxSynchronize(), cudaStreamSynchronize(), cuStreamSynchronize(),
   cudaEventSynchronize(), or cuEventSynchronize(),
3. Memory copies via any blocking version of cudaMemcpy*() or cuMemcpy*(),
4. Module loading/unloading via cuModuleLoad() or cuModuleUnload(),
5. Context destruction via cudaDeviceReset() or cuCtxDestroy().
6. Prior to executing a stream callback added by cudaStreamAddCallback or cuStreamAddCallback.

这里，添加cudaDeviceSynchronize()，就可以正常输出了：

#include <stdio.h>

__global__ void hello() {
    printf("Hello world from device");
}

int main() {
    hello<<<1, 1>>>();
    printf("Hello world from host");
    cudaDeviceSynchronize()；
    return 0;
}