CUDA编程一、基本概念和cuda向量加法

       

目录

一、cuda编程的基本概念入门

1、GPU架构和存储结构

2、cuda编程模型

3、cuda编程流程

二、cuda向量加法实践

1、代码实现

2、代码运行和结果


        有一段时间对模型加速比较感兴趣,其中的一块儿内容就是使用C++和cuda算子优化之类一起给模型推理提速。之前一直没有了解过cuda编程,也没有学习过C++相关的东西。强迫自己来学习一下cuda编程,同时也学习一下C++,扩宽一下AI相关的领域知识。主要是能够理解怎么使用cuda来提升模型的推理速度,学习的目标就是要会使用cuda编程实现基本的向量加法乘法、能使用C++和cuda混合编程实现神经网络的一些基本模块、最终能够完成C++语言和cuda混合编程(自己实现算子或者调用英伟达成熟的库)完成一个LLM模型的前向推理过程。这里是cuda编程的第一篇入门篇,了解基本概念、gpu的架构、cuda编程模型和实现cuda向量加法,对cuda编程有一个基础的了解和实践。

一、cuda编程的基本概念入门

1、GPU架构和存储结构

GPU全称图形处理器(graphics processing unit),主要是做图像和图像等涉及到并行计算的微处理器。 GPU和CPU同样有自己的架构,GPU更重计算、CPU更重逻辑控制。从硬件层面来说,GPU的内部构成如下图——详解GPU:

GPU通常包括图形显存控制器、压缩单元、BIOS、图形和计算阵列、总线接口、电源管理单元、视频管理单元、显示接口等,我们用来做深度学习就主要用到了它的图形和计算阵列模块。

GPU微架构

从微架构角度来说,GPU是有一个个SM(Streaming Multiprocessors)构成的。如下图:

这是A100 安培架构显卡的SM内部结构图,SM由L1缓存、指令缓存、寄存器(Register)和Wrap scheduler等构成。图中的绿色部分是tensor core 也可以称作SP(Streaming processor),用于浮点数的计算,它可以支持一个时钟周期完成两个16×16矩阵的乘法操作,其他版本如Volta完成两个两个4×4半精度浮点矩阵的计算、Turing完成64个半精度浮点的乘加操作,总之计算速度更慢。

内存模型

GPU的内存也是多层级结构的,具体结构如下:

通用内存DRAM目前最好的显卡采用了HBM(High Bandwidth Memory 高带宽内存);更近一级的是L2缓存,所有的SM共享;L2之上就是L1缓存,SM独有,所有的显存共享;L1之上的就是寄存器,线程独有的。如图,英伟达的cuda编程指南中给出示意图:

2、cuda编程模型

GPU其实可以看做一个超多线程处理器,一个运行多次使用不同数据执行的程序,可以使用很多不同的线程来执行,在GPU上就是把这个函数编译为设备的指令集——kernel核函数。cuda就是实现这样功能的一个代码库,可以让开发者使用高级语言来实现上述GPU的多线程并行执行,加速计算速度。

图中显示一个kernel会被grid中的线程块一起执行。这里就有几个概念,gird、block和thread。一个grid有多个block构成;一个block有多个thread构成。其中grid中的block有x/y/z三个维度,总数有最大值,每个维度上有各自的最大值,需要查阅当前的cuda规范。同时block中的线程也分x/y/z三个维度,总数有最大值,每个维度上有各自的最大值。一般来说,block中的线程数最大为1024个。线程的序号由block数目和线程在block中的位置,对于上述kernel1,thread(4,2)来说,线程Id

threadId

=(threadIdx.x+threadIdx.y*blockDim.x)+(blockIdx.x+blockIdx.y*gridDim.x)*(blockDim.x*blockDim.y)

=(4+2*5)+(1+1*3)*(5*3)=14+60=74

内存模型

线程在内存的使用是什么样的?GPU的内存模型如下:

block中的线程共用shared Mem,线程独立拥有寄存器和本地内存,其它的内存都是所有的block共享的。

3、cuda编程流程

cuda编程流程其实有点像我们使用GPU进行模型训练,模型训练中首先是模型加载和数据的处理;然后是把模型参数和数据都从CPU内存移动到GPU内存(显存)上;最后进行模型训练。那cuda编程宏观上也是这么个逻辑,这里摘抄一段知乎博主小小将博文《CUDA编程入门极简教程》总结的流程如下:

  1. 分配host内存,并进行数据初始化;
  2. 分配device内存,并从host将数据拷贝到device上;
  3. 调用CUDA的核函数在device上完成指定的运算;
  4. 将device上的运算结果拷贝到host上;
  5. 释放device和host上分配的内存。

cuda编程的重点和难点也在于第三步cuda核函数的设计实现(设计一个跑通的可能不难但是设计一个高效率的可能就很难了),核函数的定义如下:

