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端,需要采用下面任一操作:
- Kernel launch via <<<>>> or cuLaunchKernel()
- Synchronization via cudaDeviceSynchronize(), cuCtxSynchronize(), cudaStreamSynchronize(), cuStreamSynchronize(),
cudaEventSynchronize(), or cuEventSynchronize(), - Memory copies via any blocking version of cudaMemcpy*() or cuMemcpy*(),
- Module loading/unloading via cuModuleLoad() or cuModuleUnload(),
- Context destruction via cudaDeviceReset() or cuCtxDestroy().
- 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;
}