【CUDA】学习记录（7）- Global Memory

作者: 不会code的程序猿 | 来源:发表于2017-05-05 16:56 被阅读251次

Professional CUDA C Programing
代码下载：http://www.wrox.com/WileyCDA/WroxTitle/Professional-CUDA-C-Programming.productCd-1118739329,descCd-DOWNLOAD.html

Memory

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

Benefits of a Memory Hierarchy

一般来说，应用程序不会在任何时间点访问任意数据或运行任意代码。程序获取资源是有规律的，也就是计算机体系结构经常提到的局部原则：时间局部性和空间局部性。
时间局部性：如果在某时刻访问了某数据，很小可能在段时间内还会访问该数据。
空间局部性：如果某时刻访问了某数据，则下一时刻很可能访问与之相邻的数据。
总体：如果速度越快，容量越小。

Screenshot from 2017-05-05 16:03:13.png

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

CUDA Memory Model

对于编程人员来讲，memory分为两类：
➤ Programmable: 我们可以灵活操作的部分。
➤ Non-programmable: 不能控制的部分。
对CPU而言，L1和L2缓存对我们而言是non-programmable memory.
CUDA将memory完全暴露给了用户：
➤ Registers
➤ Shared memory
➤ Local memory
➤ Constant memory
➤ Texture memory
➤ Global memory

Screenshot from 2017-05-05 16:11:45.png

每个thread有自己独立的registers和local memory，每个block中的所有threads共享share memory，所有的线程都可以访问global memory，其中constant和texture是只读内存。

Registers

寄存器是GPU最快的memory，kernel中没有什么特殊声明的自动变量都是放在寄存器中的。当数组的索引是constant类型且在编译期能被确定的话，就是内置类型，数组也是放在寄存器中。
寄存器是每个thread的私有变量，一旦thread执行结束，寄存器变量就会失效。寄存器也是稀缺资源，在Fermi每个thread最多63个registers, Kepler(255).在每个kernel中使用较少的寄存器，可以使更多的block 驻留在SM上，实现更多并发的blocks，进而提高occupy和性能。
如果kernel使用的register超过硬件限制，这部分会使用local memory来代替register，即所谓的register spilling，我们应该尽量避免这种情况。编译器有相应策略来最小化register的使用并且避免register spilling。
-Xptxas -v,-abi=no选项可以查看每个thread使用的寄存器数量，shared memory和constant memory的大小。

__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：
➤ 编译期间无法确定值的本地数组
➤消耗太多寄存器的较大的结构体或数组
➤ 任何超过寄存器限制的变量
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为单位分配的，如果每个block占用的share memory过多，那么每个SM上驻留的blocks就少，active warp的数目也会减少。
不同于register，shared memory尽管在kernel里声明的，但是他的生命周期是伴随整个block，而不是单个thread。当该block执行完毕，所拥有的资源就会被释放，重新分配给别的block。
shared memory是同一个block 中thread交流的基本方式。同一个block中的thread通过shared memory中的数据来相互合作。获取shared memory的数据前必须先用__syncthreads()同步。L1 cache和shared memory使用相同的64KB on-chip memory，我们也可以使用API来动态配置二者的大小。

Constant memory

Constant Memory驻留在device Memory，并且使用专用的constant cache（per-SM）。该Memory的声明应该以__connstant__修饰。constant的范围是全局的，针对所有kernel，对于所有GPU其大小都是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实际上也是global Memory在一块，但是他有自己专有的只读cache。这个cache在浮点运算很有用，texture Memory是针对2D空间局部性的优化策略，所以thread要获取2D数据就可以使用texture Memory来达到很高的性能。

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数量依赖于下面两个因素：
➤ Distribution of memory addresses across the threads of that warp.（没有读明白？？？）
➤ Alignment of memory addresses per transaction.地址对齐
一般来说，所需求的transaction越多，潜在的不必要数据传输就越多，从而导致throughput efficiency降低。
对于一个既定的warp memory请求，transaction的数量和throughput efficiency是由CC版本决定的。对于CC1.0和1.1来说，对于global memory的获取是非常严格的。而1.1以上，由于cache的存在，获取要轻松的多。

GPU Caches

跟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

Screenshot from 2017-05-05 18:55:42.png

Static Global Memory

