[每日一氵] Nsight Systems (nsys) 使用记录以及cuda程序优化

2023-11-21 01:40
文章标签 使用 程序 优化 每日 记录 cuda systems nsight nsys

本文主要是介绍[每日一氵] Nsight Systems (nsys) 使用记录以及cuda程序优化，希望对大家解决编程问题提供一定的参考价值，需要的开发者们随着小编来一起学习吧！

主要内容来自这个课程界面：
https://www.nvidia.cn/training/instructor-led-workshops/

Nsight Systems (nsys) 原来这么有用啊，我每次安装cuda的时候，都不安装他，不过配置DL环境确实不需要[手动狗头]

这样运行文件

nvcc -o xx xx.cu -run

Nsys这么用

# 运用 nsys profile 分析刚编译好的可执行文件
nsys profile --stats=true ./xx

nsys profile将生成一个qdrep报告文件，该文件可以以多种方式使用。 我们在这里使用--stats = true标志表示我们希望打印输出摘要统计信息。 输出的信息有很多，包括：

  • 配置文件配置详细信息
  • 报告文件的生成详细信息
  • CUDA API统计信息
  • CUDA核函数的统计信息
  • CUDA内存操作统计信息（时间和大小）
  • 操作系统内核调用接口的统计信息

如下：

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-a10b-4a77-7d5c-f462.qdstrm" file to disk...
Creating final output files...Processing [==============================================================100%]
Saved report file to "/tmp/nsys-report-a10b-4a77-7d5c-f462.qdrep"
Exporting 22723 events: [=================================================100%]Exported successfully to
/tmp/nsys-report-a10b-4a77-7d5c-f462.sqliteCUDA API Statistics: # CUDA API统计信息Time(%)  Total Time (ns)  Num Calls    Average     Minimum    Maximum           Name         -------  ---------------  ---------  -----------  ---------  ---------  ---------------------55.9        220024635          3   73341545.0      35564  219942207  cudaMallocManaged    39.1        154081013          1  154081013.0  154081013  154081013  cudaDeviceSynchronize5.0         19599393          3    6533131.0    5868170    7536695  cudaFree             0.0            54357          1      54357.0      54357      54357  cudaLaunchKernel     CUDA Kernel Statistics: # CUDA核函数的统计信息Time(%)  Total Time (ns)  Instances    Average     Minimum    Maximum                      Name                    -------  ---------------  ---------  -----------  ---------  ---------  -------------------------------------------100.0        154061080          1  154061080.0  154061080  154061080  addVectorsInto(float*, float*, float*, int)CUDA Memory Operation Statistics (by time): # CUDA内存操作统计信息（时间）Time(%)  Total Time (ns)  Operations  Average  Minimum  Maximum              Operation            -------  ---------------  ----------  -------  -------  -------  ---------------------------------82.6         99842969       20879   4782.0     1823   169216  [CUDA Unified Memory memcpy HtoD]17.4         21020960         768  27371.0     1375   159872  [CUDA Unified Memory memcpy DtoH]CUDA Memory Operation Statistics (by size in KiB): # CUDA内存操作统计信息（大小）Total     Operations  Average  Minimum  Maximum               Operation            ----------  ----------  -------  -------  --------  ---------------------------------393216.000       20879   18.833    4.000  1012.000  [CUDA Unified Memory memcpy HtoD]131072.000         768  170.667    4.000  1020.000  [CUDA Unified Memory memcpy DtoH]Operating System Runtime API Statistics: # 操作系统内核调用接口的统计信息Time(%)  Total Time (ns)  Num Calls   Average    Minimum   Maximum        Name     -------  ---------------  ---------  ----------  -------  ---------  --------------53.9       1349784189         74  18240326.9    24368  100131135  poll          41.7       1042453633         74  14087211.3    15428  100074482  sem_timedwait 3.5         87328279        587    148770.5     1023   16811695  ioctl         0.9         21850661         90    242785.1     1235    7474212  mmap          0.0           624849         77      8114.9     2460      18975  open64        0.0           113233          4     28308.3    23925      32553  pthread_create0.0           107072         23      4655.3     1296      13371  fopen         0.0            86168          3     28722.7    20436      43529  fgets         0.0            85314         11      7755.8     4313      13945  write         0.0            40344         14      2881.7     1294       4315  munmap        0.0            29311         16      1831.9     1057       3519  fclose        0.0            27759          5      5551.8     2789       8032  open          0.0            26388         13      2029.8     1114       3558  read          0.0            16141          3      5380.3     3831       6160  pipe2         0.0             8240          2      4120.0     3544       4696  socket        0.0             7423          2      3711.5     1435       5988  fgetc         0.0             6363          4      1590.8     1442       1841  mprotect      0.0             6290          2      3145.0     2664       3626  fread         0.0             5900          1      5900.0     5900       5900  connect       0.0             4790          2      2395.0     1221       3569  fcntl         0.0             1913          1      1913.0     1913       1913  bind          0.0             1418          1      1418.0     1418       1418  listen        Report file moved to "/xxx/task/report3.qdrep"
Report file moved to "/xxx/task/report3.sqlite"

