CUDA C 编程权威指南（7、8、9、10）

第 7 章 调整指令级原语 · 学习 CUDA 指令及其在应用程序行为中的作用 · 单浮点数和双浮点数的精确度对比 · 有关标准函数及 CUDA 内部函数的性能和精确度的实验 · 从不安全的内存访问中发现未定义行为 · 理解运算指令的意义和使用不当所产生的后果 CUDA 和其他遵循 IEEE——754 标准的编程模式支持两种浮点精度：32 位和 64 位。这些不同的格式也分别被称为单精度和双精度 7.1.2 内部函数和标准函数 CUDA 还将所有算数函数分成内部函数和标准函 数。标准函数用于支持可对主机和设备进行访问并标准化主机和设备的操作。标准函数包含来自于C 标准数学库的数学运算，如 sqrt 、 exp 和 sin 。单指令运算如乘法和加法，也包含在标准函数中 CUDA内置函数只能对设备代码进行访问。 在 CUDA 中，许多内部函数与标准函数是有关联的，这意味着存在与内部函数功能相 同的标准函数。举个例子，标准函数中的双精度浮点平方根函数也就是 sqrt 。有相同功能 的内部函数是 __dsqrt_rn 。还有执行单精度浮点除法运算的内部函数： __fdividef 。 7.1.3 原子操作指令 每个原子函数都能实现一个基本的数学运算，如加法、乘法或减法。不同于目前介绍过的其他指类型，当原子操作指令在两个竞争线程共享的 内存空间进行操作时，会有一个定义好的行为。 举例

__global__ void incr(int *ptr)
{
	int temp=*ptr;
	temp=temp+1;
	*ptr=temp;

}

这个内核从一个内存位置上读取一个数据，同时将其值加一，然后将得到的值写回到 相同位置。注意，这里没有使用线程 ID 来改变正在被访问的内存位置，内核启动时每个线 程都会从相同地址读写。如果启用一个含 32 个线程的线程块来运行这个核函数，那么会得 到什么样的输出？你可能会说对于 32 个线程，每个线程都会增加 1 。事实上，结果是不确 定的。这是因为不止一个线程对同一个内存位置进行写操作，这叫作数据竞争，或者称为 对内存的不安全访问。数据竞争的定义是两个或者多个独立的正在执行的线程访问同一个 地址，并且至少其中一个访问会修改该地址。直到程序真正被执行时，才能知道在这个过 程中哪一个线程赢得了胜利。因此，对于这个例子或任何会发生数据竞争的应用程序来 说，其结果是不能事先确定的。 使用原子操作指令可以避免这种事情的发生。

int atomicAdd(int *M,int V)

__global__ void incr(__global__ int *ptr)
{
	int temp=atomicAdd(ptr,1);
}

如果启动 32 个线程，存储在*ptr 所指位置中的值应该是 32 。 7.2 程序优化指令 7.2.1 单精度与双精度的比较 单精度和双精度浮点运算在通信和计算上的性能差异是 不可忽略的。在这种情况下，使用双精度数值能够使总的程序运行时间增加近一倍（虽然 这个结果可能取决于应用程序是计算密集型还是I/O 密集型）。在设备端进行数据通信的 时间也是使用单精度数值的两倍，这是由双精度数值长度是单精度数值长度的两倍造成 的。随着全局内存输入/ 输出数量和每条指令执行的位操作数量的增加，设备上的计算时间也会增加 单精度与双精度的结果有较大的数值差异，这些结果可能在迭代过程中不断被积累，即第一次迭代产生的不精确的结果作为下一次迭代的输入继续参与运算，导致最终结果偏差很大。因此，考虑到数值精确度，在迭代应用中可能更需要使用双精度变量。  

 7.2.2.2 操纵指令生成

7.2.3 了解原子指令

