GPU并行计算CUDA

一、CUDA 和 GPU 简介

  • CUDA 是显卡厂商 NVIDIA 推出的运算平台,是一种通用并行计算架构,使得 GPU 能够解决复杂的计算问题。
  • 开发人员可以使用 C 语言来为 CUDA 架构编写程序,可以在支持 CUDA 的处理器上以超高性能运行,CUDA 3.0 已经开始支持 C++。
  • GPU 是图形处理器,显卡的处理核心
  • 电脑显示器上显示的图像,在显示在显示器上之前,要经过一系列的图形计算,这个过程叫渲染,针对图形计算的这些操作设计了一种处理器,也就是 GPU。

二、GPU 工作原理与结构

  • GPU 采用流式并行计算模式,可对每个数据行进行独立的并行计算。
  • GPU 和 CPU 的区别:
    • CPU 基于低延时设计,由运算器(ALU:算数逻辑单元)和控制器(CU),以及若干个寄存器和高速缓冲存储器组成,功能模块较多,擅长逻辑控制,串行运算。
    • GPU 基于大吞吐量设计,拥有更多的 ALU 用于数据处理,适合对密集数据进行并行处理,擅长大规模并行计算。
    • GPU 为图形图像专门设计,在矩阵运算,数值计算方面具有独特优势,特别是浮点和并行计算上远远优于 CPU,GPU 的优势在于快。

在这里插入图片描述


三、GPU 编程模型

  • 异构计算 = CPU + GPU
  • 一个 GPU 包含多个 SM(Streaming Multiprocessor),而每个 SM 又包含多个 core
  • 一个 block 只能调度到一个 SM 上运行,直到 block 运行完毕。
  • 一个 SM 可以同时运行多个 block,因为有多个 core
  • 每个 block 以 warp(一般为 32 个线程或 64 个线程)作为一次执行的单位(真正的同时执行)
    • 在具体的硬件执行中,一个 core 会同时执行 warp,一个 block 会被绑定到一个 core 上,即使这个 block 内部可能有 1024 个线程,这些线程组会被相应的调度器来进行调度,在逻辑层面上可以认为 1024 个线程同时执行,但在硬件上是 warp 同时执行,这一点其实和操作系统的线程调度是一样的
    • 假如一个 core 同时能执行 64 个线程,但一个 block 有 1024 个线程,那这 1024 个线程会分 16 次执行。
  • 显存层面:一个 block 内的 thread 共享一块 share memory(一般是 SM 的一级缓存)。GPU 和 CPU 一样有着多级 cache寄存器的架构,把全局显存的数据加载到共享显存上再处理可以有效地加速。

在这里插入图片描述


四、Grid、Block、Thread 的关系

