linux - CUDA - 「branch」と「divergent branch」の Visual Profiler の結果に関する混乱 (2)

Question

NVIDIA Visual Profiler を使用してコードを分析しています。テストカーネルは次のとおりです。

//////////////////////////////////////////////////////////////// Group 1
static __global__ void gpu_test_divergency_0(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < 0)
    {
         a[tid] = tid;
    }
    else
    {
         b[tid] = tid;
    }
}
static __global__ void gpu_test_divergency_1(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid == 0)
    {
         a[tid] = tid;
    }
    else
    {
         b[tid] = tid;
    }
}
static __global__ void gpu_test_divergency_2(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid >= 0)
    {
         a[tid] = tid;
    }
    else
    {
         b[tid] = tid;
    }
}
static __global__ void gpu_test_divergency_3(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid > 0)
    {
         a[tid] = tid;
    }
    else
    {
         b[tid] = tid;
    }
}
//////////////////////////////////////////////////////////////// Group 2
static __global__ void gpu_test_divergency_4(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < 0)
    {
         a[tid] = tid + 1;
    }
    else
    {
         b[tid] = tid + 2;
    }
}
static __global__ void gpu_test_divergency_5(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid == 0)
    {
         a[tid] = tid + 1;
    }
    else
    {
         b[tid] = tid + 2;
    }
}
static __global__ void gpu_test_divergency_6(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid >= 0)
    {
         a[tid] = tid + 1;
    }
    else
    {
         b[tid] = tid + 2;
    }
}
static __global__ void gpu_test_divergency_7(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid > 0)
    {
         a[tid] = tid + 1;
    }
    else
    {
         b[tid] = tid + 2;
    }
}
//////////////////////////////////////////////////////////////// Group 3
static __global__ void gpu_test_divergency_8(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < 0)
    {
         a[tid] = tid + 1.0;
    }
    else
    {
         b[tid] = tid + 2.0;
    }
}
static __global__ void gpu_test_divergency_9(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid == 0)
    {
         a[tid] = tid + 1.0;
    }
    else
    {
         b[tid] = tid + 2.0;
    }
}
static __global__ void gpu_test_divergency_10(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid >= 0)
    {
         a[tid] = tid + 1.0;
    }
    else
    {
         b[tid] = tid + 2.0;
    }
}
static __global__ void gpu_test_divergency_11(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid > 0)
    {
         a[tid] = tid + 1.0;
    }
    else
    {
         b[tid] = tid + 2.0;
    }
}

<<< 1, 32 >>> でテストカーネルを起動すると、プロファイラーから次のような結果が得られました。

gpu_test_divergency_0 :  Branch Efficiency = 100% branch = 1 divergent branch = 0
gpu_test_divergency_1 :  Branch Efficiency = 100% branch = 1 divergent branch = 0
gpu_test_divergency_2 :  Branch Efficiency = 100% branch = 1 divergent branch = 0
gpu_test_divergency_3 :  Branch Efficiency = 100% branch = 1 divergent branch = 0

gpu_test_divergency_4 :  Branch Efficiency = 100% branch = 3 divergent branch = 0
gpu_test_divergency_5 :  Branch Efficiency = 100% branch = 3 divergent branch = 0
gpu_test_divergency_6 :  Branch Efficiency = 100% branch = 2 divergent branch = 0
gpu_test_divergency_7 :  Branch Efficiency = 100% branch = 3 divergent branch = 0

gpu_test_divergency_8 :  Branch Efficiency = 100% branch = 3 divergent branch = 0
gpu_test_divergency_9 :  Branch Efficiency = 75%  branch = 4 divergent branch = 1
gpu_test_divergency_10 : Branch Efficiency = 100% branch = 2 divergent branch = 0
gpu_test_divergency_11 : Branch Efficiency = 75%  branch = 4 divergent branch = 1

<<< 1, 64 >>> でテストカーネルを起動すると、プロファイラーから次のような結果が得られました。

