cuda调试(一)vs2019-windows-Nsight system--nvtx使用,添加nvToolsExt.h文件

news2024/11/17 19:51:11

cuda调试

由于在编程过程中发现不同的网格块的结构,对最后的代码结果有影响,所以想记录一下解决办法。

CUDA的Context、Stream、Warp、SM、SP、Kernel、Block、Grid

cuda context (上下文)

context类似于CPU进程上下,表示由管理层 Drive 层分配的资源的生命周期,多线程分配调用的GPU资源同属一个context下,通常与CPU的一个进程对应。

CUDA Stream

CUDA Stream是指一堆异步的CUDA操作,他们按照host代码调用的顺序执行在device上。
Stream维护了这些操作的顺序,并在所有预处理完成后允许这些操作进入工作队列,同时也可以对这些操作进行一些查询操作。
这些操作包括host到device的数据传输,launch kernel以及其他的由host发起由device执行的动作。
这些操作的执行总是异步的,CUDA runtime会决定这些操作合适的执行时机。我们则可以使用相应的cuda api来保证所取得结果是在所有操作完成后获得的。同一个stream里的操作有严格的执行顺序,不同的stream则没有此限制。

CUDA API可分为同步和异步两类,同步函数会阻塞host端的线程执行,异步函数会立刻将控制权返还给host从而继续执行之后的动作。

当我们使用CUDA异步函数与多流(Multi Stream)时,多线程间既可以实现并行进行数据传输与计算,如下图所示。不过需要注意的是, CUDA runtime API默认的default stream是同步串行的,且一个进程内的所有线程都在default stream下,需要显式声明default之外的Stream才可以实现多流并发。
在这里插入图片描述

显卡硬件架构:SM、SP、Warp

具体到nvidia硬件架构上,有以下两个重要概念:

SP(streaming processor):最基本的处理单元,也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算,也就是很多个SP同时做处理。

SM(streaming multiprocessor):多个SP加上其他的一些资源组成一个SM,也叫GPU大核,其他资源如:warp scheduler,register,shared memory等。SM可以看做GPU的心脏(对比CPU核心),register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此,这些有限的资源就使每个SM中active warps有非常严格的限制,也就限制了并行能力。如下图是一个SM的基本组成,其中每个绿色小块代表一个SP。
在这里插入图片描述
每个SM包含的SP数量依据GPU架构而不同,Fermi架构GF100是32个,GF10X是48个,Kepler架构都是192个,Maxwell都是128个。当一个kernel启动后,thread会被分配到很多SM中执行。大量的thread可能会被分配到不同的SM,但是同一个block中的thread必然在同一个SM中并行执行。

Warp调度

一个SP可以执行一个thread,但是实际上并不是所有的thread能够在同一时刻执行。Nvidia把32个threads组成一个warp,warp是调度和运行的基本单元。warp中所有threads并行的执行相同的指令。一个warp需要占用一个SM运行,多个warps需要轮流进入SM。由SM的硬件warp scheduler负责调度。目前每个warp包含32个threads(Nvidia保留修改数量的权利)。所以,一个GPU上resident thread最多只有 SM*warp个

同一个warp中的thread可以以任意顺序执行,active warps被SM资源限制。当一个warp空闲时,SM就可以调度驻留在该SM中另一个可用warp。在并发的warp之间切换是没什么消耗的,因为硬件资源早就被分配到所有thread和block,所以新调度的warp的状态已经存储在SM中了。

每个SM有一个32位register集合放在register file中,还有固定数量的shared memory,这些资源都被thread瓜分了,由于资源是有限的,所以,如果thread比较多,那么每个thread占用资源就叫少,thread较少,占用资源就较多,这需要根据自己的要求作出一个平衡。

软件架构:Kernel、Grid、Block

上面的context与stream类似进程、线程的概念,具体到我们如何调用GPU上的线程实现我们的算法,则是通过Kernel实现的。在GPU上调用的函数成为CUDA核函数(Kernel function),核函数会被GPU上的多个线程执行。我们可以通过如下方式来定义一个kernel:

func_name<<<grid, block>>>(param1, param2, param3....);

在这里插入图片描述
Grid:由一个单独的kernel启动的所有线程组成一个grid,grid中所有线程共享global memory。Grid由很多Block组成,可以是一维二维或三维。

Block:一个grid由许多block组成,block由许多线程组成,同样可以有一维、二维或者三维。block内部的多个线程可以同步(synchronize),可访问共享内存(share memory)。

CUDA中可以创建的网格数量跟GPU的计算能力有关,可创建的Grid、Block和Thread的最大数量如下所示:
在这里插入图片描述
以上是引用:

https://zhuanlan.zhihu.com/p/266633373

在这里插入图片描述
在这里插入图片描述
通过deviceQuery可以知道,每个block中只有1024个thread,而32 × 32× 32 = 32768 > 1024:会出现以下错误,运行时参数传递得太大了,出现这类问题后,cuda仍可继续提供服务,仅单纯拒绝了启动核函数

cudaErrorInvalidConfiguration = 9,"invalid configuration argument"

所有CUDA kernel的启动都是异步的,当CUDA kernel被调用时,控制权会立即返回给CPU。在分配Grid、Block大小时,我们可以遵循这几点原则

  1. 保证block中thread数目是32的倍数。这是因为同一个block必须在一个SM内,而SM的Warp调度是32个线程一组进行的。
  2. 避免block太小:每个blcok最少128或256个thread。
  3. 根据kernel需要的资源调整block,多做实验来挖掘最佳配置。
  4. 保证block的数目远大于SM的数目。

配置Nsight system

低开销的性能分析工具,Nvidia nsight Systems旨在提供开发人员优化其软件所需的洞察力。在工具中可视化活动数据,以帮助用户调查瓶颈,避免推断误报,并以更高的性能提高概率进行优化。用户将能够识别问题,例如GPU不足、不必要的GPU同步、CPU并行化不足,甚至目标平台的CPU和GPU之间的算法异常。

  1. 打开该程序
  2. 上面的 command line with arguments中填写 要调试的cuda项目的exe文件我的目录是F:\E_cuda\test1\testCUDA\x64\Debug\testCUDA.exe
    注意这个路径需要是全英文的
    3.下面的 working directory 填写 cuda中的Debug文件,我的cuda中的默认目录是
    C:\ProgramData\NVIDIA Corporation\CUDA Samples\v11.0\bin\win64\Debug

在这里插入图片描述

F:\E_cuda\test1\testCUDA\x64\Debug\testCUDA.exe

在这里插入图片描述
然后start就行了
在这里插入图片描述
但是还不知道怎么去看具体的kernel

在这里插入图片描述
这个Report包含5部分内容:

  1. Analysis Summary (分析总结,内容非常全面,包含了Target的详细信息,Process summary, Module summary, Thread summary, Environment Variables, CPU info, GPU info等等)
  2. Timeline View (展示CPU/GPU各个核的工作时间线,一般用来来勘察模型训练或者推理的瓶颈在哪里)
  3. Diagnostics Summary (顾名思义,诊断总结。就是程序在运行中做了什么,有什么warning , error,或者message的,都在这里汇总)
  4. Symbol Resolution Logs(暂时不知道是干嘛的)
    Files (执行结果的log 文件:pid_stdout.log,& 执行出错的log 文件pid_stderr.log)

简单看看Timeline view。如下,这里有三个CPU核在工作,它们启动和停止的时间可以从timeline上看到。还可以看到下面有三个Thread的时间线。
在这里插入图片描述

https://blog.csdn.net/NXHYD/article/details/112915968

这位博主的图和说明

我在b站上找的视频发现别人用的是 nsight system profiles 这个,(还没找到是啥,只有一个visoul profiler—需要java的环境)但是他讲的时候又说是NVTX,有点晕

https://www.bilibili.com/video/BV13w411o7cu/?spm_id_from=333.337.search-card.all.click&vd_source=0a4d8c47345ce0df71cd9cdb01575134

