一、简介

1. 兼容性信息

  1. 对于RTX 4050，Driver版本是12.1，所以 CUDA Runtime应小于12.1，对于4050显卡来说，已经查到 CUDA Runtime版本可以使用11.5–11.8（对于40系显卡来说目前都是这样），所以下载11.8版本，总之不论如何，一定要有 CUDA Driver 版本 >= CUDA Runtime 版本。但高太多可能会出问题。
  2. 40系目前兼容的最早cuda版本为11.8. 下载的pytorch也需要与这个版本对应.
  3. CUDA Toolkit可以认为是一个软件安装包，它可以安装cuda driver，nvcc(编译器)，libraries, CUDA Samples。传统上，安装 NVIDIA Driver 和 CUDA Toolkit 的步骤是分开的，但实际上我们可以直接安装 CUDA Toolkit，系统将自动安装与其版本匹配的 NVIDIA Driver。下面我们讲述安装 CUDA Toolkit 的方法。
  4. 为什么nvcc和nvidia-smi显示的CUDA版本不同 (参看stackoverflow，很清晰)CUDA有两个主要的API：runtime(运行时) API和driver API，这两个API都有对应的CUDA版本。driver API是通过GPU driver installer安装的，可理解为系统出厂安装的默认驱动，nvidia-sim显示的是 driver API。runtime API是通过CUDA toolkit安装的，可理解为用户自己安装的驱动，nvcc显示的是runtime API。
  5. CUDA软件主要包括三部分：
    • CUDA Toolkit： 库文件、运行环境和开发工具， 主要是面向开发者CUDA编译环境。
    • CUDA Driver：用户驱动组件，用于运行CUDA程序，可以理解为CUDA运行环境。
    • Nvidia GPU 驱动：显卡核心驱动，就是硬件驱动

CUDA 5.0后有一个工具叫做nvprof的命令行分析工具

核函数支持 c++ 的重载机制。

版本和兼容性(Versioning and Compatibility)

CUDA driver API的版本定义在驱动的头文件中的CUDA_VERSION宏内。可以在程序中将该宏调出，以检查程序是否可以在目标设备上运行。

针对新版本的CUDA driver API编译的程序、插件、库，并不能在旧版本的驱动上运行。

关于CUDA driver API：

  1. 由于一个系统只能装一个版本的驱动。因此驱动版本要足够高。
  2. 默认情况下，nvcc编译程序时，库和插件是静态编译的。静态编译不要求库和插件的驱动版本和CUDA运行库保持一致。但是动态链接则要求版本一致。

Pascal及以上架构，支持指令级的优先级调度。不再是以线程块为SM的最小调度单位，而是以指令为最小调度单位，且具有优先级。 这意味着具有冗长kernel的线程块不再会占据太多的计算资源，或是发生timeout。但是这也有缺点：当多个进程创建了上下文时，以往基于线程块的调度不会造成太多的上下文切换，但现在的指令级调度则会造成很多的上下文切换，降低效率。(注意跟GPU内线程的上下文切换不同，GPU内线程上下文切换几乎不浪费时间，直接换一个指针就好)。因此最好设置为Exclusive-process，这样只有一个进程会使用设备。

设备函数可以有返回值。

cudaThreadSynchronize()在cuda10.0以后被弃用了，用 cudaDeviceSynchronize() 来代替。与父块的子内核显式同步（即在设备代码中使用 cudaDeviceSynchronize()）在 CUDA 11.6 中已弃用，并计划在未来的 CUDA 版本中删除，Note that this is different from host-side cudaDeviceSynchronize, which is still supported。

那些可以到达__syncthreads()的线程只需要等待其他可以到达该点的线程，而不是等待块内所有其他线程。

NVVP是可视化的非常实用的工具。

限制内核性能的主要包括但不限于以下因素

  • 存储带宽
  • 计算资源
  • 指令和内存延迟

解决虚假依赖的最好办法就是多个工作队列，这样就从根本上解决了虚假依赖关系，Hyper-Q就是这种技术，32个硬件工作队列同时执行多个流，这就可以实现所有流的并发，最小化虚假依赖。

对于2D或3D数组，可以使用cudaMallocPitch()和cudaMalloc3D()来分配内存。这两个函数会自动padding，以满足内存对齐的要求，提高内存读写效率。

语法

