4.3 内存访问模式

发表于 2018-05-03 | 分类于 CUDA，Freshman | 评论数： 0 | 阅读次数：

Abstract: 本文介绍内存的访问过程，也就是从应用发起请求到硬件实现的完整操作过程，这里是优化内存瓶颈的关键之处，也是CUDA程序优化的基础。

Keywords: 内存访问模式，对齐，合并，缓存，结构体数组，数组结构体

内存访问模式

"物有本末，事有终始，知所先后，则近道矣" ——《大学·大学之道章》

这句话出自《大学》，对于技术学习也有启发意义。在学习过程中，理解事物的本质和发展规律，把握学习的先后顺序，才能更好地掌握知识。

对于CUDA编程来说，内存访问模式是性能优化的基础，是"本"；而各种具体的优化技巧是"末"。只有深入理解内存访问的本质规律，才能在实际编程中运用各种优化技术。

今天我们要学习的是CUDA中最重要的课程之一，本文将通过一些简单通俗、但足够恰当的比喻和实例，让大家更容易理解。

多数GPU程序容易受到内存带宽的限制，所以最大程度地利用全局内存带宽，提高全局加载效率，是调控内核函数性能的基本条件。如果不能正确调控全局内存使用，那么优化方案可能收效甚微。

CUDA执行模型告诉我们，CUDA执行的基本单位是线程束，所以，内存访问也是以线程束为基本单位发布和执行的。我们本文研究的就是一个线程束的内存访问，不同线程的内存请求，其目标位置的不同，可以产生非常多种情况。所以本篇就是研究这些不同情况，以及如何实现最佳的全局内存访问。

注意：访问可以是加载，也可以是存储。

对齐与合并访问

全局内存通过缓存实现加载和存储的过程如下图：

1-1

全局内存是一个逻辑层面的模型，我们编程的时候有两种模型考虑：一种是逻辑层面的，也就是我们在写程序的时候（包括串行程序和并行程序），写的一维（多维）数组、结构体、定义的变量，这些都是在逻辑层面的；一种是硬件角度，就是一块DRAM上的电信号，以及最底层内存驱动代码所完成的数字信号处理。

L1表示一级缓存，每个SM都有自己的L1，但是L2是所有SM公用的，除了L1缓存外，还有只读缓存和常量缓存，这个我们后面会详细介绍。

核函数运行时需要从全局内存（DRAM）中读取数据，只有两种粒度，这个是关键的：

128字节
32字节

解释下"粒度"，可以理解为最小单位，也就是核函数运行时每次读内存，哪怕是读一个字节的变量，也要读128字节，或者32字节，而具体是32还是128还是要看访问方式：

使用一级缓存
不使用一级缓存

对于CPU来说，一级缓存或者二级缓存是不能被编程的，但是CUDA是支持通过编译指令停用一级缓存的。如果启用一级缓存，那么每次从DRAM上加载数据的粒度是128字节，如果不使用一级缓存，只使用二级缓存，那么粒度是32字节。

还要强调一下CUDA内存模型的内存读写，我们现在讨论的都是单个SM上的情况，多个SM只是下面我们描述的情形的复制：SM执行的基础是线程束，也就是说，当一个SM中正在被执行的某个线程需要访问内存，那么，和它同线程束的其他31个线程也要访问内存，这个基础就表示，即使每个线程只访问一个字节，那么在执行的时候，只要有内存请求，至少是32个字节，所以不使用一级缓存的内存加载，一次粒度是32字节而不是更小。

在优化内存的时候，我们要最关注的是以下两个特性：

对齐内存访问
合并内存访问

我们把一次内存请求——也就是从内核函数发起请求，到硬件响应返回数据这个过程称为一个内存事务（加载和存储都可以）。

当一个内存事务的首个访问地址是缓存粒度（32或128字节）的偶数倍的时候：比如二级缓存32字节的偶数倍64，128字节的偶数倍256的时候，这个时候被称为对齐内存访问，非对齐访问就是除上述的其他情况，非对齐的内存访问会造成带宽浪费。

