占有量を増やしてカーネルのパフォーマンスを改善しますか？

Question

ここ

はGT 440の私のカーネルの計算ビジュアルプロファイラの出力です：

• カーネルの詳細：グリッド・サイズ：[100 1 1]、ブロックサイズ：[256 1 1]
• 登録比： 0.84375（32768分の27648）[スレッド当たり35個のレジスタ]
• 共有記憶率：0.336914（49152分の16560） ブロック当たり5520バイト] SMあたり
• アクティブなブロック：SMあたり3（最大アクティブブロック： 8）
0 SMあたりの
• アクティブスレッド：768（SMあたり最大アクティブスレッド：1536）
• ポテンシャル人数：0.5（48分の24）
• 入居制限要因は：、にご注意ください

登録太字の箇条書き。カーネル実行時間は121195 usです。

私はいくつかのローカル変数を共有メモリに移動することによってスレッドあたりのレジスタ数を減らしました。計算ビジュアルプロファイラ出力になった：

• カーネルの詳細：グリッドサイズ：[100 1]、ブロックサイズ：[256 1]
• レジスタ比：1（32768分の32768）当たりの[30個のレジスタスレッド]
• 共有メモリ比：0.451823（49152分の22208）[ブロック] SMあたり
• アクティブなブロックあたり5552バイト：4（SMあたりの最大アクティブなブロック：8）SMあたり
• アクティブスレッド： 1024（SMごとの最大アクティブスレッド数：153 6）
• 潜在人数：0.666667（48分の32）
• 人数制限要因は：したがって、今4ブロックが同時に以前のバージョンで3ブロックに対して単一SM上で実行され

登録します。ただし、実行時間は115756 usでほぼ同じです！どうして？ブロックが完全に独立していて、異なるCUDAコアで実行されていませんか？

Answer 1

高い占有率が自動的により高いパフォーマンスに変換されることを暗黙のうちに前提としています。それはほとんどの場合そうではありません。

NVIDIAアーキテクチャでは、GPUの命令パイプラインのレイテンシを隠すために、MPごとに特定の数のアクティブワーピングが必要です。あなたのFermiベースのカードでは、その要件は約30％の最小占有率に換算されます。レイテンシのボトルネックがGPUの別の部分に移動する可能性があるため、その最小値より高い占有率を目指しても、必ずしもスループットが向上するとは限りません。あなたのエントリーレベルのGPUには多くのメモリ帯域幅がありません。そして、MPあたり3ブロックでコードメモリ帯域幅を制限するのに十分な可能性があります。その場合、ブロック数を増やすことはパフォーマンスに影響しません（メモリーコントローラーの競合やキャッシュミスのためにダウンすることさえあります）。さらに、カーネルのレジスタフットプリントを減らすために、変数を共有メモリに書き込んだとします。フェルミでは、共用メモリは約1000Gb/sの帯域幅しか持たないのに対し、レジスタは約8000Gb/sです（これを示すマイクロベンチマーキングの結果は下のリンクを参照してください）。したがって、変数をより低速のメモリに移動しました。これはパフォーマンスに悪影響を及ぼし、高い占有率がもたらす利益を相殺します。

まだご覧になっていない方は、GTC 2010の「Vasily Volkov」のプレゼンテーションをお勧めします。「より低い人数でパフォーマンスが向上します」(pdf)ここでは、命令レベルの並列性を利用することで、非常に低いレベルの占有でGPUスループットを非常に高いレベルにまで高めることができる方法を示しています。

Answer 2

talonmiesは、すでにあなたの質問に答えたので、私はちょうど上記の回答で述べたV.ボルコフ氏のプレゼンテーションの最初の部分に触発されたコードを共有したいです。

これは、コードされています、結果は命令レベルの並列性（あれば下の占有とカーネルはまだ、高いパフォーマンスを発揮することができることを意味

