问题标签 [ptx]

我想在 CUDA C 代码中使用汇编代码，以减少昂贵的执行，就像我们在 c 编程中使用asm一样。

是否可以？

我了解 Fermi GPU 支持预取到 L1 或 L2 缓存。但是，在 CUDA 参考手册中我找不到任何关于它的信息。

Dues CUDA 允许我的内核代码将特定数据预取到特定级别的缓存？

我知道如何从 a 生成.ptx文件.cu以及如何从 a 生成.cubin文件.ptx.但我不知道如何获得最终的可执行文件。

更具体地说，我有一个sample.cu文件，它被编译为sample.ptx. 然后我使用 nvcc 编译sample.ptx为sample.cubin. 但是，.cubin如果没有主机代码，则无法直接执行此文件。如何将.cubin文件链接到我的原始.cu文件以生成最终的可执行文件？

假设我有一个代码可以让用户通过threads_per_block调用内核。然后我想检查输入是否有效（例如 <=512 表示计算能力 CC <2.0 和 1024 表示 CC >=2.0）。

现在我想知道如果我nvcc -arch=sm_13在我的计算机中使用 CC2.0 的图形卡编译代码时会发生什么，当用户通过时threads_per_block == 1024？这是：

• 一个有效的输入 - 因为我运行的卡有 CC2.0，或者......
• 无效，因为我为 CC1.3 编译它？

还是nvcc -arch=sm_13仅仅意味着 CC1.3 至少是必要的，但是在更高的 CC 上运行它时，尽管可以使用那些更高的功能？

我安装了 CUDA 4.0，以及具有 Compute Capability 2.0 的设备（GTX 460 卡）。

“cubin”和“ptx”文件有什么区别？

我认为 cubin 是 gpu 的本机代码，因此这是特定于微架构的，而 ptx 是通过 JIT 编译在 Fermi 设备（例如 Geforce GTX 460）上运行的中间语言。当我编译.cu源文件时，我可以在 ptx 或 cubin 目标之间进行选择。如果我想要 cubin 文件，我选择code=sm_20. 但是如果我想要一个 ptx 文件，我会使用code=compute_20.

这是对的吗？

您是否建议阅读内核的 PTX 代码以进一步优化内核？

一个例子：我读到，可以从 PTX 代码中找出自动循环展开是否有效。如果不是这种情况，则必须在内核代码中手动展开循环。

• PTX 代码还有其他用例吗？
• 您查看您的 PTX 代码吗？
• 在哪里可以找到如何读取 CUDA 为我的内核生成的 PTX 代码？

有人知道如何使用新的 LLVM 后端使用 C/C++ 代码注释 PTX 汇编器吗？

可以使用 CUDA 4.0 或更早版本轻松获得它，但 NVCC 在将 CUDA 工具包升级到 4.2 版后拒绝了我的所有标志。

正如Nvidia的Inline PTX Assembly文档中所写，使用内联汇编的语法是： asm("temp_string" : "constraint"(output) : "constraint"(input));
这里有两个例子：
asm("vadd.s32.s32.s32 %0, %1.h0, %2.h0;" : "=r"(v) : "r"(a), "r"(b));
asm("vadd.u32.u32.u32 %0.b0, %1, %2, %3;" : "=r"(v) : "r"(a), "r"(b), "r"(z));
在这两个例子中，都有参数比如：h0或者b0跟在%n. 我浏览了 CUDA 的官方文档，并没有发现对h0or的含义有任何顾虑b0。我见过h0, h1, b0, b1, b2. b3我猜h0还是h1代表一个 16 位的值，而bn代表一个字节的值。有人知道这些的确切含义吗？

感谢罗杰达尔的帮助。我阅读了 PTX ISA 3.0 并找到了答案。
“h”表示半字。h0表示 32 位字的低半字。h1表示 32 位字的高半字。“b”表示整数字节。b0, b1,b2分别b3代表一个 32 位字的前 8 位、第二个 8 位、第三个 8 位和最高 8 位。

编辑：这个问题是原版的重做版本，所以前几个回答可能不再相关。

我很好奇强制非内联的设备函数调用对设备函数内的同步有何影响。我有一个简单的测试内核来说明有问题的行为。

内核获取一个缓冲区并将其传递给设备函数，以及一个共享缓冲区和一个指示变量，该变量将单个线程标识为“boss”线程。设备函数有不同的代码：boss线程首先花时间在共享缓冲区上做一些琐碎的操作，然后写入全局缓冲区。在同步调用之后，所有线程都写入全局缓冲区。在内核调用之后，主机打印全局缓冲区的内容。这是代码：

CUDA 代码：

test_main.cu

test_kernel.cu

我从 CUDA SDK 中编译了这段代码，以利用 test_main.cu 中的“cutilsafecall()”函数，但如果您想在 SDK 之外编译，当然可以取出这些代码。我使用 CUDA 驱动程序/工具包 4.0 版、计算能力 2.0 版进行编译，代码在具有 Fermi 架构的 GeForce GTX 480 上运行。

预期的输出是

0 1 2 3 ... blockDim.x-1

但是，我得到的输出是

1 1 2 3 ... blockDim.x-1

这似乎表明boss线程执行了条件“scratchBuffer[0] = 1;” 语句 AFTER 所有线程执行“scratchBuffer[threadIdx.x] = threadIdx.x;” 声明，即使它们被 __syncthreads() 屏障隔开。

即使老板线程被指示将标记值写入同一线程中线程的缓冲区位置，也会发生这种情况；sentinel 是缓冲区中存在的最终值，而不是适当的 threadIdx.x 。

导致代码产生预期输出的一种修改是更改条件语句

如果（是老板）{

到

如果（IS_BOSS（））{

; 即，将散度控制变量从存储在参数寄存器中更改为在宏函数中计算。（请注意源代码中相应行的注释。）我一直在关注这个特殊的变化，试图找出问题所在。在查看带有“isBoss”条件（即损坏的代码）和“IS_BOSS()”条件（即工作代码）的内核的反汇编 .cubin 时，指令中最显着的区别似乎是没有反汇编的损坏代码中的 SSY 指令。

以下是使用 "cuobjdump -sass test_kernel.cubin" 反汇编 .cubin 文件生成的反汇编内核。第一个“EXIT”之前的所有内容都是内核，之后的所有内容都是设备功能。唯一的区别在于设备功能。

分解的对象代码：

“破”代码

“工作”代码

“SSY”指令存在于工作代码中，但不存在于损坏的代码中。cuobjdump 手册用“设置同步点；在可能不同的指令之前使用”来描述该指令。这让我认为，由于某种原因，编译器无法识别损坏代码中存在分歧的可能性。

我还发现，如果我注释掉 __noinline__ 指令，那么代码会产生预期的输出，实际上由其他“损坏”和“工作”版本产生的程序集完全相同。所以，这让我觉得当一个变量通过调用堆栈传递时，该变量不能用于控制分歧和随后的同步调用；在这种情况下，编译器似乎没有认识到分歧的可能性，因此没有插入“SSY”指令。有谁知道这是否确实是 CUDA 的合法限制，如果是，是否在任何地方都有记录？

提前致谢。

CUDA 2.0 设备上是否有任何方法可以仅为一个特定变量禁用 L1 缓存？我知道可以在编译时禁用 L1 缓存，为所有内存操作添加-Xptxas -dlcm=cg标志nvcc。但是，我只想对特定全局变量的内存读取禁用缓存，以便所有其余的内存读取都通过 L1 缓存。

根据我在网络上进行的搜索，可能的解决方案是通过 PTX 汇编代码。