【GPU】Nvidia CUDA 编程基础教程——利用基本的 CUDA 内存管理技术来优化加速应用程序
博主未授权任何人或组织机构转载博主任何原创文章，感谢各位对原创的支持！
博主链接
本人就职于国际知名终端厂商，负责modem芯片研发。
在5G早期负责终端数据业务层、核心网相关的开发工作，目前牵头6G算力网络技术标准研究。博客内容主要围绕：
5G/6G协议讲解
算力网络讲解（云计算，边缘计算，端计算）
高级C语言讲解
Rust语言讲解
利用基本的 CUDA 内存管理技术来优化加速应用程序
使用nsys性能分析器帮助应用程序迭代地进行优化
如要确保优化加速代码库的尝试真正取得成功，唯一方法便是分析应用程序以获取有关其性能的定量信息。nsys 是指 NVIDIA 的Nsight System命令行分析器。该分析器附带于CUDA工具包中，提供分析被加速的应用程序性能的强大功能。
nsys 使用起来十分简单，最基本用法是向其传递使用 nvcc 编译的可执行文件的路径。随后 nsys 会继续执行应用程序，并在此之后打印应用程序 GPU 活动的摘要输出、CUDA API 调用以及统一内存活动的相关信息。
在加速应用程序或优化已经加速的应用程序时，我们应该采用科学的迭代方法。作出更改后需分析应用程序、做好记录并记录任何重构可能会对性能造成何种影响。尽早且经常进行此类观察通常会让您轻松获得足够的性能提升，以助您发布加速应用程序。此外，经常分析应用程序将使您了解到对 CUDA 代码库作出的特定更改会对其实际性能造成何种影响：而当只在代码库中进行多种更改后再分析应用程序时，将很难得知这一点。
使用nsys分析应用程序
#include <stdio.h>/** Host function to initialize vector elements. This function* simply initializes each element to equal its index in the* vector.*/void initWith(float num, float *a, int N)
{for(int i = 0; i < N; ++i){a[i] = num;}
}/** Device kernel stores into `result` the sum of each* same-indexed value of `a` and `b`.*/__global__
void addVectorsInto(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];}
}/** Host function to confirm values in `vector`. This function* assumes all values are the same `target` value.*/void checkElementsAre(float target, float *vector, int N)
{for(int i = 0; i < N; i++){if(vector[i] != target){printf("FAIL: vector[%d] - %0.0f does not equal %0.0f\n", i, vector[i], target);exit(1);}}printf("Success! All values calculated correctly.\n");
}int main()
{const int N = 2<<24;size_t size = N * sizeof(float);float *a;float *b;float *c;cudaMallocManaged(&a, size);cudaMallocManaged(&b, size);cudaMallocManaged(&c, size);initWith(3, a, N);initWith(4, b, N);initWith(0, c, N);size_t threadsPerBlock;size_t numberOfBlocks;/** nsys should register performance changes when execution configuration* is updated.*/threadsPerBlock = 1;numberOfBlocks = 1;cudaError_t addVectorsErr;cudaError_t asyncErr;addVectorsInto<<<numberOfBlocks, threadsPerBlock>>>(c, a, b, N);addVectorsErr = cudaGetLastError();if(addVectorsErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(addVectorsErr));asyncErr = cudaDeviceSynchronize();if(asyncErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(asyncErr));checkElementsAre(7, c, N);cudaFree(a);cudaFree(b);cudaFree(c);
}运行以下命令，编译.cu
nvcc -o single-thread-vector-add 01-vector-add/01-vector-add.cu -run
使用nsys分析程序
nsys profile --stats=true ./single-thread-vector-add
输出的结果如下：
Warning: LBR backtrace method is not supported on this platform. DWARF backtrace method will be used.
Collecting data...
Success! All values calculated correctly.
Processing events...
Capturing symbol files...
Saving temporary "/tmp/nsys-report-2314-afad-f037-6b95.qdstrm" file to disk...
Creating final output files...Processing [==============================================================100%]
Saved report file to "/tmp/nsys-report-2314-afad-f037-6b95.qdrep"
Exporting 1080 events: [==================================================100%]Exported successfully to
/tmp/nsys-report-2314-afad-f037-6b95.sqliteCUDA API Statistics:Time(%) Total Time (ns) Num Calls Average Minimum Maximum Name ------- --------------- --------- ------------ ---------- ---------- ---------------------90.8 2323752043 1 2323752043.0 2323752043 2323752043 cudaDeviceSynchronize8.4 213955285 3 71318428.3 18101 213912843 cudaMallocManaged 0.8 20304172 3 6768057.3 6068829 7925314 cudaFree 0.0 49268 1 49268.0 49268 49268 cudaLaunchKernel CUDA Kernel Statistics:Time(%) Total Time (ns) Instances Average Minimum Maximum Name ------- --------------- --------- ------------ ---------- ---------- -------------------------------------------100.0 2323741800 1 2323741800.0 2323741800 2323741800 addVectorsInto(float*, float*, float*, int)Operating System Runtime API Statistics:Time(%) Total Time (ns) Num Calls Average Minimum Maximum Name ------- --------------- --------- ---------- ------- --------- --------------49.7 2990336371 40 74758409.3 22644 100075885 sem_timedwait 48.2 2903562111 40 72589052.8 29203 100134978 poll 1.7 100969145 658 153448.5 1022 17198753 ioctl 0.4 22367169 89 251316.5 1341 7861414 mmap 0.0 1620737 77 21048.5 4498 41568 open64 0.0 130635 3 43545.0 39445 47538 pthread_create0.0 107175 23 4659.8 1465 15258 fopen 0.0 100180 3 33393.3 11897 75802 fgets 0.0 91575 11 8325.0 4469 13643 write 0.0 47228 13 3632.9 1536 5613 munmap 0.0 40520 13 3116.9 1486 6730 read 0.0 27315 16 1707.2 1096 3773 fclose 0.0 24081 4 6020.3 3120 8757 open 0.0 18900 3 6300.0 5804 6932 pipe2 0.0 14374 3 4791.3 1042 11424 fgetc 0.0 13291 2 6645.5 6007 7284 socket 0.0 10612 7 1516.0 1018 3996 fcntl 0.0 7702 2 3851.0 3795 3907 fread 0.0 6937 1 6937.0 6937 6937 connect 0.0 6393 3 2131.0 2053 2238 mprotect 0.0 2253 1 2253.0 2253 2253 bind 0.0 1660 1 1660.0 1660 1660 listen Report file moved to "/dli/task/report1.qdrep"
Report file moved to "/dli/task/report1.sqlite"
流多处理器（Streaming Multiprocessors）及查询GPU的设备配置
NVIDIA GPU 包含称为流多处理器或 SM 的功能单元，线程块均可安排在 SM 上运行，如下图：
根据 GPU 上的 SM 数量以及线程块要求，可在 SM 上安排运行多个线程块，如下：
如果网格维度能被 GPU 上的 SM 数量整除，则可充分提高 SM 的利用率。以下是闲置的 SM，
流多处理器和Warps
运行 CUDA 应用程序的 GPU 具有称为流多处理器（或 SM）的处理单元。在核函数执行期间，将线程块提供给 SM 以供其执行。为支持 GPU 执行尽可能多的并行操作，您通常可以选择线程块数量数倍于指定 GPU 上 SM 数量的网格大小来提升性能。
此外，SM 会在一个名为warp的线程块内创建、管理、调度和执行包含 32 个线程的线程组。本课程将不会更深入探讨 SM 和warp，但值得注意的是，您也可选择线程数量数倍于 32 的线程块大小来提升性能。
以编程方式查询GPU设备属性
由于 GPU 上的 SM 数量会因所用的特定 GPU 而异，因此为支持可移植性，您不得将 SM 数量硬编码到代码库中。相反，应该以编程方式获取此信息。
以下所示为在 CUDA C/C++ 中获取 C 结构的方法，该结构包含当前处于活动状态的 GPU 设备的多个属性，其中包括设备的 SM 数量：
int deviceId;
cudaGetDevice(&deviceId); // `deviceId` now points to the id of the currently active GPU.cudaDeviceProp props;
cudaGetDeviceProperties(&props, deviceId); // `props` now has many useful properties about// the active GPU device.
获得统一内存的细节
您一直使用 cudaMallocManaged 分配旨在供主机或设备代码使用的内存，并且现在仍在享受这种方法的便利之处，即在实现自动内存迁移且简化编程的同时，而无需深入了解 cudaMallocManaged 所分配统一内存 (UM) 实际工作原理的详细信息。nsys profile 提供有关加速应用程序中 UM 管理的详细信息，并在利用这些信息的同时结合对 UM 工作原理的更深入理解，进而为优化加速应用程序创造更多机会。
分配 UM 时，它最初可能并未驻留在 CPU 或 GPU 上，当某些工作首次请求内存时，将会发生分页错误。分页错误将触发所请求的内存发生迁移，如下图：
只要在系统中并未驻留内存的位置请求内存，此过程便会重复，如下：
如果已知将在未驻留内存的位置访问内存，则可使用异步预取，异步预取能以更大批量移动内存，并会防止发生分页错误。如下：
统一内存(UM)的迁移
分配 UM 时，内存尚未驻留在主机或设备上。主机或设备尝试访问内存时会发生页错误，此时主机或设备会批量迁移所需的数据。同理，当 CPU 或加速系统中的任何 GPU 尝试访问尚未驻留在其上的内存时，会发生页错误并触发迁移。
能够执行页错误并按需迁移内存对于在加速应用程序中简化开发流程大有助益。此外，在处理展示稀疏访问模式的数据时（例如，在应用程序实际运行之前无法得知需要处理的数据时），以及在具有多个 GPU 的加速系统中，数据可能由多个 GPU 设备访问时，按需迁移内存将会带来显著优势。
有些情况下（例如，在运行时之前需要得知数据，以及需要大量连续的内存块时），我们还能有效规避页错误和按需数据迁移所产生的开销。
异步内存预取
在主机到设备和设备到主机的内存传输过程中，我们使用一种技术来减少页错误和按需内存迁移成本，此强大技术称为异步内存预取。通过此技术，程序员可以在应用程序代码使用统一内存 (UM) 之前，在后台将其异步迁移至系统中的任何 CPU 或 GPU 设备。此举可以减少页错误和按需数据迁移所带来的成本，并进而提高 GPU 核函数和 CPU 函数的性能。
此外，预取往往会以更大的数据块来迁移数据，因此其迁移次数要低于按需迁移。此技术非常适用于以下情况：在运行时之前已知数据访问需求且数据访问并未采用稀疏模式。
CUDA 可通过 cudaMemPrefetchAsync 函数，轻松将托管内存异步预取到 GPU 设备或 CPU。以下所示为如何使用该函数将数据预取到当前处于活动状态的 GPU 设备，然后再预取到 CPU：
int deviceId;
cudaGetDevice(&deviceId); // The ID of the currently active GPU device.cudaMemPrefetchAsync(pointerToSomeUMData, size, deviceId); // Prefetch to GPU device.
cudaMemPrefetchAsync(pointerToSomeUMData, size, cudaCpuDeviceId); // Prefetch to host. `cudaCpuDeviceId` is a// built-in CUDA variable.
UM内存预取的例子
#include <stdio.h>void initWith(float num, float *a, int N)
{for(int i = 0; i < N; ++i){a[i] = num;}
}__global__
void addVectorsInto(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];}
}void checkElementsAre(float target, float *vector, int N)
{for(int i = 0; i < N; i++){if(vector[i] != target){printf("FAIL: vector[%d] - %0.0f does not equal %0.0f\n", i, vector[i], target);exit(1);}}printf("Success! All values calculated correctly.\n");
}int main()
{int deviceId;int numberOfSMs;cudaGetDevice(&deviceId);cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);printf("Device ID: %d\tNumber of SMs: %d\n", deviceId, numberOfSMs);const int N = 2<<24;size_t size = N * sizeof(float);float *a;float *b;float *c;cudaMallocManaged(&a, size);cudaMallocManaged(&b, size);cudaMallocManaged(&c, size);/** Prefetching can also be used to prevent CPU page faults.*/cudaMemPrefetchAsync(a, size, cudaCpuDeviceId);cudaMemPrefetchAsync(b, size, cudaCpuDeviceId);cudaMemPrefetchAsync(c, size, cudaCpuDeviceId);initWith(3, a, N);initWith(4, b, N);initWith(0, c, N);cudaMemPrefetchAsync(a, size, deviceId);cudaMemPrefetchAsync(b, size, deviceId);cudaMemPrefetchAsync(c, size, deviceId);size_t threadsPerBlock;size_t numberOfBlocks;threadsPerBlock = 256;numberOfBlocks = 32 * numberOfSMs;cudaError_t addVectorsErr;cudaError_t asyncErr;addVectorsInto<<<numberOfBlocks, threadsPerBlock>>>(c, a, b, N);addVectorsErr = cudaGetLastError();if(addVectorsErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(addVectorsErr));asyncErr = cudaDeviceSynchronize();if(asyncErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(asyncErr));/** Prefetching can also be used to prevent CPU page faults.*/cudaMemPrefetchAsync(c, size, cudaCpuDeviceId);checkElementsAre(7, c, N);cudaFree(a);cudaFree(b);cudaFree(c);
}使用一下命令编译：
nvcc -o prefetch-to-cpu 01-vector-add/01-vector-add.cu -run
使用nsys分析代码：
nsys profile --stats=true ./prefetch-to-cpu
分析数据如下：
在这里插入代码片Warning: LBR backtrace method is not supported on this platform. DWARF backtrace method will be used.
Collecting data...
Success! All values calculated correctly.
Processing events...
Capturing symbol files...
Saving temporary "/tmp/nsys-report-9279-82c3-782a-d763.qdstrm" file to disk...
Creating final output files...Processing [==============================================================100%]
Saved report file to "/tmp/nsys-report-9279-82c3-782a-d763.qdrep"
Exporting 1104 events: [==================================================100%]Exported successfully to
/tmp/nsys-report-9279-82c3-782a-d763.sqliteCUDA API Statistics:Time(%) Total Time (ns) Num Calls Average Minimum Maximum Name ------- --------------- --------- ------------ ---------- ---------- ---------------------89.4 2340459996 1 2340459996.0 2340459996 2340459996 cudaDeviceSynchronize9.8 255349510 3 85116503.3 36248 255274065 cudaMallocManaged 0.8 20709879 3 6903293.0 6217307 8071714 cudaFree 0.0 53112 1 53112.0 53112 53112 cudaLaunchKernel CUDA Kernel Statistics:Time(%) Total Time (ns) Instances Average Minimum Maximum Name ------- --------------- --------- ------------ ---------- ---------- -------------------------------------------100.0 2340444055 1 2340444055.0 2340444055 2340444055 addVectorsInto(float*, float*, float*, int)Operating System Runtime API Statistics:Time(%) Total Time (ns) Num Calls Average Minimum Maximum Name ------- --------------- --------- ---------- ------- --------- --------------49.4 3184377039 42 75818500.9 8752 100074730 sem_timedwait 48.2 3103601506 42 73895274.0 29641 100134382 poll 2.0 129742838 659 196878.4 1044 30576740 ioctl 0.4 23901468 89 268555.8 2229 8000483 mmap 0.0 1865799 77 24231.2 8604 48789 open64 0.0 217260 23 9446.1 3057 36618 fopen 0.0 208945 3 69648.3 21948 162061 fgets 0.0 167026 3 55675.3 37588 66053 pthread_create0.0 95862 11 8714.7 4929 13843 write 0.0 73579 14 5255.6 1672 12079 munmap 0.0 47937 16 2996.1 1634 8320 fclose 0.0 45543 4 11385.8 5203 17390 open 0.0 41010 25 1640.4 1004 8487 fcntl 0.0 38115 13 2931.9 1341 6055 read 0.0 27299 2 13649.5 11105 16194 socket 0.0 25534 3 8511.3 1914 20151 fgetc 0.0 24055 3 8018.3 5628 11625 pipe2 0.0 14576 1 14576.0 14576 14576 connect 0.0 13798 2 6899.0 6526 7272 fread 0.0 9455 3 3151.7 1811 3897 mprotect 0.0 4318 1 4318.0 4318 4318 bind 0.0 2986 1 2986.0 2986 2986 listen Report file moved to "/dli/task/report2.qdrep"
Report file moved to "/dli/task/report2.sqlite"
【GPU】Nvidia CUDA 编程基础教程——利用基本的 CUDA 内存管理技术来优化加速应用程序相关推荐
- 【GPU】Nvidia CUDA 编程高级教程——利用蒙特卡罗法求解近似值（NVSHMEM）
博主未授权任何人或组织机构转载博主任何原创文章,感谢各位对原创的支持! 博主链接 本人就职于国际知名终端厂商,负责modem芯片研发. 在5G早期负责终端数据业务层.核心网相关的开发工作,目前牵头6G ...
- 2023年的深度学习入门指南(10) - CUDA编程基础
2023年的深度学习入门指南(10) - CUDA编程基础 上一篇我们走马观花地看了下SIMD和GPGPU的编程.不过线条太粗了,在开发大模型时遇到问题了肯定还会晕. 所以我们还是需要深入到CUDA中 ...
- CUDA编程基础知识整理
CUDA编程基础知识整理 CUDA编程和GPU架构基本概念 CUDA核函数调度方式 前言:最近在尝试通过CUDA加速Swin Transformer,第一次接触CUDA的相关知识,将这些天学习到的内容 ...
- Windows驱动编程基础教程
前言 本书非常适合熟悉Windows应用编程的读者转向驱动开发.所有的内容都从最基础的编程方法入手.介绍相关的内核API,然后举出示范的例子.这本书只有不到70页,是一本非常精简的小册子.所以 ...
- Windows驱动编程基础教程 （转）
Windows驱动编程基础教程(转) 我经常在网上遇到心如火燎的提问者.他们碰到很多工作中的技术问题,是关于驱动开发的.其实绝大部分他们碰到的"巨大困难"是被老牛们看成初级得 ...
- 楚狂人Windows驱动编程基础教程
版权声明 本书是免费电子书.作者保留一切权利.但在保证本书完整性(包括版权声明.前言.正文内容.后记.以及作者的信息),并不增删.改变其中任何文字内容的前提下,欢迎任何读者以任何形式(包括各种 ...
- CUDA 编程简单入门 Advance CUDA 编程基础 (C++ programming)
Advance CUDA编程基础 (C++ programming) GPU 架构 CUDA 编程基础 基本代码框架 CUDA Execution Model Case Study : Vector ...
- CV：NVIDIA驱动程序安装图文教程(根据Anaconda的CUDA版本去安装对应匹配的NVIDIA)之详细攻略
CV:NVIDIA驱动程序安装图文教程(根据Anaconda的CUDA版本去安装对应匹配的NVIDIA)之详细攻略 目录 根据Anaconda的CUDA版本去安装对应匹配的NVIDIA驱动程序图文教程 ...
- java程序设计之网络编程基础教程_Java程序设计之网络编程基础教程
基本信息 书名:Java程序设计之网络编程基础教程(21世纪高等学校计算机基础实用规划教材) :43.50元 作者:李芝兴 主编 出版社:清华大学出版社 出版日期:2012-12-1 ISBN:978 ...
最新文章
- java主窗体设计代码_java窗体设计+GUI经典代码全放送
- lncRNA是什么？
- ITK:将所有像素的总和缩放为常数
- java中document解析jsp,JSP基于dom解析xml实例详解
- oracle 一个实例创建多个数据库_Oracle闪回，为你的数据库上一个安全防线
- 第一类边界条件，三角单元剖分，线性插值的位场延拓，LDLT高效求解
- 28. (附加)字符串的组合(C++版本)
- python中strptime函数_python datetime中strptime用法详解
- css文字覆盖线性渐变,利用css使文字渐变
- 一款真正可以操作的数据中心可视化管理软件
- c语言total用法,C语言 这个表达式怎么理解 新手请大神详述total += isalnum(ch[i])!=0;...
- 支付宝支付php开发demo下载地址,tp使用支付宝接口demo
- 百度短网址 php,PHP实现百度、网易、新浪短网址服务的API接口调用
- 中国大学moocpython笔记_中国大学MOOC —— 学习笔记（二）
- 开机后黑屏看不到桌面_电脑开机黑屏只有鼠标怎么办？电脑开机后不显示桌面的多种解决方法...
- 微信自动邀请加群！！！
- 证明四元数表示旋转的过程实部为0，虚部为罗德里格斯公式结果
- android crash监控系统,Android-Crash监控
- JVM内存模型和性能调优：JVM调优工具详解及调优实战：jstat调优小实战- 第40篇
- python爬取短视频,Python爬虫一键下载yy全站短视频详细步骤（附源码）