何时与寄存器/局部变量一起使用易失性

Question

在CUDA中用易失性限定符声明寄存器数组的意义是什么？

当我尝试使用可变关键字和寄存器数组时，它将溢出的寄存器内存数量移到本地内存中。(即强迫CUDA使用寄存器而不是本地内存)这是预期的行为吗？

我没有在CUDA文档中找到任何关于登记数组的使用情况的信息。

以下是这两个版本的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]

Answer 1

volatile限定符向编译器指定，对变量(读或写)的所有引用都应导致内存引用，并且这些引用必须按照程序中指定的顺序进行。volatile限定符的使用在Shane书的第12章"CUDA编程“中得到了说明。

使用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)。