gpu_test_divergency_0 :  Branch Efficiency = 100% branch = 2 divergent branch = 0
gpu_test_divergency_1 :  Branch Efficiency = 100% branch = 2 divergent branch = 0
gpu_test_divergency_2 :  Branch Efficiency = 100% branch = 2 divergent branch = 0
gpu_test_divergency_3 :  Branch Efficiency = 100% branch = 2 divergent branch = 0

gpu_test_divergency_4 :  Branch Efficiency = 100% branch = 6 divergent branch = 0
gpu_test_divergency_5 :  Branch Efficiency = 100% branch = 6 divergent branch = 0
gpu_test_divergency_6 :  Branch Efficiency = 100% branch = 4 divergent branch = 0
gpu_test_divergency_7 :  Branch Efficiency = 100% branch = 5 divergent branch = 0

gpu_test_divergency_8 :  Branch Efficiency = 100%  branch = 6 divergent branch = 0
gpu_test_divergency_9 :  Branch Efficiency = 85.7% branch = 7 divergent branch = 1
gpu_test_divergency_10 : Branch Efficiency = 100%  branch = 4 divergent branch = 0
gpu_test_divergency_11 : Branch Efficiency = 83.3% branch = 6 divergent branch = 1

Linux で CUDA 機能 2.0 と NVIDIA Visual Profiler v4.2 を備えた「GeForce GTX 570」を使用しています。ドキュメントによると：

"branch" - "カーネルを実行するスレッドが実行する分岐の数。このカウンターは、ワープ内の少なくとも 1 つのスレッドが分岐を実行する場合に 1 ずつ増加します。"

"分岐分岐" - "ワープ内の分岐分岐の数。このカウンターは、ワープ内の少なくとも 1 つのトレッドがデータ依存の条件付き分岐を介して分岐する (つまり、異なる実行パスに従う) 場合に 1 ずつ増加します。"

しかし、私は結果について本当に混乱しています。各テストグループの「枝」の数が異なるのはなぜですか? そして、なぜ 3 番目のテストグループだけが正しい「分岐分岐」を持っているように見えるのでしょうか?

@JackOLantern: リリースモードでコンパイルしました。私はあなたのやり方でそれを分解しました。「gpu_test_divergency_4」の結果はあなたのものとまったく同じですが、「gpu_test_divergency_0」の結果は異なります。

    Function : _Z21gpu_test_divergency_0PfS_
/*0000*/     /*0x00005de428004404*/     MOV R1, c [0x1] [0x100];
/*0008*/     /*0x94001c042c000000*/     S2R R0, SR_CTAid_X;
/*0010*/     /*0x84009c042c000000*/     S2R R2, SR_Tid_X;
/*0018*/     /*0x20009ca320044000*/     IMAD R2, R0, c [0x0] [0x8], R2;
/*0020*/     /*0xfc21dc23188e0000*/     ISETP.LT.AND P0, pt, R2, RZ, pt;
/*0028*/     /*0x0920de0418000000*/     I2F.F32.S32 R3, R2;
/*0030*/     /*0x9020204340004000*/     @!P0 ISCADD R0, R2, c [0x0] [0x24], 0x2;
/*0038*/     /*0x8020804340004000*/     @P0 ISCADD R2, R2, c [0x0] [0x20], 0x2;
/*0040*/     /*0x0000e08590000000*/     @!P0 ST [R0], R3;
/*0048*/     /*0x0020c08590000000*/     @P0 ST [R2], R3;
/*0050*/     /*0x00001de780000000*/     EXIT;

あなたが言ったように、変換命令（この場合はI2F）は余分なブランチを追加しないと思います。

