博主未授权任何人或组织机构转载博主任何原创文章，感谢各位对原创的支持！

博主链接

本人就职于国际知名终端厂商，负责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 AverageMinimumMaximum Name ------- --------------- --------- ------------ ---------- ---------- ---------------------90.8 23237520431 2323752043.0 2323752043 2323752043 cudaDeviceSynchronize8.4 2139552853 71318428.3 18101 213912843 cudaMallocManaged 0.8 2030417236768057.360688297925314 cudaFree 0.0 492681 49268.0 49268 49268 cudaLaunchKernelCUDA Kernel Statistics:Time(%) Total Time (ns) Instances AverageMinimumMaximum Name------- --------------- --------- ------------ ---------- ---------- -------------------------------------------100.0 23237418001 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 poll1.7 100969145 658 153448.51022 17198753 ioctl 0.4 22367169 89 251316.51341 7861414 mmap0.01620737 7721048.5449841568 open64 0.0 130635343545.0 3944547538 pthread_create0.0 107175 234659.8146515258 fopen 0.0 100180333393.3 1189775802 fgets 0.0 91575 118325.0446913643 write 0.0 47228 133632.91536 5613 munmap 0.0 40520 133116.91486 6730 read0.0 27315 161707.21096 3773 fclose 0.0 2408146020.33120 8757 open0.0 1890036300.05804 6932 pipe2 0.0 1437434791.3104211424 fgetc 0.0 1329126645.56007 7284 socket 0.0 1061271516.01018 3996 fcntl 0.0 770223851.03795 3907 fread 0.0 693716937.06937 6937 connect 0.0 639332131.02053 2238 mprotect0.0 225312253.02253 2253 bind0.0 166011660.01660 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 AverageMinimumMaximum Name ------- --------------- --------- ------------ ---------- ---------- ---------------------89.4 23404599961 2340459996.0 2340459996 2340459996 cudaDeviceSynchronize9.8 2553495103 85116503.3 36248 255274065 cudaMallocManaged 0.8 2070987936903293.062173078071714 cudaFree 0.0 531121 53112.0 53112 53112 cudaLaunchKernelCUDA Kernel Statistics:Time(%) Total Time (ns) Instances AverageMinimumMaximum Name------- --------------- --------- ------------ ---------- ---------- -------------------------------------------100.0 23404440551 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.98752 100074730 sem_timedwait 48.2 3103601506 42 73895274.0 29641 100134382 poll2.0 129742838 659 196878.41044 30576740 ioctl 0.4 23901468 89 268555.82229 8000483 mmap0.01865799 7724231.2860448789 open64 0.0 217260 239446.1305736618 fopen 0.0 208945369648.3 21948162061 fgets 0.0 167026355675.3 3758866053 pthread_create0.0 95862 118714.7492913843 write 0.0 73579 145255.6167212079 munmap 0.0 47937 162996.11634 8320 fclose 0.0 45543411385.8520317390 open0.0 41010 251640.41004 8487 fcntl 0.0 38115 132931.91341 6055 read0.0 27299213649.5 1110516194 socket 0.0 2553438511.319141 fgetc 0.0 2405538018.3562811625 pipe2 0.0 14576114576.0 1457614576 connect 0.0 1379826899.06526 7272 fread 0.0 945533151.71811 3897 mprotect0.0 431814318.04318 4318 bind0.0 298612986.02986 2986 listen Report file moved to "/dli/task/report2.qdrep"Report file moved to "/dli/task/report2.sqlite"