GPUNvidia CUDA 编程基础教程——利用基本的 CUDA 内存管理技术来优化加速应用程序
Posted 从善若水
tags:
篇首语：本文由小常识网(cha138.com)小编为大家整理，主要介绍了GPUNvidia 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.sqlite
CUDA API Statistics:
Time(%) Total Time (ns) Num Calls Average Minimum Maximum Name
------- --------------- --------- ------------ ---------- ---------- ---------------------
90.8 2323752043 1 2323752043.0 2323752043 2323752043 cudaDeviceSynchronize
8.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_create
0.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-82c3GPUNvidia CUDA 编程高级教程——利用蒙特卡罗法求解近似值（CUDA-Aware MPI）
GPUNvidia CUDA 编程高级教程——利用蒙特卡罗法求解近似值（CUDA-Aware MPI）
GPUNvidia CUDA 编程高级教程——利用蒙特卡罗法求解 的近似值
GPUNvidia CUDA 编程高级教程——利用蒙特卡罗法求解 的近似值