（因为这章内容比较碎，不好提炼，大部分为原文翻译，人工翻译，不是机器翻译。抵制不负责任的机翻从你我做起！翻译不易，转载贴上出处。——630056108@qq.com）

 

由于CUDA对C的拓展，所有包含CUDA代码的源文件都必须由NVCC编译一遍。

当提到 编译工作流 时runtime提供了主机和设备机之间的交互。

Runtime必须建立在低层次的API函数、CUDA 驱动函数之上。驱动函数通过包含底层函数、主机对设备机的类函数、CUDA模型、设备机动态库 提供了对控制的拓展。

大多数程序并不需要使用驱动函数，

 

3.1 用NVCC编译

CUDA内核函数可由CUDA推荐的固定结构书写——PTX（并行线程执行结构），也可由更高效（指编写高效）的高层次API如C编写。但都需要用NVCC编译成 能在设备上运行的二进制代码。

NVCC是一个编译器，它简化了C或PTX代码编写。

3.1.1 编译工作流

3.1.1.1 离线编译

NVCC主要编译流程如下

1. 将设备代码编译成汇编或二进制形式
2. 通过查询runtime加载函数和kernel调用函数修改主机代码，将<<<…>>>语法写入内核中

修改过的代码被输出成C的形式，可以用别的工具编译成目标代码。

此时程序可以

1. 链接成编译过的主机代码
2. 或 忽略修改过的主机代码并且使用CUDA驱动函数来 加载和执行PTX代码或CUBIN项目

3.1.1.2 运行时编译

任何PTX代码在程序执行时被调用，然后才由设备驱动（解释器）编译的这种编译叫做“运行时编译”。运行时编译增加了程序加载时间，但使PTX代码获得了因设备驱动器（解释器）优化而带来的性能优化。并且这也是运行一个未编译的PTX代码的唯一的方法。

当设备驱动为程序运行时编译了一些PTX代码时，它自动缓存了一份所编译的代码，防止接下来再编译，但当驱动升级时这些缓存将自动失效。

 

3.1.2 二进制兼容性

二进制代码与架构一一对应。一个CUDA可执行程序使用编译器选择-code编译成对应特异性架构的代码。举个栗子，使用选择-code=sm_35编译出的二进制代码能在计算能力为3.5的设备上运行。这种编译有向下兼容性，但不向上兼容（旧设备跑不了新代码）。

值得注意的是，二进制兼容性只支持pc端，不支持嵌入式的Tegra芯片。并且，这两者编译出的代码不能共用。

 

3.1.3 PTX（线程并行执行能力）兼容性

一些PTX特性只支持高等级的设备。例如，Warp Shuffle（瓦片乱序执行）只能在计算能力在3.0以上的设备上运行。编译设置”-arch”指明了目标设备的计算能力，如Warp Shuffle需要设置“-arch=compute_30”(或更高)

含有计算性能要求的PTX代码 编译成的目标代码总要求不小于这个计算性能要求。需要注意的是，从早期代码编译来的程序可能并没有完全利用硬件性能。如，编译为Pascal架构写的PTX代码，当它的目标设备是Volta架构时，可能没有用到Tensor加速核，而导致性能浪费。因此，旧代码编译到新设备上，可能表现会变糟。

 

3.1.4 程序兼容性

想要在特定设备上运行，程序必须加载二进制代码或与设备性能相符的PTX代码。

编译参数举例：

nvcc x.cu

              -gencode arch=compute_35,code=sm_35

              -gencode arch=compute_60,code=sm_35

译者注：我的设备是6.1能力的，经实践得出，这样编译超过计算能力的代码，将导致程序运行不成功，即使没有使用高等级设备的特性。

宏__CUDA_ARCH__标注了设备等级。如编译选择arch=compute_35,__CUDA_ARCH__就等于350

使用驱动API的程序，必须编译成分立文件，并且在执行时要详尽地加载和执行对应的文件。（这在说个啥，还没碰到过，先略过）