当一个线程束内的线程访问的内存都在一个内存块里的时候，就会出现合并访问。

对齐合并访问的状态是理想化的，也是最高效的访问方式，当线程束内的所有线程访问的数据在一个内存块，并且数据是从内存块的首地址开始被需要的，那么对齐合并访问出现了。为了最大化全局内存访问的理想状态，尽量将线程束访问内存组织成对齐合并的方式，这样的效率是最高的。下面看一个例子。

一个线程束加载数据，使用一级缓存，并且这个事务所请求的所有数据在一个128字节的对齐的地址段上，具体形式如下图，这里请求的数据是连续的，其实可以不连续，但是不要越界就好。

上面蓝色表示全局内存，下面橙色是线程束要的数据，绿色就是对齐的地址段。
如果一个事务加载的数据分布在不同的对齐地址段上，就会有以下两种情况：
1. 连续的，但是不在一个对齐的段上，比如，请求访问的数据分布在内存地址1~~128，那么0~~127和128~255这两段数据要传递两次到SM
2. 不连续的，也不在一个对齐的段上，比如，请求访问的数据分布在内存地址0~~63和128~~191上，明显这也需要两次加载。
上图就是典型的一个线程束，数据分散开了，thread0的请求在128之前，后面还有请求在256之后，所以需要三个内存事务，而利用率，也就是从主存取回来的数据被使用到的比例，只有128/(128×3)的比例。这个比例低会造成带宽的浪费，最极端的表现，就是如果每个线程的请求都在不同的段，也就是一个128字节的事务只有1个字节是有用的，那么利用率只有1/128。

这里总结一下内存事务的优化关键：用最少的事务次数满足最多的内存请求。事务数量和吞吐量的需求随设备的计算能力变化。

全局内存读取

注意我们说的都是读取，也就是加载过程，写或者叫做存储是另外一回事！

SM加载数据，根据不同的设备和类型分为三种路径：

一级和二级缓存
常量缓存
只读缓存

常规的路径是一级和二级缓存，需要使用常量和只读缓存的需要在代码中显式声明。但是提高性能，主要还是要取决于访问模式。

控制全局加载操作是否通过一级缓存可以通过编译选项来控制，当然比较老的设备可能就没有一级缓存。

编译器禁用一级缓存的选项是：

-Xptxas -dlcm=cg

编译器启用一级缓存的选项是：

-Xptxas -dlcm=ca

当一级缓存被禁用的时候，对全局内存的加载请求直接进入二级缓存，如果二级缓存缺失，则由DRAM完成请求。

每次内存事务可由一个、两个或者四个部分执行，每个部分有32个字节，也就是32，64或者128字节一次（注意前面我们讲到是否使用一级缓存决定了读取粒度是128还是32字节，这里增加的64并不在此情况，所以需要注意）。

启用一级缓存后，当SM有全局加载请求会首先尝试一级缓存，如果一级缓存缺失，则尝试二级缓存，如果二级缓存也没有，那么直接访问DRAM。

在有些设备上一级缓存不用来缓存全局内存访问，而是只用来存储寄存器溢出的本地数据，比如Kepler的K10,K20。

内存加载可以分为两类：

缓存加载
没有缓存的加载

内存访问有以下特点：

是否使用缓存：一级缓存是否介入加载过程
对齐与非对齐：如果访问的第一个地址是32的倍数
合并与非合并：访问连续数据块则是合并的

缓存加载

下面是使用一级缓存的加载过程，图片表达很清楚，我们只用少量文字进行说明：