下面绍了怎样声明一个静态的global variable。先声明了一个float全局变量，在checkGlobal-Variable中，该值被打印出来，随后，其值便被改变。在main中，这个值使用最终当全局变量被改变后，将值拷贝回host。使用cudaMemcpyToSymbol拷贝数据（这个应该是常量内存的拷贝函数，适合于1个warp中所有threads读取同一个地址的数据，但是是只读的，但是在这个例子中变量明明被改变了？？？）。

#include <cuda_runtime.h>
#include <stdio.h>

/*
 * An example of using a statically declared global variable (devData) to store
 * a floating-point value on the device.
 */

__device__ float devData;

__global__ void checkGlobalVariable()
{
    // display the original value
    printf("Device: the value of the global variable is %f\n", devData);

    // alter the value
    devData += 2.0f;
}
int main(void)
{
    // initialize the global variable
    float value = 3.14f;
    CHECK(cudaMemcpyToSymbol(devData, &value, sizeof(float)));
    printf("Host:   copied %f to the global variable\n", value);

    // invoke the kernel
    checkGlobalVariable<<<1, 1>>>();

    // copy the global variable back to the host
    CHECK(cudaMemcpyFromSymbol(&value, devData, sizeof(float)));
    printf("Host:   the value changed by the kernel to %f\n", value);

    CHECK(cudaDeviceReset());
    return EXIT_SUCCESS;
}

实验结果:ge740m(在我的博客中有时候显卡是tesla k80，有时候是gt750，由于服务器经常掉所以有些结果是自己的电脑测试结果，hym-gt740m,ccit-Tesla K80).

hym@hym-ThinkPad-Edge-E440:~/CodeSamples/chapter04$ nvcc globalVariable.cu  -o globalVariable
nvcc warning : The 'compute_20', 'sm_20', and 'sm_21' architectures are deprecated, and may be removed in a future release (Use -Wno-deprecated-gpu-targets to suppress warning).
hym@hym-ThinkPad-Edge-E440:~/CodeSamples/chapter04$ ./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
hym@hym-ThinkPad-Edge-E440:~/CodeSamples/chapter04$

尽管host和device的代码保存在同一个文件中，但是他们是完全不同的，host的代码不能访问device的变量，device的代码也不能访问host的变量。
我们可能会反驳说，用下面的代码就能在host端获得device的全局变量：

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

➤1.cudaMemcpyToSymbol 是CUDA的runtime API，是GPU的实现。
➤ 在该处devData表示的是一个符号，而不是表示该变量的地址。
➤ 在kernel函数中，devData用来表示global memory中的一个变量。
错误的方式：
cudaMemcpy(&devData, &value, sizeof(float),cudaMemcpyHostToDevice);
cudaMemcpy不能用&devData这种方式来传递变量,正如上面所说，devData只是个符号，取址这种操作本身就是错误的：
正确的解决方法：

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

CUDA pinned memory：host和device端的代码都能够访问。
注意：在文件范围内可见，不代表可以访问。

__host__cudaError_t cudaMemcpyToSymbol (const void*symbol, const void *src, size_t count, size_t offset,cudaMemcpyKind kind)

symbol:Device symbol address
src:Source memory address
count:Size in bytes to copy
offset:Offset from start of symbol in bytes
kind:Type of transfer

__host__cudaError_t cudaMemcpyFromSymbol (void*dst, const void *symbol, size_t count, size_t offset, cudaMemcpyKind kind)

Example:使用全局的GPU global数组

// includes, system
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>