__global__ void kernelFunction(float *result, float *a, float *b){
    doSomething
}

使用__global__对核函数进行限定,表示该函数是一个GPU核函数,在GPU的线程中被执行。

使用一个核函数整体的代码流程如下:

__global__ void kernelFunction(float *result, float *a, float *b){
    doSomething
}
int main(){
    ......
    // 分配内存和显存
    cudaMallocManaged();
    //数据初始化
    initWith();
    // 每一个gird有多少个block 最大2^31-1   x方向最大2^31-1  y,z 方向65535
    dim3 gridDim(x,y,z);
    // 每一个block有多少个线程  最大1024  x,y方向最大1024  z最大64
    dim3 blockDim(x,y,z);
    //执行核函数
    kernelFunction <<< gridDim, blockDim >>>();
    cudaDeviceSynchronize(); // 同步
    ......
}

使用kernelFunction<<<gridDim,blockDim>>>()来指定对应的gridDim和blockDim并且启动和函数。根据wiki的数据显示:每一个gird有多少个block 最大2^31-1   x方向最大2^31-1  y,z 方向65535; 每一个block有多少个线程  最大1024  x,y方向最大1024  z最大64。

二、cuda向量加法实践

1、代码实现

       接下来基于cuda来实现两个一维矩阵(一维向量)的加法。按照上述cuda编程流程,首先需要进行数据初始化,然后把数据传输到GPU上,然后进行cuda核函数的计算,最后得到结果释放资源。首先看一下数据怎么在CPU和GPU上灵活的传输,新版本的cuda有如下API:

cudaError_t cudaMallocManaged(void** ptr, size_t size)

该函数运行我们在内存和显存开辟size_t大小的空间,并智能的进行数据在CPU和GPU上的移动。

现在需要设计核函数,简单起见我们设置grid和block都为一维的,核函数逻辑就可以按照如下设计:

__global__ void addVectorskernel(float *result, float *a, float *b, int N){
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    for (int i=index; i<N; i+=stride){
        result[i] = a[i] + b[i]; // 元素a[i] + 元素 b[i]
    }
    
}

其中blockDim就表示一个block中有多少个线程,gridDim表示一个grid(一个gpu)中有多个block,那么总线程数就是blockDim.x * gridDim.x,每个线程处理的向量元素就是N/(blockDim.x * gridDim.x),因此就会有内部的循环,循环的步长也是总线程数blockDim.x * gridDim.x。当N=102400000,blockDim.x= 256,gridDim.x = 10,a矩阵的值全为3.0,b矩阵的值全为4.0,那么就可以得到如下代码nvcc_vector_add.cu:

#include<stdio.h>
#include<assert.h>
#include<cstdio>
#include<sys/time.h>
#include<iostream>
// 编译加链接
// nvcc -o nvcc_vector_add.cu  nvcc_vector_add.o
// 直接运行即可
// 向量加法核函数
__global__ void addVectorskernel(float *result, float *a, float *b, int N){
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    for (int i=index; i<N; i+=stride){
        result[i] = a[i] + b[i]; // 元素a[i] + 元素 b[i]
    }
}
// 初始化数组 a
void initWith(float num, float *a, int N) {
    for(int i = 0; i < N; ++i) {
      a[i] = num;
    }
  };

int main(){
    const int N = 102400000;
    const int M = 10;
    size_t Mem = N * sizeof(float);
    float *a;
    float *b;
    float *c;

    cudaMallocManaged(&a, Mem);
    cudaMallocManaged(&b, Mem);
    cudaMallocManaged(&c, Mem);

    initWith(3.0, a, N); // 将数组a中所有的元素初始化为3
    initWith(4.0, b, N); // 将数组b中所有的元素初始化为4
    initWith(0.0, c, N); // 将数组c中所有的元素初始化为0,数组c是结果向量

    for(int i=0;i<M;i++){
        printf("%f ",a[i]);
    }
    printf("\n");
    printf("******************\n");
    for(int i=0;i<M;i++){
        printf("%f ",b[i]);
    }
    printf("\n");
    printf("******************\n");
    for(int i=0;i<M;i++){
        printf("%f ",c[i]);
    }
    printf("\n");
    printf("******************\n");

    // 配置参数
    size_t threadsPerBlock = 256;
    // size_t numberOfBlocks = (N + threadsPerBlock - 1) / threadsPerBlock;
    size_t numberOfBlocks = 10;

    struct timeval start;
    struct timeval end;
    gettimeofday(&start,NULL);
    addVectorskernel <<< numberOfBlocks, threadsPerBlock >>> (c, a, b, N); // 执行核函数
    cudaDeviceSynchronize(); // 同步,且检查执行期间发生的错误
    gettimeofday(&end,NULL);
    float time_use;
    time_use=(end.tv_sec-start.tv_sec)*1000000+(end.tv_usec-start.tv_usec);//微秒
    std::cout <<"vector_add gpu time cost is "<<time_use/1000/100<< " ms"<< std::endl;


    for(int i=0;i<M;i++){
        printf("%f ",a[i]);
    }
    printf("\n");
    printf("******************\n");
    for(int i=0;i<M;i++){
        printf("%f ",b[i]);
    }
    printf("\n");
    printf("******************\n");
    for(int i=0;i<M;i++){
        printf("%f ",c[i]);
    }
    printf("\n");
    printf("******************\n");
    return 0;
}

