问题标签 [ptx]

Well, I have a program which generates, JIT-compiles and runs PTX subprograms on GPU. Programs are running just fine and the run-times are pretty good - around 500x speedup vs CPU. Problem is that the compilation takes way too long, erasing all GPU speedup turning it to slowdown :)

Question is, is there a faster, more efficient way to do this? Can I reuse some resources, make the process more stream-like?

Edit: Each PTX program is run only once, and they are all very different, so JIT caching gains no benefit

This is my code, which is pretty much the same as nvidia-provided example JIT app:

我正在尝试在 .cu 文件上使用 nvcc 构建 .ptx 代码。对于某些它可以工作，而对于某些它会产生“找不到文件”错误：

但是当我尝试另一个文件时它可以工作：

也许当我跑

在其中一些我的 PATH 变量设置不同？知道如何修复那些不起作用的吗？谢谢，bb

我正在编写一个代码来使用幂级数逼近一个函数，并希望利用#pragma unroll 和 FMA 指令，如下所示：

代码将被编译成这样的程序集：

我想避免使用常量内存并使用像这样的立即值：

我知道我可以使用#define宏来做到这一点，但是当系数很多时很不方便。

是否有任何 C 数据类型修饰符（或编译器选项）可以将我的系数数组转换为立即值而不是使用常量内存？

我试过了，它不适用于static double,static __constant__ double和static __device__ double。

我的最后一个问题是：我猜使用立即值应该比使用常量内存更快？

使用 64 位整数时，CUDA 是否会在内部在 SM35 上生成漏斗移位指令？我的内核使用 SM35 目标编译为 PTX，它显示了一条常规的shl.b64指令。我知道这个 PTX 输出还没有完全优化，一旦加载模块（cuModuleLoad）就会生成本机架构代码。

我正在使用 Nvidia 驱动程序做一个小型 OpenCL 基准测试，我的内核执行 1024 保险丝乘加并将结果存储在一个数组中：

但是当我查看生成的 PTX 时，我看到了这个：

如果我是正确的，PTX 使用1026个浮点寄存器来执行 1024 个操作，并且永远不会重复使用一个寄存器两次，即使它可以只使用 2 个寄存器来执行所有乘加操作。1026 远高于线程允许拥有的最大寄存器数（根据规范），所以我猜这最终会导致内存溢出。

它是编译器错误还是我完全错过了什么？

我在 Quadro K2000 GPU 上使用 nvcc 6.5 版。

编辑

实际上我确实错过了规格中的一些内容：

“由于 PTX 支持虚拟寄存器，编译器前端生成大量寄存器名称是很常见的。PTX 不需要显式声明每个名称，而是支持创建一组具有公共前缀字符串的变量的语法附加整数后缀。例如，假设一个程序使用大量的 .b32 变量，比如一百个，命名为 %r0, %r1, ..., %r99"

我正在尝试减少内核中使用寄存器的数量，因此我决定尝试内联 PTX。

这个内核：

使用 14 个寄存器，我认为这超出了需要，所以我编写了 Inline PTX：

但是这个内核也使用了 14 个寄存器！我有点困惑。我只声明了 10 个寄存器，在 ptx 文件中没有其他变量。我该如何解决这种情况？

这里在文档中指出，prefetchptxprefetchu指令“在指定状态空间中预取包含指定内存层次结构级别的通用地址的行”。还提到语法是

我想知道uniform cache这里指的是什么；而语法（在第 2 行）指定数据将被预取到 L1 中？prefetchu当存在prefetch允许预取到 L1 的指令时，这不是多余的吗？例如，下面的代码行有什么区别？

我正在研究一个粒子代码，其中广泛使用刷新为零来提取性能。但是，有一个浮点比较语句我不希望被刷新。一种解决方案是使用内联 PTX，但它引入了不必要的指令，因为在 PTX 中没有布尔类型，而只是谓词寄存器： C++ 代码：

PTX：

有没有一种方法可以用一条指令进行非 FTZ 比较，而无需在 PTX/SASS 中对整个内容进行编码？

在 CUDA 程序的 NVPTX(LLVM IR) 中，内存地址空间的标识符从 0 到 5（见下表）。

我在同一个 LLVM IR 程序中看到，内存地址被标识为“通用”或其他类型，如图所示。

对于“通用”（默认情况下，没有标识符）：

用以分享'：

我的问题是，对于通用内存地址空间，数据实际上位于硬件、片外、片上存储器或本地寄存器中的什么位置？有人可以解释最终如何管理通用类型的地址空间吗？

有没有办法查看 CUDA 源代码的基本块信息？我正在尝试编译一个 CUDA 代码，以便它给我的输出是一个列出所有基本块的文件。但到目前为止，我还没有在 nvcc 中找到任何允许我这样做的选项。

如果我们可以从 IR (PTX) 中提取基本块信息，那也会有所帮助。