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》

相关实践学习
部署Stable Diffusion玩转AI绘画(GPU云服务器)
本实验通过在ECS上从零开始部署Stable Diffusion来进行AI绘画创作,开启AIGC盲盒。
目录
相关文章
|
3月前
|
缓存 安全 Java
Java并发编程进阶:深入理解Java内存模型
Java并发编程进阶:深入理解Java内存模型
46 0
|
21天前
|
并行计算 算法 IDE
【灵码助力Cuda算法分析】分析共享内存的矩阵乘法优化
本文介绍了如何利用通义灵码在Visual Studio 2022中对基于CUDA的共享内存矩阵乘法优化代码进行深入分析。文章从整体程序结构入手,逐步深入到线程调度、矩阵分块、循环展开等关键细节,最后通过带入具体值的方式进一步解析复杂循环逻辑,展示了通义灵码在辅助理解和优化CUDA编程中的强大功能。
|
22天前
|
缓存 算法 数据处理
如何选择合适的内存访问模式
【10月更文挑战第20天】如何选择合适的内存访问模式
38 1
|
2月前
|
监控 算法 Java
深入理解Java中的垃圾回收机制在Java编程中,垃圾回收(Garbage Collection, GC)是一个核心概念,它自动管理内存,帮助开发者避免内存泄漏和溢出问题。本文将探讨Java中的垃圾回收机制,包括其基本原理、不同类型的垃圾收集器以及如何调优垃圾回收性能。通过深入浅出的方式,让读者对Java的垃圾回收有一个全面的认识。
本文详细介绍了Java中的垃圾回收机制,从基本原理到不同类型垃圾收集器的工作原理,再到实际调优策略。通过通俗易懂的语言和条理清晰的解释,帮助读者更好地理解和应用Java的垃圾回收技术,从而编写出更高效、稳定的Java应用程序。
|
2月前
|
存储 并行计算 算法
CUDA统一内存:简化GPU编程的内存管理
在GPU编程中,内存管理是关键挑战之一。NVIDIA CUDA 6.0引入了统一内存,简化了CPU与GPU之间的数据传输。统一内存允许在单个地址空间内分配可被两者访问的内存,自动迁移数据,从而简化内存管理、提高性能并增强代码可扩展性。本文将详细介绍统一内存的工作原理、优势及其使用方法,帮助开发者更高效地开发CUDA应用程序。
|
2月前
|
缓存 Linux C语言
C语言 多进程编程(六)共享内存
本文介绍了Linux系统下的多进程通信机制——共享内存的使用方法。首先详细讲解了如何通过`shmget()`函数创建共享内存,并提供了示例代码。接着介绍了如何利用`shmctl()`函数删除共享内存。随后,文章解释了共享内存映射的概念及其实现方法,包括使用`shmat()`函数进行映射以及使用`shmdt()`函数解除映射,并给出了相应的示例代码。最后,展示了如何在共享内存中读写数据的具体操作流程。
|
3月前
|
开发框架 监控 .NET
|
4月前
|
SQL 资源调度 关系型数据库
实时计算 Flink版产品使用问题之在使用Flink on yarn模式进行内存资源调优时,如何进行优化
实时计算Flink版作为一种强大的流处理和批处理统一的计算框架,广泛应用于各种需要实时数据处理和分析的场景。实时计算Flink版通常结合SQL接口、DataStream API、以及与上下游数据源和存储系统的丰富连接器,提供了一套全面的解决方案,以应对各种实时计算需求。其低延迟、高吞吐、容错性强的特点,使其成为众多企业和组织实时数据处理首选的技术平台。以下是实时计算Flink版的一些典型使用合集。
|
4月前
|
设计模式 安全 NoSQL
Java面试题:设计一个线程安全的单例模式,并解释其内存占用和垃圾回收机制;使用生产者消费者模式实现一个并发安全的队列;设计一个支持高并发的分布式锁
Java面试题:设计一个线程安全的单例模式,并解释其内存占用和垃圾回收机制;使用生产者消费者模式实现一个并发安全的队列;设计一个支持高并发的分布式锁
68 0
|
4月前
|
设计模式 缓存 安全
Java面试题:设计模式在并发编程中的创新应用,Java内存管理与多线程工具类的综合应用,Java并发工具包与并发框架的创新应用
Java面试题:设计模式在并发编程中的创新应用,Java内存管理与多线程工具类的综合应用,Java并发工具包与并发框架的创新应用
41 0