更简单地介绍 CUDA

在这里插入图片描述这篇文章是对 CUDA 的超级简单介绍,CUDA 是 NVIDIA 流行的并行计算平台和编程模型。我之前在2013年写过一篇文章《CUDA简单介绍》,多年来一直很受欢迎。但 CUDA 编程变得更加容易,GPU 也变得更快,所以是时候进行更新(甚至更简单)的介绍了。
CUDA C++ 只是使用 CUDA 创建大规模并行应用程序的方法之一。它允许您使用强大的 C++ 编程语言来开发由 GPU 上运行的数千个并行线程加速的高性能算法。许多开发人员通过这种方式加速了计算和带宽需求大的应用程序,包括支持正在进行的人工智能革命(称为深度学习)的库和框架。
因此,您已经听说过 CUDA,并且有兴趣学习如何在自己的应用程序中使用它。如果您是 C 或 C++ 程序员,这篇博文应该会给您一个良好的开端。要继续进行操作,您需要一台具有支持 CUDA 的 GPU(Windows、Mac 或 Linux,以及任何 NVIDIA GPU 都可以)的计算机,或者具有 GPU 的云实例(AWS、Azure、IBM SoftLayer 和其他云服务)供应商有它们)。您还需要安装免费的 CUDA 工具包。您还可以使用在云中的 GPU 上运行的 Jupyter Notebook 进行操作。
让我们开始吧!!!
在这里插入图片描述

Starting Simple

我们将从一个简单的 C++ 程序开始,该程序将两个数组的元素相加,每个数组有 100 万个元素。

#include <iostream>
#include <math.h>// function to add the elements of two arrays
void add(int n, float *x, float *y)
{for (int i = 0; i < n; i++)y[i] = x[i] + y[i];
}int main(void)
{int N = 1<<20; // 1M elementsfloat *x = new float[N];float *y = new float[N];// initialize x and y arrays on the hostfor (int i = 0; i < N; i++) {x[i] = 1.0f;y[i] = 2.0f;}// Run kernel on 1M elements on the CPUadd(N, x, y);// Check for errors (all values should be 3.0f)float maxError = 0.0f;for (int i = 0; i < N; i++)maxError = fmax(maxError, fabs(y[i]-3.0f));std::cout << "Max error: " << maxError << std::endl;// Free memorydelete [] x;delete [] y;return 0;
}

首先,编译并运行这个 C++ 程序。将上面的代码放入一个文件中并将其另存为 add.cpp,然后使用 C++ 编译器进行编译。我使用的是 Mac,所以我使用 clang++,但您可以在 Linux 上使用 g++,或者在 Windows 上使用 MSVC。

clang++ add.cpp -o add

然后运行它:

> ./add
Max error: 0.000000

(在 Windows 上,您可能需要将可执行文件命名为 add.exe 并使用 .\add 运行它。)
正如所料,它打印出求和没有错误,然后退出。现在我想让这个计算在 GPU 的多个核心上(并行)运行。事实上,迈出第一步非常容易。
首先,我只需要把我们的add函数变成GPU可以运行的函数,在CUDA中称为内核。为此,我所要做的就是向函数添加说明符 global ,它告诉 CUDA C++ 编译器这是一个在 GPU 上运行的函数,可以从 CPU 代码中调用。

// CUDA Kernel function to add the elements of two arrays on the GPU
__global__
void add(int n, float *x, float *y)
{for (int i = 0; i < n; i++)y[i] = x[i] + y[i];
}

这些 global 函数称为内核,在 GPU 上运行的代码通常称为设备代码,而在 CPU 上运行的代码称为主机代码。

Memory Allocation in CUDA