由于 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.

具体的属性名称可以参考这里：
https://docs.nvidia.com/cuda/cuda-runtime-api/structcudaDeviceProp.html

查询信息的时候这样用：

#include <stdio.h>int main()
{/** Assign values to these variables so that the output string below prints the* requested properties of the currently active GPU.*/int deviceId;int computeCapabilityMajor;int computeCapabilityMinor;int multiProcessorCount;int warpSize;cudaGetDevice(&deviceId);cudaDeviceProp prop;cudaGetDeviceProperties(&prop, deviceId);warpSize = prop.warpSize;multiProcessorCount = prop.multiProcessorCount;computeCapabilityMajor = prop.major;computeCapabilityMinor = prop.minor;/** There should be no need to modify the output string below.*/printf("Device ID: %d\nNumber of SMs: %d\nCompute Capability Major: %d\nCompute Capability Minor: %d\nWarp Size: %d\n", deviceId, multiProcessorCount, computeCapabilityMajor, computeCapabilityMinor, warpSize);
}

另外，重点理解这句话：
在这里插入图片描述
优化一个cuda程序，大概有以下几个方向：

  • 将变量初始化在GPU上，或者用 cudaMemPrefetchAsync 异步搬运，搬到GPU上，然后再搬回来
  • 修改核函数，加上stride 如上边那个截图
  • 读取GPU上SM数量，cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);根据这个数量给CUDA核函数传入block/grid数和thread/block数
  • 除了节省GPU跑的时间，也得节省开发者的时间，加上这个，时间不会消耗多少的：
cudaError_t kernelFuncErrs;
cudaError_t asyncErr;kernelFunc<<<numberOfBlocks, threadsPerBlock>>>(c, a, b, N);kernelFuncErrs= cudaGetLastError();
if(kernelFuncErrs != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(kernelFuncErrs));asyncErr = cudaDeviceSynchronize();
if(asyncErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(asyncErr));
  • 然后就是，nsys profile --stats=true ./xx 迭代优化程序，详见：
    https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#memory-optimizations

这篇关于[每日一氵] Nsight Systems (nsys) 使用记录以及cuda程序优化的文章就介绍到这儿，希望我们推荐的文章对编程师们有所帮助！


http://www.chinasem.cn/article/398927。

相关文章

Vue3 的 shallowRef 和 shallowReactive：优化性能

Vue3 的 shallowRef 和 shallowReactive：优化性能

大家对 Vue3 的 ref 和 reactive 都很熟悉，那么对 shallowRef 和 shallowReactive 是否了解呢？ 在编程和数据结构中，“shallow”（浅层）通常指对数据结构的最外层进行操作，而不递归地处理其内部或嵌套的数据。这种处理方式关注的是数据结构的第一层属性或元素，而忽略更深层次的嵌套内容。 1. 浅层与深层的对比 1.1 浅层（Shallow） 定义
阅读更多...
中文分词jieba库的使用与实景应用（一）

中文分词jieba库的使用与实景应用（一）

知识星球：https://articles.zsxq.com/id_fxvgc803qmr2.html 目录 一.定义： 精确模式（默认模式）： 全模式： 搜索引擎模式： paddle 模式（基于深度学习的分词模式）： 二 自定义词典 三.文本解析   调整词出现的频率 四. 关键词提取 A. 基于TF-IDF算法的关键词提取 B. 基于TextRank算法的关键词提取
阅读更多...
使用SecondaryNameNode恢复NameNode的数据