特别需要注意，Volta架构引入了独立线程调度，这会使得以前依赖于SIMT调度机制的代码运行混乱，得到错误的结果。所以，编译那些程序时，要特定-arch=compute_60 -code=sm_60而不能用默认的70

可以使用-arch=sm_35 来简化-arch=compute_35 -code=sm_35

 

3.1.5 C/C++兼容性

主机端支持全部C/C++，设备端只支持一部分，需要查表“C/C++ Language Support”

 

3.1.6 64位兼容性

64位的NVCC使用64位模式编译。设备机只在主机端代码以64位模式编译时编译成64位。32位的于此相同。

32位的NVCC也可以用-m64来编译64位设备代码

64位的NVCC也可以用-m32来编译32位设备代码。

 

3.2 CUDA C 运行时

运行时是调用cudart库，它与程序链接，通过cudart.lib or libcudaer.a静态链接，或通过cudart.dll or libcudart.so动态调用。动态调用需要调用库作为程序安装的一个部分被安装。这些链接的入口指针都被CUDA预先固定了。

设备内存：总览运行时使用设备内存情况。

共享内存：说明共享内存的使用，可以作为性能最大化评估点。

页锁定内存：使内核执行时，主机和设备之间内存映射重叠。

异步并行执行：说明异步并行执行在不同等级系统内的需求的概念和API。

多设备系统：说明如何将编程模型拓展到一个由单主机多设备的系统中去。

错误校对：描述如何正确地检测错误和在运行时生成错误提示。

调用栈：提到用于管理CUDA C调用的运行时函数。

纹理和表面内存：展示了纹理和表面内存这种使用内存的另一种方式。这也展示了GPU纹理化的内在操作。

图形（函数库）协同性：介绍了不同 运行时 可协同工作的函数，主要有图形api，OpenGL，Direct3D

 

3.2.1 初始化

第一步：创建CUDA上下文。这个上下文是设备的第一个上下文，它在主机程序中所有线程中共享。创建的同时，那些“运行时编译的设备代码”编译并且装入设备内存中。这些操作都是在底层，并且首要上下文对程序是透明的。

当主机调用cudaDeviceRestart()时，销毁当前主机线程创建的首要上下文。当其他线程新的运行时函数调用时，将创建新的首要上下文。

 

3.2.2 设备内存

设备内存可以像 访问“线性存储”或“CUDA数组”一样访问。

当使用纹理内存时，CUDA数组是不透明的内存结构（意思是要懂CUDA数组）

线性内存在设备的40位地址空间中，所以可以通过指针来访问那些分开存储的空间，就像二叉树那样。

线性内存的申请使用cudaMalloc(),用cudaFree()释放，主机和设备间搬运内存用cudaMemcpy()；

线性内存也可以通过cudaMallocPitch()和cudaMalloc3D()来申请。这要求数据符合2D、3D结构要求。对应的内存拷贝是cudaMemcpy2D()、cudaMemcpy3D();

下面的代码说明了一个width*height 2D数组，并且如何在设备中循环

__global__ void My2DArraySampleKernel(float *devPtr, size_t pitch, int width, int height) {

    for (int r = 0; r < height; r++) {

        float *row = (float*)((char*)devPtr + r * pitch);

        for (int c = 0; c < width; c++) {

            float element = row[c];

        }

    }

}

void host2DArraySample() {

    int width = 64, height = 64;

    float *devPtr;//设备存储空间的指针

    size_t pitch;

    cudaMallocPitch(&devPtr, &pitch, width * sizeof(float), height);

    My2DArraySampleKernel<<<100, 512 >>> (devPtr, pitch, width, height);

 

}

译者注：

 pitch的意思是

 

 

标签：代码,编程,内存,编译,CUDA,PTX,CUDA10.0,设备
来源： https://blog.csdn.net/lvxiangyu11/article/details/87906670