为了在 GPU 上计算,我需要分配 GPU 可访问的内存。 CUDA 中的统一内存通过提供可供系统中所有 GPU 和 CPU 访问的单一内存空间,使这一过程变得简单。要在统一内存中分配数据,请调用 cudaMallocManaged(),它返回一个可以从主机 (CPU) 代码或设备 (GPU) 代码访问的指针。要释放数据,只需将指针传递给 cudaFree()。
我只需要将上面代码中对 new 的调用替换为对 cudaMallocManaged() 的调用,并将对 delete [] 的调用替换为对 cudaFree 的调用。

  float *x, *y;cudaMallocManaged(&x, N*sizeof(float));cudaMallocManaged(&y, N*sizeof(float));...// Free memorycudaFree(x);cudaFree(y);

最后,我需要启动 add() 内核,它会在 GPU 上调用它。 CUDA 内核启动是使用三尖括号语法 <<< >>> 指定的。我只需将它添加到参数列表之前的调用中即可。

add<<<1, 1>>>(N, x, y);

简单的!我很快就会详细介绍尖括号内的内容;现在您需要知道的是这一行启动一个 GPU 线程来运行 add()。
还有一件事:我需要 CPU 等待内核完成后再访问结果(因为 CUDA 内核启动不会阻塞调用 CPU 线程)。为此,我只需在对 CPU 进行最终错误检查之前调用 cudaDeviceSynchronize() 即可。
这是完整的代码:

#include <iostream>
#include <math.h>
// Kernel function to add the elements of two arrays
__global__
void add(int n, float *x, float *y)
{for (int i = 0; i < n; i++)y[i] = x[i] + y[i];
}int main(void)
{int N = 1<<20;float *x, *y;// Allocate Unified Memory – accessible from CPU or GPUcudaMallocManaged(&x, N*sizeof(float));cudaMallocManaged(&y, N*sizeof(float));// initialize x and y arrays on the hostfor (int i = 0; i < N; i++) {x[i] = 1.0f;y[i] = 2.0f;}// Run kernel on 1M elements on the GPUadd<<<1, 1>>>(N, x, y);// Wait for GPU to finish before accessing on hostcudaDeviceSynchronize();// Check for errors (all values should be 3.0f)float maxError = 0.0f;for (int i = 0; i < N; i++)maxError = fmax(maxError, fabs(y[i]-3.0f));std::cout << "Max error: " << maxError << std::endl;// Free memorycudaFree(x);cudaFree(y);return 0;
}

CUDA 文件的文件扩展名为 .cu。因此,将此代码保存在名为 add.cu 的文件中,并使用 CUDA C++ 编译器 nvcc 进行编译。

> nvcc add.cu -o add_cuda
> ./add_cuda
Max error: 0.000000

这只是第一步,因为正如所写的,该内核仅对于单个线程是正确的,因为运行它的每个线程都会对整个数组执行添加操作。此外,由于多个并行线程都会读取和写入相同的位置,因此存在竞争条件。
注意:在 Windows 上,您需要确保在 Microsoft Visual Studio 项目的配置属性中将平台设置为 x64。

Profile it!

我认为了解内核运行时间的最简单方法是使用 nvprof(CUDA 工具包附带的命令行 GPU 分析器)运行它。只需在命令行中输入 nvprof ./add_cuda:

上面是 nvprof 的截断输出,显示了对 add 的单个调用。在 NVIDIA Tesla K80 加速器上大约需要半秒,在我用了 3 年的 Macbook Pro 上的 NVIDIA GeForce GT 740M 上大约需要同样的时间。
让我们通过并行性使其更快。

Picking up the Threads

现在您已经运行了一个带有一个线程来执行一些计算的内核,如何使其并行?关键在于 CUDA 的 <<<1, 1>>> 语法。这称为执行配置,它告诉 CUDA 运行时要使用多少个并行线程来在 GPU 上启动。这里有两个参数,但让我们从更改第二个参数开始:线程块中的线程数。 CUDA GPU 使用大小为 32 倍数的线程块运行内核,因此选择 256 个线程是合理的大小。

add<<<1, 256>>>(N, x, y);