在这里插入图片描述

  • CUDA 中线程分成三个层次:线程、线程块、线程网格。
    • 线程:CUDA 中的基本执行单元,由硬件支持、开销很小,每个线程执行相同的代码
    • 线程块(Block):若干线程的分组,Block 内的一个块至多 512 个线程,或 1024 个线程(根据不同的 GPU 规格),线程块可以是一维二维或者三维的,同一个 block 中的 threads 可以同步,也可以通过 share memory 通信。
    • 线程网格(Grid):若干线程块的网格。
  • CUDA 中每一个线程都有一个唯一的标识 IDThreadIdx
    • threadIdx 是一个 uint3 类型,表示一个线程的索引。
    • blockIdx 是一个 uint3 类型,表示一个线程块的索引,一个线程块中通常有多个线程。
    • blockDim 是一个 dim3 类型,表示线程块的大小。
    • gridDim 是一个 dim3 类型,表示网格的大小,一个网格中通常有多个线程块。
  • grid 划分成 1 维,block 划分为 1 维:
    int threadId = blockIdx.x * blockDim.x + threadIdx.x;
    
  • grid 划分成 1 维,block 划分为 2 维:
    int threadId = blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
    
  • grid 划分成 1 维,block 划分为 3 维:
    int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z+ threadIdx.z * blockDim.y * blockDim.x+ threadIdx.y * blockDim.x + threadIdx.x;
    
  • grid 划分成 2 维,block 划分为 1 维:
    int blockId = blockIdx.y * gridDim.x + blockIdx.x;
    int threadId = blockId * blockDim.x + threadIdx.x;
    
  • grid 划分成 2 维,block 划分为 2 维:
    int blockId = blockIdx.x + blockIdx.y * gridDim.x;
    int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;
    
  • grid 划分成 2 维,block 划分为 3 维:
    int blockId = blockIdx.x + blockIdx.y * gridDim.x;
    int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)+ (threadIdx.z * (blockDim.x * blockDim.y))+ (threadIdx.y * blockDim.x) + threadIdx.x;
    
  • grid 划分成 3 维,block 划分为 1 维:
    int blockId = blockIdx.x + blockIdx.y * gridDim.x+ gridDim.x * gridDim.y * blockIdx.z;
    int threadId = blockId * blockDim.x + threadIdx.x;
    
  • grid 划分成 3 维,block 划分为 2 维:
    int blockId = blockIdx.x + blockIdx.y * gridDim.x+ gridDim.x * gridDim.y * blockIdx.z;
    int threadId = blockId * (blockDim.x * blockDim.y)+ (threadIdx.y * blockDim.x) + threadIdx.x;
    
  • grid 划分成 3 维,block 划分为 3 维
    int blockId = blockIdx.x + blockIdx.y * gridDim.x+ gridDim.x * gridDim.y * blockIdx.z;
    int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)+ (threadIdx.z * (blockDim.x * blockDim.y))+ (threadIdx.y * blockDim.x) + threadIdx.x;
    

五、CUDA 程序结构

  • CUDA 程序的结构大体是:{主机串行 → GPU 并行} + → 主机串行,这样的串并交叉结构。
  • 主机串行过渡到 GPU 并行时需要将数据从主机内存上拷贝到 GPU 设备显存上,GPU 执行完毕时也需要把数据拷贝回来。

六、CUDA 核函数和配置

  • 主机调用设备代码的唯一接口就是 kernel 函数,使用限定词:__global__
  • 使用核函数需要在核函数名后面加 <<<>>> 指定核函数配置。
  • <<<>>> 运算符完整的执行配置参数形式是 <<< Dg, Db, Ns, S>>>
    • 参数 Dg:
      • 用于定义整个 grid 的维度和尺寸,即一个 grid 有多少个 block。
      • Dg 为 dim3 类型。
      • Dim3 Dg(Dg.x, Dg.y, 1) 表示 grid 中每行有 Dg.x 个 block、每列有 Dg.y 个 block、并且只有一个 grid。
      • 整个 grid 中共有 Dg.x * Dg.y 个 block,其中 Dg.xDg.y 最大值为 65535
    • 参数 Db:
      • 用于定义一个 block 的维度和尺寸,即一个 block 有多少个 thread。
      • Db 为 dim3 类型。
      • Dim3 Db(Db.x, Db.y, Db.z) 表示整个 block 中每行有 Db.x 个 thread,每列有 Db.y 个 thread,高度为 Db.z
      • Db.xDb.y 最大值为 512Db.z 最大值为 62
      • 一个 block 中共有 Db.x * Db.y * Db.z 个 thread。
    • 参数 Ns 是一个可选参数,用于设置每个 block 除了静态分配的 shared Memory 以外,最多能动态分配的 shared memory大小,单位为 byte。不需要动态分配时该值为 0 或省略不写。
    • 参数 S 是一个 cudaStream_t 类型的可选参数,初始值为零,表示该核函数处在哪个流之中。
  • <<<DimGrid, DimBlock>>> 指定线程网格和线程块维度,若当前硬件无法满足用户配置,则核函数不会被执行,直接返回错误。

