1

我注意到 PTX 代码允许一些具有复杂语义的指令，例如位字段提取 ( bfe)、查找最重要的非符号位 ( bfind) 和人口计数 ( popc)。

显式使用它们比在 C/C++ 中使用其预期语义编写代码更有效吗？

例如：“人口计数”，或popc，表示计数一位。所以我应该写：

__device__ int popc(int a) {
  int d = 0;
  while (a != 0) {
    if (a & 0x1)  d++;
    a = a >> 1;
  }   
  return d;
}

对于该功能，或者我应该使用：

__device__ int popc(int a) {
    int d;
    asm("popc.u32 %1 %2;"："=r"(d): "r"(a));
    return d;
}

? 内联 PTX 会更高效吗？我们是否应该编写内联 PTX 以获得最佳性能？

还有 - GPU 是否有一些与 PTX 指令相对应的额外魔法指令？

4

2 回答 2

2

编译器可能会识别您正在做什么并使用花哨的指令来执行它，或者它可能不会。在一般情况下，唯一知道的方法是通过使用添加到nvcc-ptx的标志来查看 ptx 程序集中的编译输出。如果编译器为您生成它，则无需自己手动编码内联程序集（或使用 instrinsic）。

此外，它在一般情况下是否会产生性能差异取决于代码路径是否以重要方式使用，以及其他因素，例如内核的当前性能限制器（例如计算受限或内存-边界）。

于 2013-09-02T12:18:58.810 回答
0

除了@RobertCrovella 的回答之外，还有几点：

  • 即使您确实将 PTX 用于某些事情 - 这应该很少发生。将其限制为不超过几行 PTX 行的小功能 - 然后您可以在您认为合适的情况下将其重新用于多种用途，您的大部分编码都在 C/C++ 中。
  • 这个原则的一个例子是 @njuffa 提到的内在函数，在（我认为这不是该文件的正式副本）。请通读以了解您可以使用哪些内在函数。当然，这并不意味着您应该全部使用它们。
  • 对于您的具体示例-您确实希望 PTX 超过第一个版本；它当然不会造成任何伤害。但是，同样，它也是您不需要实际编写 PTX 的一个示例，因为popc它具有相应的__popc内在函数（同样，正如 @njuffa 所指出的那样）。
  • 您可能还想查看一些基于 CUDA 的库的源代码，以了解他们选择使用哪种 PTX 片段。
于 2017-01-09T18:47:05.917 回答