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\temp，Linux 使用环境变量 TMPDIR 或默认目录 /tmp

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

■ 自己在电脑上实验分离编译，VS中能通过，gcc 中没有成功，编译过程没有问题，运行时提示 couldn't get the symbol addr，留坑。

//---------- b.h ----------
#define N 8
extern __device__ int g[N];
extern __device__ void bar(void);

//---------- b.cu ----------
#include "b.h"
__device__ int g[N];
__device__ void bar(void)
{
    g[threadIdx.x]++;
}

//---------- a.cu ----------
#include <stdio.h>
#include "b.h"
__global__ void foo(void) 
{
    __shared__ int a[N];
    a[threadIdx.x] = threadIdx.x;
    __syncthreads();
    g[threadIdx.x] = a[blockDim.x - threadIdx.x - 1];
    bar();
}

int main(void) 
{
    unsigned int i;
    int *dg, hg[N];
    int sum = 0;
    foo << <1, N >> >();
    if (cudaGetSymbolAddress((void**)&dg, g)) 
    {
        printf("couldn't get the symbol addr\n");
        return 1;
    }
    if (cudaMemcpy(hg, dg, N * sizeof(int), cudaMemcpyDeviceToHost)) 
    {
        printf("couldn't memcpy\n");
        return 1;
    }
    for (i = 0; i < N; i++) 
        sum += hg[i];
    if (sum == 36)
        printf("PASSED\n");
    else 
        printf("FAILED (%d)\n", sum);
    return 0;
}

■ 书上用到的编译命令

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