5.2 全局内存

全局内存是CUDA中的一个主要抽象，内核通过其读写设备内存[1]。由于设备内存直接附加在GPU上，由集成在GPU中的内存控制器读写，所以峰值带宽会非常高：在高端的CUDA卡上通常会超过100G/s。

设备内存可以被CUDA内核使用设备指针访问，下面是简单用于内存设置的内核的一个例子。

template<class T> global void
GPUmemset(int *base, int value, size_t N) {
    for (size_t i = BlockIdx.x*blockDim.x + threadIdx.x; i < N; i += gridDim.x*blockDim.x)
        base[i] = value;
}

设备指针base驻留在设备地址空间，独立于CUDA程序中主机代码使用的CPU地址空间。因此，CUDA程序中的主机代码可以在设备指针上执行指针算术运算，但它们不能解引用这些指针[2]。

这一内核写入整型值到由base和N给定的地址范围。在指定的线程块和网格参数来启动内核时，对blockIdx、blockDim和gridDim的引用使内核能够正确地操作。

5.2.1 指针

当使用CUDA运行时时，设备指针与主机指针类型均为void*。驱动程序API则使用一个整型值类型的typedef定义——CUdeviceptr，它与主机指针有相同宽度（即，32位操作系统上32位宽，64位同理），如下：

if defined(_x86_64) || defined(AMD64) || defined(_M_AMD64)  
typedef unsigned long long CUdeviceptr;  
#else  
typedef unsigned int CUdeviceptr;  
#endif

类型uintptr_t是在头文件<stdio.h>中定义的，在C++0x中被引入。开发者可使用它在主机指针（void*）和设备指针（CDeviceptr）之间方便的转换，如下：

CUDeviceptr devicePtr;  
void *p;  
p = (void *) (uintptr_t) devicePtr;  
devicePtr = (CUDeviceptr) (uintptr_t) p;

主机可在设备指针上进行指针的算术运算，作为参数传入一个内核或被内存复制调用，但是主机不能使用这些指针读或写设备内存。

驱动程序API中的32位和64位指针

因为驱动程序API中对指针的原始定义是32位的，在CUDA中添加对64位指针的支持需要定义CUdeviceptr，并且，所有使用CUdeviceptr

作为一个参数的驱动程序API函数，需要有所变化来改变支持位数。

[3] 举个例子，函数cuMemAlloc()，从

CResult CUDAAPI cuMemAlloc(CUdeviceptr *dptr, unsigned int bytesize);

改变为

Cuiresult CUDAAPI cuMemAlloc(CUdeviceptr *dptr, size_t bytesize);

为了同时支持旧的应用程序（函数cuMemAlloc()使用32位CUdeviceptr链接的）和新的应用程序，CUDA.h头文件包含了两个代码块，在开发者更新到新的API后，使用预处理技术不用改变函数名而改变绑定即可。

首先，一个代码块暗中改变函数名称来映射新的有不同语义的函数。

if defined(_CUDA_API_VERSION_INTERNAL) || _CUDA_API_VERSION >= 3020
    #define cuDeviceTotalMem cuDeviceTotalMem_v2
    ...
    #define cuTexRefGetAddress cuTexRefGetAddress_v2
    #endif /* _CUDA_API_VERSION Internal */ _CUDA_API_VERSION >= 3020 */

使用这一方式，客户端代码使用同样的旧函数名称，但是编译的代码会生成对新函数的引用，新函数名称后缀添加了_v2。

在之后的头文件中，旧函数依旧使用以前的定义方式。因此，开发者应使用最新的CUDA版本编译程序以得到最新的函数定义与语义。

CUDA.h使用相同的策略来处理函数语义因版本改变的问题，例如cuStreamDestroy()。

5.2.2 动态内存分配

大多数CUDA中的全局内存通过动态分配得到。使用CUDA运行时，函数

cudaeError_t CUDAAlloc(void **, size_t);  
cudaeError_t CUDAFree(void);

分别分配与释放全局内存。相关的驱动程序API函数是

CResult CUDAAPI cuMemAlloc(CUdeviceptr *dptr, size_t bytesize);  
CResult CUDAAPI cuMemFree(CUdeviceptr dptr);

分配全局内存成本非常高昂。CUDA驱动程序实现了一个CUDA小型内存请求的子分配器（suballocation），但是如果这个子分配器必须创建一个新的内存块，这需要调用操作系统的一个成本很高的内核模式驱动程序。如果这种情况发生，CUDA驱动程序必须与GPU同步，这可能会中断CPU/GPU的并发。因此，在性能要求很高的代码中避免分配或释放全局内存是一个较好的做法。

1. 等步长内存分配