// includes CUDA
#include <cuda_runtime.h>
#include<device_launch_parameters.h>
// includes, project
#include <helper_cuda.h>
#include <helper_functions.h> // helper functions for SDK examples
#define N (3)
__device__ int d_attrSelect[3];
__global__ void checkGlobalVariable()
{
    for (int i = 0; i < 3; i++)
    {
    
        d_attrSelect[i] += i;
        printf("Device: the value of the global variable is %d\n", d_attrSelect[i]);
    }
}
int main()
{
    int h_attrSelect[N] = { 1, 1, 1 };
    for (int i = 0; i <3; i++)
    {
        printf("Host: the value of the global variable is %d\n", h_attrSelect[i]);
    }
    size_t sz = (size_t)(N) * sizeof(int);
    //checkCudaErrors(cudaMemcpyToSymbol(d_attrSelect, &h_attrSelect[0], sz,size_t(0),cudaMemcpyHostToDevice));
    checkCudaErrors(cudaMemcpyToSymbol(d_attrSelect, h_attrSelect, sz, size_t(0), cudaMemcpyHostToDevice));
    int blocksize = 1;
    int gridsize = 1;
    dim3 block(blocksize);
    dim3 grid(gridsize);
    checkGlobalVariable <<<grid, block >>>();
    cudaDeviceSynchronize();
    //checkCudaErrors(cudaMemcpy(&h_attrSelect[0], d_attrSelect, sz, cudaMemcpyDeviceToHost));
    //checkCudaErrors(cudaMemcpyFromSymbol(&h_attrSelect[0], d_attrSelect, sz));
    checkCudaErrors(cudaMemcpyFromSymbol(h_attrSelect, d_attrSelect, sz,size_t(0),cudaMemcpyDeviceToHost));
    for (int i = 0; i < 3; i++)
    {
        printf("Host: the value of the global variable is %d\n", h_attrSelect[i]);
    }
    checkCudaErrors(cudaDeviceReset());
    getchar();
    return 0;
}

注意：有时候程序会报找不到helper_cuda.h文件，我们要记得include相关的helper等文件。比如windows一般保存在：C:\ProgramData\NVIDIA Corporation\CUDA Samples\v7.5\common\inc。在vs2013下有时候blockIdx等内置变量报未定义的错误，这时要添加头文件”#include<device_launch_parameters.h>。

Unifiled Memory

// includes, system
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>

// includes CUDA
#include <cuda_runtime.h>
//#include<device_launch_parameters.h>
// includes, project
#include <helper_cuda.h>
#include <helper_functions.h> // helper functions for SDK examples
#define N (3)
__device__ __managed__ int d_attrSelect[3];
__global__ void checkGlobalVariable()
{
    for (int i = 0; i < 3; i++)
    {
    
        d_attrSelect[i] += i;
        printf("Device: the value of the global variable is %d\n", d_attrSelect[i]);
    }
}
int main()
{
    d_attrSelect[0] = 1;
    d_attrSelect[1] = 1;
    d_attrSelect[2] = 1;
    for (int i = 0; i <3; i++)
    {
        printf("Host: the value of the global variable is %d\n", d_attrSelect[i]);
    }
    size_t sz = (size_t)(N) * sizeof(int);
    //checkCudaErrors(cudaMemcpyToSymbol(d_attrSelect, &h_attrSelect[0], sz,size_t(0),cudaMemcpyHostToDevice));
    //checkCudaErrors(cudaMemcpyToSymbol(d_attrSelect, h_attrSelect, sz, size_t(0), cudaMemcpyHostToDevice));
    int blocksize = 1;
    int gridsize = 1;
    dim3 block(blocksize);
    dim3 grid(gridsize);
    checkGlobalVariable <<<grid, block >>>();
    cudaDeviceSynchronize();
    //checkCudaErrors(cudaMemcpy(&h_attrSelect[0], d_attrSelect, sz, cudaMemcpyDeviceToHost));
    //checkCudaErrors(cudaMemcpyFromSymbol(&h_attrSelect[0], d_attrSelect, sz));
    //checkCudaErrors(cudaMemcpyFromSymbol(h_attrSelect, d_attrSelect, sz,size_t(0),cudaMemcpyDeviceToHost));
    for (int i = 0; i < 3; i++)
    {
        printf("Host: the value of the global variable is %d\n", d_attrSelect[i]);
    }
    checkCudaErrors(cudaDeviceReset());
    getchar();
    return 0;
}

在使用unifiled memory时一定要考虑到同步的问题：

 checkGlobalVariable <<<grid, block >>>();
 cudaDeviceSynchronize();

该处如果没有device的同步，cpu端再读取数据就会出错。

【CUDA】学习记录（7）- Global Memory

Memory

Benefits of a Memory Hierarchy

CUDA Memory Model

Registers

Local Memory

Shared Memory

Constant memory

Texture Memory

Global Memory

GPU Caches

CUDA Variable Declaration Summary

Static Global Memory

Unifiled Memory

相关文章

网友评论

延伸阅读

深度阅读

栏目导航

热点阅读

CUDA