当使用内联PTX asm（）指令时，'volatile'会做什么？

当我们在一般的C/C++ CUDA代码，例如编写内联PTX汇编：当使用内联PTX asm（）指令时，'volatile'会做什么？

__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()语句没有副作用，除了改变输出操作数。为确保在生成PTX期间不删除或移动asm，应使用volatile关键字

我不明白这意味着什么。所以，

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

来源

2017-04-23 einpoklum

为什么我的asm（）被删除？或者说，如果编译器注意到它没有效果，为什么我应该介意它被删除？

如果编译器检测到您的内联PTX不会改变状态而不是在线程本地作用域以外的任何状态，则可以自由删除它作为优化。 一般来说，这正是你想要发生的事情。但有时候，事实并非如此。您的意图和编译器的优化策略可能并不总是以您希望或期望的方式相交。警惕和所有。

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

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

怎么会多了一个什么时候面向非挥发性及挥发性ASM（）指令分别表征编译器的行为？

与标准CUDA内核代码，易失性确保了编译器荣誉发射在其输出给定的内联PTX操作，而不是将其暴露于被优化由代码分析。

来源

2017-04-23 10:53:16 talonmies

当使用内联PTX asm（）指令时，'volatile'会做什么？

回答

相关问题