CUDA编程（五）

关注内存的存取模式

上一篇博客我们使用Thread完成了简单的并行加速，虽然我们的程序运行速度有了50甚至上百倍的提升，但是根据内存带宽来评估的话我们的程序还远远不够，
除了通过Block继续提高线程数量来优化性能，这次想给大家先介绍一个访存方面非常重要的优化，同样可以大幅提高程序的性能~

什么样的存取模式是高效的？

大家知道一般显卡上的内存是 DRAM，因此最有效率的存取方式，是以连续的方式存取，单纯说连续存取可能比较抽象，我们还是通过例子来看这个问题。

之前的程序，大家可以看到我们非常重要的核函数部分：

// __global__ 函数 (GPU上执行) 计算立方和
__global__ static void sumOfSquares(int *num, int* result, clock_t* time)
{

    //表示目前的 thread 是第几个 thread（由 0 开始计算）
    const int tid = threadIdx.x;

    //计算每个线程需要完成的量
    const int size = DATA_SIZE / THREAD_NUM;

    int sum = 0;

    int i;

    //记录运算开始的时间
    clock_t start;

    //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录
    if (tid == 0) start = clock();

    for (i = tid * size; i < (tid + 1) * size; i++) {

        sum += num[i] * num[i] * num[i];

    }

    result[tid] = sum;

    //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行
    if (tid == 0) *time = clock() - start;

}

在计算立方和的部分，虽然看起来是连续存取内存位置（每个 thread 对一块连续的数字计算平方和），但是实际上并不是这样的，我们要考虑到实际上 thread 的执行方式。

前面提过，当一个 thread 在等待内存的数据时，GPU 会切换到下一个 thread。也就是说，实际上线程执行的顺序是类似

thread 0 -> thread 1 -> thread 2 -> thread 3 -> thread 4 ->...

因此，在同一个 thread 中连续存取内存，在实际执行时反而不是连续了，下图很明显的反应了这个问题，我们的存取是跳跃式的。

要让实际执行结果是连续的存取，我们应该要让 thread 0 读取第一个数字，thread 1 读取第二个数字…依此类推，很容易可以想象，通过这种存储方式，我们取数字的时候就变成了连续存取。

改进存取模式

根据我们上面的分析，我们原本的核函数并不是连续存取的，读取数字完全是跳跃式的读取，这会非常影响内存的存取效率，因此我们下一步要将取数字的过程变成：

thread 0 读取第一个数字，thread 1 读取第二个数字…

这点通过对核函数的for循环进行一个小修改就可以达到了~

// __global__ 函数 (GPU上执行) 计算立方和
__global__ static void sumOfSquares(int *num, int* result, clock_t* time)
{

    //表示目前的 thread 是第几个 thread（由 0 开始计算）
    const int tid = threadIdx.x;

    int sum = 0;

    int i;

    //记录运算开始的时间
    clock_t start;

    //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录
    if (tid == 0) start = clock();

    //改为连续存取（thread 0 读取第一个数字，thread 1 读取第二个数字 …）
    for (i = tid; i < DATA_SIZE; i += THREAD_NUM) {

        sum += num[i] * num[i] * num[i];

    }

    result[tid] = sum;

    //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行
    if (tid == 0) *time = clock() - start;

}

通过上面对for循环的一个小修改就可以达到目的了，那么这么一个微小的修改到底有多大作用呢？

完整代码 :

#include <stdio.h>
#include <stdlib.h>
#include <time.h>

//CUDA RunTime API
#include <cuda_runtime.h>

//1M
#define DATA_SIZE 1048576

#define THREAD_NUM 1024

int data[DATA_SIZE];

//产生大量0-9之间的随机数
void GenerateNumbers(int *number, int size)
{
    for (int i = 0; i < size; i++) {
        number[i] = rand() % 10;
    }
}

//打印设备信息
void printDeviceProp(const cudaDeviceProp &prop)
{
    printf("Device Name : %s.\n", prop.name);
    printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);
    printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);
    printf("regsPerBlock : %d.\n", prop.regsPerBlock);
    printf("warpSize : %d.\n", prop.warpSize);
    printf("memPitch : %d.\n", prop.memPitch);
    printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);
    printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
    printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
    printf("totalConstMem : %d.\n", prop.totalConstMem);
    printf("major.minor : %d.%d.\n", prop.major, prop.minor);
    printf("clockRate : %d.\n", prop.clockRate);
    printf("textureAlignment : %d.\n", prop.textureAlignment);
    printf("deviceOverlap : %d.\n", prop.deviceOverlap);
    printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);
}

