nvcc编译过程第一篇:https://blog.csdn.net/shungry/article/details/89715468
nvcc编译过程第二篇:https://blog.csdn.net/shungry/article/details/89788342
正文:
nvcc的分散编译必须在CUDA5.0版本之后才支持,在此之前是不支持的。也就是说以前的版本一贯支持单独的主机代码的编译,只是设备CUDA需要的代码嵌入在其中。从CUDA 5.0,支持单独的设备代码的编译。
如上图,分散编译的过程就是通过nvcc的 --relocatable-device-code 命令生成可重定位可执行设备代码,--complie 命令用于控制在主机对象上停止编译,以上两个命令结合起来就是 --device-c 命令(就是 --relocatable device code=true --compile的意思)。以上编译由Device Linker (即前篇说的 nvlink),可以使用命令 --device-link来调用设备链接器(不写默认也会链接),进行设备的链接。
附上官方nvcc手册的例子,a.obj 和 b.obj 就是分离编译的结果,foo.lib是一个外部的静态库,通过--library-path=<path>来进行加载链接。
nvcc --gpu-architecture=sm_50 --device-c a.cu b.cu
nvcc --gpu-architecture=sm_50 a.obj b.obj foo.lib --library-path=<path>
注意,能否分散编译还会有以下的变化
对于5.0以前的版本,设备代码中的extern是不能定义为外部变量的,只能用来定义动态分配共享内存,如:
extern __shared__ int array[];
而在5.0版本之后,由于加入了分散编译,支持多设备文件代码进行独立编译,所以extern 是可以作为定义外部变量的。
想了解更多可以直接看官方nvcc文档:https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html