可以同时用 __host__ 和 __device__ 修饰函数，从而减少代码冗余，此时编译器将分别在主机和设备上编译该函数. 可以通过 __noinline__ 建议编译器不要将一个设备函数当作内联函数。

二、线程组织

网格和块

网格和块的维度一般是二维和三维的，也就是说一个网格通常被分成二维的块，而每个块常被分成三维的线程。

线程块被分配到某一个SM上以后，将分为多个线程束，每个线程束一般是32个线程。当SM内的资源没办法处理一个完整块，程序将无法启动或者返回错误值。

GPU中每个SM都能支持数百个线程并发执行，每个GPU通常有多个SM，当一个核函数的网格被启动的时候，多个block会被同时分配给可用的SM上执行。同一个SM上可以有不止一个常驻的线程束，有些在执行，有些在等待，他们之间状态的转换是不需要开销的。

线程束

由于硬件设计，调度会以一整个线程束为单位进行。因此应该充分利用线程束中的所有线程进行计算，以尽可能少的线程束的数量，同时提高硬件的使用率(现在基本上是Volta及更新的架构)。

设置线程块大小时，为了补齐32的倍数而创建的线程，如线程块定义为31个线程，则会补1个线程，这个线程是inactive的。

线程束一旦被激活来到片上，那么他就不会再离开SM直到执行结束。线程束一旦被激活来到片上，那么他就不会再离开SM直到执行结束。

线程束调度器(warp scheduler)会选择ready状态的线程束，将其调度到SP上执行。

分支分化

从Volta架构开始，Independent Thread Scheduling被引入，线程束内的线程不再完全同步。每个线程都会有自己独立的PC。遇到分支时，不再像之前的架构一样，只有(半)线程束内的线程条件一致时，才会跳过分支；Volta架构的调度优化器会将线程束中的线程，按照分支条件是否满足，重新组合成SIMT单元，从而跳过分支。

Volta架构的Independent Thread Scheduling无疑是高效的，但是这是一个跟旧架构完全不同的特性。在编写旧架构的CUDA程序时，程序员会默认线程束内的线程一定会同步执行。Volta架构的新特性破坏了这一假设，无疑会给程序带来一些问题，需要注意。

并行规约问题

动态并行

动态并行，相当于串行编程的中的递归调用。

父线程块启动子网格需要显示的同步。父网格和子网格共享相同的全局和常量内存。父网格子网格有不同的局部内存。

三、内存组织

程序具有局部性特点。时间局部性，就是一个内存位置的数据某时刻被引用，那么在此时刻附近也很有可能被引用，随时间流逝，该数据被引用的可能性逐渐降低。空间局部性，如果某一内存位置的数据被使用，那么附近的数据也有可能被使用。

在CPU编程中，我们学过空间局部性与时间局部性。对全局内存的访问要尽量进行合并访问与存储，让对内存的访问更加集中(访问的内存地址尽量连续)，这样才能达到最大的带宽。

GPU上的内存设备有：寄存器、共享内存、本地内存、常量内存、纹理内存、全局内存。

CUDA中每个线程都有自己的私有的本地内存；线程块有自己的共享内存，对线程块内所有线程可见；所有线程都能访问读取常量内存和纹理内存，但是不能写，因为他们是只读的。

不同块内线程不能相互影响，他们是物理隔离的。

寄存器

在核函数内不加修饰的声明一个变量，此变量就存储在寄存器中。在核函数中定义的有常数长度的数组也是在寄存器中分配地址的。寄存器对于每个线程是私有的，寄存器变量的声明周期和核函数一致。寄存器是SM中的稀缺资源，一个线程如果使用更少的寄存器，那么就会有更多的常驻线程块，SM上并发的线程块越多。

为了避免寄存器溢出，可以在核函数的代码中配置额外的信息来辅助编译器优化

// (线程块内包含的最大线程数, 每个SM中预期的最小的常驻内存块参数)
__global__ void __lauch_bounds__(maxThreadaPerBlock,minBlocksPerMultiprocessor) kernel(...) 
{
    /* kernel code */
}

可以在编译选项中加入-maxrregcount=32来控制一个编译单元里所有核函数使用的最大数量。

本地内存

核函数中符合存储在寄存器中但不能进入被核函数分配的寄存器空间中的变量将存储在本地内存中。本地内存实质上是和全局内存一样在同一块存储区域当中的，其访问特点——高延迟，低带宽。对于2.0以上的设备，本地内存存储在每个SM的一级缓存，或者设备的二级缓存上。