七、CUDA 限定符

  • 函数限定符。

    函数限定符在何处执行从何处调用特性
    __device__设备设备函数的地址无法获取
    __global__设备主机返回类型必须为空
    __host__主机主机等同于不使用任何限定符
  • 变量限定符。

    变量限定符位于何处可以访问的线程主机访问
    __device__全局存储器线程网格内的所有线程通过运行时库访问
    __constant__固定存储器线程网格内的所有线程通过运行时库访问
    __shared__共享存储器线程块内的所有线程不可从主机访问

八、同步

  • CPU 启动 kernel 函数是异步的,它并不会阻塞等到 GPU 执行完 kernel 函数才执行后面的 CPU 部
    分。
  • 因此如果后续程序立即需要用到上一个 kernel 函数的结果我们需要显式设置同步障来阻塞 CPU 程序。
  • 一个线程块内需要同步共享存储器的共享变量 __shared__时,需要在使用前显式调用 __syncthreads() 同步块内所有线程。
  • 同一个 Grid 中不同 Block 之间无法设置同步。

九、CUDA 运行时 API

  • cudaMemcpy:用于在主机和设备之间拷贝数据,其中 cudaMemcpyKind 枚举类型常用的有:cudaMemcpyHostToDevice 表示把主机数据拷贝到显存,以及逆向的 cudaMemcpyDeviceToHostcudaMemcpyHostToHostcudaMemcpyDeviceToDevice
    __host__ cudaError_t cudaMemcpy( void* dst, const void* src, size_t count, cudaMemcpyKind kind)
    
  • cudaMalloc:在设备上分配动态显存,两个限定符表示可以在主机或设备上调用。
    __host__ __device__ cudaError_t cudaMalloc( void** devPtr, size_t size )
    
  • cudaFree:释放回收在设备上分配的动态显存。
    __host__ __device__ cudaError_t cudaFree( void* devPtr )
    
  • cudaThreadSynchronize:等待 GPU 代码运行结束。

十、CUDA 向量加法

cmake_minimum_required(VERSION 3.16.3)project(CUDATest LANGUAGES CXX CUDA)add_executable(vector_add vector_add.cu)
#include <iostream>
#include <random>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"void initDevice(int devNum) {int dev = devNum;cudaDeviceProp deviceProp;cudaGetDeviceProperties(&deviceProp, dev);printf("Using device: %d: %s \n", dev, deviceProp.name);cudaSetDevice(dev);
}// 全局随机数引擎
std::default_random_engine engine(static_cast<unsigned>(time(nullptr)));
std::uniform_real_distribution<float> distribution(0.0, 100.0);
void initData(float* arr, int nElem) {for (int i = 0; i < nElem; i++) {arr[i] = distribution(engine);}
}__global__ void vector_add(float* arr1, float * arr2, float* res, int n) {int i = blockIdx.x * blockDim.x + threadIdx.x;if (i >= n) return;res[i] = arr1[i] + arr2[i];
}void vector_add_cpu(float* arr1, float * arr2, float* res, int n) {for (int i = 0;i < n;i++) {res[i] = arr1[i] + arr2[i];}   
}void checkRes(float* cpuRes, float* gpuRes, const int n) {double epsilon = 1.0E-8;for (int i = 0;i < n;i++) {if (abs(cpuRes[i] - gpuRes[i]) > epsilon) {printf("Result don't match\n");return;} }printf("match!\n");
}int main() {initDevice(0);int nElem = 2048 * 2048;int nByte = sizeof(float) * nElem;// 主机内存float* h_arr1 = (float*)malloc(nByte);float* h_arr2 = (float*)malloc(nByte);float* h_res = (float*)malloc(nByte);float* res = (float*)malloc(nByte);// 初始化数据initData(h_arr1, nElem);initData(h_arr2, nElem);memset(h_res, 0, nByte);memset(res, 0, nByte);// GPU 内存申请float* d_arr1;float* d_arr2;float* d_res;// 设备显存cudaMalloc(&d_arr1, nByte);cudaMalloc(&d_arr2, nByte);cudaMalloc(&d_res, nByte);// 数据拷贝cudaMemcpy(d_arr1, h_arr1, nByte, cudaMemcpyHostToDevice);cudaMemcpy(d_arr2, h_arr2, nByte, cudaMemcpyHostToDevice);// 执行向量加法int threads = 32;int blocks = (nElem + threads - 1) / threads; vector_add<<<blocks, threads>>>(d_arr1, d_arr2, d_res, nElem);// 数据拷贝cudaMemcpy(h_res, d_res, nByte, cudaMemcpyDeviceToHost);// cpu 计算vector_add_cpu(h_arr1, h_arr2, res, nElem);cudaThreadSynchronize();// 对比结果checkRes(h_res, res, nElem);// 释放显存cudaFree(d_arr1);cudaFree(d_arr2);cudaFree(d_res);// 释放内存free(h_arr1);free(h_arr2);free(h_res);free(res);return 0;
}