2、代码运行和结果

以上是一个.cu单文件,怎么运行它呢?需要使用英伟达的cuda环境进行编译和链接。在安装cuda toolkit的环境下,使用nvcc编译器进行编译:

//编译和链接
nvcc nvcc_vector_add.cu -o nvcc_vector_add

//运行
./nvcc_vector_add

参数-o表示生成可执行文件,编译结果如下:

编译后生成一个可执行文件,直接运行该可执行文件,得到如下结果:

可以看到结果正确,耗时为2ms,显存使用384M。

调整一下blockDim、gridDim的大小看看耗时的变化情况。

blockDim=256,gridDim=5

vector_add gpu time cost is 2.23614 ms

blockDim=256,gridDim=10

vector_add gpu time cost is 2.01713 ms

blockDim=256,gridDim=20

vector_add gpu time cost is 2.05501 ms

blockDim=128,gridDim=10

vector_add gpu time cost is 2.1049 ms

blockDim=512,gridDim=10

vector_add gpu time cost is 2.027 ms

以上结果具有误差因为都只跑了一次,没有多次求平均值,但是可以说明gridDim和blockDim对性能是有影响的。一般来说blockDim选择为32的倍数,因为一个wrap的线程束是32,blockDim这样设置可以减少bank conflict。

CUDA编程入门极简教程

Cuda Core VS Tensor Core

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:/a/162808.html

如若内容造成侵权/违法违规/事实不符,请联系我们进行投诉反馈qq邮箱809451989@qq.com,一经查实,立即删除!

相关文章

2023年最新软件安装管家目录

最新软件安装管家目录 软件目录 ①【电脑办公】电脑系统&#xff08;直接安装&#xff09;Win7Win8Win10OfficeOffice激活office2003office2007office2010office2013office2016office2019office365office2021wps2021Projectproject2007project2010project2016project2019projec…

python基础练习题库实验3

题目1 编写一个程序&#xff0c;根据以下定价计算成本。 Number of itemsCost1-50每件3美元 邮费: 10美元超过50每件2美元 邮寄&#xff1a;免费 举个例子&#xff1a; 代码 items_num input("Enter the number of items: ") items_num_i int(items_num) ite…

JVM虚拟机:通过日志学习PS+PO垃圾回收器

我们刚才设置参数的时候看到了-XXPrintGCDetails表示输出详细的GC处理日志&#xff0c;那么我们如何理解这个日志呢&#xff1f;日志是有规则的&#xff0c;我们需要按照这个规则来理解日志中的内容&#xff0c;它有两个格式&#xff0c;一个格式是GC的格式&#xff08;新生代&…

Leetcode—206.反转链表【简单】

2023每日刷题&#xff08;三十三&#xff09; Leetcode—206.反转链表 头插法实现代码 /*** Definition for singly-linked list.* struct ListNode {* int val;* struct ListNode *next;* };*/ struct ListNode* reverseList(struct ListNode* head) {if(head NULL…

PaddlePaddle:开源深度学习平台

深度学习作为人工智能领域的重要分支&#xff0c;正在全球范围内得到广泛应用。而在构建和训练深度学习模型时&#xff0c;选择一个高效、易用且功能强大的开源平台是至关重要的。PaddlePaddle&#xff08;即飞桨&#xff09;作为国内领先的深度学习平台&#xff0c;一直以来都…

13.真刀实枪做项目---博客系统(页面设计)

文章目录 1.预期效果1.1博客列表页效果1.2博客详情页效果1.3博客登陆页效果1.4博客编辑页效果 2.实现博客列表页2.1实现导航栏2.2实现版心2.3实现个人信息2.4实现博客列表2.5博客列表页完整代码 3.实现博客正文页3.1引入导航栏3.2引入版心3.3引入个人信息3.4实现博客正文3.5博客…

C++17中std::variant的使用

可变参数模板类std::variant表示类型安全联合体(type-safe union)。std::variant的实例在任何给定时间要么保存其替代类型之一的值&#xff0c;要么在错误的情况下无值。 与union一样&#xff0c;如果std::variant保存某个对象类型T的值&#xff0c;则T的对象表示形式将直…

【音视频基础】AVI文件格式