共享内存

每个SM都有一定数量的由线程块分配的共享内存，共享内存是片上内存，跟主存相比，速度要快很多，也即是延迟低，带宽高。其类似于一级缓存，但是可以被编程。

共享内存在核函数内使用修饰符__share__声明，生命周期和线程块一致。因为共享内存是块内线程可见的，使用__syncthreads可缓解竞争。

声明一个二维浮点数共享内存数组的方法是: __shared__ float a[size_x][size_y];

填充静态声明的共享内存: _shared__ int tile[BDIMY][BDIMX+IPAD];。唯一要注意的就是索引，当填充了以后我们的行不变，列加宽了，所以索引的时候要相应改变。

一个线程块使用的共享内存过多，导致更多的线程块没办法被SM启动，影响活跃的线程束数量。

动态声明一个共享内存数组，可以使用extern关键字，并在核函数启动时添加第三个参数。注意，动态声明只支持一维数组。

extern __shared__ int tile[];
kernel<<<grid,block,isize*sizeof(int)>>>(...); // isize就是共享内存要存储的数组的大小。比如一个十个元素的int数组，isize就是10.

当多个线程要访问一个存储体的时候，冲突就发生了，注意这里是说访问同一个存储体，而不是同一个地址。

访问模式查询: 可以通过以下语句，查询是4字节还是8字节：

cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig * pConfig);

在可以配置的设备上，可以用下面函数来配置新的存储体大小：

cudaError_t cudaDeviceSetShareMemConfig(cudaSharedMemConfig config);

每个SM上有64KB的片上内存，共享内存和L1共享这64KB，并且可以配置。CUDA为配置一级缓存和共享内存提供以下两种方法： 按设备进行配置(cudaError_t cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig);) 按核函数进行配置(cudaError_t cudaFuncSetCacheConfig(const void* func,enum cudaFuncCacheca cheConfig);)

一级缓存和共享内存都在同一个片上，但是行为大不相同，共享内存靠的的是存储体来管理数据，而L1则是通过缓存行进行访问。我们对共享内存有绝对的控制权，但是L1的删除工作是硬件完成的。

共享内存有个特殊的形式是，分为32个同样大小的内存模型，称为存储体，可以同时访问。32个存储体的目的是对应一个线程束中有32个线程，这些线程在访问共享内存的时候，如果都访问不同存储体（无冲突），那么一个事务就能够完成，否则（有冲突）需要多个内存事务了，这样带宽利用率降低。

同步基本方法：障碍、内存栅栏。障碍是所有调用线程等待其余调用线程达到障碍点。内存栅栏，所有调用线程必须等到全部内存修改对其余线程可见时才继续进行。

弱排序内存模型: 核函数内连续两个内存访问指令，如果独立，其不一定哪个先被执行。Volatile修饰符声明一个变量，则对该变量同一时间只能执行一个指令。

共享内存的数据布局

在CPU中，如果用循环遍历二维数组，尤其是双层循环的方式，我们倾向于内层循环对应x，因为这样的访问方式在内存中是连续的。但是GPU的共享内存并不是线性的，而是二维的，分成不同存储体的，并且，并行也不是循环。

我们的数据是按照行放进存储体中的。所以x[threadIdx.y][threadIdx.x]这种访问方式是最优的，threadIdx.x在线程束中体现为连续变化的，而对应到共享内存中也是遍历共享内存的同一行的不同列。

内存填充
配置共享内存

内存栅栏???

  1. 线程块内 void __threadfence_block(); 保证同一块中的其他线程对于栅栏前的内存写操作可见

  2. 网格级内存栅栏 void __threadfence(); 挂起调用线程，直到全局内存中所有写操作对相同的网格内的所有线程可见

  3. 系统级栅栏，夸系统，包括主机和设备， void __threadfence_system(); 挂起调用线程，以保证该线程对全局内存，锁页主机内存和其他设备内存中的所有写操作对全部设备中的线程和主机线程可见。

volatile修饰符声明一个变量，防止编译器优化，volatile声明的变量始终在全局内存中。

常量内存

常量内存驻留在设备内存中，每个SM都有专用的常量内存缓存，常量内存使用关键字__constant__。

常量内存在核函数外，全局范围内声明，并对同一编译单元中的所有核函数可见。声明时就申请了一块常量内存。主机端代码可以初始化常量内存的,不能被核函数修改。初始化函数如下：