十一、CUDA 矩阵乘法

#include <iostream>
#include <random>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"void initDevice(int devNum) {int dev = devNum;cudaDeviceProp deviceProp;cudaGetDeviceProperties(&deviceProp, dev);printf("Using device: %d: %s \n", dev, deviceProp.name);cudaSetDevice(dev);
}// 全局随机数引擎
std::default_random_engine engine(static_cast<unsigned>(time(nullptr)));
std::uniform_real_distribution<float> distribution(0.0, 100.0);
void initData(float* arr, int nElem) {for (int i = 0; i < nElem; i++) {arr[i] = distribution(engine);}
}__global__ void matrix_mul(float* A, float * B, float* C, int n) {int col = blockIdx.x * blockDim.x + threadIdx.x;int row = blockIdx.y * blockDim.y + threadIdx.y;if (row >= n || col >= n) return;float val = 0.0;for (int i = 0;i < n;i++) {val += A[row * n + i] * B[i * n + col];}C[row * n + col] = val;
}void matrix_mul_cpu(float* A, float * B, float* C, int n) {int sum = 0;for (int i = 0;i < n;i++) {for (int j = 0;j < n;j++) {for (int k = 0;k < n;k++) {sum += A[i * n + k] * B[k * n + j];}C[i * n + j] = sum;sum = 0;}}
}void checkRes(float* cpuRes, float* gpuRes, const int nElem) {double epsilon = 1.0E-8;for (int i = 0;i < nElem;i++) {if (abs(cpuRes[i] - gpuRes[i] > epsilon)) {printf("Result don't match\n");return;} }printf("match!\n");
}int main() {initDevice(0);int n = 64;int nElem = n * n;int nByte = sizeof(float) * nElem;// 主机内存float* h_A = (float*)malloc(nByte);float* h_B = (float*)malloc(nByte);float* h_C = (float*)malloc(nByte);float* C = (float*)malloc(nByte);// 初始化数据initData(h_A, nElem);initData(h_B, nElem);memset(h_C, 0, nByte);memset(C, 0, nByte);// GPU 内存申请float* d_A;float* d_B;float* d_C;// 设备显存cudaMalloc(&d_A, nByte);cudaMalloc(&d_B, nByte);cudaMalloc(&d_C, nByte);// 数据拷贝cudaMemcpy(d_A, h_A, nByte, cudaMemcpyHostToDevice);cudaMemcpy(d_B, h_B, nByte, cudaMemcpyHostToDevice);// 执行矩阵乘法dim3 threads(32, 32);int blockX = (threads.x + n - 1) / threads.x;int blockY = (threads.y + n - 1) / threads.y;dim3 blocks(blockX, blockY);matrix_mul<<<blocks, threads>>>(d_A, d_B, d_C, n);// 数据拷贝cudaMemcpy(h_C, d_C, nByte, cudaMemcpyDeviceToHost);// cpu 计算matrix_mul_cpu(h_A, h_B, C, n);cudaThreadSynchronize();// 对比结果checkRes(C, h_C, nElem);cudaFree(d_A);cudaFree(d_B);cudaFree(d_C);free(h_C);free(C);return 0;
}

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

如若内容造成侵权/违法违规/事实不符,请联系长河编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

【Docker学习】docker version查看版本信息