kernel0 GFlops = 21.069281 Occupancy = 66% 
kernel1 GFlops = 21.183354 Occupancy = 33% 
kernel2 GFlops = 21.224517 Occupancy = 16.7% 

私のGeForce GT540Mで

#include<stdio.h> 

#define N_ITERATIONS 8192 

//#define DEBUG 

/********************/ 
/* CUDA ERROR CHECK */ 
/********************/ 
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) exit(code); 
    } 
} 

/********************************************************/ 
/* KERNEL0 - NO INSTRUCTION LEVEL PARALLELISM (ILP = 0) */ 
/********************************************************/ 
__global__ void kernel0(int *d_a, int *d_b, int *d_c, unsigned int N) { 

    const int tid = threadIdx.x + blockIdx.x * blockDim.x ; 

    if (tid < N) { 

     int a = d_a[tid]; 
     int b = d_b[tid]; 
     int c = d_c[tid]; 

     for(unsigned int i = 0; i < N_ITERATIONS; i++) { 
      a = a * b + c; 
     } 

     d_a[tid] = a; 
    } 

} 

/*****************************************************/ 
/* KERNEL1 - INSTRUCTION LEVEL PARALLELISM (ILP = 2) */ 
/*****************************************************/ 
__global__ void kernel1(int *d_a, int *d_b, int *d_c, unsigned int N) { 

    const int tid = threadIdx.x + blockIdx.x * blockDim.x; 

    if (tid < N/2) { 

     int a1 = d_a[tid]; 
     int b1 = d_b[tid]; 
     int c1 = d_c[tid]; 

     int a2 = d_a[tid+N/2]; 
     int b2 = d_b[tid+N/2]; 
     int c2 = d_c[tid+N/2]; 

     for(unsigned int i = 0; i < N_ITERATIONS; i++) { 
      a1 = a1 * b1 + c1; 
      a2 = a2 * b2 + c2; 
     } 

     d_a[tid]  = a1; 
     d_a[tid+N/2] = a2; 
    } 

} 

/*****************************************************/ 
/* KERNEL2 - INSTRUCTION LEVEL PARALLELISM (ILP = 4) */ 
/*****************************************************/ 
__global__ void kernel2(int *d_a, int *d_b, int *d_c, unsigned int N) { 

    const int tid = threadIdx.x + blockIdx.x * blockDim.x; 

    if (tid < N/4) { 

     int a1 = d_a[tid]; 
     int b1 = d_b[tid]; 
     int c1 = d_c[tid]; 

     int a2 = d_a[tid+N/4]; 
     int b2 = d_b[tid+N/4]; 
     int c2 = d_c[tid+N/4]; 

     int a3 = d_a[tid+N/2]; 
     int b3 = d_b[tid+N/2]; 
     int c3 = d_c[tid+N/2]; 

     int a4 = d_a[tid+3*N/4]; 
     int b4 = d_b[tid+3*N/4]; 
     int c4 = d_c[tid+3*N/4]; 

     for(unsigned int i = 0; i < N_ITERATIONS; i++) { 
      a1 = a1 * b1 + c1; 
      a2 = a2 * b2 + c2; 
      a3 = a3 * b3 + c3; 
      a4 = a4 * b4 + c4; 
     } 

     d_a[tid]  = a1; 
     d_a[tid+N/4] = a2; 
     d_a[tid+N/2] = a3; 
     d_a[tid+3*N/4] = a4; 
    } 

} 