cudaError_t cudaMemcpyToSymbol(const void* symbol,const void *src,size_t count);

当线程束中所有线程都从相同的地址取数据时，常量内存表现较好，但是如果不同的线程取不同地址的数据，常量内存就不那么好了，因为常量内存的读取机制是：一次读取会广播给所有线程束内的线程。

纹理内存

???

全局内存

一般在主机端代码里定义，也可以在设备端定义，不过需要加修饰符，只要不销毁，是和应用程序同生命周期的。全局内存访问是对齐。

GPU缓存

GPU缓存不可编程，其行为出厂是时已经设定好了。GPU上有4种缓存：一级缓存、二级缓存、只读常量缓存、只读纹理缓存。每个SM都有一个一级缓存，所有SM公用一个二级缓存。CUDA允许我们配置读操作的数据是使用一级缓存和二级缓存，还是只使用二级缓存。每个SM有一个只读常量缓存，只读纹理缓存，它们用于设备内存中提高来自于各自内存空间内的读取性能。

只读缓存对于分散访问的更好。当所有线程读取同一地址的时候常量缓存最好，只读缓存这时候效果并不好，只读缓存粒度为32。使用函数 _ldg从只读缓存读取。

__global__ void copyKernel(float * in,float* out)
{
    int idx=blockDim*blockIdx.x+threadIdx.x;
    out[idx]=__ldg(&in[idx]);

}

静态全局内存

数据传输

数据传输的异步版本API如下。

cudaError_t cudaMemcpyAsync(void* dst, const void* src, size_t count,cudaMemcpyKind kind, cudaStream_t stream = 0);

执行异步数据传输时，主机端的内存必须是固定的，非分页的。

内存管理

在数据传输之前，CUDA驱动会锁定页面，或者直接分配固定的主机内存，将主机源数据复制到固定内存上，然后从固定内存传输数据到设备上。下面函数用来分配固定内存：

cudaError_t cudaMallocHost(void ** devPtr,size_t count); // 分配锁页内存

分配count字节的固定内存，这些内存是页面锁定的，可以直接传输到设备的. 固定的主机内存释放使用：

cudaError_t cudaFreeHost(void *ptr);

固定内存的释放和分配成本比可分页内存要高很多，但是传输速度更快，所以对于大规模数据，固定内存效率更高。

使用锁页内存后，可以将锁页内存映射到设备内存上。使用PCI-E的DMA与设备内存进行数据传输，而不需要CPU来搬运数据。可以使用流和内存映射，来让CPU程序、数据传输和内核执行并行。

合并写内存(Write-Combining Memory)???

零拷贝内存/内存映射(Mapped Memory)

零拷贝内存是固定内存，不可分页。 通过以下函数创建零拷贝内存：

cudaError_t cudaHostAlloc(void ** pHost,size_t count,unsigned int flags);

cudaHostAllocDefalt和cudaMallocHost函数一致，cudaHostAllocPortable函数返回能被所有CUDA上下文使用的固定内存，cudaHostAllocWriteCombined返回写结合内存，在某些设备上这种内存传输效率更高。cudaHostAllocMapped产生零拷贝内存。

如果要使用内存映射，必须在其他CUDA函数执行前，执行cudaSetDeviceFlags()并传入cudaDeviceMapHost，来使能设备的内存映射功能。否则cudaHostGetDevicePointer()函数会返回error。

零拷贝内存虽然不需要显式的传递到设备上，但是设备还不能通过pHost直接访问对应的内存地址，设备需要访问主机上的零拷贝内存，需要先获得另一个地址，这个地址帮助设备访问到主机对应的内存，方法是：

cudaError_t cudaHostGetDevicePointer(void ** pDevice,void * pHost,unsigned flags);

pDevice就是设备上访问主机零拷贝内存的指针了。此处flag必须设置为0。使用内存映射，可以让CPU/GPU之间的数据传输隐式执行，而不需要显示的分配GPU内存并传输数据。

统一虚拟寻址UVA

设备架构2.0以后，cudaHostGetDevicePointer基本上可以不再使用。

统一内存寻址

CUDA6.0出现统一内存寻址。

cudaError_t cudaMallocManaged(void ** devPtr,size_t size,unsigned int flags=0);