就像很多应用一样&#xff0c;docker也使用version来查看版本信息。但因为docker包含有不少独立组件&#xff0c;version的作用范围会更广一些。 用法1&#xff1a; docker --version 描述&#xff1a; 输出安装的Docker CLI 的版本号。关于Docker CLI&#xff0c;请访问。 实操…

ROS2专栏(三) | 理解ROS2的动作

​ 1. 创建一个动作 目标&#xff1a; 在ROS 2软件包中定义一个动作。 1.1 新建包 设置一个 workspace 并创建一个名为 action_tutorials_interfaces 的包&#xff1a; mkdir -p ros2_ws/src #you can reuse existing workspace with this naming convention cd ros2_ws/s…

【C++】:类和对象(下)

目录 一&#xff0c;再谈构造函数1.初始化列表2. 隐式类型转换的过程及其优化3. 隐式类型转换的使用4. explcit关键字5. 单参数和多参数构造函数的隐式类型转换 二&#xff0c;static成员1.静态成员变量2.静态成员函数3. static 成员的应用 三&#xff0c;友元3.1 友元函数3.2 …

SQL注入漏洞扫描---sqlmap

what SQLMap是一款先进的自动执行SQL注入的审计工具。当给定一个URL时&#xff0c;SQLMap会执行以下操作&#xff1a; 判断可注入的参数。判断可以用哪种SQL注入技术来注入。识别出目标使用哪种数据库。根据用户的选择&#xff0c;读取哪些数据库中的数据。 更详细语法请参考…

Mac 安装 JDK21 流程

一、下载JDK21 访问Oracle官方网站或选择OpenJDK作为替代品。Oracle JDK从11版本开始是商业的&#xff0c;可能需要支付费用。OpenJDK是一个免费开源选项。 Oracle JDK官方网站&#xff1a;Oracle JDK Downloads OpenJDK官方网站&#xff1a;OpenJDK Downloads 这里以JDK21为…

Nginx实现端口转发与负载均衡配置

前言&#xff1a;当我们的软件体系结构较为庞大的时候&#xff0c;访问量往往是巨大的&#xff0c;所以我们这里可以使用nginx的均衡负载 一、配置nginx实现端口转发 本地tomcat服务端口为8082 本地nginx端口为8080 目的&#xff1a;将nginx的8080转发到tomcat的8082端口上…

如何从Mac电脑恢复任何删除的视频

Microsoft Office是包括Mac用户在内的人们在世界各地创建文档时使用的最佳软件之一。该软件允许您创建任何类型的文件&#xff0c;如演示文稿、帐户文件和书面文件。您可以使用 MS Office 来完成。所有Microsoft文档都可以在Mac上使用。大多数情况下&#xff0c;您处理文档&…

网络安全审计

一、什么叫网络安全审计 网络安全审计是按照一定的安全策略&#xff0c;利用记录、系统活动和用户活动等信息&#xff0c;检查、审查和检验操作时间的环境及活动&#xff0c;从而发现系统漏洞、入侵行为或改善系统性能的过程&#xff0c;它是提高系统安全性的重要手段。 系统…

springboot整合mybatis配置多数据源(mysql/oracle)

目录 前言导入依赖坐标创建mysql/oracle数据源配置类MySQLDataSourceConfigOracleDataSourceConfig application.yml配置文件配置mysql/oracle数据源编写Mapper接口编写Book实体类编写测试类 前言 springboot整合mybatis配置多数据源&#xff0c;可以都是mysql数据源&#xff…

2024 年 16 款最佳屏幕录制软件榜单

屏幕录制软件是一个经常被忽视的方便工具。 无论您是想要制作教程的内容创建者还是想要录制演示文稿的业务人员&#xff0c;屏幕录像机都是必须的&#xff01;但是&#xff0c;在使用您能找到的第一台录音机之前&#xff0c;我们建议您花些时间找到符合您需求的录音机。有这么…

pytorch笔记:ModuleDict