合并限制以及为纹理操作和2D内存复制而设的对齐限制，激发了等步长内存分配的应用。这一想法是当创建一个2D数组时，指向数组的指针每移动到不同的行后必须保持相同的对齐特性。数组步长是数组中每一行的字节数 [4] 。等步长内存分配使用一个宽度（单位为字节）和高度，填补宽度到一个合适的硬件相关的宽度，传回基指针和步长给驱动程序。通过使用这些分配函数，把选择步长的任务委托给驱动程序，开发者可以让他们的代码适应未来增大对齐宽度的架构变化 [5] 。

CUDA应用程序经常需要满足硬件强制的对齐约束。不仅仅针对基地址，对内存复制宽度（字节）和绑定到纹理的线性内存也是同样的。因为对齐约束是由机器决定的，CUDA提供API使开发者为驱动程序选择适当的对齐。这些API使CUDA应用程序实现了独立于硬件的代码，并且可以在未来适应现在尚未出现的架构。

图5-1 步长与宽度

图5-1展示了一个在数组上执行的等步长内存分配，这个数组的宽度为352字节宽。在内存分配之前，宽度会填补到最近的64的倍数。对给定的数组宽度，除了行和列，一个数组元素的地址可以按如下代码计算：

inline T *  
getElement(T *base, size_t Pitch, int row, int col)  
{  
return (T *) ((char *) base + row * Pitch) + col;  
}

执行等步长内存分配的CUDA运行时函数如下：

template<class T> inline __host__udaError_tudaMallocPitch(T **devPtr, size_t *pitch, size_t widthInBytes, size_t height);

CUDA运行时同样包含了CUDAAlloc3D()函数，这个函数会使用CUDAPitchPtr和CUDAExtent结构体分配3D内存。

extern _host__ CUDAError_t CUDAAPI CUDAAlloc3D(struct CUDAPitchedPtr* pitchedDevPtr, struct CUDAExtent extent);

CUDAPitchedPtr结构接收分配的内存，定义如下：

structudaPitchedPtr {
    void *ptr;
    size_t pitch;
    size_t xsize;
    size_t ysize;
};

cudaPitchedPtr::ptr指定指针，cudaPitchedPtr::pitch指定了分配的步长（单位为字节），cudaPitchedPtr::xsize和cudaPitchedPtr::ysize分别指定了分配的逻辑宽度和高度。cudaExtent如下定义：

struct CUDAExtent {
    size_t width;
    size_t height;
    size_t depth;
};

CUDAExtent::width在数组与线性设备内存中分别做不同处理。在数组中，它指定数组成员的宽度；对线性设备内存，它会指定步长（用字节描述的宽度）。

驱动程序API使用的等步长内存分配函数如下：

CResult CUDAAPI cuMemAllocPitch(CuDeviceptr *dptr, size_t *pPitch, size_t WidthInBytes, size_t Height, unsigned int ElementSizeBytes);

ElementSizeBytes参数可能为4、8或16字节，而且它会引发分配的等步长内存分配分别填补到64、128或256字节边界。这些“对齐”在SM 1.0和SM 1.1硬件上的4、8、16字节内存事务合并处理中需要。那些不关心性能的应用程序可以指定为4。

函数CUDAAllocPitch()/cuMemAllocPitch()返回的步长是在调用者传入的一个以字节为单位的宽度基础上，补齐到满足全局加载/存储

操作的合并限制和纹理操作绑定API的对齐限制。内存分配数量是高*宽。

对3D数组，开发者可以在执行内存分配之前用深度乘上高度。这一考虑只适用于会通过全局加载/存储访问的数组，这是由于3D纹理操作不能被绑定到全局内存。

2. 内核内分配

费米架构硬件可以使用malloc()函数动态的分配全局内存。由于这可能需要GPU来中断CPU，可能会很慢。样例程序mallocSpeed.cu测量了内核中malloc()和free()函数的性能。

代码清单5-3展示了mallocSpeed.cu中的关键内核和计时子程序。一个重要的提示，CUDASetDeviceLimit()函数必须连同CUDALimitMallocHeapSize参数在内核可能调用的malloc()函数之前调用。mallocSpeed.cu中的调用需要一个完整的10亿字节（2 $^{30}$ ）。

CUDART_CHECK(udaDeviceSetLimit(cudaLimitMallocHeapSize, 1<<30));

当调用CUDADeviceSetLimit()时，所请求数量的内存被分配，并且不能再被其他目的使用。

代码清单5-3 MallocSpeed函数和内核

.global__void
AllocateBuffers( void **out, size_t N )
\{
	sizet i  $=$  blockIdx.x\*blockDim.x + threadIdx.x;
_out[i]  $=$  malloc(N);
\}
_global__void
FreeBuffers( void **in )
\{
	sizet i  $=$  blockIdx.x\*blockDim.x + threadIdx.x;
free(in[i]);
\}
 JudaError_t