如果我仅使用此更改运行代码,它将为每个线程执行一次计算,而不是将计算分散到并行线程中。为了正确地做到这一点,我需要修改内核。 CUDA C++ 提供了关键字,让内核可以获取正在运行的线程的索引。具体来说,threadIdx.x 包含当前线程在其块中的索引,blockDim.x 包含块中线程的数量。我将修改循环以使用并行线程跨过数组。

__global__
void add(int n, float *x, float *y)
{int index = threadIdx.x;int stride = blockDim.x;for (int i = index; i < n; i += stride)y[i] = x[i] + y[i];
}

add 功能没有太大变化。事实上,将索引设置为 0 并将步长设置为 1 使其在语义上与第一个版本相同。
将文件保存为 add_block.cu 并再次在 nvprof 中编译并运行。在本文的其余部分中,我将仅显示输出中的相关行。

在这里插入图片描述
这是一个很大的加速(从 463 毫秒降至 2.7 毫秒),但并不奇怪,因为我从 1 个线程增加到 256 个线程。 K80 比我的小型 Macbook Pro GPU 更快(3.2 毫秒)。让我们继续努力以获得更好的性能。

Out of the Blocks

CUDA GPU 具有许多并行处理器,分为流式多处理器 (SM)。每个SM可以运行多个并发线程块。例如,基于 Pascal GPU 架构的 Tesla P100 GPU 有 56 个 SM,每个 SM 最多能够支持 2048 个活动线程。为了充分利用所有这些线程,我应该启动具有多个线程块的内核。
现在您可能已经猜到执行配置的第一个参数指定了线程块的数量。并行线程块一起构成了所谓的网格。由于我有 N 个元素要处理,每个块有 256 个线程,因此我只需要计算块数即可获得至少 N 个线程。我只是将 N 除以块大小(小心向上舍入,以防 N 不是 blockSize 的倍数)。
在这里插入图片描述
在这里插入图片描述
我还需要更新内核代码以考虑整个线程块网格。 CUDA提供了gridDim.x,它包含网格中块的数量,以及blockIdx.x,它包含当前线程块在网格中的索引。图 1 说明了使用 blockDim.xgridDim.x threadIdx.xCUDA 中对数组(一维)进行索引的方法。这个想法是,每个线程通过计算到其块开头的偏移量(块索引乘以块大小:blockIdx.x * blockDim.x)并添加块内线程的索引(threadIdx.x)来获取其索引。代码 blockIdx.x * blockDim.x + threadIdx.x 是惯用的 CUDA

__global__
void add(int n, float *x, float *y)
{int index = blockIdx.x * blockDim.x + threadIdx.x;int stride = blockDim.x * gridDim.x;for (int i = index; i < n; i += stride)y[i] = x[i] + y[i];
}

更新后的内核还将步幅设置为网格中的线程总数 (blockDim.x * gridDim.x)。 CUDA 内核中的这种类型的循环通常称为网格跨度循环。
将文件另存为 add_grid.cu 并再次在 nvprof 中编译并运行。
在这里插入图片描述
在 K80 的所有 SM 上运行多个块,这又是 28 倍的加速!我们只使用 K80 上 2 个 GPU 之一,但每个 GPU 有 13 个 SM。请注意,我笔记本电脑中的 GeForce 有 2 个(较弱的)SM,运行内核需要 680us。

Summing Up

以下是 Tesla K80 和 GeForce GT 750M 上三个版本的 add() 内核的性能概要。
在这里插入图片描述正如您所看到的,我们可以在 GPU 上实现非常高的带宽。本文中的计算非常受带宽限制,但 GPU 也擅长计算密集型计算,例如密集矩阵线性代数、深度学习、图像和信号处理、物理模拟等。

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

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

相关文章

【Vue实现参数传递:查询参数 vs. 动态路由】

文章目录 查询参数传递1. 什么是查询参数&#xff1f;2. 在Vue中使用查询参数步骤 1&#xff1a;在路由配置中定义查询参数步骤 2&#xff1a;在组件中使用查询参数步骤 3&#xff1a;在页面中生成链接 3. 查询参数传递的优势 动态路由传递1. 什么是动态路由&#xff1f;2. 在V…

