问题标签 [ptx]

我具有以下内联汇编功能，可在 32 位 Visual Studio 2008 的调试模式下正常工作：

pa 和 pb 在设备上全局分配，例如

但是，此代码在发布模式下失败，在线asm("ld.global.b32 r1, [s0+8];"::); 如何在发布模式下使用内联 ptx 正确加载函数参数？

PS 使用 -G 标志（生成 GPU 调试信息）构建发布模式会导致代码在发布模式下正确运行。谢谢，

我在转换.cu为.ptx. 我使用nvcc如下：

"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\bin\nvcc" -ptx -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\bin" -o foo.ptx foo.cu

作为回报显示以下内容：

foo.cu位于\CUDA\v5.0\bin.

:) 当我试图管理我的内核资源时，我决定研究 PTX，但有几件事我不明白。这是我写的一个非常简单的内核：

然后我使用编译它：nvcc --ptxas-options=-v -keep main.cu我在控制台上得到了这个输出：

结果 ptx 如下：

现在有一些我不明白的事情：

• 根据 ptx 程序集，使用了 4+5+8+5=22 个寄存器。那为什么它used 2 registers在编译期间说？
• 查看程序集，我意识到 threadId、blockId 等的数据类型是u16. 这是在 CUDA 规范中定义的吗？或者这可能在不同版本的 CUDA 驱动程序之间有所不同？
• 有人可以向我解释这一行：mul.wide.u16 %r1, %rh1, %rh2;？%r1是u32，为什么使用wide而不是u32？
• 寄存器的名称是如何选择的？在我的花瓶中，我理解%r部分，但我不理解h,(null),d部分。它是根据数据类型长度选择的吗？即：h对于 16 位，对于 32 位为空，d对于 64 位？
• 如果我用 this 替换内核的最后两行out[idx] = in[idx];，那么当我编译程序时，它会说使用了 3 个寄存器！现在怎么可能使用更多的寄存器？

请忽略我的测试内核不检查数组索引是否超出范围这一事实。

非常感谢。

这个问题类似于cuModuleLoadDataEx 选项，但我想再次提出这个主题并提供更多信息。

通过 cuModuleLoadDataEx 使用 NV 驱动程序加载 PTX 字符串时，它似乎忽略了所有选项。我提供了完整的工作示例，以便任何有兴趣的人都可以直接轻松地复制它。首先是一个小的 PTX 内核（将其保存为 small.ptx），然后是加载 PTX 内核的 C++ 程序。

主文件

构建（假设 CUDA 安装在 /usr/local/cuda 下，我使用 CUDA 5.0）：

如果有人能够从编译过程中提取任何有意义的信息，那就太好了！解释 cuModuleLoadDataEx 的 CUDA 驱动程序 API 的文档（以及它应该接受的选项）http://docs.nvidia.com/cuda/cuda-driver-api/index.html

如果我运行它，日志是空的，jitTime甚至没有被 NV 驱动程序触及：

编辑：

我设法获得了 JIT 编译时间。然而，驱动程序似乎需要一个 32 位值的数组作为 OptVals。void *不像手册中所说的那样，我的系统上的指针数组 ( ) 是 64 位。所以，这有效：

我相信不可能对数组做同样的事情void *。以下代码不起作用：

编辑

查看 JIT 编译时间jitOptVals[0]会产生误导。如评论中所述，JIT 编译器缓存以前的翻译，如果找到缓存的编译，则不会更新 JIT 编译时间。由于我正在查看此值是否已更改，因此我假设该调用忽略了所有选项。它没有。它工作正常。

两个事实：CUDA 5.0 允许您在不同的对象文件中编译 CUDA 代码，以便稍后链接。CUDA 架构 2.x 不再自动内联函数。

像往常一样在 C/C++ 中，我已经实现了一个函数并将其头文件__device__ int foo()放在. 该函数在其他 CUDA 源文件中调用。functions.cufunctions.hufoo

当我检查时functions.ptx，我看到foo()溢出到本地内存。出于测试目的，我评论了所有的内容，foo()并return 1; 根据.ptx. （我无法想象它是什么，因为该函数什么都不做！）

但是，当我将实现移动foo()到头文件 functions.hu 并添加__forceinline__限定符时，没有任何内容写入本地内存！

这里发生了什么？ 为什么 CUDA 不自动内联这么简单的函数？

单独的头文件和实现文件的重点是让我的生活更容易维护代码。但是，如果我必须在标头和它们中粘贴一堆函数（或所有函数）__forceinline__，那么它有点违背了 CUDA 5.0 不同编译单元的目的......

有没有办法解决？

简单，真实的例子：

函数.cu：

上述函数溢出到本地内存。

函数.ptx：

考虑这 3 个微不足道的最小内核。他们的寄存器使用率比我预期的要高得多。为什么？

A：

对应的ptx：

乙：

对应的ptx：

C：

对应的ptx：

问题：

为什么空内核 A 和 B 使用 2 个寄存器？CUDA 总是使用一个隐式寄存器，但为什么要使用 2 个额外的显式寄存器？

内核 C 更令人沮丧。10个寄存器？但是只有2个指针。这为指针提供了 2*2 = 4 个寄存器。即使有另外 2 个神秘的寄存器（由内核 A 和内核 B 建议），这将给出 6 个总数。 还是不到10个！

如果您有兴趣，这里是ptx内核 A 的ptx代码。内核 B 的代码完全相同，以整数值和变量名为模。

而对于内核 C...

1. 为什么它首先声明一个本地内存变量 ( .local) ？
2. 为什么将两个指针（作为函数参数给出）存储在寄存器中？他们没有特殊的参数空间吗？
3. 也许这两个函数参数指针属于寄存器 - 这解释了这两.reg .b64 行。但.reg .s64线是什么？为什么会在那里？

情况变得更糟：

丁：

给

那么操作参数（指针）从 10 个寄存器减少到 6 个寄存器？

我想在http://docs.nvidia.com/cuda/pdf/ptx_isa_3.1.pdf中使用 SIMD 视频指令（vadd4、vmax4 等）第 8.7.13 节

我在我的代码中尝试了以下内容

其中 i,j,k,l 是 int 变量。我使用了“r”，因为它是 .u32 reg 的约束

但是在编译时，我收到以下错误

错误：未知的寄存器名称“r”

我应该在这里用什么代替“r”？还是代码中还有其他问题？（我使用的是计算能力为 3.5 的 Tesla 卡）

LLVM 的 NVPTX 后端（由 NVIDIA 提供）是否支持 CUDA5 / Compute Capability 3.5 设备中的新动态并行功能？

我想开始学习 PTX，我应该从哪里开始？有什么好书/资源可以做到这一点吗？

如果这可能有帮助，我已经知道 x86/x64 ASM（或多或少）

我不明白 mov 指令在 PTX 中是如何工作的。

如果 a 是寄存器或立即数，这会将 a 移动到 d 中。顺便说一句，如果 a 是全局、本地或共享状态空间中的变量，则可以将 a 的地址移入 d。

假设 a 是全局内存中的一个变量，指向一个值为 0x1 的 u64……我如何将 0x1 存储到 d 中，因为我只能得到 a 的地址？

我不确定如何获取值而不是地址.. 像英特尔 ASM 的 mov eax，ebx 用于地址和 mov eax，[ebx] 用于值（取消引用 ebx）