CUDA学习笔记(三)——CUDA内存

时间:2024-03-10 22:00:28

转自:http://blog.sina.com.cn/s/blog_48b9e1f90100fm5f.html

结合lec07_intro_cuda.pptx学习

内存类型

       CGMA: Compute to Global Memory Access ratio

 

       Constant memory只允许device只读,比global memory 能够提供更快更多的并行数据访问路径给kernel。

       Register和local memory是线程私有的。Shared memory是同一个block中的线程共享的。

       Table 1显示了cuda声明变量的语法。Scope表示变量能够被访问的线程范围。包括thread:线程单独访问,每个thread都有一个变量,如果kernel声明一个scope为thread的变量y,在启动x个线程后,就会有x个版本的变量y。block:被block中的所有thread访问,grid:被grid中的所有线程访问。

       Lifetime是变量的生存期。注意:如果生存期为kernel,那么在kernel不同的启动之间,变量的值是不会被保存下来的。每次启动一次kernel都要对变量进行初始化。生存期为application的变量,必须在所有函数体外进行声明,变量可以在程序执行中保存下来并可以被所有kernel访问。

 

       非数组自动变量:除了在kernel和device函数中声明的数组外,其他所有自动变量都在寄存器中。这些变量称为scalar变量,scope是单独的线程。当一个kernel声明了一个自动变量,系统会为执行这个kernel函数的所有线程copy这个变量。线程终止后,所有变量也就不存在了。

       自动数组变量:存在global memory中,对它们的访问需要长延迟。他们的scope也是单独的线程。因此,对这种变量尽量避免使用。

(__device__)__shared__修饰的变量,表示CUDA中的共享变量。共享变量的scope是block,block中的所有线程都可以看到共享变量的同一个版本。Lifetime是kernel,kernel结束,共享变量内存也就不存在了。对共享内存的访问非常快而且是高度并行的。CUDA编程者通常用共享内存来保留一部分在kernel中用的多的全局内存数据。

       (__device__) __constant__修饰的变量表示常数变量constant variable。Constant variable必须在函数体外进行声明。Scope是grids, lifetime是整个应用程序的执行。Constant variable常用于为kernel function提供输入值,存储在global memory中但被cached。一个程序constant variable最大可以使65536个字节。

       __device__修饰的变量是global variable,存储在global memory中。对global memory的访问非常慢。由于global variable对所有kernel中的所有线程都是可见的。因此,global variable可以作为跨block的线程之间的协同方法。但是,如果不终止目前的kernel,无法保证线程之间数据的一致性。因此,global variable通常作为kernel function之间的信息传递。

       指针只能用于指向global memory的数据对象,不用于device memory。指针有两种典型用法:第一,如果一个对象由host function分配,指向此对象的指针被cudaMalloc()初始化并能够作为参数传递给kernel function。第二,在global memory中声明的变量的地址可以分配给一个指针变量。例如,

float * ptr=&GlobalVar。

减少全局内存通信的策略

       由于全局内存大而慢,共享内存小而快。常用的策略是把数据划分成片tile,每一片适合共享内存的使用。对这些tile的kernel计算可以独立的进行。

__global__ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width)

{

1. __shared__float Mds[TILE_WIDTH][TILE_WIDTH];

2. __shared__float Nds[TILE_WIDTH][TILE_WIDTH];

3. int bx = blockIdx.x; int by = blockIdx.y;

4. int tx = threadIdx.x; int ty = threadIdx.y;

// Identify the row and column of the Pd element to work on

5. int Row = by * TILE_WIDTH + ty;

6. int Col = bx * TILE_WIDTH + tx;

7. float Pvalue = 0;

// Loop over the Md and Nd tiles required to compute the Pd element

8. for (int m = 0; m < Width/TILE_WIDTH; ++m) {

// Coolaborative loading of Md and Nd tiles into shared memory

9.           Mds[ty][tx] = Md[Row][m*TILE_WIDTH + tx];

10.          Nds[ty][tx] = Nd[m*TILE_WIDTH + ty][Col];

11.         __Syncthreads();

12.         for (int k = 0; k < TILE_WIDTH; ++k)

13.                Pvalue += Mds[ty][k] * Nds[k][tx];

14. }

15. Pd[Row][Col] = Pvalue;

}

 

硬件限制:

       GeForce 8800GTX每个SM有8K个寄存器,整个处理器有128K个寄存器。一个SM最多有768个线程。如果要达到这个线程最大数,每个线程只能用8K/768=10个寄存器。如果每个线程要用11个寄存器,那么线程数就会减少。例如,如果一个block有256个线程,那么每个SM中只有1/3的线程同时存在。

       共享内存也会限制线程数目。在GeForce 8800 GTX中,每个SM有16K bytes大小的共享内存。而共享内存是block使用的。每个SM最多有8个block,所以,如果一个SM中有8个block,那么每个block最多能够使用2K字节的共享内存。以矩阵乘为例,若tile大小为16*16,那么,每个block需要16*16*4=1K字节存储Mds,需要1K字节存储Nds。因此,一个block需要2K字节的共享内存。根据共享内存16K 字节的限制,最多有8个block可以同时存在于一个SM中,这也是硬件限制上的最大block数目了。若tile size是32*32,那么每个block需要8K 字节共享内存,那么一个SM只能有2个block。

 

注意:不时关注CUDA主页,关注CUDA相关的最新研究

http://www.nvidia.com/object/cuda_home.html#