SQL-Labs46关order by注入姿势

君衍. 四十六关 ORDER BY数字型注入1、源码分析2、rand()盲注3、if语句盲注4、时间盲注5、报错注入6、Limit注入7、盲注脚本 四十六关 ORDER BY数字型注入 请求方式注入类型拼接方式GET报错、布尔盲注、延时盲注ORDER BY $id 我们直接可以从界面中得知传参的参数为SORT&#x…

Linux内核处理并发与竞争的一种方法:信号量

一. 简介 本文来学习Linux内核处理并发与竞争的一种方法:信号量。 本文主要对Linux内核提供的信号量进行简单的介绍。 二. Linux内核处理并发与竞争的一种方法:信号量 1. 信号量简介 大家如果有学习过 FreeRTOS 或者 UCOS 的话就应该对信号量很熟悉,因为信号量是同步…

Yolo v9 “Silence”模块结构及作用!

论文链接&#xff1a;&#x1f47f; YOLOv9: Learning What You Want to Learn Using Programmable Gradient Information 代码链接&#xff1a;&#x1f47f; https://github.com/WongKinYiu/yolov9/tree/main Silence代码 class Silence(nn.Module):def __init__(self):supe…

vue2和vue3对比(语法层面)

阅读文章你将收获&#xff1a; 1 了解不使用组件化工具时&#xff0c;vue在html是如何使用的 2 知道vue2的生命周期函数有哪些 3 知道如何在组件化开发中使用vue 4 大致了解了vue2和vue3在使用上什么不同 最后&#xff1a;vue2和vue3除了下面我列出的有差异化的地方&…

高防服务器的价格受到哪些因素的影响?

高防服务器可以帮助网站拒绝服务攻击&#xff0c;能够进行定时扫描网络节点&#xff0c;并且查找可能会存在的安全漏洞的服务器类型&#xff0c;对于经常遭受到DDOS攻击的大型游戏网络企业会选择高防服务器。 那面对价格多样的高防服务器我们应该怎样进行选择&#xff0c;高防服…

day41打卡

day41打卡 46. 携带研究材料&#xff08;第六期模拟笔试&#xff09; 状态表示 ​ 二维&#xff1a;dp[i] [j] 表示从下标为[0-i]的物品里任意取&#xff0c;放进容量为j的背包&#xff0c;价值总和最大是多少。 一维&#xff1a; ​ dp[j]表示&#xff1a;容量为j的背包&a…

模型 HBG(品牌增长)

系列文章 分享 模型&#xff0c;了解更多&#x1f449; 模型_总纲目录。品牌增长法。 1 HBG(品牌增长)模型的应用 1.1 江小白使用HBG模型提高品牌知名度和销售额 选择受众市场&#xff1a;江小白的目标客户是年轻人&#xff0c;他们喜欢简单、时尚的产品。因此&#xff0c;江…

面试前端性能优化八股文十问十答第三期

面试前端性能优化八股文十问十答第三期 作者&#xff1a;程序员小白条&#xff0c;个人博客 相信看了本文后&#xff0c;对你的面试是有一定帮助的&#xff01;关注专栏后就能收到持续更新&#xff01; ⭐点赞⭐收藏⭐不迷路&#xff01;⭐ 1&#xff09;如何⽤webpack来优化…

数据结构D4作业

1.实现单向循环链表的功能 loop.c #include "loop.h" loop_p create_loop() { loop_p H(loop_p)malloc(sizeof(loop)); if(HNULL) { printf("创建失败\n"); return NULL; } H->len0; H->nextH; ret…

一文彻底搞懂SQL优化