第 8 章 GPU 加速库和 OpenACC 利用 CUDA 库（ GPU 加速库）探索并行编程的新高度 · 了解多种 CUDA 库共享的通用工作流 · 尝试在线性代数、傅里叶变换和随机数中使用 CUDA 库 ·CUDA 6.0 中新的库特性 · 使用 OpenACC 指令在 GPU 上进行应用程序加速

8.1.2 通用的CUDA库工作流

 NVIDIA函数库的通用工作流如下所示：

1. 在库操作中创建一个特定的库句柄来管理上下文信息。 2. 为库函数的输入输出分配设备内存。 3. 如果输入格式不是函数库支持的格式则需要进行转换。 4. 将输入以支持的格式填入预先分配的设备内存中。 5. 配置要执行的库运算。 6. 执行一个将计算部分交付给 GPU 的库函数调用。 7. 取回设备内存中的计算结果，它可能是库设定的格式。 8. 如有必要，则将取回的数据转换成应用程序的原始格式。 9. 释放 CUDA 资源。 10. 继续完成应用程序的其他工作。 各种cuda库的细节，不详细解释，需要时看书 8.8 OpenACC 的使用 OpenACC 的线程模型与 CUDA 的线程模型类似，但添加了一个并行的维度。 Open- ACC可以分为 gang 、 worker 和 vector 3 个并行层次。在上层， gang 类似于 CUDA 线程块。一 个gang 可包含一个或多个执行的线程，在每个 gang 内部每个 gang 都包含一个或多个 worker。在 CUDA 中，一个 worker 类似于线程中的一个线程束。每个 worker 都有一个向量 宽度，由一个或多个同时执行相同指令的向量元素组成。每个向量元素都类似于一个 CUDA线程，因为它是一个单一的执行流。 OpenACC 和 CUDA 线程模型之间的主要区别在 于，OpenACC 在编程模型中直接指出了 worker 的概念（即线程束），而在 CUDA 中并没有明确建立线程束。 根据任务是否通过 gang 、 worker 、 vector 并行执行， OpenACC 执行被分成几种模式。现在，假设在一个OpenACC 程序的并行计算区域中，创建了 G 个 gang ，其中每个 gang 包含 W个 worker ，每个 worker 的向量宽度为 V 。那么，总共有 G×W×V 个执行线程处理这个并行 区域。 本章作了解 第 9 章 多 GPU 编程 到目前为止，本书中的大部分示例使用的都是单一的 GPU 。在本章中，会介绍多 GPU 编程的内容：在一个计算节点内或者跨多个GPU 加速节点实现跨 GPU 扩展应用。 CUDA 提 供了大量实现多GPU 编程的功能，包括：在一个或多个进程中管理多设备，使用统一的虚 拟寻址（Unified Virtual Addressing ， UVA ）直接访问其他设备内存， GPUDirect ，以及使用 流和异步函数实现的多设备计算通信重叠。在本章中需要掌握的内容有以下几个方面： · 在多 GPU 上管理和执行内核 · 跨 GPU 的重叠计算和通信 · 使用流和事件实现多 GPU 同步执行 · 在 GPU 加速集群上扩展 CUDA-aware MPI 应用程序 为了设计一个利用多 GPU 的程序，需要跨设备分配工作负载。根据应用程序，这种分 配会导致两种常见的 GPU 间通信模式： · 问题分区之间没有必要进行数据交换，因此在各 GPU 间没有数据共享 · 问题分区之间有部分数据交换，在各 GPU 间需要冗余数据存储 9.1.1 在多 GPU 上执行 下面的代码说明了如何确定使能 CUDA 设备的数量，对其进行遍历，并查询性能。