对齐合并的访问，利用率100%
对齐的，但是不是连续的，每个线程访问的数据都在一个块内，但是位置是交叉的，利用率100%
连续非对齐的，线程束请求一个连续的非对齐的32个4字节数据，那么会出现，数据横跨两个块，但是没有对齐，当启用一级缓存的时候，就要两个128字节的事务来完成
线程束所有线程请求同一个地址，那么肯定落在一个缓存行范围（缓存行就是主存上一个可以被一次读到缓存中的一段数据），那么如果按照请求的是4字节数据来说，使用一级缓存的利用率是4/128=3.125%
比较坏的情况，前面提到过最坏的，就是每个线程束内的线程请求的都是不同的缓存行内，这里比较坏的情况就是，所有数据分布在N个缓存行上，其中1≤N≤32，那么请求32个4字节的数据，就需要N个事务来完成，利用率也是1/N

CPU和GPU的一级缓存有显著的差异，GPU的一级缓存可以通过编译选项等控制，CPU不可以，而且CPU的一级缓存的替换算法是有使用频率和时间局部性的，GPU则没有。

没有缓存的加载

没有缓存的加载是指没有通过一级缓存，二级缓存则是不得不经过的。

当不使用一级缓存的时候，内存事务的粒度变为32字节，更细粒度的好处是提高利用率，这个很好理解，比如你每次喝水只能选择一瓶大瓶500ml的或者一个小瓶的250ml，当你非常渴的时候需要400ml水分，喝大瓶的比较方便，因为如果喝小瓶的一瓶不够，还需要再喝一瓶，此时大瓶的方便。但如果你需要200ml的水分的时候，小瓶的利用率就高很多。细粒度的访问就是用小瓶喝水，虽然体积小，但是每次的利用率都高了不少，针对上面使用缓存的情况5，可能效果会更好。

继续我们的图解：

对齐合并访问128字节，不用说，还是最理想的情况，使用4个段，利用率100%
对齐不连续访问128字节，都在四个段内，且互不相同，这样的利用率也是100%
连续不对齐，一个段32字节，所以，一个连续的128字节的请求，即使不对齐，最多也不会超过五个段，所以利用率是4/5=80%，如果不明白为什么不能超过5个段，请注意前提是连续的，这个时候不可能超过五段
所有线程访问一个4字节的数据，那么此时的利用率是4/32=12.5%
最坏的情况，所有目标数据分散在内存的各个角落，那么需要N个内存段，此时与使用一级缓存的对比也是有优势的，因为N×128还是要比N×32大不少，这里假设N不会因为128还是32而变化，而实际情况，当使用大粒度的缓存行的时候，N有可能会减小

非对齐读取示例

下面就非对齐读取进行演示，代码如下：

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

void sumArrays(float * a, float * b, float * res, int offset, const int size)
{
    for(int i = 0, k = offset; k < size; i++, k++)
    {
        res[i] = a[k] + b[k];
    }
}

__global__ void sumArraysGPU(float*a, float*b, float*res, int offset, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int k = i + offset;
    if(k < n)
        res[i] = a[k] + b[k];
}

int main(int argc, char **argv)
{
    int dev = 0;
    cudaSetDevice(dev);

    int nElem = 1 << 18;
    int offset = 0;
    if(argc >= 2)
        offset = atoi(argv[1]);
    printf("Vector size:%d\n", nElem);
    int nByte = sizeof(float) * nElem;
    float *a_h = (float*)malloc(nByte);
    float *b_h = (float*)malloc(nByte);
    float *res_h = (float*)malloc(nByte);
    float *res_from_gpu_h = (float*)malloc(nByte);
    memset(res_h, 0, nByte);
    memset(res_from_gpu_h, 0, nByte);

    float *a_d, *b_d, *res_d;
    CHECK(cudaMalloc((void**)&a_d, nByte));
    CHECK(cudaMalloc((void**)&b_d, nByte));
    CHECK(cudaMalloc((void**)&res_d, nByte));
    CHECK(cudaMemset(res_d, 0, nByte));
    initialData(a_h, nElem);
    initialData(b_h, nElem);

    CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));

    dim3 block(1024);
    dim3 grid(nElem/block.x);
    double iStart, iElaps;
    iStart = cpuSecond();
    sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d, offset, nElem);
    cudaDeviceSynchronize();
    iElaps = cpuSecond() - iStart;
    CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost));
    printf("Execution configuration<<<%d,%d>>> Time elapsed %f sec --offset:%d \n", grid.x, block.x, iElaps, offset);

    sumArrays(a_h, b_h, res_h, offset, nElem);

    checkResult(res_h, res_from_gpu_h, nElem);
    cudaFree(a_d);
    cudaFree(b_d);
    cudaFree(res_d);

    free(a_h);
    free(b_h);
    free(res_h);
    free(res_from_gpu_h);

    return 0;
}

