CUDA compiler driver nvcc 散点 part 2

Posted cuancuancuanhao

tags:

篇首语:本文由小常识网(cha138.com)小编为大家整理,主要介绍了CUDA compiler driver nvcc 散点 part 2相关的知识,希望对你有一定的参考价值。

● nvcc 编译流程图

技术图片

● sm 是向前兼容的,高的版本号是在低版本号的基础上添加了新功能得到的,同一 compute_XY 编译的 .cu 文件仅能向后 sm_ZW 的实 GPU 版本(Z > X)

技术图片

技术图片

 

● 虚拟 GPU 完全由它提供给应用程序的一组功能或特征来定义

● PTX 可以视为虚拟 GPU 的 汇编,以文本格式表示,便于进一步编译为各格式的二进制机器码

● 编译时应尽量降低虚拟 GPU 版本(增加兼容性),同时尽量提高实际 GPU 版本(在知道运行 GPU 的情况下)

● 即时编译(JIT)模式下驱动知道 runtime GPU 的信息,可以编译最佳版本的代码,离线编译和 JIT 的流程图分别如下:

技术图片技术图片

 

● 仅指定虚 GPU 版本而不指定实 GPU 版本时(如 nvcc x.cu -arch=compute_50 [-code=compute_50]),PTX 将延迟到运行时才进行编译,有启动延迟

● 消灭启动延迟的方法:

  ■ CUDA 驱动编译缓存

  ■ 编译时指定多个实 GPU 版本(如 nvcc x.cu -arch=compute_50 -code=sm_50,sm_52),设备函数的多个版本存储在 x.fatbin 中,运行时由驱动自动识别和调用

● 关于 -arch 和 -code 的要点

  ■ 仅指定 -arch 为虚 GPU 版本,-code 自动匹配最接近的版本(如 nvcc x.cu -arch=compute_50 等价于 nvcc x.cu -arch=compute_50 -code=compute_50)

  ■ 仅指定 -arch 为实 GPU 版本,-code 自动匹配最接近的版本(如 nvcc x.cu -arch=sm_50 等价于 nvcc x.cu -arch=compute_50 -code=sm_50,compute_50)

  ■ 同时指定 -arch 和 -code 为虚 GPU 版本,必须一致

  ■ 均不指定,使用默认值(如 nvcc x.cu 等价于 nvcc x.cu -arch=compute_30 -code=)

  ■ 默认 -arch 值就是 sm_XX

  ■ 编译第一阶段中有宏定义 __CUDA_ARCH__ 代表虚 GPU 版本,可用于 __device__ 函数中,用于指明该函数所用的虚 GPU 版本

● 没有指定 --keep 时 nvcc 使用临时目录来保存中间文件,编译完成后立即删除,Windows 中使用环境变量 TEMP 或默认目录 C:Windows emp,Linux 使用环境变量 TMPDIR 或默认目录 /tmp

● CUDA 5.0 开始支持分离编译,流程图如下。

技术图片

■ 自己在电脑上实验分离编译没有成功,先把书上的代码和编译指令堆在这里,以后再来看,留坑

 1 //---------- b.h ----------
 2 #define N 8
 3 extern __device__ int g[N];
 4 extern __device__ void bar(void);
 5 
 6 //---------- b.cu ----------
 7 #include "b.h"
 8 __device__ int g[N];
 9 __device__ void bar(void)
10 {
11     g[threadIdx.x]++;
12 }
13 
14 //---------- a.cu ----------
15 #include <stdio.h>
16 #include "b.h"
17 __global__ void foo(void) 
18 {
19     __shared__ int a[N];
20     a[threadIdx.x] = threadIdx.x;
21     __syncthreads();
22     g[threadIdx.x] = a[blockDim.x - threadIdx.x - 1];
23     bar();
24 }
25 
26 int main(void) 
27 {
28     unsigned int i;
29     int *dg, hg[N];
30     int sum = 0;
31     foo << <1, N >> >();
32     if (cudaGetSymbolAddress((void**)&dg, g)) 
33     {
34         printf("couldn‘t get the symbol addr
");
35         return 1;
36     }
37     if (cudaMemcpy(hg, dg, N * sizeof(int), cudaMemcpyDeviceToHost)) 
38     {
39         printf("couldn‘t memcpy
");
40         return 1;
41     }
42     for (i = 0; i < N; i++) 
43         sum += hg[i];
44     if (sum == 36)
45         printf("PASSED
");
46     else 
47         printf("FAILED (%d)
", sum);
48     return 0;
49 }

■ 可能用到的编译指令

nvcc --gpu - architecture = sm_50 --device - c a.cu b.cu
nvcc --gpu - architecture = sm_50 a.o b.o

nvcc --gpu - architecture = sm_50 --device - c a.cu b.cu
nvcc --gpu - architecture = sm_50 --device - link a.o b.o --output - file link.o
g++ a.o b.o link.o --library - path = <path> --library = cudart

nvcc --gpu - architecture = sm_50 --device - link a.o b.o --cubin --output - file link.cubin

nvcc --gpu - architecture = sm_50 --device - c a.cu b.cu
nvcc --lib a.o b.o --output - file test.a
nvcc --gpu - architecture = sm_50 test.a

nvcc --gpu - architecture = sm_50 --device - c a.ptx

nvcc --gpu - architecture = sm_50 --device - c a.cu b.cu
nvcc --gpu - architecture = sm_50 --device - link a.o b.o --output - file link.o
nvcc --lib --output - file libgpu.a a.o b.o link.o
g++ host.o --library = gpu --library - path = <path> --library = cudadevrt --library = cudart

 

以上是关于CUDA compiler driver nvcc 散点 part 2的主要内容,如果未能解决你的问题,请参考以下文章

CUDA主机代码的互斥锁? boost::mutex 由于使用 nvcc 重新定义了 BOOST_COMPILER 没有解决方案

cuda报错: nvcc fatal : Host compiler targets unsupported OS

MatConvNet报错:The NVCC compiler ‘/usr/local/cuda-11.1/bin/nvcc‘ does not appear to be valid. 解决办法

“不知道如何处理' nvcc 致命错误

如何在 ns3 中使用 nvcc 作为编译器

如何为 nvcc 添加预定义宏?