问题标签 [ptx]

在 CUDA PTX 中，有一个特殊的寄存器保存线程的 warp 索引：%warpid. 现在，规范说：

请注意，它%warpid 是易失的，并在读取时返回线程的位置，但其值可能会在执行期间发生变化，例如，由于抢占后线程的重新调度。

嗯，那是什么位置？它不应该是块内的位置，例如一维网格%tid.x / warpSize吗？它是 SM 中的一些 slot-for-a-warp（例如 warp 调度程序或一些内部队列）吗？我很困惑。

动机：我想%tid.x / warpSize通过使用这个特殊的寄存器来省去计算和释放寄存器的麻烦。

我正在通过 c++filt 过滤我编译的 PTX，但它只会破坏一些名称/标签并保留一些原样。例如，这个：

被分解为：

而不是至少这个：

我意识到 c++filt 没有明确支持 CUDA PTX；但请注意，未分解的名称与示例中的已分解名称的区别仅在于添加_param_0等_param_1后缀（还有一个问题是应该如何分解这些名称的前缀，但让我们忘记这一点）。

• 我可以做些什么来强制 c++filt 也适用于参数名称/标签吗？更一般地说，对于 PTX 文件中所有损坏的 C++ 名称？
• 除了它已经具有的“格式”（[-s|--format {none,auto,gnu,lucid,arm,hp,edg,gnu-v3,java,gnat,dlang}]）之外，是否有可能/容易通过对 CUDA“格式”的认识来增强 c++filt？
• 如果 c++filt 在这种情况下不能使用或不能适应使用，我应该如何进行解构？

我正在尝试使用 CUDA SDK 中的 ptxjit 示例作为检测与 GPU 设备交互的基础。

我已经成功地编译了检测代码，并控制设备加载和执行具有 CUDA 功能 2.0 的 Geforce GT440 的 PTX 模块。

在具有 CUDA 功能 5.0 的 Geforce 830M 系统（笔记本电脑使用 bumblebee 控制离散 GPU）上编译相同的仪器代码时，代码编译但给了我 209（CUDA_ERROR_NO_BINARY_FOR_GPU）。

我试图编译内核以与 CUDA 功能 5.0 兼容，但没有成功，仍然是同样的错误。

有任何想法吗？

运行设备端 CUDA 代码是否有可能知道为运行内核网格的每个块分配了多少（静态和/或动态）共享内存？

在主机端，您知道启动的内核拥有（或将拥有）多少共享内存，因为您自己设置了该值；但是设备端呢？在该大小的上限中很容易编译，但该信息不可用（除非明确传递）给设备。是否有获取它的 on-GPU 机制？CUDA C Programming Guide似乎没有讨论这个问题（在共享内存部分之内或之外）。

我正在查看PTX ISA 规范（针对 PTX v2）中的存储指令缓存模式表。它提供了有关四种缓存模式的详细信息：

• .wb：缓存回写所有相干级别

• .cg：在全局级别缓存（缓存在 L2，而不是 L1）

• .cs: 缓存流，可能被访问一次

• .wt：缓存直写（到系统内存）

解释性文字让我有些困惑。

对于正在写入的行（相对于需要驱逐的其他行），哪些模式会导致立即更改 L1 缓存、L2 缓存和全局内存？我会将我从文本中理解的内容发布为答案，如果我错了，请纠正我。

注意：为了这个问题，让我们忽略本地和共享内存。

PTX 是一种中间表示，用于将 C/C++ GPU 代码最终编译为单个微架构的 SASS 汇编语言。因此，它不应该受到特定 nVIDIA GPU 微架构的实际指令集中的特定孔/gaffs/flukes/特质的阻碍。

现在，PTX 有一条指令用于计算寄存器中前导零的数量：clz. 然而 - 它缺少相应的ctz指令，该指令计算尾随零的数量。这些操作是“对称的”，人们肯定会期望在指令集中看到两者或一个都没有——同样，特别是如果它是抽象的并且不受特定硬件上可用的东西的约束。流行的 CPU 架构已经拥有这两种架构多年了。

奇怪的是，CUDA 标头device_functions.h声明了函数

这个功能：

• 具有与计数尾随零几乎相同的语义 - 仅在全零输入上有所不同。
• 不会转换为单个 PTX 指令，而是两个：按位否定，然后是clz.
• 也错过了它的潜在对手，__fls- 找到最后一组。

那么，这是为什么呢？为什么 PTX 中缺少一个明显显而易见的指令，以及一个与标题中几乎相同的“假内置”？

当我们在我们通常的 C/C++ CUDA 代码中编写内联 PTX 程序集时，例如：

我们可以在volatile之后添加关键字asm，例如：

内联 PTX 程序集的 CUDA文档说：

编译器假定asm()语句除了更改输出操作数外没有副作用。为确保asm在 PTX 生成过程中不被删除或移动，您应该使用 volatile 关键字

我不明白那是什么意思。所以，

• 为什么我的asm()会被删除？或者更确切地说，如果编译器注意到它没有效果，我为什么要介意它被删除？
• 为什么asm()在 PTX 生成过程中移动 my 会出现问题？这是优化过程的一部分，不是吗？
• 当分别面对非易失性和易失性asm()指令时，如何更准确地描述编译器的行为？

随着 gcc 7.1 的发布，我们现在可以为 openmp 4.5 配置 gcc，卸载到 Nvidia PTX GPGPU。这就是他们在发行说明中所说的（大约）。

所以我的问题是，在编译 openmp 4.5 以针对 nvidia 设备时，是否有任何特殊标志来激活此配置？

In CUDA, each thread knows its block index in the grid and thread index within the block. But two important values do not seem to be explicitly available to it:

• Its index as a lane within its warp (its "lane id")
• The index of the warp of which it is a lane within the block (its "warp id")

Assuming the grid is 1-dimensional(a.k.a. linear, i.e. blockDim.y and blockDim.z are 1), one can obviously obtain these as follows:

and if you don't trust the compiler to optimize that, you could rewrite it as:

is that the most efficient thing to do? It still seems like a lot of waste for every thread to have to compute this.

_{(inspired by this question.)}

在研究使用方式时，我正在摆弄一些 SASS %laneid。在浪费了某人生命的一分钟失态之后（对不起 - 你知道你是谁），我现在有以下几点：

CUDA 代码：

SASS（对于 SM 6.1）：

因此，STG 指令——我猜想存储到全局内存——不会立即使用 SR_LANEID，而是使用内联 PTX 将其放入的寄存器。这是因为（Pascal）GPU 无法从特殊寄存器中存储，还是错过了优化机会？