15.8_延伸阅读

15.8 延伸阅读

《数字图像处理》一书中对归一化相关（pp. 583∼586583\sim 586 ）和对数变换（pp. 168∼169168\sim 169 ）有相关讨论，其中对数变换部分在我们的示例程序中被用来计算输出像素：

Gonzalez, Rafael C., and Richard E. Woods. Digital image processing. Addison-Wesley, Reading, MA, 1992. www(imageprocessingplace.com/root_files_V3/publications.htm

我们的示例程序需要模板去匹配输入图像的每一个像素，而J.P.Lewis曾经讨论过一个更加渐近有效的方式来加速这种程序中各种相关操作类型的实现。他通过快速傅里叶变换计算分子，并用区域求和表来计算系数的分母：

Lewis, J.P. Fast template matching. Vision Interface 10, 1995, pp. 120~123。其扩充版名为“Fast Normalized Correlation”，可以在http://bit.ly/NJnZPI找到。

附录A CUDA专家手册库

正如第1章所述，本书附带的源代码遵循（简化的）2句版BSD开源许可证。源代码可在www.cudahandbook.com下载，开发人员可以在https://github.com/ArchaeaSoftware/cudahandbook找到对应的Git存储库。

本附录简要介绍CUDA专家手册库（chLib）。该库包含一套位于源代码项目的chLib/子目录的可移植头文件。chLib不可重用于软件产品。它提供了最小功能，使用尽可能少量的源代码来说明本书涵盖的概念。chLib可以移植到任何目标操作系统的CUDA环境，所以它往往必须支持这些操作系统的共有特性。

A. 1 计时操作

CUDA专家手册库包括一个使用QueryPerformanceCounter()（在Windows系统）和gettimeofday()（非Windows系统）的可移植计时库。一个示例用法如下：

float   
TimeNULLKernelLaunches(int cIterations  $= 1000000$  ） { chTimerTimestamp start,stop; chTimerGetTime( &start ); for（int  $\mathrm{i} = 0$  ；i<cIterations;  $\mathrm{i + + }$  ）{ NullKernel<<<1,1>>(); } CUDAThreadSynchronize(); chTimerGetTime( &stop); returnle6\*chTimerElapsedTime( &start,&stop)/ (float)cIterations;   
}

这个函数测定启动指定次数的内核所需的时间，并返回每次启动消耗的微秒值。chTimerTimestamp是一个高分辨率的时间戳。通常它是一个64位的计数器，它随时间单调递增，所以需要两个时间戳来计算时间间隔。

chTimerGetTime()函数取得当前时间的快照。

chTimerElapsedTime()函数返回两个时间戳间隔的秒数。这些计时器的分辨率是非常精细的（也许是微秒），所以chTimerElapsedTime()返回双精度浮点值。

#ifdef WIN32
#include <windows.h>
typedef LARGE_integer chTimerTimestamp;
#else
typedef struct timeval chTimerTimestamp;
#endif
void chTimerGetTime(chTimerTimestamp *p);
double chTimerElapsedTime(chTimerTimestamp *pStart, chTimerTimestamp *pEnd);
double chTimerBandwidth(chTimerTimestamp *pStart, chTimerTimestamp *pEnd, double cBytes);

在支持CUDA的GPU上对性能进行隔离测量时，我们可以使用CUDA事件，例如，测量一个内核的设备内存带宽。使用CUDA事件进行计时是

一把双刃剑：它们较少受伪系统级事件影响，如网络流量，但有时会导致过于乐观的计时结果。

A.2 线程操作

chLib包括最低限度的线程操作库，支持创建“工作”CPU线程，并拥有允许一个父线程把工作委派给工作线程的工具。线程操作是一个特别难以进行抽象的功能，因为不同的操作系统需要启用不同的工具。有些操作系统甚至有“线程池”，使线程容易被回收，因此应用程序不必让线程挂起以等待一个同步事件（当一些工作出现，即会发出信号）。

代码清单A-1给出了来自chLib/chThread.h的抽象线程操作。它包括一个processorCount()函数和一个C++类WorkerThread。前者返回可用的CPU核的数量（很多应用程序使用多线程以充分利用多个CPU核，例如第14章的多线程N-体实现，意在让每个核上运行一个线程），后者支持一些简单的线程操作。

创建和销毁

· delegateSynchronous(): 父线程指定一个工作线程要执行的函数指针，而且该函数直到工作线程完成才返回。

· delegateAsynchronous(): 父线程指定一个工作线程要异步执行的函数指针；workerThread::waitForAll必须被调用以便该父线程与它的子线程同步。
· 成员函数waitAll()进入等待，直到所有指定的工作线程完成其被委派的工作。

代码清单A-1 workerThread类

//  
// Return the number of execution cores on the platform.  
//  
unsigned int processorCount();  
//  
// workerThread class - includes a thread ID (specified to constructor)  
//  
class workerThread  
{  
public:  
    workerThread(int cpuThreadId = 0);  
    virtual *workerThread();  
    bool initialize();  
    // thread routine (platform specific)  
    static void threadRoutine(LPVOID);  
    // call this from your app thread to delegate to the worker.  
    // it will not return until your pointer-to-function has been called with the given parameter.  
    bool delegateSynchronous(void (*pfn)(void *), void *parameter);  
    // call this from your app thread to delegate to the worker asynchronously. Since it returns immediately, you must call waitAll later  
    bool delegateAsynchronous(void (*pfn)(void *), void *parameter);  
    static bool waitAll( workerThread *p, size_t N);  
};

A. 3 驱动程序API工具

chDrv.h包含为驱动程序API开发者提供的一些有用工具：chCUDADevice类，如代码清单A-2所示，它简化了设备和上下文的管理。其loadModuleFromFile方法简化了从一个.cubin或.ptx文件创建一个模块。

此外，chGetErrorString()函数传回一个对应于错误值的只读字符串。不仅针对驱动程序API的CResult类型实现了这一函数，chGetErrorString()的一个特化也包装了CUDA运行时的CUDAGetErrorString()函数。

代码清单A-2 chCUDADevicel类

class chCUDADevic   
{ public: chCUDADevic(); virtual -chCUDADevic(); CUresult Initialize( int ordinal, list<string>& moduleList, unsigned int Flags  $= 0$  unsigned int numOptions  $= 0$  CUjit_option \*options  $=$  NULL, void \*\*optionValues  $=$  NULL); CUresult loadModuleFromFileCUmodule \*pModule, string fileName, unsigned int numOptions  $= 0$  CUjit_option \*options  $=$  NULL, void \*\*optionValues  $=$  NULL); CUdevice device() const { return m_device; } CUcontext context() const { return m_context;} CUmodule module (string s） const { return (*m Modules.find(s)). second; } private: CUdevice m_device; CUcontext m_context; map<string,CUmodule> m Modules; };

A. 4 Shmoo工具

“shmoo图”是测试电路图性能的可视化曲线，它随两个输入的变化而变化（例如，电压和时钟频率）。在编写代码以确定不同内核的最佳线程块配置参数时，它有类似的用处：改变如线程块大小和循环展开因子等输入，观察输出性能。代码清单A-3显示了chShmooRange类和chShmooIterator类。前者封装了参数范围，后者使for循环可以方便地在给定范围内迭代。

代码清单A-3 chShmooRange类和chShmooIterator类

class chShmooRange {
public:
    chShmooRange() {}
    void Initialize(int value);
    bool Initialize(int min, int max, int step);
    bool isStatic() const{return m_min=m_max;} 
    friend class chShmooIterator;
    int min() const{return m_min;} 
    int max() const{return m_max;} 
private:
    bool m_initializer;
    int m_min, m_max, m_step;
}; 
class chShmooIterator
{
public:
    chShmooIterator(const chShmooRange& range);
    int operator *() const{return m_i;} 
    operator bool() const{return m_i<=m_max;} 
    void operator++(int) {m_i=m_step;};
private:
    int m_i;
    int m_max;
    int m_step;
};

命令行分析器还包括一个特化，用来基于命令行参数创建chShmooRange：在关键词之前追加“min”、“max”和“step”，而返回的结果为对应的范围。如果这三项均未提供，函数返回false。例如，concurrencyKernelKerne样例（在本书附带源代码的concurrency/子目录下）对流的数量和时钟周期数这两个参数在一定范围内进行测量。在命令行提取这些值的代码如下所示：

chShmooRange streamsRange;   
const int numStreams  $= 8$  .   
if（！chCommandLineGet(&streamsRange，"Streams"，argc，argv））{ streamsRange.Initialize(numStreams）;   
}   
chShmooRange cyclesRange;   
{ const int minCycles  $= 8$  · const int maxCycles  $= 512$  · const int stepCycles  $= 8$  · cyclesRange.Initialize(minCycles,maxCycles,stepCycles）； chCommandLineGet(&cyclesRange，"Cycles"，argc，argv）;

并且用户可以按照如下方式为应用程序指定参数。

concurrencyKernelKernel -- minStreams 2 --maxStreams 16 stepStreams 2

A. 5 命令行分析工具

chCommandLine.h给出了一个可移植的命令行分析库（只有100行的C++代码）。它包括模板函数chCommandLineGet()和chCommandLineGetBool()。前者传回给定类型的变量，而后者返回命令行中是否包括给定关键字的判断。

template<typename T> T
chCommandLineGet(T *p, const char *keyword, int argc, char *argv[])

正如上一节所述，chCommandLineGet()的一个特化将传回chShmooRange的一个实例。为了保证这一特化被编译，chShmoo.h必须在chCommandLine.h之前被包含。

A. 6 错误处理

chError.h实现了一组基于goto语句的错误处理宏，这一错误处理机制已在本书1.2.3小节提到过。这些宏的执行步骤如下：

·把返回值赋给一个名为status的变量；

  • 检查status是否成功，如果在调试模式下，错误报告输出到stderr;
    如果status包含一个错误，则goto到名为Error的语句标签。

CUDA运行时版本的错误处理如下:

ifdef DEBUG
define CUDART_CHECK( fn ) do {
(status) = (fn);
if ( CUDASuccess != (status)) {
fprintf( stderr, "CUDA Runtime Failure (line %d of file %s):\n\t" \)
"\\(s returned 0x\\){%}\)\n", \)
LINE , FILE , #fn, status,udaGetErrorString(status) ;
goto Error;
}
while (0);
else
#define CUDART_CHECK( fn ) do {
status = (fn);
if (CUDASuccess != (status)) {
goto Error;
}
while (0);
endif

do…while语句是一个C语言编程里的惯用组合，经常用在宏中。该语句让每次宏调用执行一条语句。如果变量status和标签Error：中的任何一个没有定义，这些宏将产生编译错误。

使用goto语句的一个隐含条件是所有的变量必须在代码块的顶部进行声明。否则，一些编译器会产生错误，因为goto语句可以绕过初始化。出现这种情况时，待初始化的变量必须移到第一个goto语句之前，或者移到一个基本代码块以使goto语句超出其范围。

代码清单A-4给出了一个遵循这一惯例的示例函数。返回值和中间资源被初始化为能够由清理代码进行处理的值。在这种情况下，所有由该函数分配的资源也由该函数释放，所以清理代码和错误处理代码是相同的。那些只会释放它们分配的一些资源的函数，必须在不同的代码块中实现成功和失败两种情形。

代码清单A-4 goto风格的错误处理示例

double   
TimedReduction( int \*answer, const int \*deviceIn,size_t N, int cBlocks,int cThreads, pfnReduction hostReduction   
1 double ret  $= 0.0$  int \*deviceAnswer  $= 0$  int \*partialSums  $= 0$  .. CUDAEvent_t start  $= 0$  .. CUDAEvent_t stop  $= 0$  .. CUDAError_t status; CUDART_CHECK(cudaMalloc(&deviceAnswer,sizeof(int))); CUDART_CHECK(cudaMalloc(&partialSums.cBlocks\*sizeof(int))
CUDART_CHECK(udaEventCreate(&start));  
CUDART_CHECK(udaEventCreate(&stop));  
CUDART_CHECK(udaThreadSynchronize());  
CUDART_CHECK(udaEventRecord(start, 0));  
hostReduction(
    deviceAnswer,
    partialSums,
    deviceIn,
    N,
    cBlocks,
    cThreads);
  
CUDART_CHECK(udaEventRecord(stop, 0));
CUDART_CHECK(udaMemcpy(
    answer,
    deviceAnswer,
    sizeof(int),
   udaMemcpyDeviceToHost));
ret = chEventBandwidth(start, stop, N*sizeof(int)) / powf(2.0f, 30.0f);  
// fall through to free resources before returning  
Error:
   udaFree(deviceAnswer);
   udaFree(partialSums);
   udaEventDestroy(start);
   udaEventDestroy(stop);
   return ret;

术语表

赋别名（aliasing）针对同一个内存指针，建立多个访问方式。例如，CUDA中一个映射锁页缓冲会在主机指针和设备指针之间互为别名；一个纹理引用绑定到设备内存，成为编程设备内存的别名。

AOS array of structures的缩写，参见结构体数组。

应用程序编程接口（API） application programming interface的缩写。

结构体数组（array of structures）一种内存结构，描述对象的元素在内存上是连续存放的（就好像在一个结构体上声明）。与数组结构相对。

异步的（asynchronous）函数调用在所请求的操作执行完毕之前返回。为保证得到正确的结果，基于异步操作的CUDA应用程序必须使用CUDA流或事件进行CPU/GPU同步。

计算密度（computational density）计算开销相对于外部内存传输开销的比例。

常量内存（constant memory）只读内存，当执行同一内存位置的读操作时，为广播机制作了优化。

中央处理器（CPU） central processing unit的缩写。当代计算机（无论是x86、x86-64，还是ARM）的大脑。

CUDA数组（CUDA array）结构细节对开发者未公开的一维、二维或三维数组。应用程序可用内存复制函数读写CUDA数组。CUDA内核可通过纹理读取CUDA数组，或者利用表面加载/存储内置函数对它进行读/写。

CUDA运行时（CUDART） CUDA runtime的缩写。带有语言集成特性的高级API。

设备驱动程序接口（DDI） device driver interface的缩写。设备驱动程序接口的例子包括XPDDM和WDDM。

请求式换页（demand paging）操作系统可把页面标记为“非驻留”，当应用程序试图访问一个非驻留页面时，硬件发出中断信号。操作系统可以使用该机制，依据一些启发式规则，对“一段时间”没有被访问的页面标记为非驻留，使这些页面的内容换回到磁盘以释放更多的物理内存供更活跃的虚拟页面使用。[1]如果应用程序再次访问某换出的页面，该页面会根据“请求”重新加载到内存（可能在与原来不同的物理页面上）。迄今为止，GPU实现了一个比较强大

的虚拟内存系统，可以让虚拟地址和物理地址分离，但未实现硬件上的请求式换页。

设备内存（device memory）适宜GPU访问的内存。CUDA数组、全局内存、常量内存和本地内存都是设备内存的不同形式。

直接内存访问（DMA） direct memory access的缩写。外设异步于、独立于CPU进行CPU内存的读或写。

驱动程序（driver） 使用操作系统的工具对外设的硬件功能予以暴露的软件。

驱动程序API（driver API）低级API，允许对CUDA工具进行全面访问。

动态指令总数（dynamic instruction count）一个程序实际执行的机器指令数目。与静态指令总数相对。

错误纠正码（ECC） error correction code的缩写。一些CUDA硬件保护GPU外部内存接口的方式，它预留12.5%的设备内存（可用内存每8位外配1位纠错码），并用它来检测和（有时）纠正内存事务的错误。可用nvidia-smi或英伟达管理库查询是发生了可纠正（1位）的错误，还是无法纠正（2位）的错误。

前端总线（FSB）在非NUMA系统配置下，芯片组与内存的接口。

全局内存（global memory） CUDA内核使用指针进行读写的设备内存。

图形处理器（GPU） graphics processing unit的缩写。

GPU时间（GPU time）通过CUDA事件测定的时间，与系统计数器时间不同。这一时间可以指导优化，但它无法提供系统整体性能的准确图景。与系统时钟时间相对。

高性能计算（HPC） high performance computing的缩写。

ILP instruction level parallelism的缩写。参见指令级并行。

指令级并行（instruction level parallelism）程序执行过程中，不同操作之间的细粒度并行。

内置函数（intrinsic function）直接对应于一个低级机器指令的函数。

即时编译（JIT） just-in-time compilation场合下的缩写。另请参阅在线编译。

内核模式（kernel mode）可以执行如编辑页表等敏感操作的特权执行模式。

内核转换（kernel think）从用户模式到内核模式的转换。该操作需要几千个时钟周期，所以操作系统上运行的驱动程序在需要内核转换操作以提交命令给硬件时，必须在执行内核转换之前在用户模式下排队等候硬件指令。

束内线程（lane）一个线程束内的线程。束内线程编号可以使用threadIdx.x&31计算。

内存管理单元（MMU）memory management unit的缩写。负责把虚拟地址转换到物理地址，并在所指定地址无效时发出错误信号的CPU或GPU硬件。

节点（node）NUMA系统上产生内存带宽的一个单元。在廉价的NUMA系统中，节点通常对应于物理CPU。

非一致内存访问（NUMA） nonuniform memory access的缩写。指的是如AMD皓龙或者英特尔Nehalem处理器的内存架构，其中的内存控制器集成到CPU以得到更低的延迟和更高的性能。

占用率（occupancy）一个SM上执行的线程束数量与理论最大值的比值。

在线编译（online compilation）在运行之际（而不是在开发人员生成应用程序时）进行编译。

事前允诺（opt-in）一个开发者必须在接口级别请求行为改变的API机制。例如，创建一个阻塞事件是一个“事前允诺”，因为开发者必须为创建事件API传入一个特殊的标志。由于现有应用程序依赖于旧行为，事前允诺是一种公开新功能而无须冒让步风险的途径。

事后拒绝（opt-out）一个禁止合法行为的API机制，例如，创建一个禁用计时的事件。

可换页内存（pageable memory）可以被VMM换出的内存。操作系统的设计者更喜欢可换页内存，因为它能使操作系统把页面“换出”到磁盘并把可用物理内存用于其他目的。

页面故障（page fault）当应用程序访问到被操作系统标记为非驻留的虚拟内存时发生的执行故障。如果访问有效，操作系统更新其数据结构（可能是把页面读入物理内存并更新指向该内存的物理地址）并恢复执行。如果访问无效，操作系统发送一个异常信号给应用程序。

页面锁定内存（page-locked memory）已经由操作系统在物理上予以分配并被标记为非分页的内存。通常这是为了支持硬件通过DMA访问内存。

PCIe PCI Express总线，CUDA用来在主机和设备内存之间交换数据。

锁页内存（pinned memory）参见页面锁定内存（page-locked memory）。

等步长内存分配（pitched memory allocation）一种内存分配方式，其中每行字节数需要另外指定，不是由每行元素数乘以每个元素大小得到。用于满足对齐约束，数组的每一行都必须对齐在同一位置。

等步长线性结构（pitch-linear layout）用于等步长内存分配的内存结构，由一个“基地址和每行字节数”（即“步长”）构成的“元组”指定。

断定（predicate）布尔型的1个位，值为“真”或“假”。在C语言中，整数可以通过评估它是否非零（真）或零（假）转换为一个断定。

进程（process）多任务操作系统的执行单位，它拥有自己的地址空间和资源生命周期的管理权（如文件句柄）。当进程退出时，与它相关联的所有资源均由操作系统“清理”。

页表项（PTE） page table entry的缩写。

并行线程执行（PTX） parallel thread execution的缩写，一种中间汇编语言和字节码，它会被作为驱动程序JIT执行的输入，编译成针对目标GPU的二进制代码。

SASS 支持CUDA的GPU的汇编语言级原生指令集。该缩写的确切全称已不可考，但着色器汇编（shader ASsembly）语言似乎是一个合理的猜测！

SBIOS 系统BIOS（基本输入/输出系统）。控制计算机系统最基本的I/O子系统（例如，是否启用可能不被某些操作系统支持的CPU或芯片组功能）的固件。SBIOS是较操作系统更低层级的。

共享内存（shared memory） CUDA内核可访问的快速板载GPU内存，用于保存临时结果。

单指令多数据（SIMD）single instruction multiple data的缩写，一个并行编程原语，涉及针对不同数据并行执行一个统一操作。在CUDA硬件中，流处理器簇以SIMD方式运行32个线程。x86硬件上的SSE指令也以SIMD方式处理通过宽寄存器操作包装的数据。

流处理器簇（SM） streaming multiprocessor的缩写，GPU的核心执行单元之一。一个GPU包含的SM数目范围可以从2到几十个。此外，一个GPU的指令集可以被指定一个版本号，例如，SM 2.0。

扩展的流处理器簇（SMX） SM在开普勒架构（SM 3.x）硬件上的扩展实现。

单指令多数据流扩展（SSE）在20世纪90年代末期添加到x86的指令集扩展，它可以使用单条指令执行4个单精度浮点运算。后来增加的特性包括支持整型SIMD运算，并把操作的字宽从128位拓展到256位。

静态指令总数（static instruction count）程序的机器指令数目；程序占用空间与静态指令总数成正比。与动态指令总数相对。

数组结构（SOA）使用一个数组来描述对象的每个元素的内存结构。与结构体数组（AOS）相对。

同步的（synchronous）该形容词用来描述那些直到所请求的操作完成才返回的函数。

特斯拉计算集群驱动程序（TCC） Tesla compute cluster driver的缩写。它是可以在Windows Vista和更高版本上运行的XPDDM类驱动程序。它没有WDDM的优势（Windows桌面管理器加速、图形互操作性、模拟分页），但可以无须执行内核转换就可提交命令到硬件，并实现64位统一地址空间。

Thrust 为提高CUDA生产效率开发的C++库，参考了STL的部分特性。

三字母缩写（TLA） 为three-letter acronym的缩略语。

线程本地存储（TLS）为thread local storage的缩略语。

最小精度单位（ulp） unit of last precision的缩写，意指浮点数尾数的最低有效数字。

用户模式（user mode）非特权执行模式，在此模式下，内存一般是分页的，而硬件资源只能通过API与操作系统的内核模式软件进行交互。

统一虚拟寻址（UVA）为unified virtual addressing的缩略语。（参见2.4.5小节）

虚拟内存管理器（VMM）为virtual memory manager的缩略语。操作系统中用于管理内存的部分：分配、锁页、管理缺页，等等。

系统时钟时间（wall clock time）依据系统时钟测定一组操作执行前和执行后的时间。系统时钟时间包括了系统的所有影响并提供整体性能的最精确测量。与GPU时间相对。

线程束（warp）流处理器簇的基本执行单位。对于前三代的CUDA硬件，线程束正好有32个线程，所以在一维线程块的线程束ID可以使用threadIdx.x>>5计算。另请参阅束内线程编号（lane）。

Windows显示驱动程序模型（WDDM）为Windows display driver model的缩略语。此驱动程序模型，是新加入Windows Vista的，它把多数显示驱动程序逻辑从内核模式移入用户模式。

Windows XP显示驱动程序模型（XPDDM）为Windows XP display driver model的缩略语。从架构上来说，这个驱动程序模型可以追溯到Windows NT 4.0（即1996年）。这个缩写是为了与“WDDM”对照而一起发明的。

[1] 请求式换页硬件可以用来实现很多其他功能, 例如, 写时复制, 映射文件z/o等。可以查阅操作系统方面的教科书获取更详细的内容。

前の章15.7_性能评价
15.8_延伸阅读 - CUDA专家手册 | OpenTech