/********/ 
/* MAIN */ 
/********/ 
void main() { 

    const int N = 1024; 

    int *h_a    = (int*)malloc(N*sizeof(int)); 
    int *h_a_result_host = (int*)malloc(N*sizeof(int)); 
    int *h_a_result_device = (int*)malloc(N*sizeof(int)); 
    int *h_b    = (int*)malloc(N*sizeof(int)); 
    int *h_c    = (int*)malloc(N*sizeof(int)); 

    for (int i=0; i<N; i++) { 
     h_a[i] = 2; 
     h_b[i] = 1; 
     h_c[i] = 2; 
     h_a_result_host[i] = h_a[i]; 
     for(unsigned int k = 0; k < N_ITERATIONS; k++) { 
      h_a_result_host[i] = h_a_result_host[i] * h_b[i] + h_c[i]; 
     } 
    } 

    int *d_a; gpuErrchk(cudaMalloc((void**)&d_a, N*sizeof(int))); 
    int *d_b; gpuErrchk(cudaMalloc((void**)&d_b, N*sizeof(int))); 
    int *d_c; gpuErrchk(cudaMalloc((void**)&d_c, N*sizeof(int))); 

    gpuErrchk(cudaMemcpy(d_a, h_a, N*sizeof(int), cudaMemcpyHostToDevice)); 
    gpuErrchk(cudaMemcpy(d_b, h_b, N*sizeof(int), cudaMemcpyHostToDevice)); 
    gpuErrchk(cudaMemcpy(d_c, h_c, N*sizeof(int), cudaMemcpyHostToDevice)); 

    // --- Creating events for timing 
    float time; 
    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 

    /***********/ 
    /* KERNEL0 */ 
    /***********/ 
    cudaEventRecord(start, 0); 
    kernel0<<<1, N>>>(d_a, d_b, d_c, N); 
#ifdef DEBUG 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 
#endif 
    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 
    cudaEventElapsedTime(&time, start, stop); 
    printf("GFlops = %f\n", (1.e-6)*(float)(N*N_ITERATIONS)/time); 
    gpuErrchk(cudaMemcpy(h_a_result_device, d_a, N*sizeof(int), cudaMemcpyDeviceToHost)); 
    for (int i=0; i<N; i++) if(h_a_result_device[i] != h_a_result_host[i]) { printf("Error at i=%i! Host = %i; Device = %i\n", i, h_a_result_host[i], h_a_result_device[i]); return; } 

    /***********/ 
    /* KERNEL1 */ 
    /***********/ 
    gpuErrchk(cudaMemcpy(d_a, h_a, N*sizeof(int), cudaMemcpyHostToDevice)); 
    cudaEventRecord(start, 0); 
    kernel1<<<1, N/2>>>(d_a, d_b, d_c, N); 
#ifdef DEBUG 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 
#endif 
    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 
    cudaEventElapsedTime(&time, start, stop); 
    printf("GFlops = %f\n", (1.e-6)*(float)(N*N_ITERATIONS)/time); 
    gpuErrchk(cudaMemcpy(h_a_result_device, d_a, N*sizeof(int), cudaMemcpyDeviceToHost)); 
    for (int i=0; i<N; i++) if(h_a_result_device[i] != h_a_result_host[i]) { printf("Error at i=%i! Host = %i; Device = %i\n", i, h_a_result_host[i], h_a_result_device[i]); return; } 

    /***********/ 
    /* KERNEL2 */ 
    /***********/ 
    gpuErrchk(cudaMemcpy(d_a, h_a, N*sizeof(int), cudaMemcpyHostToDevice)); 
    cudaEventRecord(start, 0); 
    kernel2<<<1, N/4>>>(d_a, d_b, d_c, N); 
#ifdef DEBUG 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 
#endif 
    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 
    cudaEventElapsedTime(&time, start, stop); 
    printf("GFlops = %f\n", (1.e-6)*(float)(N*N_ITERATIONS)/time); 
    gpuErrchk(cudaMemcpy(h_a_result_device, d_a, N*sizeof(int), cudaMemcpyDeviceToHost)); 
    for (int i=0; i<N; i++) if(h_a_result_device[i] != h_a_result_host[i]) { printf("Error at i=%i! Host = %i; Device = %i\n", i, h_a_result_host[i], h_a_result_device[i]); return; } 

    cudaDeviceReset(); 

} 

ですILP）が悪用されます。