深度学习模型部署(十四)CUDA编程-流&sharememory

于 2024-08-13 10:52:31 发布
阅读量807 收藏 22
点赞数 19
分类专栏： 并行编程 模型部署实战 文章标签： c++ 边缘计算 CUDA
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/qq_45983373/article/details/140956361
版权

流stream

前面提到，CUDA的线程结构为：grid，block，threads，每个线程之间是并行执行的。我们到目前为止的代码流程为：读取内存，进行计算，导出结果。如果我们要处理一个非常大的数据时，读取内存就需要占用很长时间的时候，那这段时间岂不是浪费了？能否并行读取数据？或者在读取到一部分数据后就开始部分运算？答案是可以的，CUDA提供了流(stream)技术，stream之间是并行的，stream内是串行的。方便程序员编写异步执行代码。
在这里插入图片描述

#include <iostream>

#define SPLIT 4

__global__ void vector_add(float *a, float *b, float *c, int width){
    int index = (blockIdx.x * gridDim.y + blockIdx.y)*blockDim.x * blockDim.y + threadIdx.x * blockDim.y + threadIdx.y;
    // 计算线程的全局索引，第blockIdx.x行第blockIdx.y列的block中的第threadIdx.x行第threadIdx.y列的线程
    if (index >= width) return;
    c[index] = a[index] + b[index];
    
}

void single_stream(float *a, float *b, float *c, int width){
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start,0); //0是指默认流
    cudaEventSynchronize(start);

    float *d_a, *d_b, *d_c;
    cudaMalloc((void **)&d_a, width * sizeof(float));
    cudaMalloc((void **)&d_b, width * sizeof(float));
    cudaMalloc((void **)&d_c, width * sizeof(float));

    cudaMemcpy(d_a, a, width * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, width * sizeof(float), cudaMemcpyHostToDevice);

    dim3 dimGrid(1, width >> 10);
    dim3 dimBlock(256, 4);
    vector_add<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, width);
    cudaMemcpy(c, d_c, width * sizeof(float), cudaMemcpyDeviceToHost);

    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    float time;
    cudaEventElapsedTime(&time, start, stop);
    std::cout << "Single stream time: " << time << "ms" << std::endl;

    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
}