CUDA6.0中设备代码不能调用cudaMallocManaged，只能主机调用，所有托管内存必须在主机代码上动态声明，或者全局静态声明. 页面故障???

线程束洗牌指令

一个线程束可以看做是32个内核并行执行这32个核函数中寄存器变量在硬件上相邻，为相互访问提供了物理基础。线程束内线程相互访问数据不通过共享内存或者全局内存，使得通信效率高很多，线程束洗牌指令传递数据，延迟极低，且不消耗内存。

// 
int __shfl(int var,int srcLane,int width=warpSize);
// 从与调用线程相关的线程中复制数据: 调用线程得到当前束内线程编号减去delta的编号的线程内的var值
int __shfl_up(int var,unsigned int delta,int with=warpSize);
// 从与调用线程相关的线程中复制数据: 调用线程得到当前束内线程编号加上delta的编号的线程内的var值
int __shfl_down(int var,unsigned int delta,int with=warpSize);
// 
int __shfl_xor(int var,int laneMask,int with=warpSize);

流和事件

事件通过下面指令添加到CUDA流：

cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0);

CDUA提供了一种控制事件行为和性能的函数：

cudaError_t cudaEventCreateWithFlags(cudaEvent_t* event, unsigned int flags);

空流不需要显式声明，而是隐式的，他是阻塞的，跟所有阻塞流同步。默认创建的非空流是阻塞版本。

隐式同步

忽略隐式同步会造成性能下降。隐式同步常出现在内存操作上，比如锁页主机内存分布、设备内存分配、设备内存初始化、同一设备两地址之间的内存复制、一级缓存和共享内存配置修改。

显式同步

常见的显式同步有同步设备、同步流、同步流中的事件、使用事件跨流同步。

  1. 阻塞主机线程，直到设备完成所有操作： 
    cudaError_t cudaDeviceSynchronize(void);

    这个函数我们前面常用，但是尽量少用，这个会拖慢效率。

  2. 使用流的同步

流

没有声明流的那些GPU操作指令，核函数是在空流上执行的，空流是阻塞流，同时我们声明定义的流如果没有特别指出，声明的也是阻塞流。 流的使用：

cudaStream_t stream; // 声明
cudaError_t cudaStreamCreate(cudaStream_t stream); // 创建默认的阻塞流
cudaError_t cudaStreamCreateWithFlags(cudaStream_t* pStream, unsigned int flags); // 创建流并设置是否阻塞
cudaError_t cudaStreamCreateWithPriority(cudaStream_t* stream, unsigned int flags,int priority); // 创建流并设定优先级
kernel_name<<<grid, block, sharedMemSize, stream>>>(argument list); // 添加
cudaError_t cudaStreamQuery(cudaStream_t stream); // 查询
cudaError_t cudaStreamSynchronize(cudaStream_t stream); // 阻塞
cudaError_t cudaStreamDestroy(cudaStream_t stream); // 销毁

cudaStreamSynchronize会阻塞主机，直到流完成。cudaStreamQuery则是立即返回，如果查询的流执行完了，那么返回cudaSuccess否则返回cudaErrorNotReady。

事件

事件的作用就是在流中设定一些标记用来同步，和检查是否执行到关键点位（事件位置），也是用类似的函数

cudaEvent_t event; // 声明
cudaError_t cudaEventCreate(cudaEvent_t* event); // 创建默认事件
cudaError_t cudaEventCreateWithFlags(cudaEvent_t* event, unsigned int flags); // 创建事件并设定事件的行为和性能
cudaError_t cudaEventQuery(cudaEvent_t event); // 异步查询
cudaError_t cudaEventSynchronize(cudaEvent_t event); // 同步查询
cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t start, cudaEvent_t stop); // 记录两个事件之间的时间间隔
cudaError_t cudaEventDestroy(cudaEvent_t event); // 销毁

事件有四种配置

  1. cudaEventDefault:
  2. cudaEventBlockingSync: 使用cudaEventSynchronize同步,阻塞调用线程
  3. cudaEventDisableTiming: 表示事件不用于计时，可以减少系统不必要的开支
  4. cudaEventInterprocess: 可能被用于进程之间的事件

事件提供了一个流之间同步的方法：

cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event);

这条命令的含义是，指定的流要等待指定的事件，事件完成后流才能继续，这个事件可以在这个流中，也可以不在，当在不同的流的时候，这个就是实现了跨流同步。