就是这个
以下是我使用visoul profiler—的报错
在这里插入图片描述
英伟达的官方文档关于nsight-system的

https://docs.nvidia.com/nsight-systems/2020.3/profiling/index.html

NVTX是一种CUDA Profiler的工具

可以用于在CUDA程序中进行标记和注释,以便更好地理解和优化程序的性能。以下是使用NVTX的一些示例代码和步骤:

  1. 在CUDA代码中包含nvToolsExt.h文件。
    #include <nvToolsExt.h>
    解决:右键cuda项目——属性——配置属性——C/C++——常规——附加包含目录
    可能遇到的问题:无法在 Visual Studio 中打开 nvToolsExt.h 文件,可能是因为 Visual Studio 找不到 CUDA 的 include 目录。

错误示范
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\include
but我这里没有这个头文件
在这里插入图片描述
C:\ProgramData\NVIDIA Corporation\CUDA Samples\v11.0\common\inc
这个目录下也没有
在这里插入图片描述

正确示范 可以通过在vs项目属性里面进行设置

第一步:在项目属性页,打开vc++目录(c程序是c目录会略微有些不同),打开外部包含目录,输入:$(NVTOOLSEXT_PATH)\include
在这里插入图片描述
第二步,打开链接器–常规----附加库目录,输入:$(NVTOOLSEXT_PATH)\lib\$(Platform)
在这里插入图片描述
第三步,打开链接器–输入—附加依赖项,输入:nvToolsExt64_1.lib
在这里插入图片描述
在 .cu 文件中注释代码,还应该设置:打开CUDA C/C+±-common—附加依赖项,输入:$(NVTOOLSEXT_PATH)\include
在这里插入图片描述
在这里插入图片描述
英伟达官方NVTX文档

https://docs.nvidia.com/nsight-visual-studio-edition/nvtx/index.html

  1. 在需要标记的代码块前后插入标记函数。
    nvtxRangePushA("My Code Block"); // 开始标记代码块 nvtxRangePop(); // 结束标记

  2. 可以使用nvtxEventAttributes_t结构体自定义标记的颜色和描述等信息。

nvtxEventAttributes_t eventAttrib = {0};
eventAttrib.version = NVTX_VERSION;
eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE;
eventAttrib.colorType = NVTX_COLOR_ARGB;
eventAttrib.color = 0xFF00FF00;
eventAttrib.messageType = NVTX_MESSAGE_TYPE_ASCII;
eventAttrib.message.ascii = "My Custom Event";
nvtxRangePushEx(&eventAttrib);
// 代码块
nvtxRangePop();

  1. 还可以使用nvtxMarkA()函数在代码中插入注释。
nvtxMarkA("My Comment");

  1. 在运行程序时,可以使用nvprof或Nvvp等工具来查看标记和注释。
    注意:使用NVTX可能会对程序的性能产生一定的影响

cuda运行时间记录

在python里面
异步计算的结果是没有同步的时间测量是不准确的。要获得精确的测量值,应该在测量前调用 torch.cuda.synchronize(),或者使用 torch.cuda.Event 记录时间,如下所示:

 
 
start_event = torch.cuda.Event(enable_timing=True)
end_event = torch.cuda.Event(enable_timing=True)
start_event.record()
 
# 在这里运行一些东西
 
end_event.record()
torch.cuda.synchronize() # 等待事件被记录!
elapsed_time_ms = start_event.elapsed_time(end_event)
 

https://blog.csdn.net/dedell/article/details/121574306

先这样,后续会补充,老师让我去搞二维的共享内存的程序,祝我好运!!!

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

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

相关文章

信息系统项目管理师——OSI模型、协议

OSI七层网络模型&#xff08;掌握&#xff09; 应用层&#xff1a;对软件提供接口以使程序能使用网络服务&#xff0c;如事务处理程序、文件传送协议和网络管理等。&#xff08;HTTP、Telnet、FTP、SMTP&#xff09;表示层&#xff1a;程序和网络之间的翻译官&#xff0c;管理…

可选链运算符(?.)与空值合并运算符(??)