void multi_stream(float *a, float *b, float *c, int width){
    cudaStream_t stream1, stream2, stream3,stream4;
    cudaStreamCreate(&stream1);
    cudaStreamCreate(&stream2);
    cudaStreamCreate(&stream3);
    cudaStreamCreate(&stream4);

    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start, 0);
    cudaEventSynchronize(start);

    float *d_a, *d_b, *d_c_1, *d_c_2, *d_c_3, *d_c_4;

    cudaMalloc((void **)&d_a, width * sizeof(float));
    cudaMalloc((void **)&d_b, width * sizeof(float));
    cudaMalloc((void **)&d_c_1, width / SPLIT * sizeof(float));
    cudaMalloc((void **)&d_c_2, width / SPLIT * sizeof(float));
    cudaMalloc((void **)&d_c_3, width / SPLIT * sizeof(float));
    cudaMalloc((void **)&d_c_4, width / SPLIT * sizeof(float));

    dim3 dimGrid(1, 1);
    dim3 dimBlock(width / SPLIT, 1);
    
    cudaMemcpyAsync(d_a, a, width / SPLIT * sizeof(float), cudaMemcpyHostToDevice, stream1);
    // 注意这里要用异步拷贝
    cudaMemcpyAsync(d_b, b, width / SPLIT * sizeof(float), cudaMemcpyHostToDevice, stream1);
    vector_add<<<dimGrid, dimBlock, 0, stream1>>>(d_a, d_b, d_c_1, width / SPLIT);
    cudaMemcpyAsync(c, d_c_1, width / SPLIT * sizeof(float), cudaMemcpyDeviceToHost, stream1);

    cudaMemcpyAsync(d_a + width / SPLIT, a + width / SPLIT, width / SPLIT * sizeof(float), cudaMemcpyHostToDevice, stream2);
    cudaMemcpyAsync(d_b + width / SPLIT, b + width / SPLIT, width / SPLIT * sizeof(float), cudaMemcpyHostToDevice, stream2);
    vector_add<<<dimGrid, dimBlock, 0, stream2>>>(d_a + width / SPLIT, d_b + width / SPLIT, d_c_2, width / SPLIT);
    cudaMemcpyAsync(c + width / SPLIT, d_c_2, width / SPLIT * sizeof(float), cudaMemcpyDeviceToHost, stream2);

    cudaMemcpyAsync(d_a + 2 * width / SPLIT, a + 2 * width / SPLIT, width / SPLIT * sizeof(float), cudaMemcpyHostToDevice, stream3);
    cudaMemcpyAsync(d_b + 2 * width / SPLIT, b + 2 * width / SPLIT, width / SPLIT * sizeof(float), cudaMemcpyHostToDevice, stream3);
    vector_add<<<dimGrid, dimBlock, 0, stream3>>>(d_a + 2 * width / SPLIT, d_b + 2 * width / SPLIT, d_c_3, width / SPLIT);
    cudaMemcpyAsync(c + 2 * width / SPLIT, d_c_3, width / SPLIT * sizeof(float), cudaMemcpyDeviceToHost, stream3);

    cudaMemcpyAsync(d_a + 3 * width / SPLIT, a + 3 * width / SPLIT, width / SPLIT * sizeof(float), cudaMemcpyHostToDevice, stream4);
    cudaMemcpyAsync(d_b + 3 * width / SPLIT, b + 3 * width / SPLIT, width / SPLIT * sizeof(float), cudaMemcpyHostToDevice, stream4);
    vector_add<<<dimGrid, dimBlock, 0, stream4>>>(d_a + 3 * width / SPLIT, d_b + 3 * width / SPLIT, d_c_4, width / SPLIT);
    cudaMemcpyAsync(c + 3 * width / SPLIT, d_c_4, width / SPLIT * sizeof(float), cudaMemcpyDeviceToHost, stream4);

    cudaStreamSynchronize(stream1);
    cudaStreamSynchronize(stream2);
    cudaStreamSynchronize(stream3);
    cudaStreamSynchronize(stream4);

    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);

    float time;
    cudaEventElapsedTime(&time, start, stop);
    std::cout << "Multi stream time: " << time << "ms" << std::endl;

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c_1);
    cudaFree(d_c_2);
    cudaFree(d_c_3);
    cudaFree(d_c_4);
}

bool if_right(float *c, int width){
    for(int i = 0; i < width; i++){
        if(c[i] != 3.0f){
            std::cout << "Error at: " << i << std::endl;
            return false;
        }
    }
    return true;
}

int main(){
    int width = 1 << 12;
    float *a = new float[width];
    float *b = new float[width];
    float *c = new float[width];

    for(int i = 0; i < width; i++){
        a[i] = 1.0f;
        b[i] = 2.0f;
    }

    single_stream(a, b, c, width);
    if_right(c, width);
    multi_stream(a, b, c, width);
    if_right(c, width);

    return 0;
}

sharememory

之前讨论过，GPU中的内存模型分为几个层次，读取速度由快到慢，大小由小到大：

  • register：寄存器
  • share memory
  • global memory

其中，register是线程自己占有的，share memory是block占有的，一个block内的线程可以共享，global memory是全局的，也就是我们常说的显存，存储介质是片外存储，DRAM存储。

DRAM不能一次读取一个字节，只能一次读取一行，这是其实现原理决定的

合理使用share memory能够有效降低访存延迟，提高整体效率。
例如进行向量元素求和。
如果使用最简单的实现，对于一个长度为D的向量，需要顺序执行D次访存，时间复杂度也是 O ( n ) O(n) O(n)：

