當使用內聯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'會做什麼？

回答

相關問題