CUDA - Memory Model

CUDA ---- Memory Model

Memory

kernel性能高低是不能单纯的从warp的执行上来解释的。比如之前博文涉及到的,将block的维度设置为warp大小的一半会导致load efficiency降低,这个问题无法用warp的调度或者并行性来解释。根本原因是获取global memory的方式很差劲。

众所周知,memory的操作在讲求效率的语言中占有极重的地位。low-latency和high-bandwidth是高性能的理想情况。但是购买拥有大容量,高性能的memory是不现实的,或者不经济的。因此,我们就要尽量依靠软件层面来获取最优latency和bandwidth。CUDA将memory model unit分为device和host两个系统,充分暴露了其内存结构以供我们操作,给予用户充足的使用灵活性。

Benefits of a Memory Hierarchy

一般来说,程序获取资源是有规律的,也就是计算机体系结构经常提到的局部原则。其又分为时间局部性和空间局部性。 相信大家对计算机内存方面的知识都很熟悉了,这里就不多说了,只简单提下。

 

GPU和CPU的主存都是用DRAM实现,cache则是用lower-latency的SRAM来实现。GPU和CPU的存储结构基本一样。而且CUDA将memory结构更好的呈现给用户,从而能更灵活的控制程序行为。

CUDA Memory Model

对于程序员来说,memory可以分为下面两类:

  • Programmable:我们可以灵活操作的部分。
  • Non-programmable:不能操作,由一套自动机制来达到很好的性能。

在CPU的存储结构中,L1和L2 cache都是non-programmable的。对于CUDA来说,programmable的类型很丰富:

  • Registers
  • Shared memory
  • Local memory
  • Constant memory
  • Texture memory
  • Global memory

下图展示了memory的结构,他们各自都有不用的空间、生命期和cache。

 

其中constant和texture是只读的。最下面这三个global、constant和texture拥有相同的生命周期。

Registers

寄存器是GPU最快的memory,kernel中没有什么特殊声明的自动变量都是放在寄存器中的。当数组的索引是constant类型且在编译期能被确定的话,就是内置类型,数组也是放在寄存器中。

寄存器变量是每个线程私有的,一旦thread执行结束,寄存器变量就会失效。寄存器是稀有资源。在Fermi上,每个thread限制最多拥有63个register,Kepler则是255个。让自己的kernel使用较少的register就能够允许更多的block驻留在SM中,也就增加了Occupancy,提升了性能。

使用nvcc的-Xptxas -v,-abi=no(这里Xptxas表示这个是要传给ptx的参数,不是nvcc的,v是verbose,abi忘了,好像是application by interface)选项可以查看每个thread使用的寄存器数量,shared memory和constant memory的大小。如果kernel使用的register超过硬件限制,这部分会使用local memory来代替register,即所谓的register spilling,我们应该尽量避免这种情况。编译器有相应策略来最小化register的使用并且避免register spilling。我们也可以在代码中显式的加上额外的信息来帮助编译器做优化:

__global__ void
__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)
kernel(...) {
    // your kernel body
}

maxThreadsPerBlock指明每个block可以包含的最大thread数目。minBlocksPerMultiprocessor是可选的参数,指明必要的最少的block数目。

我们也可以使用-maxrregcount=32来指定kernel使用的register最大数目。如果使用了__launch_bounds__,则这里指定的32将失效。

Local Memory

有时候,如果register不够用了,那么就会使用local memory来代替这部分寄存器空间。除此外,下面几种情况,编译器可能会把变量放置在local memory:

  • 编译期无法决定确切值的本地数组。
  • 较大的结构体或者数组,也就是那些可能会消耗大量register的变量。
  • 任何超过寄存器限制的变量。

local memory这个名字是有歧义的:在local memory中的变量本质上跟global memory在同一块存储区。所以,local memory有很高的latency和较低的bandwidth。在CC2.0以上,GPU针对local memory会有L1(per-SM)和L2(per-device)两级cache。

Shared Memory

__shared__修饰符修饰的变量存放在shared memory。因为shared memory是on-chip的,他相比localMemory和global memory来说,拥有高的多bandwidth和低很多的latency。他的使用和CPU的L1cache非常类似,但是他是programmable的。

按惯例,像这类性能这么好的memory都是有限制的,shared memory是以block为单位分配的。我们必须非常小心的使用shared memory,否则会无意识的限制了active warp的数目。

不同于register,shared memory尽管在kernel里声明的,但是他的生命周期是伴随整个block,而不是单个thread。当该block执行完毕,他所拥有的资源就会被释放,重新分配给别的block。

shared memory是thread交流的基本方式。同一个block中的thread通过shared memory中的数据来相互合作。获取shared memory的数据前必须先用__syncthreads()同步。L1 cache和shared memory使用相同的64KB on-chip memory,我们也可以使用下面的API来动态配置二者:

cudaError_t cudaFuncSetCacheConfig(const void* func, enum cudaFuncCachecacheConfig);

func是分配策略,可以使用下面几种:

cudaFuncCachePreferNone: no preference (default)

cudaFuncCachePreferShared: prefer 48KB shared memory and 16KB L1 cache

cudaFuncCachePreferL1: prefer 48KB L1 cache and 16KB shared memory

cudaFuncCachePreferEqual: Prefer equal size of L1 cache and shared memory, both 32KB

Fermi仅支持前三种配置,Kepler支持全部。

Constant Memory

Constant Memory驻留在device Memory,并且使用专用的constant cache(per-SM)。该Memory的声明应该以__connstant__修饰。constant的范围是全局的,针对所有kernel,对于所有CC其大小都是64KB。在同一个编译单元,constant对所有kernel可见。