__global__ void vectorSum(float *input ,float *result， int N){
	float temp_sum = 0.0;
	for(int i = 0; i < N; ++i){
		temp_sum += input[i];
	}
	result[0] = temp_sum;
}

如果使用二分法，先计算一半元素的和，再将得到的结果再求和，访存变成了2D次，但是由于是二分法，同时读取，所以时间复杂度减小为 O ( log ⁡ n ) O(\log^n) O(logn)。这种将一个向量经过计算得到一个标量的过程，被称为规约，这种二分法的规约方式被称为相邻规约：

__global__ void bisectionVectorSum(float *input, float* result, int N){
    int tid = threadIdx.x;
    if (tid >= N) return;
    int step = 1;
    while(step < N){
        if(tid % (2*step) == 0){
            input[tid] += input[tid + step];
        }
        __syncthreads();
        step *= 2;
    }
    if(tid == 0){
        *result = input[0];
    }

}

这种相邻规约会导致读取时多个线程同时读取相邻位置，导致访存冲突，CUDA中还有一种规约方式，称为交错规约：

__global__ void bisectVectorSum(float *input, float *result, int N){
    int tid = threadIdx.x;
    for(int step = N/2; step > 0; step /= 2){
        if(tid < step){
            input[tid] += input[tid + step];
        }
        __syncthreads();
    }
    if(tid == 0){
        *result = input[0];
    }
}

以上都是没有使用share memory版本的，通过合理使用share memory，可以将访存降低至D次，同时保持时间成本为 O ( log ⁡ n ) O(\log^n) O(logn)

__global__ void shareMemoryVectorSum(float *input, float *result, int N){
    int tid = threadIdx.x;
    __shared__ float temp[1024];
    temp[tid] = input[tid];
    __syncthreads();
    for(int step = N/2; step > 0; step /= 2){
        if(tid < step){
            temp[tid] += temp[tid + step];
        }
        __syncthreads();
    }
    if(tid == 0){
        *result = temp[0];
    }

}

share memory由32个bank构成，根据设备计算能力的不同，bank的大小也不同，计算能力3.x的，bank的宽度是8个字节64位，即share memory一共大小为32*8 = 256个byte。
需要注意的是，访问share memory时，每个bank同时只能有一个线程在读写，否则就会造成bank conflict，降低读写效率。所以像上面提到的相邻规约，就不如交错规约，因为交错规约每个线程访问的bank都相隔较远，不会产生bank conflict。

实际中还涉及到当多个线程访问同一个时会触发broadcast等具体情况，看官方手册。