int ngpus;
cudaGetDeviceCount(&ngpus);
for(int i=0;i 
        
         显式地指定哪个
         GPU
         是当前所有 CUDA运算的目标。使用以下函数设置当前设备： 
         
         
         
cudaError_t cudaSetDevice(int id);
 
         
          没有显式地调用cudaSetDevice
          函数，那么当前设备会被自动设置设备
          0
          。 
          
          
          
          
           下面的代码准确展示了如何执行内核和在单一的主机线程中进行内存拷贝，使用循环 遍历设备： 
           
           
           
for(int i=0;i>>();
	cudaMemcpyAsync();//异步传输
}
 
           
            在单一节点内获取
            GPU
            的数量和它们的性能，可以使用下述函数： 
            
cudaError_t cudaGetDeviceCount(int *count);
cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp *prop,int id);
cudaError_t cudaDeviceReset(int id);
 
            
             9.1.2.1 
             启用点对点访问 
             
             
             
             
              点对点访问允许各
              GPU
              连接到同一个
              PCIe
              根节点上，使其直接引用存储在其他
              GPU
              设备内存上的数据。
              
             
              检查当前设备是否支持点对点访问
              
              
              
cudaError_t cudaDeviceCanAccessPeer(int * canAccessPeer,int device,int peerDevice);
 
              
               如果设备
               device
               能够直接访问对等设备
               peerDevice
               的全局内存，那么函数变量
               can-
               AccessPeer
               返回值为整型
               1
               ，否则返回
               0
               。 
               
              
               在两个设备间，必须用以下函数显式地启用点对点内存访问：
               
               
               
cudaError_t cudaDeviceEnablePeerAccess(int peerDevice,unsigned int flags);
 
               
                这个函数允许从当前设备到
                peerDevice
                进行点对点访问。
                flag
                参数被保留以备将来使 
                
               
                用，目前必须将其设置为
                0
                。一旦成功，该对等设备的内存将立即由当前设备进行访问。
                
                
                
                
                 这个函数允许从当前设备到
                 peerDevice
                 的访问，但 不允许从peerDevice
                 到当前设备的访问。如果希望对等设备能直接访问当前设备的内存，则需要另一个方向单独的匹配调用。 
                 
                 
                
                 显示禁用点对点访问
                 
                 
                 
cudaError_t cudaDeviceDisablePeerAccess(int peerDevice);
 
                 
                  32
                  位应用程序不支持点对点访问。
                  
                  
                  
                  
                   9.1.3 
                   多
                   GPU
                   间的同步（做了解）
                   
                   
                   
                    1.
                    选择这个应用程序将使用的
                    GPU
                    集。 
                    
                   
                    2.
                    为每个设备创建流和事件。 
                    
                   
                    3.
                    为每个设备分配设备资源（如设备内存）。 
                    
                   
                    4.
                    通过流在每个
                    GPU
                    上启动任务（例如，数据传输或内核执行）。 
                    
                   
                    5.
                    使用流和事件来查询和等待任务完成。 
                    
                   
                    6.
                    清空所有设备的资源。
                    
                    
                    
                    
                     9.3 
                     多
                     GPU
                     上的点对点通信
                     
                    
                     9.3.3 统一虚拟寻址的点对点内存访问 
                     
                     
                    
                     判断当前设备是否支持统一寻址
                     
                     
                     
