0

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

__device__ __inline__ uint32_t bfind(uint32_t val)
{
    uint32_t ret;
    asm ("bfind.u32 %0, %1;" : "=r"(ret): "r"(val));
    return ret;
}

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

__device__ __inline__ uint32_t bfind(uint32_t val)
{
    uint32_t ret;
    asm volatile ("bfind.u32 %0, %1;" : "=r"(ret): "r"(val));
    return ret;
}

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

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

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

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

1 回答 1

3

为什么我的 asm() 会被删除？或者更确切地说，如果编译器注意到它没有效果，我为什么要介意它被删除？

如果编译器检测到您的内联 PTX 不会在线程本地范围以外的任何地方更改状态，则可以随意将其作为优化删除。一般来说，这正是你想要发生的事情。但有时，并非如此。您的意图和编译器的优化策略可能并不总是以您想要或期望的方式相交。警告购买者和所有这些。

如果我的 asm() 在 PTX 生成期间被移动，为什么会出现问题？这是优化过程的一部分，不是吗？

这不是问题，是优化过程的一部分；但有时您可能想规避这一点。想象一下，您正在制作微基准测试，编译器决定重新排序您在内联 PTX 中编写的精心设计的指令序列（经典案例是将调用移动到发出的代码中的错误位置，从而破坏时序部分或内存事务模式设计） . 结果不会是你想要的。我想这可能会非常令人沮丧。

当分别面对非易失性和易失性 asm() 指令时，如何更准确地描述编译器的行为？

与标准 CUDA 内核代码一样，volatile 确保编译器尊重在其输出中发出给定的内联 PTX 操作，而不是通过代码分析将其暴露在优化之外。

于 2017-04-23T10:53:16.717 回答