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

このコードは、test_main.cu の「cutilsafecall()」関数を利用するために CUDA SDK 内からコンパイルしましたが、SDK の外でコンパイルしたい場合はもちろん、これらを取り除くことができます。CUDA Driver/Toolkit バージョン 4.0、コンピューティング機能 2.0 でコンパイルし、コードは Fermi アーキテクチャの GeForce GTX 480 上で実行されました。

期待される出力は

0 1 2 3 ... blockDim.x-1

しかし、私が得る出力は

1 1 2 3 ... blockDim.x-1

これは、ボス スレッドが条件 "scratchBuffer[0] = 1;" を実行したことを示しているようです。すべてのスレッドが「scratchBuffer[threadIdx.x] = threadIdx.x;」を実行した後のステートメント ステートメントは __syncthreads() バリアで区切られていますが。

これは、同じワープ内のスレッドのバッファー位置にセンチネル値を書き込むようにボス スレッドが指示された場合でも発生します。センチネルは、適切な threadIdx.x ではなく、バッファーに存在する最終的な値です。

コードが期待される出力を生成するようにする変更の 1 つは、条件ステートメントを変更することです。

if(isBoss) {

に

if(IS_BOSS()) {

; つまり、発散制御変数をパラメータ レジスタに格納することから、マクロ関数で計算するように変更します。(ソース コード内の適切な行のコメントに注意してください。) 問題を突き止めるために私が焦点を当ててきたのは、この特定の変更です。'isBoss' 条件 (つまり、壊れたコード) と 'IS_BOSS()' 条件 (つまり、動作中のコード) を持つカーネルの逆アセンブルされた .cubin を見ると、命令の最も顕著な違いは、逆アセンブルされた壊れたコードの SSY 命令。

.cubin ファイルを "cuobjdump -sass test_kernel.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の正当な制限であるかどうか、そしてもしそうなら、これがどこかに文書化されているかどうかは誰にも分かりますか?

前もって感謝します。

4

1 に答える 1

3

これは単に CUDA 4.1/4.2 で修正されたコンパイラのバグのようです。CUDA 4.2 の asker では再現しません。

于 2012-10-01T21:20:44.647 に答える