战术摸鱼大师
关注
专栏目录
【人工智能】国产开源大模型 baichuan-7B 百川聊天 BaiChat 快速开始上手实战&效果实测 —— 希望 baichuan-7B 再打磨打磨吧
AI天才研究院
06-16 1万+
baichuan-7B 是由百川智能开发的一个开源可商用的大规模预训练语言模型。基于 Transformer 结构，在大约1.2万亿 tokens 上训练的70亿参数模型，支持中英双语，上下文窗口长度为4096。在标准的中文和英文权威 benchmark（C-EVAL/MMLU）上均取得同尺寸最好的效果。
【ChatGPT模型精调训练】AI 大模型精调 Fine-Tuning （微调）训练图文代码实战详解
最新发布
AI天才研究院
03-09 1008
选择预训练模型：选择一个在类似任务上已经训练好的模型作为起点。数据准备：准备并预处理你的数据集，使其适合模型的输入格式。微调：在你的特定数据集上继续训练模型，调整模型的权重。评估：评估微调后模型的性能。应用：将微调后的模型部署到实际应用中。Fine-Tuning（精调）是指在预训练模型的基础上，对模型进行微调，使其适应特定任务。预训练模型通常在大量无标签数据上进行训练，以学习通用的语言表示。通过 Fine-Tuning，我们可以在较小的标签数据集上训练模型，以便模型能够更好地解决特定任务。
CUDA编程，指南pdf
02-27
cuda开发，指南pdf
Linux下Linpack测试GPU和CPU性能
liu_feng_zi_的专栏
07-18 3893
前面的文章已经介绍了如何搭建Linpack环境：https://blog.csdn.net/liu_feng_zi_/article/details/107429720 接下来说明如何使用 一、编写运行脚本： 新建文件：run_linpack.sh，内容如下： export HPL_DIR =/root/LHC/Linpack/hpl-2.0_FERMI_v15 export MKL_NUM_THREADS=6 export OMP_NUM_THREADS=6 export MKL_DYNAMI.
CUDA编程简介与深度学习模型部署
qzm212422的专栏
06-30 2819
1. 简介 CUDA是NVIDIA针对外设GPU（Graphic Processing Unit）的一种并行计算架构，全称是Computer Unified Device Architecture。根据厂商可将GPU分为 N 卡（NVIDIA） 和 A 卡（ATI），通常使用的都是N卡。英伟达的显卡系列如下： 1.1 为什么要用并行计算？ 在科学计算领域所要用到的计算往往不是我们熟知的普通矩阵，而是千维甚至万维的矩阵，而普通的CPU串行计算往往不能满足与科学计算所要求的性能。最好的例子就是深度学.
CUDA编程入门极简教程
小白将
03-18 9万+
码字不易，欢迎给个赞！ 欢迎交流与转载，文章会同步发布在公众号：机器学习算法全栈工程师(Jeemy110) 目录 目录 前言 CUDA编程模型基础 向量加法实例 矩阵乘法实例 小结 参考资料 前言 2006年，NVIDIA公司发布了CUDA，CUDA是建立在NVIDIA的CPUs上的一个通用并行计算平台和编程模型，基于CUDA编程可以利用GPUs的并行计算引擎...
深度学习CUDA编程干货-kernel的编写和调用
jinfagang1211的博客
06-09 3014
title: 深度学习CUDA编程干货-kernel的编写和调用 date: 2020-06-09 12:08:54 category: 默认分类 本文介绍 深度学习CUDA编程干货-kernel的编写和调用 深度学习CUDA编程干货-kernel的编写和调用 本文由林大佬原创，转载请注明出处，来自腾讯、阿里等一线AI算法工程师组成的QQ交流群欢迎你的加入: 1037662480 上一篇给大家分享了一些CUDA编程的干货，这一篇来夯实一下，我们主要看一些基础的cuda概念。 三个层级 cuda编程主.
大语言模型的学习路线和开源模型的学习材料《一》
11-01 767
【LLMs 入门实战】 ChatGLM3 模型学习与实战【LLMs 入门实战】 ChatGLM3 模型微调学习与实战【ChatGLM2-6B入门】清华大学开源中文版ChatGLM-6B模型学习与实战【关于 ChatGLM2 + LoRA 进行finetune 】那些你不知道的事【LLMs 入门实战 】基于 🤗PEFT 的高效 🤖ChatGLM2-6B 微调【LLMs 入门实战】基于 🤗QLoRA 的高效 🤖ChatGLM2-6B 微调【LLMs 入门实战】 QLoRA微调Llama2 模型学习与实战
人工智能大模型技术基础系列之：分布式模型存储与加载
AI天才研究院
11-01 317
随着人工智能技术的快速发展，大规模语言模型（Large Language Models, LLMs）如GPT-3、BERT、T5等在自然语言处理领域取得了突破性进展。这些模型通常包含数十亿甚至上千亿个参数，为了有效地存储、加载和推理如此庞大的模型，分布式技术成为了不可或缺的解决方案。本文将深入探讨大模型分布式存储与加载的核心技术，包括模型并行、张量并行、流水线并行等关键概念，以及它们在实际应用中的实现方法和优化策略。
目标检测YOLO实战应用案例100讲-智能目标检测系统在FPGA中的设计与实现
qq_36130719的博客
06-06 2425
本文结合FPGA芯片的高并行、低功耗、可重配置等优势，对基于卷积神经网络结构的目标检测算法，在端侧场景下的加速与部署进行研究与实现。本章节对本文涉及到的基本理论知识与技术手段进行介绍，为后续系统的方案分析与系统设计奠定理论基础。首先，介绍了卷积神经网络的基本组成部分及其原理。然后，对FPGA及其开发技术进行简单介绍。最后，对当前常用的部分基于FPGA的卷积神经网络加速方法进行概述。
深度学习（三十六）异构计算CUDA学习笔记（1）
hjimce的专栏
05-26 1万+
二、CUDA编程步骤： 1、设置显卡编号：cudaSetDevice； 2、为显卡开辟变量内存:cudaMalloc； 3、把cup上的数据拷贝到GPU上：cudaMemcpy； 4、调用内核函数__global__类型函数； 5、把计算结果拷贝到CPU上：cudaMemcpy； 6、释放显存空间cudaFree； 示例代码： //计算a、b相加，得到c，size输入向量的维度 c
HPL on Intel MPI+MKL
vanderong的专栏
10-12 2431
Disable CUDA 1. Compiling Make.CUDA:  delete -DCUDA 2. Set CUDA_DGEMM_SPLIT, CUDA_DTRSM_SPLIT to 0.00 (run_linpack script) Modify code so that the codes could run directly on CPU nodes /src/cu
深度学习配置CUDA8.0/9.0及对应版本cuDNN安装
热门推荐
saintchen的博客
12-25 12万+
本人为中科院测地所博士生，所研究专业为自然地理学（遥感数据分析方向），研究课题偏向于深度学习。 由于本人不是计算机专业，故有关计算机配置及操作方面相较于计算机专业人员不是那么专业。所以请各位大牛大神绕道，我这里所做的一些工作比较浅显，仅供需要的各位一起交流。 本文主要讲解CUDA8.0/cuDNN5.1（CUDA9.1/cuDNN7.0）的安装及环境配置步骤。
CUDA矩阵BLAS效率
Augusdi的专栏
10-11 5423
以前曾听说cublas的效率不是很高, 今天写了个小程序对cublas的矩阵乘法速度进行了一个测试, 发现结果并非如此. 至少就矩阵乘法来说，cublas的效率很不错, 相对CPU有非常高的加速比.测试程序是在sdk例子simpleCUBLAS的基础上修改而成, 测试内容是分别用cublas和CPU函数计算两个N阶矩阵A和B的乘积, 然后对结果进行校验，并计算各自的flops.(每个矩阵元素加乘看
HPL TEST Q&A
vanderong的专栏
10-01 3448
Background 1. mvapich2-1.9a rpm -qa | grep mvapich2 wget http://mvapich.cse.ohio-state.edu/download/mvapich2/mvapich2-1.9a.tgz tar -xzf mvapich2-1.9a.tgz ./configure 无参数，之后make install出错 ./confi
CUDA编程手册（一）
Jony0917的专栏
02-13 2336
1 简介 1.1 从图像处理到通用并行计算 由于市场对实时高清3D图形计算的强烈需求，可编程图像处理器(programmable Graphic Processor Unit)或简称GPU，已发展成高并行、多核、多线程的处理器，拥有强大算力和高内存带宽。 以下图一和图二比较了CPU和GPU之间浮点运算能力和内存带宽间方面的差异： 图一：CPU和GPU每秒浮点运算次数比较 图二：CPU与GPU之...
CUDA编程实战：从安装到深度学习模型部署全攻略
学习者将学习到如何将优化后的深度学习模型部署到生产环境中，包括模型的训练、测试、优化以及打包等环节。 整个教程不仅仅局限于理论知识的讲解，更重要的是结合了大量实战实例，通过实际案例来展示CUDA编程在解决...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值