使用内联 PTX asm() 指令时，'volatile' 有啥作用？

Posted 2023-02-16

【中文标题】使用内联 PTX asm() 指令时，\'volatile\' 有啥作用？

【英文标题】：When using inline PTX asm() instructions, what does 'volatile' do?使用内联 PTX asm() 指令时，'volatile' 有什么作用？

【发布时间】：2017-04-23 10:23:38

【问题描述】：

当我们在我们通常的 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;

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

__device__ __inline__ uint32_t bfind(uint32_t val)

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

CUDA documentation on inline PTX assembly 说：

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

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

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

【参考方案1】：

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

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

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

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

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

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