CUDA 设备堆栈和同步；SSY 指令

Question

CUDA 设备堆栈和同步；SSY 指令

8

编辑：这个问题是原问题的重新制作版本，所以前几个回答可能已经不再相关。

我想知道在设备函数中强制禁用内联调用的设备函数调用对同步的影响。我有一个简单的测试内核，可以说明所讨论的行为。

内核获取一个缓冲区，并将其传递给设备函数，同时还传递了一个共享缓冲区和一个标识变量，该变量将单个线程标识为“boss”线程。设备函数具有分歧代码：老板线程首先花费时间在共享缓冲区上执行微不足道的操作，然后写入全局缓冲区。在同步调用之后，所有线程都会写入全局缓冲区。在内核调用之后，主机打印全局缓冲区的内容。以下是代码：

CUDA 代码：

test_main.cu

#include<cutil_inline.h>
#include "test_kernel.cu"

int main()
{
  int scratchBufferLength = 100;
  int *scratchBuffer;
  int *d_scratchBuffer;

  int b = 1;
  int t = 64;

  // copy scratch buffer to device
  scratchBuffer = (int *)calloc(scratchBufferLength,sizeof(int));
  cutilSafeCall( cudaMalloc(&d_scratchBuffer,
        sizeof(int) * scratchBufferLength) );
  cutilSafeCall( cudaMemcpy(d_scratchBuffer, scratchBuffer,
        sizeof(int)*scratchBufferLength, cudaMemcpyHostToDevice) );

  // kernel call
  testKernel<<<b, t>>>(d_scratchBuffer);

  cudaThreadSynchronize();

  // copy data back to host
  cutilSafeCall( cudaMemcpy(scratchBuffer, d_scratchBuffer,
        sizeof(int) * scratchBufferLength, cudaMemcpyDeviceToHost) );

  // print results
  printf("Scratch buffer contents: \t");
  for(int i=0; i < scratchBufferLength; ++i)
  {
    if(i % 25 == 0)
      printf("\n");
    printf("%d ", scratchBuffer[i]);
  }
  printf("\n");

  //cleanup
  cudaFree(d_scratchBuffer);
  free(scratchBuffer);

  return 0;
}

test_kernel.cu

#ifndef __TEST_KERNEL_CU
#define __TEST_KERNEL_CU


#define IS_BOSS() (threadIdx.x == blockDim.x - 1)

__device__
__noinline__
void testFunc(int *sA, int *scratchBuffer, bool isBoss) {

  if(isBoss)  {   // produces unexpected output-- "broken" code
//if(IS_BOSS())  {    // produces expected output-- "working" code

    for (int c = 0; c < 10000; c++)  {
      sA[0] = 1;
    }
  }

  if(isBoss) {
    scratchBuffer[0] = 1;
  }

  __syncthreads();

  scratchBuffer[threadIdx.x ] = threadIdx.x;

  return;

}

__global__
void testKernel(int *scratchBuffer)
{
  __shared__ int sA[4];

  bool isBoss = IS_BOSS();

  testFunc(sA, scratchBuffer, isBoss);
  return;
}
#endif

我从CUDA SDK中编译了这段代码，以利用test_main.cu中的“cutilsafecall()”函数，但如果您希望在SDK之外编译，则可以将其删除。我使用CUDA驱动程序/工具包版本4.0，计算能力2.0进行编译，并在具有Fermi架构的GeForce GTX 480上运行代码。

期望输出为：0 1 2 3 ... blockDim.x-1

然而，我得到的输出是：1 1 2 3 ... blockDim.x-1

这似乎表明老板线程在所有线程执行“scratchBuffer[threadIdx.x] = threadIdx.x;”语句后执行条件“scratchBuffer[0] = 1;”语句，尽管它们被__syncthreads()屏障分开。

即使老板线程被指示将一个标记值写入其同一warp中的线程的缓冲区位置，该标记也是缓冲区中存在的最终值，而不是适当的threadIdx.x。

导致代码产生预期输出的一种修改是将条件语句

