何時使用volatile與寄存器/局部變量

在CUDA中使用volatile限定符聲明寄存器數組有什麼意義？何時使用volatile與寄存器/局部變量

當我用volatile關鍵字與寄存器數組一起嘗試時，它將溢出的寄存器內存數刪除到本地內存。（即強制CUDA使用寄存器而不是本地存儲器）這是預期的行爲嗎？

我在CUDA文檔中沒有找到關於volatile寄存器陣列的使用信息。

這裏是ptxas -v輸出兩種版本

揮發性預選賽

__volatile__ float array[32];

ptxas -v輸出

ptxas info : Compiling entry function '_Z2swPcS_PfiiiiS0_' for 'sm_20' 
ptxas info : Function properties for _Z2swPcS_PfiiiiS0_ 
88 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 47 registers, 16640 bytes smem, 80 bytes cmem[0], 8 bytes cmem[16]

不揮發性預選賽

個

float array[32];

ptxas -v輸出

ptxas info : Compiling entry function '_Z2swPcS_PfiiiiS0_' for 'sm_20' 
ptxas info : Function properties for _Z2swPcS_PfiiiiS0_ 
96 bytes stack frame, 100 bytes spill stores, 108 bytes spill loads 
ptxas info : Used 51 registers, 16640 bytes smem, 80 bytes cmem[0], 8 bytes cmem[16]

來源

2013-10-18 warunapww

'volatile'限定符指定編譯器所有對變量的引用（讀或寫）應該導致內存引用，並且這些引用必須按程序中指定的順序。 Shane Cook的書籍「CUDA編程」第12章介紹了使用'volatile'限定符。我的理解是，這種使用將避免編譯器可以做的一些優化，因此改變使用的寄存器的數量。我認爲理解它實際做什麼的最好方法是反彙編帶有和不帶限定詞的相關'__global__'函數。 – JackOLantern

的volatile限定符指定到一個變量（讀或寫）的所有引用應導致存儲器參考編譯器和這些引用必須在指定的順序該程序。在Shane Cook的書籍「CUDA編程」的第12章中說明了使用volatile限定符。

使用volatile將避免編譯器可以做的一些優化，因此改變使用的已使用寄存器的數量。瞭解volatile實際在做什麼的最好方法是在有和沒有限定符的情況下反彙編相關的__global__函數。

確實考慮下面的內核函數

__global__ void volatile_test() { 

    volatile float a[3]; 

    for (int i=0; i<3; i++) a[i] = (float)i; 
} 

__global__ void no_volatile_test() { 

    float a[3]; 

    for (int i=0; i<3; i++) a[i] = (float)i; 
}

拆卸上面的內核函數可以得到

code for sm_20 
     Function : _Z16no_volatile_testv 
.headerflags @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)" 
/*0000*/  MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */ 
/*0008*/  EXIT ;     /* 0x8000000000001de7 */ 


     Function : _Z13volatile_testv 
.headerflags @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)" 
/*0000*/  MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */ 
/*0008*/  ISUB R1, R1, 0x10;  /* 0x4800c00040105d03 */ R1 = address of a[0] 
/*0010*/  MOV32I R2, 0x3f800000; /* 0x18fe000000009de2 */ R2 = 1 
/*0018*/  MOV32I R0, 0x40000000; /* 0x1900000000001de2 */ R0 = 2 
/*0020*/  STL [R1], RZ;   /* 0xc8000000001fdc85 */ 
/*0028*/  STL [R1+0x4], R2;  /* 0xc800000010109c85 */ a[0] = 0; 
/*0030*/  STL [R1+0x8], R0;  /* 0xc800000020101c85 */ a[1] = R2 = 1; 
/*0038*/  EXIT ;     /* 0x8000000000001de7 */ a[2] = R0 = 2;

正如你所看到的，不使用volatile關鍵字時，編譯器意識到a設置但從未使用（實際上，編譯器返回以下警告：變量「a」已設置，但從未使用過）並且實際上沒有反彙編代碼。

與此相反，使用volatile關鍵字時，所有對a的引用都被轉換爲內存引用（在這種情況下編寫）。

來源

2013-10-18 07:31:17 JackOLantern

何時使用volatile與寄存器/局部變量

回答

相關問題