何时将 volatile 与寄存器/局部变量一起使用

在 CUDA 中使用 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]

没有 volatile 限定符

    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]

The volatile限定符向编译器指定对变量的所有引用（读或写）都应产生内存引用，并且这些引用必须按照程序中指定的顺序。使用volatileShane Cook 的书《CUDA 编程》第 12 章对限定符进行了说明。

指某东西的用途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被转换为内存引用（在本例中为 write）。