AVI文件采用的是RIFF文件结构方式。波形音频wave&#xff0c;MIDI和数字视频AVI都采用这种格式存储。 AVI文件的整体结构如下图所示 构造RIFF文件的基本单元叫做数据块&#xff08;Chunk&#xff09;&#xff0c;每个数据块包含3个部分 4字节的数据块标记&#xff08;或者叫…

动 态 规 划

一、&#xff08;what&#xff1f;&#xff09; 二、&#xff08;why&#xff1f;&#xff09; 三、&#xff08;how&#xff1f;&#xff09; 四、典型例题分析&#xff1a; 例题1&#xff1a;神奇的兔子序列 输入&#xff1a;月数 输出&#xff1a;兔子数 思路&#xff1…

基于机器学习的居民消费影响因子分析预测

项目视频讲解: 基于机器学习的居民消费影响因子分析预测_哔哩哔哩_bilibili 主要工作内容: 完整代码: import pandas as pd import numpy as np import matplotlib.pyplot as plt import seaborn as sns import missingno as msno import warnings warnings.filterwarnin…

Python------列表 集合 字典 推导式(本文以 集合为主)

推导式&#xff1a; 推导式comprehensions&#xff08;又称解析式&#xff09;&#xff0c;是Python的一种独有特性。推导式是可以从一个数据序列 构建 另一个 新的数据序列&#xff08;一个有规律的列表或控制一个有规律列表&#xff09;的结构体。 共有三种推导&#xff…

二十三种设计模式全面解析-当你的对象需要知道其他对象的状态变化时,观察者模式是你的救星!

在软件设计的世界中&#xff0c;有一种设计模式以其简洁而强大的特性闪耀着光芒&#xff0c;它就是——观察者模式&#xff08;Observer Pattern&#xff09;。这个模式它定义了一种一对多的依赖关系&#xff0c;让多个观察者对象同时监听某一个主题对象&#xff0c;为我们创造…

【spring】ApplicationContext的实现

目录 一、ClassPathXmlApplicationContext1.1 说明1.2 代码示例1.3 截图示例 二、FileSystemXmlApplicationContext2.1 说明2.2 代码示例2.3 加载xml的bean定义示例 三、AnnotationConfigApplicationContext3.1 说明3.2 代码示例3.3 截图示例 四、AnnotationConfigServletWebSe…

Git安装与常用命令

Git简介&#xff1a; Git是一个开源的分布式版本控制系统&#xff0c;用于敏捷高效地处理任何或大或小的项目。Git是Linus Torvalds为了帮助管理Linux内核开发而开发的一个开放源代码的版本控制软件。Git与常用的版本控制工具CVS、Subversion等不同&#xff0c;它采用了分布式…

CF1899 G. Unusual Entertainment [二维数点/二维偏序]

传送门:CF [前题提要]:没什么好说的,区域赛爆炸之后发愤加训思维题.秒了div3 A~F的脑筋急转弯,然后被G卡了,树剖dfs序的想法已经想到了,题目也已经化简为两个线段是否存在一个合法位置了.但是MD不会二维数点,用一个树剖扫描线搞来搞去最后还是Tle.果然如下图所说:科技还是十分…

Netty传输object并解决粘包拆包问题

⭐️ 前言 大家好&#xff0c;笔者之前写过一篇文章&#xff0c;《Netty中粘包拆包问题解决探讨》&#xff0c;就Netty粘包拆包问题及其解决方案进行了探讨&#xff0c;本文算是这篇博客的延续。探讨netty传输object的问题。 本文将netty结合java序列化来传输object并解决粘包…

3dMax2024新功能和工作流增强功能速览

3dMax2024新功能和工作流增强功能速览 Autodesk发布的3dMax2024引入了一系列新功能和工作流增强功能&#xff0c;如下所示&#xff1a; 更新的“指定控制器”卷展栏&#xff1a;这个现代化的功能为动画师提供了更高效的工作方式&#xff0c;简化了他们的动画流程。 布尔修饰符…

【DevOps】Git 图文详解(三):常用的 Git GUI

Git 图文详解&#xff08;三&#xff09;&#xff1a;常用的 Git GUI 1.SourceTree2.TortoiseGit3.VSCode 中的 Git 如果不想用命令行工具&#xff0c;完全可以安装一个 Git 的 GUI 工具&#xff0c;用的更简单、更舒服。不用记那么多命令了&#xff0c;极易上手&#xff0c;不…

二、ST-Link驱动的安装

1、灵动mm32单片机 (1)上海灵动微电子股份有限公司 (2)mm32单片机支持ST-Link下载程序。 2、ST-Link驱动的安装 (1)下载地址 ST-Link 官网下载地址 (2)点击获取软件下载ST-Link驱动。(需要登陆ST官网账户) (3)下载后解压&#xff0c;根据电脑位数安装 .exe 文件即可。 6…