4

在 CUDA 中使用 volatile 限定符声明寄存器数组是什么意思？

当我尝试使用带有寄存器数组的 volatile 关键字时，它会删除溢出的寄存器内存到本地内存的数量。（即强制 CUDA 使用寄存器而不是本地内存）这是预期的行为吗？

我没有在 CUDA 文档中找到关于寄存器数组使用 volatile 的任何信息。

这是两个版本的 ptxas -v 输出

带有 volatile 限定符

    __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]
4

1 回答 1

7

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关键字时，所有对 to 的引用a都被转换为内存引用（在这种情况下为 write）。

于 2013-10-18T07:31:17.360 回答