//CUDA 初始化
bool InitCUDA()
{
    int count;

    //取得支持Cuda的装置的数目
    cudaGetDeviceCount(&count);

    if (count == 0) {
        fprintf(stderr, "There is no device.\n");
        return false;
    }

    int i;

    for (i = 0; i < count; i++) {

        cudaDeviceProp prop;
        cudaGetDeviceProperties(&prop, i);
        //打印设备信息
        printDeviceProp(prop);

        if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
            if (prop.major >= 1) {
                break;
            }
        }
    }

    if (i == count) {
        fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
        return false;
    }

    cudaSetDevice(i);

    return true;
}


// __global__ 函数 (GPU上执行) 计算立方和
__global__ static void sumOfSquares(int *num, int* result, clock_t* time)
{

    //表示目前的 thread 是第几个 thread（由 0 开始计算）
    const int tid = threadIdx.x;

    int sum = 0;

    int i;

    //记录运算开始的时间
    clock_t start;

    //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录
    if (tid == 0) start = clock();

    for (i = tid; i < DATA_SIZE; i += THREAD_NUM) {

        sum += num[i] * num[i] * num[i];

    }

    result[tid] = sum;

    //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行
    if (tid == 0) *time = clock() - start;

}





int main()
{

    //CUDA 初始化
    if (!InitCUDA()) {
        return 0;
    }

    //生成随机数
    GenerateNumbers(data, DATA_SIZE);

    /*把数据复制到显卡内存中*/
    int* gpudata, *result;

    clock_t* time;

    //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果，time用来存储运行时间 )
    cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
    cudaMalloc((void**)&result, sizeof(int)*THREAD_NUM);
    cudaMalloc((void**)&time, sizeof(clock_t));

    //cudaMemcpy 将产生的随机数复制到显卡内存中
    //cudaMemcpyHostToDevice - 从内存复制到显卡内存
    //cudaMemcpyDeviceToHost - 从显卡内存复制到内存
    cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);

    // 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
    sumOfSquares << < 1, THREAD_NUM, 0 >> >(gpudata, result, time);


    /*把结果从显示芯片复制回主内存*/

    int sum[THREAD_NUM];

    clock_t time_use;

    //cudaMemcpy 将结果从显存中复制回内存
    cudaMemcpy(&sum, result, sizeof(int)* THREAD_NUM, cudaMemcpyDeviceToHost);
    cudaMemcpy(&time_use, time, sizeof(clock_t), cudaMemcpyDeviceToHost);

    //Free
    cudaFree(gpudata);
    cudaFree(result);
    cudaFree(time);

    int final_sum = 0;

    for (int i = 0; i < THREAD_NUM; i++) {

        final_sum += sum[i];

    }

    printf("GPUsum: %d  gputime: %d\n", final_sum, time_use);

    final_sum = 0;

    for (int i = 0; i < DATA_SIZE; i++) {

        final_sum += data[i] * data[i] * data[i];

    }

    printf("CPUsum: %d \n", final_sum);

    return 0;
}

运行结果 :

我们看到这次运行用了894297个时钟周期

不知道大家是否还记得上次我们用1024个线程运行的最终结果：6489302个时钟周期，现在我们只是很简单的改了一下存取模式，同样使用1024个线程最终只使用了894297个时钟周期

6489302/894297= 7.26

可以看到我们的速度居然提升了7.26倍，而我们只是单纯修改了一下存取模式罢了，所以我们可以看到连续存取这个存取优化还是十分重要的，在我们没法再单纯地从线程数量上继续优化的情况下，从存取模式上进行的这个优化是十分有效的。

我们还是从内存带宽的角度来进行一下评估：

首先计算一下使用的时间：

894297/ (797000 * 1000) =  0.0011221S

然后计算使用的带宽：

数据量仍然没有变 DATA_SIZE 1048576，也就是1024*1024 也就是 1M

1M 个 32 bits 数字的数据量是 4MB。

因此，这个程序实际上使用的内存带宽约为：

4MB / 0.0011221S = 3564.745MB/s = 3.48GB/s

注意我们没进行内存存取优化之前的内存带宽是491MB/s，可以看到，我们通过这个优化一下子就把内存带宽提升到了GB级别，不得不说这是一个非常令人满意的效果，我们在没有继续增加线程数量的情况下，通过把内存的存取模式变成连续的，取得了7倍左右的加速。

总结：

这篇博客主要讲解了通过如何尽可能的连续操作内存，减少内存存取方面的时间浪费。

通过最终的结果我们可以看到，看似不起眼的一个小改进（尽可能的去连续操作内存），竟然有这近7倍的性能提升，所以希望大家记住这个优化，在优化我们的CUDA程序的时候，一定不要忘记从内存存取角度去进行一些优化，这往往能取得出乎意料的结果。

希望我的博客能帮助到大家~

参考资料：《深入浅出谈CUDA》