1. 可选链运算符Optional chaining(?.) MDN定义 可选链运算符&#xff08;?.&#xff09;允许读取位于连接对象链深处的属性的值&#xff0c;而不必明确验证链中的每个引用是否有效。?. 运算符的功能类似于 . 链式运算符&#xff0c;不同之处在于&#xff0c;在引用为空 (n…

JavaScript事件循环及任务处理

JavaScript事件循环及任务处理## 浏览器中 JavaScript 的执行流程和 Node.js 中的流程都是基于 事件循环 的。 理解事件循环的工作方式对于代码优化、性能优化很重要&#xff0c;有时对于正确的架构也很重要。 我们首先介绍事件循环工作方式的理论细节&#xff0c;然后介绍该知…

【数据库】MySQL数据库基础

目录 1.数据库&#xff1a; 2.数据库基本操作 2.1 MySQL的运行原理 2.2显示数据库&#xff1a; 2.3创建数据库 2.4使用数据库 2.5删除数据库 3.常见的数据类型 3.1数值类型&#xff1a; 3.2字符型类型 3.3日期类型 4.表的操作 4.1创建表 4.2查看表 4.3删除表 5.汇总…

【蓝桥杯PythonB组备赛】【Acwing周赛】第93场 4867. 整除数 4868. 数字替换 python解

目录 A AcWing 4867. 整除数 1.题目描述 2.思路分析 3.代码实现 B AcWing 4868. 数字替换 1.题目描述 2.思路分析 3.代码实现 A AcWing 4867. 整除数 1.题目描述 2.思路分析 为什么不能直接暴力&#xff1f; 数据&#xff1a;1 ≤ n, k ≤ 10 ** 9 1s内最多…

openpnp - 设备开机后, 吸嘴校验失败的解决方法

文章目录openpnp - 设备开机后, 吸嘴校验失败的解决方法概述重新校验吸嘴ENDopenpnp - 设备开机后, 吸嘴校验失败的解决方法 概述 设备开机后, 默认会校验吸嘴座上已经安装的2个吸嘴. 如果开机校验吸嘴失败, 就需要用向导重新校验失败的吸嘴. 具体是哪个吸嘴校验失败, 可以看…

C++的双端队列

双端队列介绍1.双端队列知识需知2.大试牛刀1.双端队列知识需知 由于队列是一种先进先出&#xff08;FIFO&#xff09;的数据结构&#xff0c;因此无法直接从队列的底部删除元素。如果希望从队列的底部删除元素&#xff0c;可以考虑使用双端队列&#xff08;deque&#xff09;。…

SpringCloud:Ribbon

目录 SpringCloud&#xff1a;Ribbon 1、负载均衡原理 2、负载均衡策略 3、饥饿加载 SpringCloud&#xff1a;Ribbon 1、负载均衡原理 2、负载均衡策略 Ribbon的负载均衡规则是一个叫做IRule的接口来定义的&#xff0c;每一个子接口都是一种规则; 内置负载均衡celue规则描…

2023年“网络安全”赛项江苏省淮安市选拔赛 任务书

任务书 一、竞赛时间 共计3小时。 二、竞赛阶段 竞赛阶段 任务阶段 竞赛任务 竞赛时间 分值 第一阶段单兵模式系统渗透测试 任务一 服务器内部信息获取 任务二 网站渗透测试 任务三 Linux系统渗透提权 任务四 Web渗透测试 第二阶段分组对抗 备战阶段 攻防对抗准备工作 系统加…

python+嵌入式——串口通信篇(收发解包)

目录前言安装pyserialpyserial大致概括整体流程硬件连接例子(简单版)详细使用serial初始化参数发包收包收包检查包并解包python struct模块结语前言 这几年&#xff0c;自己也做了一些嵌入式机器人。在整个开发的过程中&#xff0c;调通信通常会花费一段比较长的时间&#xff…

Docker是什么,怎么用?