用环境变量调整流行为

对于Linux系统中，修改方式如下：

#For Bash or Bourne Shell:
export CUDA_DEVICE_MAX_CONNECTIONS=32
#For C-Shell:
setenv CUDA_DEVICE_MAX_CONNECTIONS 32

另一种修改方法是直接在程序里写，这种方法更好用通过底层驱动修改硬件配置：

setenv("CUDA_DEVICE_MAX_CONNECTIONS", "32", 1);

创建流间依赖关系

声明成 cudaEventDisableTiming 的同步事件：

cudaEvent_t * event=(cudaEvent_t *)malloc(n_stream*sizeof(cudaEvent_t));
for(int i=0;i<n_stream;i++)
{
  cudaEventCreateWithFlag(&event[i],cudaEventDisableTiming);
}

在流中加入指令:

for(int i=0;i<n_stream;i++)
{
    kernel_1<<<grid,block,0,stream[i]>>>();
    kernel_2<<<grid,block,0,stream[i]>>>();
    kernel_3<<<grid,block,0,stream[i]>>>();
    kernel_4<<<grid,block,0,stream[i]>>>();
    cudaEventRecord(event[i],stream[i]);
    cudaStreamWaitEvent(stream[n_stream-1],event[i],0);
}

这时候，最后一个流（第5个流）都会等到前面所有流中的事件完成，自己才会完成.

重叠内核执行和数据传输

  1. 深度优先调度重叠
  2. 广度优先调度重叠

数据传输使用异步方式，注意异步处理的数据要声明称为固定内存，不能是分页的，如果是分页的可能会出现未知错误。

流回调

流函数有特殊的参数规格，必须写成下面形式参数的函数;

void CUDART_CB my_callback(cudaStream_t stream, cudaError_t status, void *data) {
    printf("callback from stream %d\n", *((int *)data));
}

然后使用：

cudaError_t cudaStreamAddCallback(cudaStream_t stream,cudaStreamCallback_t callback, void *userData, unsigned int flags);

加入流中。

使用NVCC编译CUDA程序

离线编译

NVCC进行离线编译的操作流程是： 分离CUDA程序中的主机端代码(host code)和设备端代码(device code) 将设备端代码编译成一种虚拟汇编文件(名为PTX)，再接着编译成二进制代码(名为cubin) 将主机端代码中含有”«<»>”的代码(即内核调用)替换为CUDA运行库中的函数调用代码 之后NVCC会借助其他编译器(如gcc)将主机端代码编译出来 * 主机端代码和设备端代码被编译好后，nvcc会将两段代码链接起来。

如果希望编译出来的文件能在更多的GPU上运行，则可以同时指定多组计算能力，例如：

-gencode arch=compute_35, code=sm_35
-gencode arch=compute_50, code=sm_50
-gencode arch=compute_60, code=sm_60

此时，编译出来的可执行文件将包含3个二进制版本，称为 胖二进制文件（fatbinary）。

在线编译(JIT Compilation)

设备驱动(device driver)会负责在运行时，使用PTX代码生成二进制代码。这个过程被称作在线编译(JIT Compilation, Just-In-Time Compilation)。

当设备驱动程序有关编译的部分得到优化时，同样的PTX编出来的cubin文件同样会得到优化。

二进制代码的兼容性

二进制代码cubin是受到GPU计算能力的限制的。在编译时，需要使用-code来指定将代码编译到哪个计算能力平台上，如-code=sm_35代表生成的cubin代码是运行在计算能力为3.5的平台上的。

二进制代码若要兼容，首先架构得一致。不同架构上的二进制代码不能互相兼容。其次，若执行平台的次版本号版本比编译时指定的的次版本号高，则可以运行。例如如果在编译时指定-code=sm_35，则在计算能力3.7的平台上也可以运行。反之则不可以。

PTX代码的兼容性

若PTX代码使用了较高级别架构的特有特性，则无法在较低架构上运行。若PTX在较低架构上生成，则虽然能够在所有更高级别的架构上运行，但无法充分利用这些架构的硬件特性。

应用程序兼容性

7.0以前，都是以线程束为单位在调度，线程束内指令永远是同步的，被成为锁步。Volta架构(计算能力7.x)引入了Independent Thread Scheduling，破坏了线程束内的隐式同步。因此，如果老版本的代码里面有默认锁步的代码，在Volta架构下运行时可能会因为锁步的消失而出问题，可以指定-arch=compute_60 -code=sm_70，即将PTX编到Pascal架构下以禁用Independent Thread Scheduling特性。（当然，也可以修改代码来显示同步）