使用SecondaryNameNode恢复NameNode的数据

1）需求： NameNode进程挂了并且存储的数据也丢失了，如何恢复NameNode 此种方式恢复的数据可能存在小部分数据的丢失。 2）故障模拟 （1）kill -9 NameNode进程 [lytfly@hadoop102 current]$ kill -9 19886 （2）删除NameNode存储的数据（/opt/module/hadoop-3.1.4/data/tmp/dfs/na
阅读更多...
HDFS—存储优化（纠删码）

HDFS—存储优化（纠删码）

纠删码原理 HDFS 默认情况下，一个文件有3个副本，这样提高了数据的可靠性，但也带来了2倍的冗余开销。 Hadoop3.x 引入了纠删码，采用计算的方式，可以节省约50％左右的存储空间。 此种方式节约了空间，但是会增加 cpu 的计算。 纠删码策略是给具体一个路径设置。所有往此路径下存储的文件，都会执行此策略。 默认只开启对 RS-6-3-1024k
阅读更多...
Hadoop数据压缩使用介绍

Hadoop数据压缩使用介绍

一、压缩原则 （1）运算密集型的Job，少用压缩 （2）IO密集型的Job，多用压缩 二、压缩算法比较 三、压缩位置选择 四、压缩参数配置 1）为了支持多种压缩/解压缩算法，Hadoop引入了编码/解码器 2）要在Hadoop中启用压缩，可以配置如下参数
阅读更多...
Makefile简明使用教程

Makefile简明使用教程

文章目录 规则makefile文件的基本语法：加在命令前的特殊符号：.PHONY伪目标: Makefilev1 直观写法v2 加上中间过程v3 伪目标v4 变量 make 选项-f-n-C Make 是一种流行的构建工具，常用于将源代码转换成可执行文件或者其他形式的输出文件(如库文件、文档等)。Make 可以自动化地执行编译、链接等一系列操作。 规则 makefile文件
阅读更多...
使用opencv优化图片（画面变清晰）

使用opencv优化图片（画面变清晰）

文章目录 需求影响照片清晰度的因素 实现降噪测试代码 锐化空间锐化Unsharp Masking频率域锐化对比测试 对比度增强常用算法对比测试 需求 对图像进行优化，使其看起来更清晰，同时保持尺寸不变，通常涉及到图像处理技术如锐化、降噪、对比度增强等 影响照片清晰度的因素 影响照片清晰度的因素有很多，主要可以从以下几个方面来分析 1. 拍摄设备 相机传感器：相机传
阅读更多...
JAVA智听未来一站式有声阅读平台听书系统小程序源码

JAVA智听未来一站式有声阅读平台听书系统小程序源码

智听未来，一站式有声阅读平台听书系统 🌟&nbsp;开篇：遇见未来，从“智听”开始 在这个快节奏的时代，你是否渴望在忙碌的间隙，找到一片属于自己的宁静角落？是否梦想着能随时随地，沉浸在知识的海洋，或是故事的奇幻世界里？今天，就让我带你一起探索“智听未来”——这一站式有声阅读平台听书系统，它正悄悄改变着我们的阅读方式，让未来触手可及！ 📚&nbsp;第一站：海量资源，应有尽有 走进“智听
阅读更多...
pdfmake生成pdf的使用

pdfmake生成pdf的使用

实际项目中有时会有根据填写的表单数据或者其他格式的数据，将数据自动填充到pdf文件中根据固定模板生成pdf文件的需求 文章目录 利用pdfmake生成pdf文件1.下载安装pdfmake第三方包2.封装生成pdf文件的共用配置3.生成pdf文件的文件模板内容4.调用方法生成pdf 利用pdfmake生成pdf文件 1.下载安装pdfmake第三方包 npm i pdfma
阅读更多...
零基础学习Redis(10) -- zset类型命令使用

零基础学习Redis(10) -- zset类型命令使用

zset是有序集合，内部除了存储元素外，还会存储一个score，存储在zset中的元素会按照score的大小升序排列，不同元素的score可以重复，score相同的元素会按照元素的字典序排列。 1. zset常用命令 1.1 zadd  zadd key [NX | XX] [GT | LT]   [CH] [INCR] score member [score member ...]
阅读更多...

Copyright China编程 23002807@qq.com

沪ICP备2023033072号-2