1、docker解决了什么的问题&#xff1f; Docker是属于运维和持续集成的管理工具&#xff0c;特别是在运维方面使用Docker可以让你省去很多重复的工作&#xff0c;可以让你的应用程序永远都是在相同的环境中运行。 统一标准 ● 应用构建 ○ Java、C、JavaScript ○ 打成软件包 …

【刷题笔记】之二分查找(搜索插入位置。在排序数组中查找元素的第一个和最后一个位置、x的平方根、有效的完全平方数)

1. 二分查找题目链接 704. 二分查找 - 力扣&#xff08;LeetCode&#xff09;给定一个 n 个元素有序的&#xff08;升序&#xff09;整型数组 nums 和一个目标值 target &#xff0c;写一个函数搜索 nums 中的 target&#xff0c;如果目标值存在返回下标&#xff0c;否则返回 -…

《Linux运维实战:Mysql8.0.30安装及卸载component_validate_password插件》

一、背景 由于业务系统的特殊性&#xff0c;我们需要将MySQL8.0.30主从复制集群部署在客户机房服务器上&#xff0c;且客户对数据库的密码策略要求比较高。 因为官方的8.0.30版本的Docker镜像默认是没有安装validate_password插件的&#xff0c;所以我在主从复制集群安装完成后…

minGW-w64配置途径

文章目录1 GNU、GCC与minGW2 minGW当前下载方式3 minGW-w64配置途径Step1Step2Step31 GNU、GCC与minGW GNU这个名字是GNUs Not Unix的递归首字母缩写&#xff0c;它的发音为[gnoo]&#xff0c;只有一个音节&#xff0c;发音很像"grew"&#xff0c;但需要把其中的r音替…

aws apigateway 基础概念和入门示例

参考资料 https://docs.aws.amazon.com/zh_cn/apigateway/latest/developerguide/getting-started.html apigateway基础理解 apigateway的核心概念 apigateway&#xff0c;基础服务用来管理接口的创建&#xff0c;部署和管理restapi&#xff0c;http资源和方法的集合&#…

【LeetCode】1599. 经营摩天轮的最大利润

1599. 经营摩天轮的最大利润 题目描述 你正在经营一座摩天轮&#xff0c;该摩天轮共有 4 个座舱 &#xff0c;每个座舱 最多可以容纳 4 位游客 。你可以 逆时针 轮转座舱&#xff0c;但每次轮转都需要支付一定的运行成本 runningCost 。摩天轮每次轮转都恰好转动 1 / 4 周。 …

7 Seata简介

Seata-Server安装 分布式事务解决方案 2PC即两阶段提交协议&#xff0c;是将整个事务流程分为两个阶段&#xff0c;P是指准备阶段&#xff0c;C是指提交阶段。 1. 准备阶段&#xff08;Prepare phase&#xff09; 2. 提交阶段&#xff08;commit phase&#xff09;举例&…

模电基础(2)半导体二极管

1.二极管的组成二极管&#xff1a;将PN结封装起来&#xff0c;引出两个电极就构成了半导体二极管。二极管的常见结构包括&#xff1a;点接触型&#xff08;图a&#xff09;&#xff0c;面接触型&#xff08;图b&#xff09;&#xff0c;平面型&#xff08;图c&#xff09;。 点…

世界顶级五大女程序媛,不仅技术强还都是美女

文章目录1.计算机程序创始人&#xff1a;勒芙蕾丝伯爵夫人2.首位获得图灵奖的女性&#xff1a;法兰艾伦3.谷歌经典首页守护神&#xff1a;玛丽莎梅耶尔4.COBOL之母&#xff1a;葛丽丝穆雷霍普5.史上最强游戏程序媛-余国荔说起程序员的话&#xff0c;人们想到的都会是哪些理工科…

java基础-标识符命名规范和数据类型

标识符 1.什么是标识符&#xff1f; Java中变量、方法、类等要素命名时使用的字符序列&#xff0c;称为标识符。 技巧&#xff1a;凡是自己可以起名字的地方都叫标识符。比如&#xff1a;类名、方法名、变量名、包名、常量名等 2.标识符的命名规则 1.标识符由26个英文字母大小…