しかし、これらの逆アセンブルされたコードとプロファイラーの結果との関係がわかりません。別の投稿 ( https://devtalk.nvidia.com/default/topic/463316/branch-divergent-branches/ ) から、発散分岐は実際のスレッド (ワープ) が SM で実行されている状況で計算されることを知りました。したがって、これらの逆アセンブルされたコードだけでは、実際の実行ごとの分岐分岐を推測することはできないと思います。私は正しいですか？

score 1 · Accepted Answer

フォローアップ - VOTE 組み込み関数を使用してスレッドの相違をチェックする

__ballotワープ内のスレッドの分岐をチェックする最善の方法は、vote 組み込み関数、特におよび組み込み関数を使用することだと思います__popc。__ballotおよびに関する適切な説明は__popc、シェーンクック、CUDA プログラミング、モーガンカウフマンによる本で利用できます。

のプロトタイプ__ballotは次のとおりです

unsigned int __ballot(int predicate);

predicate がゼロ以外の場合、番目のビットが設定され__ballotた値を返します。NNthreadIdx.x

一方、__popc-bit32パラメータで設定されたビット数を返します。

__ballotしたがって、、、__popcおよびを併用するatomicAddことで、ワープが発散しているかどうかを確認できます。

この目的のために、次のコードを設定しました

#include <cuda.h>
#include <stdio.h>
#include <iostream>

#include <cuda.h>
#include <cuda_runtime.h>

__device__ unsigned int __ballot_non_atom(int predicate)
{
    if (predicate != 0) return (1 << (threadIdx.x % 32));
    else return 0;
}

__global__ void gpu_test_divergency_0(unsigned int* d_ballot, int Num_Warps_per_Block)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    const unsigned int warp_num = threadIdx.x >> 5;

    atomicAdd(&d_ballot[warp_num+blockIdx.x*Num_Warps_per_Block],__popc(__ballot_non_atom(tid > 2)));
    //  atomicAdd(&d_ballot[warp_num+blockIdx.x*Num_Warps_per_Block],__popc(__ballot(tid > 2)));

}

#include <conio.h>

int main(int argc, char *argv[])
{
    unsigned int Num_Threads_per_Block      = 64;
    unsigned int Num_Blocks_per_Grid        = 1;
    unsigned int Num_Warps_per_Block        = Num_Threads_per_Block/32;
    unsigned int Num_Warps_per_Grid         = (Num_Threads_per_Block*Num_Blocks_per_Grid)/32;

    unsigned int* h_ballot = (unsigned int*)malloc(Num_Warps_per_Grid*sizeof(unsigned int));
    unsigned int* d_ballot; cudaMalloc((void**)&d_ballot, Num_Warps_per_Grid*sizeof(unsigned int));

    for (int i=0; i<Num_Warps_per_Grid; i++) h_ballot[i] = 0;

    cudaMemcpy(d_ballot, h_ballot, Num_Warps_per_Grid*sizeof(unsigned int), cudaMemcpyHostToDevice);

    gpu_test_divergency_0<<<Num_Blocks_per_Grid,Num_Threads_per_Block>>>(d_ballot,Num_Warps_per_Block);

    cudaMemcpy(h_ballot, d_ballot, Num_Warps_per_Grid*sizeof(unsigned int), cudaMemcpyDeviceToHost);

    for (int i=0; i<Num_Warps_per_Grid; i++) { 
        if ((h_ballot[i] == 0)||(h_ballot[i] == 32)) std::cout << "Warp " << i << " IS NOT divergent- Predicate true for " << h_ballot[i] << " threads\n";
            else std::cout << "Warp " << i << " IS divergent - Predicate true for " << h_ballot[i] << " threads\n";
    }

    getch();
    return EXIT_SUCCESS;
}

私は現在、計算機能 1.2 カードでコードを実行していることに注意してください。したがって、上記の例では、非組み込みのを使用してい__ballot_non_atomます。つまり、計算能力が 2.0 以上のカードをお持ちの場合は、カーネル関数で使用する命令のコメントを外してください。__ballot__ballot__ballot

上記のコードを使用すると、カーネル関数内の関連する述語を変更するだけで、上記のすべてのカーネル関数を操作できます。

前の回答

リリースモードで2.0計算機能用にコードをコンパイルし、2 つのカーネルの逆アセンブリを生成するために中間ファイルとユーティリティを保持していました。-keepcuobjdump

static __global__ void gpu_test_divergency_0(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < 0) a[tid] = tid;
    else b[tid] = tid;
}

