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端再读取数据就会出错。