编译指令：

nvcc -O3 -arch=sm_35 -Xptxas -dlcm=cg -I ../include/ sum_array_offset.cu -o sum_array_offset

运行结果：

res-1

res-cg_nvprof

编译指令，启用一级缓存：

nvcc -O3 -arch=sm_35 -Xptxas -dlcm=ca -I ../include/ sum_array_offset.cu -o sum_array_offset

res-2

res-ca_nvprof

这里我们使用的指标是：

全局加载效率 = 请求的全局内存加载吞吐量 / 所需的全局内存加载吞吐量

只读缓存

只读缓存最初是留给纹理内存加载用的，在计算能力3.5以上的设备，只读缓存也支持使用全局内存加载代替一级缓存。也就是说3.5以后的设备，可以通过只读缓存从全局内存中读数据了。

只读缓存粒度32字节，对于分散读取，细粒度优于一级缓存。

有两种方法指导内存从只读缓存读取：

使用函数__ldg
在间接引用的指针上使用修饰符

代码：

__global__ void copyKernel(float * in, float* out)
{
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    out[idx] = __ldg(&in[idx]);
}

注意函数参数，然后就能强制使用只读缓存了。

全局内存写入

内存的写入和读取（或者叫做加载）是完全不同的，并且写入相对简单很多。一级缓存不能用在Fermi和Kepler GPU上进行存储操作，发送到设备前，只经过二级缓存，存储操作在32个字节的粒度上执行，内存事务也被分为一段、两段或者四段，如果两个地址在一个128字节的段内但不在64字节范围内，则会产生一个四段的事务，其他情况以此类推。

我们将内存写入也参考前面的加载分为下面这些情况：

对齐的，访问一个连续的128字节范围。存储操作使用一个4段事务完成：
分散在一个192字节的范围内，不连续，使用3个一段事务来完成：
对齐的，在一个64字节的范围内，使用一个两段事务完成：

非对齐写入示例

与读取情况类似，且更简单，因为始终不经过一级缓存，所以略过此实验。

结构体数组与数组结构体

写过C语言的人对结构体都应该非常了解，结构体实际上就是基础数据类型组合出来的新的数据类型，这个新的数据类型在内存中表现是：结构中的成员在内存里对齐地依次排开，然后我们就有了接下来的话题，数组的结构体，和结构体的数组。

数组结构体（AoS）就是一个数组，每个元素都是一个结构体，而结构体数组（SoA）就是结构体中的成员是数组，用代码表示：

AoS（数组结构体）：

struct A a[N];

SoA（结构体数组）：

struct A{
    int a[N];
    int b[N];
} a;

如果你分不清这两个名字，没关系，记住AoS是数组就行了。CUDA对细粒度数组是非常友好的，但是对粗粒度如结构体组成的数组就不太友好了，具体表现在，内存访问利用率低。比如当一个线程要访问结构体中的某个成员的时候，当32个线程同时访问的时候，SoA的访问就是连续的，而AoS则是不连续的：

4-22

这样看来AoS访问效率只有50%。

对比AoS和SoA的内存布局，我们能得到下面结论：

并行编程范式，尤其是SIMD（单指令多数据）对SoA更友好。CUDA中普遍倾向于SoA，因为这种内存访问可以有效地合并。