if(isBoss) {

更改为

if(IS_BOSS()) {

即将分歧控制变量从参数寄存器存储更改为在宏函数中计算。（请注意源代码中适当行的注释。）这是我一直专注于追踪问题的特定更改。通过查看具有“isBoss”条件（即损坏的代码）和“IS_BOSS()”条件（即工作代码）的内核的反汇编.cubins，指令中最显着的区别似乎是在已解组装的损坏代码中缺少SSY指令。

以下是通过使用“cuobjdump -sass test_kernel.cubin”反汇编.cubin文件生成的反汇编内核。从一开始到第一个“EXIT”都是内核，之后的所有内容都是设备函数。唯一的区别在于设备函数。

code for sm_20

    Function : _Z10testKernelPi
/*0000*/     /*0x00005de428004404*/     MOV R1, c [0x1] [0x100];
/*0008*/     /*0x20009de428004000*/     MOV R2, c [0x0] [0x8];
/*0010*/     /*0x84001c042c000000*/     S2R R0, SR_Tid_X;
/*0018*/     /*0xfc015de428000000*/     MOV R5, RZ;
/*0020*/     /*0x00011de428004000*/     MOV R4, c [0x0] [0x0];
/*0028*/     /*0xfc209c034800ffff*/     IADD R2, R2, 0xfffff;
/*0030*/     /*0x9001dde428004000*/     MOV R7, c [0x0] [0x24];
/*0038*/     /*0x80019de428004000*/     MOV R6, c [0x0] [0x20];
/*0040*/     /*0x08001c03110e0000*/     ISET.EQ.U32.AND R0, R0, R2, pt;
/*0048*/     /*0x01221f841c000000*/     I2I.S32.S32 R8, -R0;
/*0050*/     /*0x2001000750000000*/     CAL 0x60;
/*0058*/     /*0x00001de780000000*/     EXIT;
/*0060*/     /*0x20201e841c000000*/     I2I.S32.S8 R0, R8;
/*0068*/     /*0xfc01dc231a8e0000*/     ISETP.NE.AND P0, pt, R0, RZ, pt;
/*0070*/     /*0xc00021e740000000*/     @!P0 BRA 0xa8;
/*0078*/     /*0xfc001de428000000*/     MOV R0, RZ;
/*0080*/     /*0x04001c034800c000*/     IADD R0, R0, 0x1;
/*0088*/     /*0x04009de218000000*/     MOV32I R2, 0x1;
/*0090*/     /*0x4003dc231a8ec09c*/     ISETP.NE.AND P1, pt, R0, 0x2710, pt;
/*0098*/     /*0x00409c8594000000*/     ST.E [R4], R2;
/*00a0*/     /*0x600005e74003ffff*/     @P1 BRA 0x80;
/*00a8*/     /*0x040001e218000000*/     @P0 MOV32I R0, 0x1;
/*00b0*/     /*0x0060008594000000*/     @P0 ST.E [R6], R0;
/*00b8*/     /*0xffffdc0450ee0000*/     BAR.RED.POPC RZ, RZ;
/*00c0*/     /*0x84001c042c000000*/     S2R R0, SR_Tid_X;
/*00c8*/     /*0x10011c03200dc000*/     IMAD.U32.U32 R4.CC, R0, 0x4, R6;
/*00d0*/     /*0x10009c435000c000*/     IMUL.U32.U32.HI R2, R0, 0x4;
/*00d8*/     /*0x08715c4348000000*/     IADD.X R5, R7, R2;
/*00e0*/     /*0x00401c8594000000*/     ST.E [R4], R0;
/*00e8*/     /*0x00001de790000000*/     RET;
    .................................

“工作中”的代码

code for sm_20

    Function : _Z10testKernelPi
/*0000*/     /*0x00005de428004404*/     MOV R1, c [0x1] [0x100];
/*0008*/     /*0x20009de428004000*/     MOV R2, c [0x0] [0x8];
/*0010*/     /*0x84001c042c000000*/     S2R R0, SR_Tid_X;
/*0018*/     /*0xfc015de428000000*/     MOV R5, RZ;
/*0020*/     /*0x00011de428004000*/     MOV R4, c [0x0] [0x0];
/*0028*/     /*0xfc209c034800ffff*/     IADD R2, R2, 0xfffff;
/*0030*/     /*0x9001dde428004000*/     MOV R7, c [0x0] [0x24];
/*0038*/     /*0x80019de428004000*/     MOV R6, c [0x0] [0x20];
/*0040*/     /*0x08001c03110e0000*/     ISET.EQ.U32.AND R0, R0, R2, pt;
/*0048*/     /*0x01221f841c000000*/     I2I.S32.S32 R8, -R0;
/*0050*/     /*0x2001000750000000*/     CAL 0x60;
/*0058*/     /*0x00001de780000000*/     EXIT;
/*0060*/     /*0x20009de428004000*/     MOV R2, c [0x0] [0x8];
/*0068*/     /*0x8400dc042c000000*/     S2R R3, SR_Tid_X;
/*0070*/     /*0x20201e841c000000*/     I2I.S32.S8 R0, R8;
/*0078*/     /*0x4000000760000001*/     SSY 0xd0;
/*0080*/     /*0xfc209c034800ffff*/     IADD R2, R2, 0xfffff;
/*0088*/     /*0x0831dc031a8e0000*/     ISETP.NE.U32.AND P0, pt, R3, R2, pt;
/*0090*/     /*0xc00001e740000000*/     @P0 BRA 0xc8;
/*0098*/     /*0xfc009de428000000*/     MOV R2, RZ;
/*00a0*/     /*0x04209c034800c000*/     IADD R2, R2, 0x1;
/*00a8*/     /*0x04021de218000000*/     MOV32I R8, 0x1;
/*00b0*/     /*0x4021dc231a8ec09c*/     ISETP.NE.AND P0, pt, R2, 0x2710, pt;
/*00b8*/     /*0x00421c8594000000*/     ST.E [R4], R8;
/*00c0*/     /*0x600001e74003ffff*/     @P0 BRA 0xa0;
/*00c8*/     /*0xfc01dc33190e0000*/     ISETP.EQ.AND.S P0, pt, R0, RZ, pt;
/*00d0*/     /*0x040021e218000000*/     @!P0 MOV32I R0, 0x1;
/*00d8*/     /*0x0060208594000000*/     @!P0 ST.E [R6], R0;
/*00e0*/     /*0xffffdc0450ee0000*/     BAR.RED.POPC RZ, RZ;
/*00e8*/     /*0x10311c03200dc000*/     IMAD.U32.U32 R4.CC, R3, 0x4, R6;
/*00f0*/     /*0x10309c435000c000*/     IMUL.U32.U32.HI R2, R3, 0x4;
/*00f8*/     /*0x84001c042c000000*/     S2R R0, SR_Tid_X;
/*0100*/     /*0x08715c4348000000*/     IADD.X R5, R7, R2;
/*0108*/     /*0x00401c8594000000*/     ST.E [R4], R0;
/*0110*/     /*0x00001de790000000*/     RET;
    .................................

“SSY”指令在正常工作代码中存在，但在有问题的代码中不存在。cuobjdump手册描述该指令为“设置同步点；用于潜在分歧指令之前。” 这使我认为编译器不认识有问题代码中可能存在的分歧。

我还发现，如果我注释掉__noinline__指令，那么代码会产生预期的输出，事实上，“有问题”的和“正常”版本产生的汇编完全相同。因此，当通过调用堆栈传递变量时，不能使用该变量来控制分歧和随后的同步调用; 在这种情况下，编译器似乎不认识可能存在的分歧，因此不插入 “SSY” 指令。是否有人知道这是否是CUDA的合法限制，如果是，是否有文档记录？

提前感谢。

- user1663964

3

SSY是设置同步相对地址的硬件指令（请注意，这不是PTX），需要查看cuobjdump指南。 - talonmies

1

我看不出有什么问题，但代码并不完整。你能否发布一个尽可能简单但完整的（即我只需复制、粘贴、编译和运行）示例代码？我建议先尽可能简化你的示例（尝试在没有对象和成员函数的情况下进行，使主函数最小化等）。 - harrism

1

你展示的代码毫无意义。memberFunc有三个参数，但内核调用只传递了两个。我几乎敢打赌这是一个共享内存/编译器优化冲突，通过声明sA为volatile可以解决。但如果你要问为什么代码能够工作或者不能工作，请发布一个真实的可编译重现案例。你现在展示的大部分内容根本没有任何帮助。 - talonmies

3

看着发布的SASS代码，我认为这两种变体并不相等。特别是，在一种情况下，“boss index”的确定将线程索引与0x3f进行比较，在另一种“破损”的情况下，代码将线程索引与0进行比较以确定“boss index”。不能排除某种编译器问题，但在手头没有可编译、可运行的重现案例之前假设这样的问题是过早的。顺便问一下，正在使用哪个CUDA版本？使用“volatile”可能只是掩盖了可能存在的问题（无论是在用户代码中还是在编译器中）。 - njuffa

1

@harrism：好的，我已经在Fermi卡上更新了CUDA 4.2，并且我也无法再现出意外的行为。编译器似乎以不同的方式处理同步，因为在任一版本的反汇编代码中都没有“SSY”指令。因此，看起来CUDA团队决定更改他们的模型并允许参数控制分歧/同步，或者这是一个被修复的编译器故障，或者编译器的其他更改作为副作用修复了这个特定问题。谢谢您的帮助。 - user1663964

显示剩余10条评论

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- harrism · Accepted Answer

看起来这只是CUDA 4.1/4.2中修复的编译器错误。在CUDA 4.2上，问题无法复现。