C/C++兼容性

对于主机端代码，nvcc支持C++的全部特性；而对于设备端代码，只支持C++的部分特性。具体查阅手册。

编程接口

回调函数(Callbacks)

可以使用cudaStreamAddCallback()函数，向流中添加callback。该callback会在流中之前所有的任务完成后被调用。如果stream参数设为0，则代表之前的所有stream的任务执行完后就调用该callback。

回调函数和cudaStreamWaitEvent()一样，对于在加在callback之后的指令，必须等待callback执行完成后，才会继续执行。

流的优先级(Stream Priorities)

运行时，高优先级stream中的线程块不能打断正在执行的低优先级stream的线程块(即不是抢占式的)。

设备选择(Device Selection)

使用cudaSetDevice()选择设备，当不选择时，默认使用设备0。

(多设备下)流和事件的执行情况

  1. 内存拷贝(will success)：如果对一个不属于当前设备的流进行内存拷贝工作，内存拷贝会成功。
  2. cudaEventSynchronize() and cudaEventQuery()(will success)：即使处于不同的设备，事件同步和事件查询依然有效。
  3. cudaStreamWaitEvent()(will success)：比较特殊，即使函数输入的流和事件不在同一个设备上，也能成功执行。也就是说，可以让流等待另一个设备上(当然当前设备也可以)的事件。这个函数可以用作多个设备间的同步。

(设备间)对等内存访问(Peer-to-Peer Memory Access)

计算能力2.0及以上的设备支持设备间对等内存访问，这意味着两个GPU之间的传输和访问可以不经过主机端中转，速度会有提升。查询cudaDeviceCanAccessPeer()可以得知设备是否支持这一特性。

进程间通讯(Interprocess Communication)

通过IPC中的cudaIpcGetMemHandle()，可以得到设备内存指针的IPC句柄。该句柄可以通过标准的IPC机制(interprocess shared memory or files)传递到另一个进程，再使用cudaIpcOpenMemHandle()解码得到该进程可以使用的设备内存指针。

错误检查(Error Checking)

对于异步函数，由于在执行前就会返回，因此返回的error code仅仅代表函数启动时的错误(如参数校验)。

内核函数不会返回值，因此只能通过cudaPeekAtLastError()或cudaGetLastError()来知悉调用内核是否有错误。

性能优化

延迟隐藏

延迟(latency)指的是线程束(从上一个动作开始)到它处于ready状态的时钟数。例如线程束先提交了一个内存访问请求，然后等了400个时钟周期，内存管理系统才返回数据，线程束可以继续执行。这400个时钟周期称为延迟。当一个线程束发生延迟时，线程束调度器(warp scheduler)会将其他处于ready状态的线程束调度到SP上。等到延迟结束后，再将该线程调度回SP继续执行。这样一来，前一个线程束的延迟，就被另一个线程束的执行所隐藏了。这一过程被称作延迟的隐藏(hidden latency)。

两种延迟

  • 算数指令延迟是一个算术操作从开始，到产生结果之间的时间，这个时间段内只有某些计算单元处于工作状态，而其他逻辑计算单元处于空闲。
  • 内存指令延迟很好理解，当产生内存访问的时候，计算单元要等数据从内存拿到寄存器，这个周期是非常长的。

造成线程(束)产生延迟的原因有：

  1. 指令执行：不同指令有不同的执行延迟
  2. 内存请求：共享内存、全局内存、PCI-E(Mapped Memory)的读写请求
  3. 同步操作：如使用__syncthreads()后，先完成的线程(束)，会等待线程块中其他线程(束)达到同步点。__syncthreads()只能同步同一个块内的线程，不能同步不同块内的线程。

当使用内存映射时，需要注意，每次内存访问都会启动一次PCI-E传输。因此，尽量保证数据只被读写一次，且尽可能合并访问以提升有效内存带宽。

在设计内核时，线程束内的线程尽量连续的访问内存。

可以使用#pragma unroll宏，来进行循环展开，减少控制指令。

计算能力 __syncthreads()消耗的指令周期
3.x 128
5.x,6.1,6.2 64
6.0 32
7.x 16

参考

  1. 人工智能编程

???矩阵的转置