MallocSpeed( double *msPerAlloc, double *msPerFree,
						// devicePointers, size_t N,
						cudaEvent_t evStart, cudaEvent_t evStop,
						int cBlocks, int cThreads )
\{
		int etAlloc, etFree;
		int cDNAError_t status;
		// CUDART_CHECK( cudaEventRecord( evStart ) );
		// AllocateBuffers<<cBlocks,cThreads>>>(devicePointers,N);
		// CUDART_CHECK( cudaEventRecord( evStop ) );
		// CUDART_CHECK( cudaThreadSynchronize(   ) );
		// CUDART_CHECK( cudaGetLastError(   ) );
		// CUDART_CHECK( cudaEventElapsedTime( &etAlloc, evStart, evStop ) );
	\}
	CUDART_CHECK( cudaEventRecord( evStart ) );
FreeBuffers<<cBlocks,cThreads>>>(devicePointers);
	CUDART_CHECK( cudaEventRecord( evStop ) );
	CUDART_CHECK( cudaThreadSynchronize(   ) );
	CUDART_CHECK( cudaGetLastError(   ) );
	// CUDART_CHECK( cudaEventElapsedTime( &etFree, evStart, evStop ) );
	*msPerAlloc = etAlloc / (double) (cBlocks*cThreads);
	*msPerFree  $=$  etFree / (double) (cBlocks*cThreads);
Error:
return status;
\}

代码清单5-4展示了在Amazon cgl.4xlarge类型实例上mallocAlloc.cu的运行输出。很明确的是分配器可以为小型内存分配进行优化：64字节的内存分配平均会使用0.39微秒来执行，而分配12KB内存会消耗 $3\sim 5$ 微秒。第一个结果（每次分配155微秒）在500线程块上分配1MB缓冲区，而每线程块使用一个线程。

代码清单5-4 样例mallocSpeed.cu的输出

Microseconds per alloc/free (1 thread per block):  
alloc free  
154.93 4.57  
Microseconds per alloc/free (32-512 threads per block, 12K allocations):  
32 64 128 256 512  
alloc free alloc free alloc free alloc free alloc free  
3.53 1.18 4.27 1.17 4.89 1.14 5.48 1.14 10.38 1.11  
Microseconds per alloc/free (32-512 threads per block, 64-byte allocations):  
32 64 128 256 512  
alloc free alloc free alloc free alloc free alloc free  
0.35 0.27 0.37 0.29 0.34 0.27 0.37 0.22 0.53 0.27

重要提示内核中使用malloc()函数分配的内存，必须使用函

数free()释放，这时在主机端调用CUDAFree()将不会成功。

5.2.3 查询全局内存数量

系统中全局内存的大小可以在CUDA初始化之前查询。

1. CUDA运行时

调用函数CUDAGetDeviceProperties(), 检查 CUDADeviceProp.totalGlobalMem:

size_t totalGlobalMem; /\*\*< Global memory on device in bytes \*/.

2. 驱动程序API

调用驱动程序API函数。

WDDM和可用内存 Windows Vista中引入的Windows显示驱动模

型，改变了显示驱动程序的内存管理。它使视频内存的数据块可以在主机内存中被换入和换出，以满足渲染所需。因此，

cuDeviceTotalMem()/CUDADeviceProp::totalGlobalMem函数所返回的内存数量并不能准确地反映出物理内存数量。

5.2.4 静态内存分配

应用程序可以静态地分配全局内存，通过使用__device__关键字标记在内存声明中进行标记即可。这一内存是由CUDA驱动程序在模块加载时分配的。

1. CUDA运行时

静态内存上的内存复制可以使用CUDAMemcpyToSymbol()和CUDAMemcpyFromSymbol()函数执行：