と

static __global__ void gpu_test_divergency_4(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < 0) a[tid] = tid + 1;
    else b[tid] = tid + 2;
}

結果は次のとおりです

gpu_test_divergency_0

/*0000*/        MOV R1, c[0x1][0x100];                 /* 0x2800440400005de4 */
/*0008*/        S2R R0, SR_CTAID.X;                    /* 0x2c00000094001c04 */
/*0010*/        S2R R2, SR_TID.X;                      /* 0x2c00000084009c04 */
/*0018*/        IMAD R2, R0, c[0x0][0x8], R2;          /* 0x2004400020009ca3 */
/*0020*/        ISETP.LT.AND P0, PT, R2, RZ, PT;       /* 0x188e0000fc21dc23 */
/*0028*/        I2F.F32.S32 R0, R2;                    /* 0x1800000009201e04 */
/*0030*/   @!P0 ISCADD R3, R2, c[0x0][0x24], 0x2;      /* 0x400040009020e043 */
/*0038*/    @P0 ISCADD R2, R2, c[0x0][0x20], 0x2;      /* 0x4000400080208043 */
/*0040*/   @!P0 ST [R3], R0;                           /* 0x9000000000302085 */
/*0048*/    @P0 ST [R2], R0;                           /* 0x9000000000200085 */
/*0050*/        EXIT ;                                 /* 0x8000000000001de7 */

と

gpu_test_divergency_4

/*0000*/        MOV R1, c[0x1][0x100];                 /* 0x2800440400005de4 */
/*0008*/        S2R R0, SR_CTAID.X;                    /* 0x2c00000094001c04 */   R0 = BlockIdx.x
/*0010*/        S2R R2, SR_TID.X;                      /* 0x2c00000084009c04 */   R2 = ThreadIdx.x
/*0018*/        IMAD R0, R0, c[0x0][0x8], R2;          /* 0x2004400020001ca3 */   R0 = R0 * c + R2
/*0020*/        ISETP.LT.AND P0, PT, R0, RZ, PT;       /* 0x188e0000fc01dc23 */   If statement
/*0028*/    @P0 BRA.U 0x58;                            /* 0x40000000a00081e7 */   Branch 1 - Jump to 0x58
/*0030*/   @!P0 IADD R2, R0, 0x2;                      /* 0x4800c0000800a003 */   Branch 2 - R2 = R0 + 2
/*0038*/   @!P0 ISCADD R0, R0, c[0x0][0x24], 0x2;      /* 0x4000400090002043 */   Branch 2 - Calculate gmem address
/*0040*/   @!P0 I2F.F32.S32 R2, R2;                    /* 0x180000000920a204 */   Branch 2 - R2 = R2 after int to float cast
/*0048*/   @!P0 ST [R0], R2;                           /* 0x900000000000a085 */   Branch 2 - gmem store
/*0050*/   @!P0 BRA.U 0x78;                            /* 0x400000008000a1e7 */   Branch 2 - Jump to 0x78 (exit)
/*0058*/    @P0 IADD R2, R0, 0x1;                      /* 0x4800c00004008003 */   Branch 1 - R2 = R0 + 1
/*0060*/    @P0 ISCADD R0, R0, c[0x0][0x20], 0x2;      /* 0x4000400080000043 */   Branch 1 - Calculate gmem address
/*0068*/    @P0 I2F.F32.S32 R2, R2;                    /* 0x1800000009208204 */   Branch 1 - R2 = R2 after int to float cast
/*0070*/    @P0 ST [R0], R2;                           /* 0x9000000000008085 */   Branch 1 - gmem store
/*0078*/        EXIT ;                                 /* 0x8000000000001de7 */

上記の逆アセンブルから、分岐分岐テストの結果は同じであると予想されます。

デバッグモードまたはリリースモードでコンパイルしていますか?

linux - CUDA - 「branch」と「divergent branch」の Visual Profiler の結果に関する混乱 (2)

1 に答える 1

Related

Reference