CUDA编程（五）关注内存的存取模式

news/2024/7/10 3:58:28 标签: cuda, GPU加速, 优化, 并行, 计算机视觉

cuda编程五">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》


http://www.niftyadmin.cn/n/1716533.html

相关文章

python输入两个数求和笔试题_Python入门练习（52）-组编程TIANTI练习：N个数之和,习题,团体,程序设计,天梯,赛,求和...

python输入两个数求和笔试题_Python入门练习（52）-组编程TIANTI练习：N个数之和,习题,团体,程序设计,天梯,赛,求和...

题目描述 本题的要求很简单&#xff0c;就是求N个数字的和。麻烦的是&#xff0c;这些数字是以有理数分子/分母的形式给出的&#xff0c;你输出的和也必须是有理数的形式。 输入格式&#xff1a; 输入第一行给出一个正整数N&#xff08;≤100&#xff09;。随后一行按格式a1/b1…
阅读更多...
二叉树的遍历_一文横扫二叉树的所有遍历方法

二叉树的遍历_一文横扫二叉树的所有遍历方法

小禹禹们&#xff0c;你们好呀&#xff0c;今天周一&#xff0c;景禹开始正式上线啦&#xff01;今日的内容&#xff0c;景禹昨天和今天两天时间给大家准备&#xff0c;今天我们谈一谈二叉树的四种遍历方式&#xff0c;看完保准让你对二叉树的遍历一网打尽。二叉树的遍历(trave…
阅读更多...
python自动测试m_appium+python搭建app自动化测试框架（一）

python自动测试m_appium+python搭建app自动化测试框架（一）

作者的配置环境和版本&#xff1a; win10 python3.6 Appium v1.4.16 下载node.js 验证安装&#xff1a; node -v 2.下载jdk 配置环境变量&#xff1a; JAVA_HOME jdk安装路径 CLASSPATH 变量值为&#xff1a;%JAVA_HOME%\lib\tools.jar;%JAVA_HOME%\lib\dt.jar path 里新建如…
阅读更多...
CUDA编程（六）进一步并行

CUDA编程（六）进一步并行

CUDA编程&#xff08;六&#xff09; 进一步并行 在之前我们使用Thread完成了简单的并行加速&#xff0c;虽然我们的程序运行速度有了50甚至上百倍的提升&#xff0c;但是根据内存带宽来评估的话我们的程序还远远不够&#xff0c;在上一篇博客中给大家介绍了一个访存方面非常…
阅读更多...
华为手机asph啥机型_华为手机降价力度大，今年发布的机型，价格真香

华为手机asph啥机型_华为手机降价力度大，今年发布的机型，价格真香

目前智能手机市场研究形成了一股多超多强的激烈竞争场面&#xff0c;大家所熟知的各大品牌就已经占据了整个手机市场的80&#xff05;。例如苹果和三星&#xff0c;已经连续数年位列全球手机市场占比的前3位&#xff0c;而作为中国智能手机的品牌门面华为手机&#xff0c;这几年…
阅读更多...
CUDA编程（七）共享内存与Thread的同步

CUDA编程（七）共享内存与Thread的同步

CUDA编程&#xff08;七&#xff09; 共享内存与Thread的同步 在之前我们通过block&#xff0c;继续增大了线程的数量&#xff0c;结果还是比较令人满意的&#xff0c;但是也产生了一个新的问题&#xff0c;即&#xff0c;我们在CPU端的加和压力变得很大&#xff0c;所以我们…
阅读更多...
CUDA编程（八）树状加法

CUDA编程（八）树状加法

CUDA编程&#xff08;八&#xff09; 树状加法 上一篇博客我们介绍了ShareMemory和Thread同步&#xff0c;最后利用这些知识完成了block内部线程结果的加和&#xff0c;减轻了CPU的负担&#xff0c;结果还是比较令人满意的&#xff0c;但是block的加和工作是使用一个thread0单…
阅读更多...
asp.net 安装element ui_Vue组件库系列二：打造属于自己的 UI 库文档（老版本的方案）...

asp.net 安装element ui_Vue组件库系列二：打造属于自己的 UI 库文档（老版本的方案）...

前一篇文章实现了封装个人的组件库功能&#xff0c;有了组件库&#xff0c;当然还要有配套说明文档&#xff0c;这样使者用起来才更方便。打包完成的dist目录是最终可放到服务器中&#xff0c;直接访问到文档的哟。在项目中配置打包examples上篇文章中&#xff0c;执行打包命令…
阅读更多...
最新文章

Copyright @ 2022~2023