AoS数据布局的简单数学运算

我们看一下AoS的例子：

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

struct naiveStruct{
    float a;
    float b;
};

void sumArrays(float * a, float * b, float * res, const int size)
{
    for(int i = 0; i < size; i++)
    {
        res[i] = a[i] + b[i];
    }
}

__global__ void sumArraysGPU(float*a, float*b, struct naiveStruct* res, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if(i < n)
        res[i].a = a[i] + b[i];
}

void checkResult_struct(float* res_h, struct naiveStruct*res_from_gpu_h, int nElem)
{
    for(int i = 0; i < nElem; i++)
        if (res_h[i] != res_from_gpu_h[i].a)
        {
            printf("check fail!\n");
            exit(0);
        }
    printf("result check success!\n");
}

int main(int argc, char **argv)
{
    int dev = 0;
    cudaSetDevice(dev);

    int nElem = 1 << 18;
    int offset = 0;
    if(argc >= 2)
        offset = atoi(argv[1]);
    printf("Vector size:%d\n", nElem);
    int nByte = sizeof(float) * nElem;
    int nByte_struct = sizeof(struct naiveStruct) * nElem;
    float *a_h = (float*)malloc(nByte);
    float *b_h = (float*)malloc(nByte);
    float *res_h = (float*)malloc(nByte_struct);
    struct naiveStruct *res_from_gpu_h = (struct naiveStruct*)malloc(nByte_struct);
    memset(res_h, 0, nByte);
    memset(res_from_gpu_h, 0, nByte);

    float *a_d, *b_d;
    struct naiveStruct* res_d;
    CHECK(cudaMalloc((void**)&a_d, nByte));
    CHECK(cudaMalloc((void**)&b_d, nByte));
    CHECK(cudaMalloc((void**)&res_d, nByte_struct));
    CHECK(cudaMemset(res_d, 0, nByte_struct));
    initialData(a_h, nElem);
    initialData(b_h, nElem);

    CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));

    dim3 block(1024);
    dim3 grid(nElem/block.x);
    double iStart, iElaps;
    iStart = cpuSecond();
    sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d, nElem);
    cudaDeviceSynchronize();
    iElaps = cpuSecond() - iStart;
    CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte_struct, cudaMemcpyDeviceToHost));
    printf("Execution configuration<<<%d,%d>>> Time elapsed %f sec\n", grid.x, block.x, iElaps);

    sumArrays(a_h, b_h, res_h, nElem);

    checkResult_struct(res_h, res_from_gpu_h, nElem);
    cudaFree(a_d);
    cudaFree(b_d);
    cudaFree(res_d);

    free(a_h);
    free(b_h);
    free(res_h);
    free(res_from_gpu_h);

    return 0;
}

编译指令：

nvcc -O3 -arch=sm_35 -Xptxas -dlcm=ca -I ../include/ AoS.cu -o AoS

AoS

nvcc -O3 -arch=sm_35 -Xptxas -dlcm=cg -I ../include/ AoS.cu -o AoS

AoS2

SoA数据布局的简单数学运算

然后看SoA的例子：

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

void sumArrays(float * a, float * b, float * res, int offset, const int size)
{
    for(int i = 0, k = offset; k < size; i++, k++)
    {
        res[i] = a[k] + b[k];
    }
}

__global__ void sumArraysGPU(float*a, float*b, float*res, int offset, int n)
{
    int i = blockIdx.x * blockDim.x * 4 + threadIdx.x;
    int k = i + offset;
    if(k + 3 * blockDim.x < n)
    {
        res[i] = a[k] + b[k];
        res[i + blockDim.x] = a[k + blockDim.x] + b[k + blockDim.x];
        res[i + blockDim.x * 2] = a[k + blockDim.x * 2] + b[k + blockDim.x * 2];
        res[i + blockDim.x * 3] = a[k + blockDim.x * 3] + b[k + blockDim.x * 3];
    }
}