varaError_t CUDAMemcpyToSymbol(
    char *symbol,
    const void *src,
    size_t count,
    size_t offset = 0,
    enum CUDAMemcpyKind kind = CUDAMemcpyHostToDevice;
);  
cudaError_t CUDAMemcpyFromSymbol(
    void *dst,
    char *symbol,
    size_t count,
    size_t offset = 0,
    enum CUDAMemcpyKind kind = CUDAMemcpyDeviceToHost;

当调用函数cudaMemcpyToSymbol()或cudaMemcpyFromSymbol()时，不要将符号名称用引号括起来，即我们应该这样使用函数：

sudoMemcpyToSymbol(g_xOffset, poffsetx, Width*Height*sizeof(int));

而不是这样：

cudMemcpyToSymbol("g_xOffset",poffsetx,...);

这两种方式的调用都会工作，但是后者会为任何的符号名称编译（即使是没有定义的符号）。如果你想让编译器报告出非法符号错误，记住不要使用引号。

CUDA运行时应用程序可以通过调用函数CUDAGetSymbolAddress()查询关联到静态分配的内存上的指针。

cudaeError_t CUDAGetSymbolAddress(void **devPtr, char *symbol);

小心：传递给CUDA内核一个静态声明的设备内存分配的符号非常容易，但是这没有作用。你必须调用CUDAGetSymbolAddress()函数并使用返回的指针。

2. 驱动程序API

使用驱动程序API的开发者可以使用函数cuModuleGetGlobal()获取静态分配内存的指针。

注意cuModuleGetGlobal()函数同时传回基指针和对象大小。如果我们不需要大小，开发者可以在bytes参数中传入NULL。一旦得到这一指针，只需要传入CUdeviceptr，内存就可以被内存复制或CUDA内核调用访问。

5.2.5 内存初始化API

为了开发者方便，CUDA提供了1D和2D内存初始化函数。由于它们使用内核实现，所以即使在没有指定流参数的情况下，它们也是异步的。然而，对必须在流内依次执行内存初始化的应用程序，这里还是有*Async()形式的函数接受一个流参数的。

1. CUDA运行时

CUDA运行时只支持以字节为单位的内存初始化:

cudaError_t cudaMemset(void *devPtr, int value, size_t count);  
cudaError_t cudaMemset2D(void *devPtr, size_t pitch, int value, size_t width, size_t height);

参数pitch指定了内存初始化中每行的字节数。

2. 驱动程序API

驱动程序API支持不同大小的1D和2D内存初始化，在表5-2中我们予以总结。这些内存初始化函数接受一个复制目标指针、待赋的值和从基地址开始写入的值的数量。参数步长是每一行的字节数（非元素数）。

表5-2 内存初始化变种

CResult CUDAAPI cuMemsetD8(CUdeviceptr dstDevice, unsigned char uc, size_t N);  
CResult CUDAAPI cuMemsetD16(CUdeviceptr dstDevice, unsigned short us, size_t N);  
CResult CUDAAPI cuMemsetD32(CUdeviceptr dstDevice, unsigned int ui, size_t N);  
CResult CUDAAPI cuMemsetD2D8(CUdeviceptr dstDevice, size_t dstPitch, unsigned char uc, size_t Width, size_t Height);  
CResult CUDAAPI cuMemsetD2D16(CUdeviceptr dstDevice, size_t dstPitch, unsigned short us, size_t Width, size_t Height);  
CResult CUDAAPI cuMemsetD2D32(CUdeviceptr dstDevice, size_t dstPitch, unsigned int ui, size_t Width, size_t Height);

现在，在同一个应用程序中CUDA运行时和驱动程序API可以和平共存，CUDA运行时开发者可以按照需要使用这些函数。参数unsigned char, unsigned short, 和unsigned int只是指定了位的模式；为了使用其他类型数据填充全局内存区间，像float，使用volatile union强制转换float到unsigned int。

5.2.6 指针查询

CUDA跟踪所有的内存分配，并提供API使应用程序可以查询CUDA中的所有指针。函数库和插件可以在此基础之上使用不同的处理策略。

1. CUDA运行时

函数CUDAPointerGetAttributes()接受一个指针作为输入，并传回一个CUDAPointerAttributes结构体，此结构体包含关于指针的信息：

struct CUDAPointerAttributes {
    enum CUDAMemoryType memoryType;
    int device;
    void *devicePointer;
    void *hostPointer;
}

当UVA启用，指针在进程内是唯一的，所以不会有相同的输入指针地址空间。当UVA未启用，输入的指针被假定为在当前设备的地址空间（表5-3）。

表5-3CORD attribute 成员

2. 驱动程序API

使用函数cuMemGetAddressRange()，开发者可以查询给定设备指针驻留的地址空间区间。

CUIresult CUDAAPI cuMemGetAddressRange(CUdeviceptr *pbase, size_t *psize, CUdeviceptr dptr);

函数接受一个设备指针作为输入，返回包含这个设备指针的基地址和分配空间大小。

伴随CUDA 4.0中新加的UVA功能，开发者可以使用函数cuPointerGetAttribute()查询CUDA以获取更多的关于一个地址的信息。

CUresult CUDAAPI cuPointerGetAttribute(void *data, CUpointer\_ attribute attribute, CUdeviceptr ptr);

这一函数接受一个设备指针作为输入，并传回属性参数相关的信息，如表5-4所示。注意对统一寻址，使用

CU_POINTER_ATTRIBUTE_DEVICE_POINTER或

CU_POINTER_ATTRIBUTE_HOST_POINTER中的任意一个会传回与所传入相同的指针值。

表5-4 cuPointerAttribute使用方法

3. 内核查询

在SM 2.x（费米架构）及之后的硬件中，开发者可以查询给定的指针是否指向了全局空间。内置函数__isGlobal()，如下：

unsigned int __isGlobal(const void *p);

返回值为1时说明指针指向全局内存，返回0则没有指向。

5.2.7 点对点内存访问

在特定的情况下，SM 2.0和之后的硬件可以映射属于其他有相同能力的GPU中的内存。在以下诸情况可以做到这一点：

·UVA必须启用。

映射双方的GPU都必须是费米架构，并基于相同的芯片。
双方GPU必须在同一个I/O集线器上。

由于点对点映射本质上属于一个多GPU特性，所以我们在多GPU一章（第9.2节）提供了更多细节。

5.2.8 读写全局内存

CUDA内核可以使用标准C语言读写全局内存，例如指针间接寻址（操作符*、操作符->）与数组索引（操作符[]）。下面有一个带模板的简单内核，用于写入一个常量到内存。

template<class T> global __void GlobalWrites(T *out, T value, size_t N) {
    for (size_t i = BlockIdx.x*blockDim.x+threadIdx.x; i < N; i += blockDim.x*gridDim.x) {
        out[i] = value;
    }
}

这个内核在任何输入下都会运行正确：任意的元素大小，任意的线程块大小，任意的网格大小。这块代码着重于演示目的，而不是追求最大的性能。那些使用更多的寄存器与内层循环中操纵多个值的CUDA内核执行的较快，但对一些线程块和网格配置，性能会刚刚好。尤其是，假如基地址和线程块大小被正确指定，硬件会执行合并内存事务，最大化内存带宽。

5.2.9 合并限制

为了达到读写数据时的最佳性能，CUDA内核必须执行内存事务合并处理。任何不满足合并所需的全套标准的内存事务称为“未合并的”。未合并内存事务的性能惩罚由 $2 \sim 8$ 倍不等。在最近的硬件上，合并内存事务对性能的影响显著减少，见表5-5。

表5-5 非合并内存访问的带宽惩罚

事务在每一个线程束的基础上被合并。为了由束执行内存读写事务的合并，一些简化的标准必须被满足：

·字至少为32位。读写字节或16位字的事务总是非合并的。
·束上的线程访问的地址是连续并递增的（即，依线程ID偏移）。
·束的基地址（束中第一个线程访问的地址）需按表5-6对齐。

表5-6 合并的对齐标准

(1)8位和16位字内存访问是始终未合并的。

函数cuMemAllocPitch()的参数ElementSizeBytes旨在适应大小限制，它会指定应用程序想要访问的内存大小（单位字节），所以步长不仅会保证对分配中给定行的合并内存事务，也会保证其他行的事务合并。

本书中的大多数内核执行合并的内存事务，是建立在输入的地址恰好对齐的基础上。在全局内存事务的处理方式上，英伟达提供了更多分别针对不同架构的信息，详见下文。

1. SM 1.x（特斯拉架构）

如前所述，SM1.0和SM1.1硬件要求线程束中的每一个线程按顺序访问临近的内存地址。SM1.2和SM1.3硬件对合并限制已经不是特别严格。为了发出合并内存请求，分割32个线程构成的线程束为两个“半束”：0-15号线程与16-31号线程。为了满足在每个半束上的内存请求，硬件会执行以下算法：

1）找到具有最小线程号的活动线程，并查找包含线程请求地址的内存段。内存段的大小由字大小决定：1字节大小字请求会需要32字节的段；2字节请求需要64字节段；其他的请求会需要128字节段。
2）找到所有请求地址位于相同段的其他活动线程。
3）如果可能的话，减少段事务大小，到64字节或32字节。

4）完成事务并标记刚接受服务的那些线程为不活动的。
5）重复 $1 \sim 4$ 步，直到所有的半束中的线程请求全部满足。

尽管同SM1.0和SM1.1限制相比，这些需求变得更加灵活，但是高效的合并操作始终需要大量的局部性特性。在实践中，这一弱化的合并限制意味着束中的线程可以在必要时使用更少的内存段置换输入。

2. SM 2.x（费米架构）

SM2.x和之后的硬件包含一级和二级缓存。二级缓存为整个芯片服务；一级缓存是在一个SM内可见的，可以设置为16KB或48KB大小。缓存行为128字节宽，对应设备内存上的128字节对齐段。当访问的是同时缓存于一级与二级缓存的内存时，使用128字节内存事务处理，但是当内存访问的是只在二级缓存中被缓存的对象，该访问只使用32字节内存事务处理。使用二级缓存可以因此减少过度读取，例如，分散模式的内存访问时。

硬件可以指定每一条指令访问的全局内存的可缓存性。默认的，编译器发出同时使用一级和二级缓存的内存访问指令（-Xptexas-dlcm=ca）。这可以通过指定-Xptexas-dlcm=cg改变为只使用二级缓存。不使用一级缓存，而只使用二级缓存的内存访问只采用32位内存事务处理，这会提升执行分散内存访问的应用程序的缓存利用率。

使用声明为volatile的指针读取内存会使所有的缓存被丢弃，而且为数据重新缓存。这一用法在轮询主机内存位置的情况下非常有用。表5-7总结了如何将一个线程束的内存请求分解为128字节缓存行请求。

表5-7 SM 2.x缓存行请求

注意在SM2.x和更高的架构上，束内线程可以以任何顺序访问任何字，包括冗余字。

3. SM 3.x（开普勒）

SM 3.x的二级缓存架构与SM 2.x中的架构相同。SM 3.x不在一级缓存中缓存全局内存。在SM 3.5中，全局内存可以通过纹理缓存访问（每SM 48KB大小），使用const restricted指针访问；或者使用sm_35_intrinsics中的内置函数__ldg()。当纹理操作直接来自设备内存，不去访问可能被其他途径并发访问的内存十分重要，因为二级缓存的存在，纹理缓存不会始终保持与二级缓存的一致性。

5.2.10 验证实验：内存峰值带宽

随同本书的源代码包含了一些验证程序，旨在以操作数大小、循环展开因素和线程块大小的最好组合最大化给定GPU的带宽。重写前文的GlobalWrites代码，变为一个接受附加参数n（内层循环中执行写操作的数量）的模板，见代码清单5-5。

代码清单5-5 内核GlobalWrites
代码清单5-6 函数ReportRow()

template<class T, const int n> global void GlobalWrites(T *out, T value, size_t N) {
    size_t i;
    for (i = n*blockIdx.x*blockDim.x+threadIdx.x; i < N-n*blockDim.x*gridDim.x; i += n*blockDim.x*gridDim.x) {
        for (int j = 0; j < n; j++) {
            size_t index = i+j*blockDim.x;
            out[index] = value;
        }
    }
    // to avoid the (index<N) conditional in the inner loop,
    // we left off some work at the end
    for (int j = 0; j < n; j++) {
        size_t index = i+j*blockDim.x;
        if (index<N) out[index] = value;
    }
}

代码清单5-6中给出的函数ReportRow()，通过调用模板函数BandwidthWrites（未给出）写入输出的其中一行，给出了对给定类型、网格和线程块大小的带宽情况。

template<class T, const int n, bool boffset> double ReportRow{ size_t N, size_t threadStart, size_t threadStop, size_t cBlocks } { int maxThreads = 0; double maxBW = 0.0; printf("%d\t", n); for (int cThreads = threadStart; cThreads <= threadStop; cThreads *= 2) { double bw; bw = BandwidthWrites<T,n,bOffset>(N, cBlocks, cThreads); if (bw > maxBW) { maxBW = bw; maxThreads = cThreads; } printf("%.2f\t", bw); } printf("%.2f\t%d\n", maxBW, maxThreads); return maxBW; }

参数threadStart和threadStop通常为32和512，32为线程束大小并且是线程块中可调度到机器的最小线程数。模板参数bOffset指定BandwidthWrites是否需要在基指针上偏移，这是由于所有的内存事务会被合并。如果程序在调用时包含命令行选项--uncoalesced，程序会在偏移指针上执行带宽测量。

注意，受制于sizeof(T)，当内核的n高于一定值时，会出现一个明显的性能退化。因为内层循环的临时变量数增长太多，将会出现溢出寄存器的情况。

表5-8中总结的5个应用程序实现了这一策略，它们测量在不同操作数大小（8位、16位、32位、64位和128位）、不同线程块大小（32、64、128、256、和512）和不同展开循环（ $1\sim 16$ ）下的内存传输带宽。CUDA硬件不一定会对所有的这些参数作出反应。例如，许多

参数设置使GK104带宽通过纹理操作达到140GB/s，但是只有操作数在32位以上时这才会生效。无论如何，对于指定的工作负载和硬件，验证程序突出了起重要作用的那个参数。同样，对于小的操作数大小，验证程序突出了循环展开是怎样帮助提升性能的（不是所有的应用程序都可以被重构来读取更大的操作数）。

表5-8 内存带宽测试

代码清单5-7给出了运行在GeForce GTX 680 GPU上的 globalRead.cu的输出示例。输出按照操作数大小进行了分组，从单字节，到2、4、8、16字节。每一组最左侧的列给出了循环展开数。32到512线程数的线程块的传输带宽在每列中依次给出，maxBW和 maxThreads列分别给出了最大带宽和对应的线程块大小。

代码清单5-7 样例globalRead.cu的输出

Running globalRead.cu microbenchmark on GeForce GTX 680

Using coalesced memory transactions

Operand size: 1 byte

Input size: 16M operands

Block Size

Operand size: 2 bytes

Input size: 16M operands

Block Size

Operand size: 4 bytes

Input size: 16M operands

Block Size

GeForce GTX 680显卡最高传输速度可达140GB/s，代码清单5-7让我们更加清楚一点，当在SM3.0上读取8位或16位字时，全局加载不应该是我们选择程序执行的方式。字节最高传输速度60GB/s，16位字最高可达101GB/s [6]。对32位操作数，为了达到最大带宽，我们需要一次2倍的循环展开和至少每线程块256线程。

这一验证程序可以帮助开发者优化他们的带宽受限型应用程序。选择一个与你的应用程序最相似的内存访问模式，在目标GPU上运行验证程序，或者，可能的话，调整验证程序以与真实的工作负载相匹配，以获得最佳的参数。

5.2.11 原子操作

SM 1.x开始支持原子操作，但是这项操作离谱的慢。全局内存中的原子操作在SM 2.x（费米架构）硬件中被提升，并且在SM 3.x（开普勒架构）中被大幅度全面提升。

大多数的原子操作，例如atomicAdd()，使代码可以使用自主导引方式替换归约（归约通常需要共享内存和同步）。SM 3.x硬件发布之前，这种编程方式会招致巨大的性能退化，因为开普勒之前的架构在处理竞争内存位置的问题上不够高效（即：许多的GPU线程同时在同一内存位置做原子操作）。

注意由于原子操作由GPU内存控制器实现，它们只在本地设备内存位置工作，在本书撰写时，尝试在非本地GPU或主机内存上执行原子操作是不可行的。

原子操作与同步

除了自主引导，原子操作同样可以使用在线程块之间的同步。CUDA硬件支持同步的主要抽象：“对比和交换”（或CAS，compare and swap）。在CUDA上，对比和交换（也可为compare and exchange——即，x86中的CMPXCHG指令）被定义如下：

int atomicCAS(int *address, int expected, int value);

函数atomicCAS()对无符号和64位变量同样有效。

这一函数从address读取的值放入old，计算（old==expected?value: old），结果传回address并返回old值。换句话说，内存地址保持不变，除非它与调用者指定的期望值相等，这种情况下，内存地址会更新为value。

可以基于CAS建立一个称为“自旋锁”的简单临界区（critical section），如下所示：

void enterSpinlock(int \*address) while atomicCAS(address,0,1);

假设自旋锁的值被初始化为0，当函数atomicCAS()执行，while循环会不断迭代直到自旋锁的值为0。这执行之后，*address自动变为1（函数atomicCAS()的第三个参数），并且其他的尝试得到这一临界区自旋锁的线程都必须等到这一值再次变为0。

拥有自旋锁的线程可以放弃这个自旋锁，通过自动换回0即可：

voidleaveSpinlock(int \*address)   
{ atomicExch(m_p,0);

在CPU上，对比和交换指令用于实现所有形式的同步。操作系统使用这些指令（一些时候与内核级线程上下文转换代码联合使用）来实现高级的同步原语。CAS同样使用在直接实现“无锁”队列和其他的数据结构中。

CUDA执行模型，在使用全局内存的原子操作进行同步时，施加了限制。不同于CPU线程，一次内核启动内的一些CUDA线程可能要等到同一内核中其他线程退出后才开始执行。在CUDA硬件上，每一个SM可切换一定数量线程块的上下文，所以对多于

MaxThreadBlocksPerSM*NumSMs数量线程块的内核启动，都需要第一批线程块退出，才能让更多的线程块开始执行。因此，开发者不假定给定内核中的所有线程都是活动的是十分重要的。

除此之外，上面给出的函数enter-spinlock()在使用块内同步时很容易导致死锁[7]，说明它不适合用在这种情况下。实际上，CUDA的硬件支持很多更好的块内线程间通信和同步的方式（可以分别基于共享内存和函数__syncthreads()）。

代码清单5-8给出了CUDASpinlock类的实现，使用了上述的算法，并且也受上述的限制。

代码清单5-8 CUDASpinlock类

class CUDASpinlock{ public: CUDASpinlock(int \*p); void acquire(); void release(); private: int \*m_p; }; inline _device_ CUDASpinlock::CUDASpinlock( int \*p) { m_p = p; } inline _device_void CUDASpinlock::acquire() { while (atomicCAS(m_p,0,1)); } inline _device_void CUDASpinlock::release(); atomicExch(m_p,0);

样例spinlockReduction.cu给出了CUDASpinlock的使用。程序要计算一个double数组的和，先在每一个线程块中基于共享内存执行一次归约，之后使用自旋锁在计算总和时进行同步。代码清单5-9给出了这个样例的SumDouble函数。注意累加部分和的操作是怎样只由每个线程块的0号线程执行的。

代码清单5-9 函数SumDouble

__global__void SumDoubles( double \*pSum, int \*spinlock, const double \*in, size_t N, int \*acquireCount)   
{ SharedMemory<double> shared; CUDASpinlock globalSpinlock( spinlock); for (size_t i  $=$  blockIdx.x\*blockDim.x+threadIdx.x; i  $<  \mathbb{N}$  . i  $^{+ = }$  blockDim.x\*gridDim.x){ shared[threadIdx.x]  $=$  in[i]; _syncthreads(); double blockSum  $=$  Reduce_block<double,double>(   ); _syncthreads(); if (threadIdx.x  $= = 0$  ) { globalSpinlock.acquire(); \*pSum  $+ =$  blockSum; _threadfence(); globalSpinlock.release(); }   
}

5.2.12 全局内存的纹理操作

对于不能方便地满足合并访问限制的应用程序，纹理映射硬件是一个令人满意的选择。硬件支持来自全局内存的纹理操作（通过CUDABindTexture() / cuTexRefSetAddress() 函数），这不会拥有比合并全局读写更快的峰值执行性能，但是对不太规则的访问会有更高的性能。纹理缓存资源同样同其他缓存资源独立。在包含TEX指令的内核调用前 [8]，驱动程序会有一个软件一致性机制使纹理缓存无效。查看第10章获取更多细节。

SM 3.x硬件添加了通过纹理缓存层读取全局内存的能力，这不需要设置和绑定纹理引用。这一功能可使用标准C++语言构建：关键字

const restrict。另一方法是使用在sm_35_intrinsics.h头文件中定义的__ldg()内置函数。

5.2.13 ECC（纠错码）

SM 2.x以及之后的服务器级特斯拉系列GPU拥有运行中纠错功能。以少量的内存（一些内存被使用在记录冗余信息）和稍低的带宽作为交换，启用ECC的GPU便可以自动地纠正单个比特位（single-bit）错误并且可以报告双位（double-bit）错误。

ECC有以下特点：

·减少大约 $12.5\%$ 的可用内存量。例如，在亚马逊EC2中的cg1.4xlarge实例上，内存数量从3071MB减少到了2678MB。
·它使上下文同步的成本变得更高。
·非合并内存处理的成本同样更高。

ECC可以使用nvidia-smi命令行工具（小节4.4）或NVML（英伟达管理库）启用或禁用。

当发现一个无法纠正的ECC错误，同步的错误报告机制会返回CUDAErrorECCUncorrectable（CUDA运行时）和CUDA_ERROR_ECC_UNCORRECTABLE（驱动程序API）。

[1] 最让开发者困惑的是，CUDA使用“设备指针”术语指代驻留在全局内存上的指针（CUDA内核可访问的设备内存）。
[2] 映射锁页指针是这一规则的一个例外。它们驻留在系统内存中，但是可以被GPU访问，在非UVA系统上，指向这一内存的主机指针和设备指针并不相同：应用程序必须调用函数cuMemHostGetDevicePointer()或cudaHostGetDevicePointer()映射主机指针到相关联的设备指针上。但是当UVA有效时，这两个指针就是相同的。
[3]出于兼容性考虑，旧函数被保留。
[4] 对齐2D内存复制的想法要比CUDA出现早的多。图形API，像苹果的QuickDraw和微软的 DirectX分别暴露了“行字节”和“间距”。同时，补齐使用移位代替乘法，简化了寻址计算，甚至使用两次移位与一次与两个2的幂的加法来简化乘法，例如640（512+128）。但在今天，整数加法的速度已经很快，等宽内存分配有着其他的动机，例如避免与缓存交互带来的性能消极影响。
[5]这不是一个意想不到的趋势，在特斯拉架构上，费米已经扩展了对不同的对齐需求的支持。
[6] 纹理操作表现得更好，读者可以运行globalReadTex.cu证实一下。
[7] 推荐用法是让每个线程块中某个线程尝试去获取自旋锁。否则，分支的代码执行会趋向于死锁。
[8] TEX是一个SASS微码指令助记符，执行纹理读取。

5.2_全局内存

5.2 全局内存

5.2.1 指针

驱动程序API中的32位和64位指针

5.2.2 动态内存分配

1. 等步长内存分配

2. 内核内分配

5.2.3 查询全局内存数量

1. CUDA运行时

2. 驱动程序API

WDDM和可用内存 Windows Vista中引入的Windows显示驱动模

5.2.4 静态内存分配

1. CUDA运行时

2. 驱动程序API

5.2.5 内存初始化API

1. CUDA运行时

2. 驱动程序API

5.2.6 指针查询

1. CUDA运行时

2. 驱动程序API

表5-4 cuPointerAttribute使用方法

3. 内核查询

5.2.7 点对点内存访问

5.2.8 读写全局内存

5.2.9 合并限制

1. SM 1.x（特斯拉架构）

2. SM 2.x（费米架构）

3. SM 3.x（开普勒）

5.2.10 验证实验：内存峰值带宽

5.2.11 原子操作

原子操作与同步

5.2.12 全局内存的纹理操作

5.2.13 ECC（纠错码）