问题标签 [ptxas]
For questions regarding programming in ECMAScript (JavaScript/JS) and its various dialects/implementations (excluding ActionScript). Note JavaScript is NOT the same as Java! Please include all relevant tags on your question; e.g., [node.js], [jquery], [json], [reactjs], [angular], [ember.js], [vue.js], [typescript], [svelte], etc.
cuda - CUDA:为 sm_20 显示错误的 lmem 统计信息?
当指定 GPU 架构时,使用该选项编译的 CUDA 内核--ptxas-options=-v
似乎显示错误的 lmem(本地内存)统计信息。sm_20
这同样为架构提供了有意义的 lmem 统计信息sm_10 / sm_11 / sm_12 / sm_13
。
有人可以澄清 sm_20 lmem 统计数据是否需要以不同的方式阅读,或者它们是完全错误的?
这是内核:
--ptxas-options=-v
并sm_20
报告:
--ptxas-options=-v
并sm_10 / sm_11 / sm_12 / sm_13
报告:
sm_20 报告一个4 bytes的 lmem ,如果您看到内核中使用了 4x1000 字节数组,这根本不可能。较旧的 GPU 架构报告正确的4000 字节lmem 统计信息。
这是用CUDA 3.2试过的。我参考了NVCC 手册(v3.2) 的打印代码生成统计部分,但它无助于解释这种异常情况。
cuda - Fermi GPU (GTX 580) 中分析的执行指令和发出指令的奇怪结果
我的内核有这样的 ptx 版本:
我数了一下,我的内核中只有 13 条指令(不包括 ret 指令)。当我将工作项的数量设置为5120时,工作组大小为64。因为有16个SM,每个SM有32个标量处理器,所以上面的代码将在一个SM中执行10次。正如我预期的那样,执行指令的数量应该是10*13 = 130。但是经过我的分析,结果是:发出指令=130,执行指令=100。1. 为什么发出指令的数量与执行指令的数量不同?没有分支,所以它们不应该是平等的吗?2. 为什么执行的指令数比预期的少?至少应该执行 ptx 版本中的所有指令吗?3. 缓存未命中(L1 和 L2)对发出指令的数量和执行的指令数量有影响吗?谢谢
cuda - 避免内联 PTX 中不必要的 mov 操作
在单独的文件中写入 PTX 时,可以将内核参数加载到寄存器中:
但是,当使用内联 PTX 时,在 CUDA(版本 01)中使用内联 PTX 程序集应用说明描述了一种语法,其中加载参数与另一个操作密切相关。它提供了这个例子:
生成:
在许多情况下,有必要将这两个操作分开。例如,可能希望将参数存储在循环外的寄存器中,然后在循环内重用和修改寄存器。我发现这样做的唯一方法是使用额外的 mov 指令,将参数从隐式加载的寄存器移动到我以后可以使用的另一个寄存器。
从单独文件中的 PTX 移动到内联 PTX 时,有没有办法避免这个额外的 mov 指令?
cuda - 在内联 PTX 中设置 32 位地址大小
我正在将作为单独文件编写的 PTX 转换为内联 PTX。在单独的 PTX 文件中,我将 ISA 和目标定义如下:
在编译器生成的 PTX 文件中,内联 PTX 后,编译器指定了 ISA 和 target,如下所示:
这.address_size 64
对我来说是有问题的,因为这意味着我必须将我在内联 PTX 中执行的指针算法从 32 位更新为 64 位。
鉴于 32 位可以寻址 4GB,比我的卡更多的内存,是否可以让编译器指定 32 位地址大小,这样我就不必更新指针算法?
考虑到新的统一寻址系统,sm_20 是否支持 32 位地址?
cuda - 在 Parallel Nsight 中调试内联 PTX
在 Parallel Nsight 中单步执行内联 PTX 时是否可以查看 PTX 寄存器?
我可以在内联 PTX 上设置断点并单步执行,但将鼠标悬停在 PTX 寄存器上不会显示它们的值。我可以打开 SASS 并将鼠标悬停在这些寄存器上确实会显示它们的值,但是很难跟踪 PTX 寄存器和 SASS 寄存器之间的关系。
c++ - OpenCL:指令和地址之间的状态空间不匹配
我正在编写一个 OpenCL 程序,并在构建时收到此错误:
相应的 ptx 行(自动生成)是:
这是我编写的函数:
错误的原因是什么?如何解决?
cuda - 在 NVIDIA GPU 上编译 ptx 代码?
我想在 NVIDIA GPU 上拦截 opencl 程序的 PTX 级别。
我想例程可能看起来像这样。
首先,我编写了一个 opencl 程序(主机和设备代码),使用 NVIDIA 编译器生成各自的 ptx 代码。然后我通过修改PTX代码来写我想做的事情(请不要问我为什么没有在设备C代码上这样做——我有一些原因)。但问题是,修改后,如何将这个 PTX 代码编译成二进制代码?
if-statement - if 的额外注册使用
我正在研究一个大型 cuda 内核,我注意到内核每个线程使用 43 个寄存器。为了了解发生了什么,我编写了一个较小的程序来计算寄存器的使用情况。我注意到每当我使用 时if
,寄存器的使用率都会上升。小代码如下:
当我编译这段代码时,每个线程使用 5 个寄存器
但是,如果我取消注释if
,每个线程使用 8 个寄存器。谁能向我解释发生了什么事?
memory - 解释 ptxas 的详细输出,第一部分
我正在尝试了解我的每个 CUDA 线程的资源使用情况,以用于手写内核。
我将我kernel.cu
的文件编译成一个kernel.o
文件nvcc -arch=sm_20 -ptxas-options=-v
我得到以下输出(通过c++filt
):
看上面的输出,这样说对吗
- 每个 CUDA 线程使用 46 个寄存器?
- 没有寄存器溢出到本地内存?
我在理解输出方面也有一些问题。
我的内核正在调用很多
__device__
函数。__global__
和__device__
函数的堆栈帧的内存总和是 72 字节吗?0 byte spill stores
和有什么区别0 bytes spill loads
为什么信息
cmem
(我假设是恒定记忆)用不同的数字重复两次?在内核中,我没有使用任何常量内存。这是否意味着编译器会在后台告诉 GPU 使用一些常量内存?
这个问题在以下内容中“继续”:解释 ptxas 的详细输出,第二部分
c - __internal_trig_reduction_slowpathd 的函数属性
目前我正在尝试优化一些cuda内核......
如果使用选项--ptxas-options=-v 编译,我会得到有关寄存器%co 的信息。
就我而言,我总是得到一些额外的行,这对我来说毫无意义:
第 1 到 4 行对我来说很清楚,但最后一行是什么?
谷歌在这里没有帮助....我已经尝试过了。
有人知道这些行的含义是什么吗?我为程序中的每个编译内核获取它们