文章目录 1. 索引优化2. WHERE 子句优化3. JOIN 优化4. 查询结果优化5. 子查询优化6. 数据库统计信息和缓存优化7. 事务优化8. 数据库配置优化9. 使用适当的数据类型10. 监控和调优 SQL(Structured Query Language&#xff09;优化是指对 SQL 查询语句进行调整和改进&#xff0…

关于 Reflect 的笔记

背景&#xff1a;Reflect 为了操作对象而提供的新Api 和 Proxy对象一样 特点 将object 对象的一些明显属于语言内部的方法&#xff0c;放到Reflect 上处理&#xff1b;修改某些object返回的异常结果&#xff0c;让其变得更合理&#xff1b;让object操作都变成函数行为&#xf…

基于ElementUI封装省市区四级联动下拉选择

基于ElementUI封装的省市区下拉级联选择 效果 数据 最新省市区JSON数据获取&#xff1a;https://xiangyuecn.github.io/AreaCity-JsSpider-StatsGov/ 参数说明 参数说明inputNumShow下拉框的数量&#xff0c;最多4个defaultAddress默认显示省市区 例&#xff1a;[‘安徽’, …

按形如 a*sqrt(b) 的格式输出一个非负整数的平方根

【题目描述】 输入一个非负整数 x&#xff0c;若能完全开平方根&#xff0c;则输出其对应的整数平方根值。 否则&#xff0c;按形如 a*sqrt(b) 的格式输出其平方根值&#xff08;a 与 b 均为整数&#xff0c;且 a≠1&#xff0c;b≠1&#xff09;。【输入输出】 典型的输入输出…

【C++初阶】--类和对象(下)

目录 一.const成员 1.权限放大问题 2.权限的缩小 二.再谈构造函数 1.构造函数体赋值 2.初始化列表 (1)概念 (2)使用 ①在对象实例化过程中&#xff0c;成员变量先依次进行初始化 ②再进行函数体内二次赋值 3.explicit关键字 (1)C为什么要存在自动隐式类型转换…

Linux内核中自旋锁驱动代码举例二

一. 简介 前面学习了不考虑中断的自旋锁的代码举例,文章地址: Linux内核自旋锁驱动代码举例一-CSDN博客 但是在 Linux系统中,中断时存在的。所以,这里学习使用带保存中断状态的自旋锁API函数,实现对Led设备的互斥访问。 二. 带保存中断状态的自旋锁函数使用 1. 准备…

Linux第64步_编译移植好的虚拟机文件

最好还是认真了解linux系统移植的整个过程&#xff0c;否则&#xff0c;可能会让你误入歧途。 1、编译移植好的tf-a 1)、编译生成“tf-a-stm32mp157d-atk-trusted.stm32” 输入“cd /home/zgq/linux/atk-mp1/tfa/my-tfa/tf-a-stm32mp-2.2.r1/回车”&#xff0c;切换到“/hom…

算法打卡day1|数组篇|Leetcode 704.二分查找、27.移除元素

数组理论基础 数组是存放在连续内存空间上的相同类型数据的集合&#xff0c;可以方便的通过下标索引的方式获取到下标下对应的数据。 1.数组下标都是从0开始的。 2.数组内存空间的地址是连续的。 正是因为数组的在内存空间的地址是连续的&#xff0c;所以我们在删除或者增添…

算法题--华为od机试考试(分苹果、字符串统计及重排、高矮个子排队)

目录 分苹果 题目描述 输入描述 输出描述 示例1 输入 输出 备注 示例2 输入 输出 解析 答案 字符统计及重排 题目描述 输入描述 输出描述 示例1 输入 输出 说明 示例2 输入 输出 说明 解析 答案 高矮个子排队 题目描述 输入描述 输出描述 备注…

【深度学习笔记】3_5 图像分类数据集fashion-mnist

注&#xff1a;本文为《动手学深度学习》开源内容&#xff0c;仅为个人学习记录&#xff0c;无抄袭搬运意图 3.5 图像分类数据集&#xff08;Fashion-MNIST&#xff09; 在介绍softmax回归的实现前我们先引入一个多类图像分类数据集。它将在后面的章节中被多次使用&#xff0c…