kernel只能从constant Memory读取数据,因此其初始化必须在host端使用下面的function调用:

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

这个function拷贝src指向的count个byte到symbol的地址,symbol指向的是在device中的global或者constant Memory。

当一个warp中所有thread都从同一个Memory地址读取数据时,constant Memory表现最好。例如,计算公式中的系数。如果所有的thread从不同的地址读取数据,并且只读一次,那么constant Memory就不是很好的选择,因为一次读constant Memory操作会广播给所有thread知道。

Texture Memory

texture Memory驻留在device Memory中,并且使用一个只读cache(per-SM)。texture Memory实际上也是global Memory在一块,但是他有自己专有的只读cache。这个cache在浮点运算很有用(具体还没弄懂)。texture Memory是针对2D空间局部性的优化策略,所以thread要获取2D数据就可以使用texture Memory来达到很高的性能,D3D编程中有两种重要的基本存储空间,其中一个就是texture。

Global Memory

global Memory是空间最大,latency最高,GPU最基础的memory。“global”指明了其生命周期。任意SM都可以在整个程序的生命期中获取其状态。global中的变量既可以是静态也可以是动态声明。可以使用__device__修饰符来限定其属性。global memory的分配就是之前频繁使用的cudaMalloc,释放使用cudaFree。global memory驻留在devicememory,可以通过32-byte、64-byte或者128-byte三种格式传输。这些memory transaction必须是对齐的,也就是说首地址必须是32、64或者128的倍数。优化memory transaction对于性能提升至关重要。当warp执行memory load/store时,需要的transaction数量依赖于下面两个因素:

  1. Distribution of memory address across the thread of that warp 就是前文的连续
  2. Alignment of memory address per transaction 对齐

一般来说,所需求的transaction越多,潜在的不必要数据传输就越多,从而导致throughput efficiency降低。

对于一个既定的warp memory请求,transaction的数量和throughput efficiency是由CC版本决定的。对于CC1.0和1.1来说,对于global memory的获取是非常严格的。而1.1以上,由于cache的存在,获取要轻松的多。

GPU Cache

跟CPU的cache一样,GPU cache也是non-programmable的。在GPU上包含以下几种cache,在前文都已经提到:

  • L1
  • L2
  • Read-only constant
  • Read-only texture

每个SM都有一个L1 cache,所有SM共享一个L2 cache。二者都是用来缓存local和global memory的,当然也包括register spilling的那部分。在Fermi GPus 和 Kepler K40或者之后的GPU,CUDA允许我们配置读操作的数据是否使用L1和L2或者只使用L2。

在CPU方面,memory的load/store都可以被cache。但是在GPU上,只有load操作会被cache,store则不会。

每个SM都有一个只读constant cache和texture cache来提升性能。

CUDA Variable Declaration Summary

下表是之前介绍的几种memory的声明总结:

 

 

Static Global Memory

下面的代码介绍了怎样静态的声明global variable(之前的博文其实都是global variable)。大致过程就是,先声明了一个float全局变量,在checkGlobal-Variable中,该值被打印出来,随后,其值便被改变。在main中,这个值使用cudaMemcpyToSymbol来初始化。最终当全局变量被改变后,将值拷贝回host。

#include <cuda_runtime.h>
#include <stdio.h>
__device__ float devData;
__global__ void checkGlobalVariable() {
    // display the original value
    printf("Device: the value of the global variable is %fn",devData);
    // alter the value
    devData +=2.0f;
}

int main(void) {
    // initialize the global variable
    float value = 3.14f;
    cudaMemcpyToSymbol(devData, &value, sizeof(float));
    printf("Host: copied %f to the global variablen", value);
    // invoke the kernel
    checkGlobalVariable <<<1, 1>>>();
    // copy the global variable back to the host
    cudaMemcpyFromSymbol(&value, devData, sizeof(float));
    printf("Host: the value changed by the kernel to %fn", value);
    cudaDeviceReset();
    return EXIT_SUCCESS;
}                                    

编译运行:

$ nvcc -arch=sm_20 globalVariable.cu -o globalVariable
$ ./globalVariable

输出:

Host: copied 3.140000 to the global variable
Device: the value of the global variable is 3.140000
Host: the value changed by the kernel to 5.140000

熟悉了CUDA的基本思想后,不难明白,尽管host和device的代码是写在同一个源文件,但是他们的执行却在完全不同的两个世界,host不能直接访问device变量,反之亦然。

我们可能会反驳说,用下面的代码就能获得device的全局变量:

cudaMemcpyToSymbol(devD6ata, &value, sizeof(float));

但是,我们应该还注意到下面的几点:

  • 该函数是CUDA的runtime API,使用的GPU实现。
  • devData在这儿只是个符号,不是device的变量地址。
  • 在kernel中,devData被用作变量。

而且,cudaMemcpy不能用&devData这种方式来传递变量,正如上面所说,devData只是个符号,取址这种操作本身就是错误的:

cudaMemcpy(&devData, &value, sizeof(float),cudaMemcpyHostToDevice);         // It’s wrong!!!

不管怎样,CUDA还是为我们提供了,利用devData这种符号来获取变量地址的方式:

cudaError_t cudaGetSymbolAddress(void** devPtr, const void* symbol);

获取地址之后,就可以使用cudaMemcpy了:

float *dptr = NULL;
cudaGetSymbolAddress((void**)&dptr, devData);
cudaMemcpy(dptr, &value, sizeof(float), cudaMemcpyHostToDevice);

我们只有一种方式能够直接获取GPU memory,即使用pinned memory,下文将详细介绍。

Memory Management

Will coming soon…

相关内容推荐