1 介绍 在 PyTorch 中&#xff0c;nn.ModuleDict 是一个方便的容器&#xff0c;用于存储一组子模块&#xff08;即 nn.Module 对象&#xff09;的字典这个容器主要用于动态地管理多个模块&#xff0c;并通过键来访问它们&#xff0c;类似于 Python 的字典 2 特点 组织性 nn…

【Unity】在空物体上实现 IPointerClickHandler 不起作用

感谢Unity接口IPointerClickHandler使用说明_哔哩哔哩_bilibiliUnity接口IPointerClickHandler使用说明, 视频播放量 197、弹幕量 0、点赞数 3、投硬币枚数 2、收藏人数 2、转发人数 0, 视频作者 游戏创作大陆, 作者简介 &#xff0c;相关视频&#xff1a;在Unity多场景同时编辑…

嵌入式开发四:STM32 基础知识入门

为方便更好的学习STM32单片机&#xff0c;本篇博客主要总结STM32的入门基础知识&#xff0c;重点在于理解寄存器以及存储器映射和寄存器映射&#xff0c;深刻体会STM32是如何组织和管理庞大的寄存器&#xff0c;从而提高开发效率的&#xff0c;为后面的基于标准库的开发做好铺垫…

专业渗透测试 Phpsploit-Framework(PSF)框架软件小白入门教程(四)

本系列课程&#xff0c;将重点讲解Phpsploit-Framework框架软件的基础使用&#xff01; 本文章仅提供学习&#xff0c;切勿将其用于不法手段&#xff01; 继续接上一篇文章内容&#xff0c;讲述如何进行Phpsploit-Framework软件的基础使用和二次开发。 当我们牢记登陆账户、…

浅谈ps/2键盘

文章目录 说明基础知识操作系统中断类型工作机制优点应用 CPU对IO设备的轮询机制轮询机制的工作原理轮询机制的特点轮询机制的优、缺点与中断机制的对比 N-Key Roller&#xff08;全键无冲&#xff09;应用领域实现原理技术限制 PS/2接口简介USB设备&PS/2设备的工作机制PS/…

11个2024年热门的AI编码助手

大家好&#xff0c;人工智能&#xff08;AI&#xff09;领域的大型语言模型&#xff08;LLMs&#xff09;已经逐渐发展成熟&#xff0c;并且深入到了我们日常的工作当中。在众多AI应用中&#xff0c;编码助手尤为突出&#xff0c;是开发人员编写更高效、准确无误代码的必备辅助…

预训练模型介绍

一、什么是GPT GPT 是由人工智能研究实验室 OpenAI 在2022年11月30日发布的全新聊天机器人模型, 一款人工智能技术驱动的自然语言处理工具 它能够通过学习和理解人类的语言来进行对话, 还能根据聊天的上下文进行互动,能完成撰写邮件、视频脚本、文案、翻译、代码等任务 二、 为…

STM32单片机实战开发笔记-EXIT外部中断检测

嵌入式单片机开发实战例程合集&#xff1a; 链接&#xff1a;https://pan.baidu.com/s/11av8rV45dtHO0EHf8e_Q0Q?pwd28ab 提取码&#xff1a;28ab EXIT模块测试 功能描述 外部中断/事件控制器由19个产生事件/中断要求的边沿检测器组成。每个输入线可以独立地配置输入类型&a…

软件测试(实验五)——Jmeter的使用

目录 实验目的 一、使用JMeter演示取样器、监听器、配置元件、断言的使用&#xff1b; 1、取样器 2、监听器 3、配置元件的使用 ① 用户定义的变量 ②HTTP信息头管理器 ③HTTP请求默认值 ④CSV数据文件设置 4、断言 ①响应断言 ②JSON断言 ③断言持续时间 二、使用…

设计模式之代理模式ProxyPattern(六)

一、代理模式介绍 1、什么是代理模式&#xff1f; 代理模式是一种结构型设计模式&#xff0c;它允许为其他对象提供一个替代品或占位符&#xff0c;以控制对这个对象的访问。 2、代理模式的角色构成 抽象主题&#xff08;Subject&#xff09;&#xff1a;定义了真实主题和代…