int deviceId=0;
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop,deviceId);
printf("GPU %d: %sunified addressingn",deviceId,prop.unifiedAddressing?"supports":"does not supports");
 
                     
                      本章后续内容仅作了解
                      
                      
                     
                      第
                      10
                      章 程序实现的注意事项 
                      
                      
                      
                       本章包含了
                       CUDA C
                       项目开发的以下几个方面： 
                       
                      
                       ·CUDA C
                       的开发过程 
                       
                      
                       ·
                       配置文件驱动优化 
                       
                      
                       ·CUDA
                       开发工具 
                       
                      
                       本章结尾提供了一个案例，逐步将
                       C
                       语言移植到
                       CUDA C
                       中，这会有助于方法的理解，形象化整个过程并说明本章涉及的工具。
                       
                       
                      
                       10.1.1.3 
                       优化 
                       
                      
                       写好c程序之后，进入下一阶段， 
                       
                        基 于CUDA的优化可以体现在以下两个层次上 
                        
                         ·
                         网格级（grid-level） 
                         
                        
                         ·
                         内核级（kernel-level） 
                         
                        
                       
                       
                       
                        优化网格级性能的方法包括同时运行多个内核以及使用CUDA
                        流和事件重叠带有数据的内核执行
                        
                        
                        
                         限制内核性能的主要原因有
                         3
                         个： 
                         
                        
                         ·
                         内存带宽 
                         
                        
                         ·
                         计算资源 
                         
                        
                         ·
                         指令和内存延迟 
                         
                         
                         
                         
                          在内核级，
                          CUDA
                          采用划分方法分配计算资源：寄存器在线程间被划分，共享内存在 
                          
                         
                          线程块间被划分。因此，内核中的资源消耗可能会限制活跃线程束的数量。 
                          
                         
                          在网格级，
                          CUDA
                          使用由线程块组成的网格来组织线程的执行，通过指定如下内容， 
                          
                         
                          可以自由选择最佳的内核启动配置参数： 
                          
                         
                          ·
                          每个线程块中线程的数量 
                          
                         
                          ·
                          每个网格中线程块的数量 
                          
                          
                          
                          
                          
                           来自每一个内核的内存请求（加载或存储）都是由单个线程束发出的。线程束中的每 
                           
                          
                           个线程都提供了一个内存地址，基于提供的内存地址，
                           32
                           个线程一起访问一个设备内存
                           块。
                           
                           
                           
                           
                            对于全局内存来说，最好的访问模式是对齐和合并访问。对齐内存访问要求所需的设 
                            
                           
                            备内存的第一个地址是
                            32
                            字节的倍数。合并内存访问指的是，通过线程束中的
                            32
                            个线程来 
                            
                           
                            访问一个连续的内存块。 
                            
                            
                            
                            
                             10.1.3 CUDA
                             代码编译
                             
                            
 
                             
                            
                             
                             
                            
                           
                          
                         
                        
                       
                      
                     
                    
                   
                  
                 
                
               
              
             
            
           
          
         
        
       
      
     
    
   
  
 
10.1.3.1 独立编译
 
 

 1.
 设备编译器将可重新定位的设备代码嵌入到主机目标文件中。 
 

 2.
 设备链接器结合设备对象。 
 

 3.
 主机链接器将设备和主机对象组合成一个最终可执行的程序。 
 
 
 
 
  考虑一个简单例子，其中有
  a.cu
  ，
  b.cu
  ，
  c.cpp 3
  个文件。假设
  a.cu
  文件中的一些核函数 
  
 
  引用
  b.cu
  文件中的一些函数或变量，因为是跨文件引用的，所以就必须使用独立编译来生 
  
 
  成可执行文件。
  
  
 
 
 
  
  
 
10.1.3.2 Makefile示例文件
 
 

 需要更换
 Makefile
 示例文件中完整的路径名称并更新可执行的文件名，以与工作环境相匹配。
 
 
10.1.3.3 将CUDA文件整合到C项目中  
 

 从
 C
 代码中调用
 CUDA
 运 行时函数，需要在主机代码中包含C
 运行时头文件，如下所示： 
 
 
 
#include
 
 
为了使cuda核函数能够像c函数那样被调用，需要用以下的声明来解决C++引用混乱的问题：
 
 
extern "C" void warpper_kernel_alunch()
{
    ...
}
 
 
  以便它可以与
  C
  代码链接。
  
 
 
  
 
  
  
 
10.1.4 CUDA错误处理 
 
cudaPeekLastError 
cudaGetLastError
cudaGeErrorString
 
常用的就是CHECK或者CALL_CUDA
 
#define CHECK(call)
{
    cudaError_t err;
    if((err=(call))!=cudaSuccess)
    {
        fprintf(stderr,"Got error %s :%dn",cudaGetErrorString(err),__FILE__,__LINE__);
        exit (1);
    }
}
 

 10.2 
 配置文件驱动优化
 
 

 10.3 CUDA
 调试