int main(int argc, char **argv)
{
    int dev = 0;
    cudaSetDevice(dev);
    int block_x = 512;
    int nElem = 1 << 18;
    int offset = 0;
    if(argc == 2)
        offset = atoi(argv[1]);
    else if(argc == 3)
    {
        offset = atoi(argv[1]);
        block_x = atoi(argv[2]);
    }
    printf("Vector size:%d\n", nElem);
    int nByte = sizeof(float) * nElem;
    float *a_h = (float*)malloc(nByte);
    float *b_h = (float*)malloc(nByte);
    float *res_h = (float*)malloc(nByte);
    float *res_from_gpu_h = (float*)malloc(nByte);
    memset(res_h, 0, nByte);
    memset(res_from_gpu_h, 0, nByte);

    float *a_d, *b_d, *res_d;
    CHECK(cudaMalloc((void**)&a_d, nByte));
    CHECK(cudaMalloc((void**)&b_d, nByte));
    CHECK(cudaMalloc((void**)&res_d, nByte));
    CHECK(cudaMemset(res_d, 0, nByte));
    initialData(a_h, nElem);
    initialData(b_h, nElem);

    CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));

    dim3 block(block_x);
    dim3 grid(nElem/block.x);
    double iStart, iElaps;
    iStart = cpuSecond();
    sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d, offset, nElem);
    cudaDeviceSynchronize();
    iElaps = cpuSecond() - iStart;

    printf("warmup Time elapsed %f sec\n", iElaps);
    iStart = cpuSecond();
    sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d, offset, nElem);
    cudaDeviceSynchronize();
    iElaps = cpuSecond() - iStart;
    CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost));
    printf("Execution configuration<<<%d,%d>>> Time elapsed %f sec --offset:%d \n", grid.x, block.x, iElaps, offset);

    sumArrays(a_h, b_h, res_h, offset, nElem);

    checkResult(res_h, res_from_gpu_h, nElem - 4 * block_x);
    cudaFree(a_d);
    cudaFree(b_d);
    cudaFree(res_d);

    free(a_h);
    free(b_h);
    free(res_h);
    free(res_from_gpu_h);

    return 0;
}

编译指令：

nvcc -O3 -arch=sm_35 -Xptxas -dlcm=ca -I ../include/ SoA.cu -o SoA

SoA-ca

nvcc -O3 -arch=sm_35 -Xptxas -dlcm=cg -I ../include/ SoA.cu -o SoA

SoA-cg

性能调整

优化设备内存带宽利用率有两个目标：

对齐合并内存访问，以减少带宽的浪费
足够的并发内存操作，以隐藏内存延迟

第三章我们讲过优化指令吞吐量的核函数，实现并发内存访问量最大化是通过以下方式得到的：

增加每个线程中执行独立内存操作的数量
对核函数启动的执行配置进行试验，以充分体现每个SM的并行性

接下来我们就按照这个思路对程序进行优化试验：展开技术和增大并行性。

展开技术

把前面讲到的展开技术用到向量加法上，我们来看看其对内存效率的影响：

代码：

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

void sumArrays(float * a, float * b, float * res, int offset, const int size)
{
    for(int i = 0, k = offset; k < size; i++, k++)
    {
        res[i] = a[k] + b[k];
    }
}

__global__ void sumArraysGPU(float*a, float*b, float*res, int offset, int n)
{
    int i = blockIdx.x * blockDim.x * 4 + threadIdx.x;
    int k = i + offset;
    if(k + 3 * blockDim.x < n)
    {
        res[i] = a[k] + b[k];
        res[i + blockDim.x] = a[k + blockDim.x] + b[k + blockDim.x];
        res[i + blockDim.x * 2] = a[k + blockDim.x * 2] + b[k + blockDim.x * 2];
        res[i + blockDim.x * 3] = a[k + blockDim.x * 3] + b[k + blockDim.x * 3];
    }
}

int main(int argc, char **argv)
{
    int dev = 0;
    cudaSetDevice(dev);
    int block_x = 512;
    int nElem = 1 << 18;
    int offset = 0;
    if(argc == 2)
        offset = atoi(argv[1]);
    else if(argc == 3)
    {
        offset = atoi(argv[1]);
        block_x = atoi(argv[2]);
    }
    printf("Vector size:%d\n", nElem);
    int nByte = sizeof(float) * nElem;
    float *a_h = (float*)malloc(nByte);
    float *b_h = (float*)malloc(nByte);
    float *res_h = (float*)malloc(nByte);
    float *res_from_gpu_h = (float*)malloc(nByte);
    memset(res_h, 0, nByte);
    memset(res_from_gpu_h, 0, nByte);

    float *a_d, *b_d, *res_d;
    CHECK(cudaMalloc((void**)&a_d, nByte));
    CHECK(cudaMalloc((void**)&b_d, nByte));
    CHECK(cudaMalloc((void**)&res_d, nByte));
    CHECK(cudaMemset(res_d, 0, nByte));
    initialData(a_h, nElem);
    initialData(b_h, nElem);

    CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));

    dim3 block(block_x);
    dim3 grid(nElem/block.x);
    double iStart, iElaps;
    iStart = cpuSecond();
    sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d, offset, nElem);
    cudaDeviceSynchronize();
    iElaps = cpuSecond() - iStart;

    printf("warmup Time elapsed %f sec\n", iElaps);
    iStart = cpuSecond();
    sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d, offset, nElem);
    cudaDeviceSynchronize();
    iElaps = cpuSecond() - iStart;
    CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost));
    printf("Execution configuration<<<%d,%d>>> Time elapsed %f sec --offset:%d \n", grid.x, block.x, iElaps, offset);

    sumArrays(a_h, b_h, res_h, offset, nElem);

    checkResult(res_h, res_from_gpu_h, nElem - 4 * block_x);
    cudaFree(a_d);
    cudaFree(b_d);
    cudaFree(res_d);

    free(a_h);
    free(b_h);
    free(res_h);
    free(res_from_gpu_h);

    return 0;
}

编译指令：

nvcc -O3 sum_array_offset_unrolling.cu -o sum_array_offset_unrolling -arch=sm_35 -Xptxas -dlcm=cg -I ../include/

unrolling-1

nvprof内存效率：

unrolling-nv

增大并行性

通过调整块的大小来实现并行性调整，也是前面讲过的套路，我们关注的还是内存利用效率。

代码同上面的展开技术。

res-block

offset=11的时候：

res-off-set-11

由于数据量相对较小，所以时间差距不大，512有最佳速度，不仅因为内存，还有并行性等多方面因素，这个前面我们也曾提到过。要看综合能力。

本文全部代码都在Github上有完整版，请访问：https://github.com/Tony-Tan/CUDA_Freshman

总结

这是我今年写作时间最长的一篇博客，写了三天，主要是代码比较多，结果也比较多。这里我们没用CMake，而是用的指令，原因是方便修改编译选项，试验时间结果不明显的原因是数据量小，部分结果和书上不一致，主要是书的时间比较久了，GPU换代太快。

全局内存本篇算是比较完整了，后面还有其他内存知识，我们继续。

内存访问模式​

对齐与合并访问​

全局内存读取​

缓存加载​

没有缓存的加载​

非对齐读取示例​

只读缓存​

全局内存写入​

非对齐写入示例​

结构体数组与数组结构体​

AoS数据布局的简单数学运算​

SoA数据布局的简单数学运算​

性能调整​

展开技术​

增大并行性​

总结​

内存访问模式

对齐与合并访问

全局内存读取

缓存加载

没有缓存的加载

非对齐读取示例

只读缓存

全局内存写入

非对齐写入示例

结构体数组与数组结构体

AoS数据布局的简单数学运算

SoA数据布局的简单数学运算

性能调整

展开技术

增大并行性

总结