なぜクワのブロック数を増やすと時間が増えますか？

Question

私が理解していることは、CUDAではブロック数が並列化されているので時間が増えることはないが、私のコードではブロック数を倍にすれば時間も倍増するということです。

#include <cuda.h> 
#include <curand.h> 
#include <curand_kernel.h> 
#include <stdio.h> 
#include <stdlib.h> 
#include <iostream> 

#define num_of_blocks 500 
#define num_of_threads 512 

__constant__ double y = 1.1; 

// set seed for random number generator 
__global__ void initcuRand(curandState* globalState, unsigned long seed){ 
    int idx = threadIdx.x + blockIdx.x * blockDim.x; 
    curand_init(seed, idx, 0, &globalState[idx]); 
} 

// kernel function for SIR 
__global__ void test(curandState* globalState, double *dev_data){ 
    // global threads id 
    int idx  = threadIdx.x + blockIdx.x * blockDim.x; 

    // local threads id 
    int lidx = threadIdx.x; 

    // creat shared memory to store seeds 
    __shared__ curandState localState[num_of_threads]; 

    // shared memory to store samples 
    __shared__ double sample[num_of_threads]; 

    // copy global seed to local 
    localState[lidx] = globalState[idx]; 
    __syncthreads(); 

    sample[lidx] = y + curand_normal_double(&localState[lidx]); 

    if(lidx == 0){ 
     // save the first sample to dev_data; 
     dev_data[blockIdx.x] = sample[0]; 
    } 

    globalState[idx] = localState[lidx]; 
} 

int main(){ 
    // creat random number seeds; 
    curandState *globalState; 
    cudaMalloc((void**)&globalState, num_of_blocks*num_of_threads*sizeof(curandState)); 
    initcuRand<<<num_of_blocks, num_of_threads>>>(globalState, 1); 

    double *dev_data; 
    cudaMalloc((double**)&dev_data, num_of_blocks*sizeof(double)); 

    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 
    // Start record 
    cudaEventRecord(start, 0); 

    test<<<num_of_blocks, num_of_threads>>>(globalState, dev_data); 

    // Stop event 
    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 
    float elapsedTime; 
    cudaEventElapsedTime(&elapsedTime, start, stop); // that's our time! 
    // Clean up: 
    cudaEventDestroy(start); 
    cudaEventDestroy(stop); 

    std::cout << "Time ellapsed: " << elapsedTime << std::endl; 

    cudaFree(dev_data); 
    cudaFree(globalState); 
    return 0; 
} 

テスト結果は次のとおりです。

number of blocks: 500, Time ellapsed: 0.39136. 
number of blocks: 1000, Time ellapsed: 0.618656. 

ので時間が増加することの理由は何ですか？それは定数メモリにアクセスするか、または共有メモリからグローバルメモリにデータをコピーするためですか？それはそれを最適化するいくつかの方法ですか？

Answer 1

パラレルで実行できるブロックの数は大きくなる可能性がありますが、オンチップリソースが限られているため、依然として有限です。カーネルの起動で要求されたブロックの数がその制限を超えた場合、それ以上のブロックは以前のブロックが終了してリソースを解放するのを待たなければなりません。

1つの制限付きリソースは、共有メモリで、カーネルは28キロバイトを使用します。 CUDA 8.0互換Nvidia GPUは、ストリーミングマルチプロセッサ（SM）あたり48〜112キロバイトの共有メモリを提供するため、一度に実行されるブロックの最大数は、GPU上のSMの数の1倍から3倍です。

その他の制限されたリソースは、スケジューラ内のレジスタおよびワード単位のさまざまなリソースです。 CUDA occupancy calculatorは、便利なExcelスプレッドシート（​​OpenOffice/LibreOfficeでも動作します）で、これらのリソースが特定のカーネルのSMあたりのブロック数を制限する方法を示しています。 nvccコマンドラインにオプション--ptxas-options="-v"を追加し、カーネルをコンパイルし、ラインがptxas情報」と言って探し：使用XXレジスタ、YYバイトSMEM、ZZバイトcmemの[0]、WWバイトcmemの[2] XX,YY、ブロックしようとしているスレッドの数、GPUの計算能力をスプレッドシートに入力してください。 1つのSM上で並列に実行できるブロックの最大数が表示されます。

テストを実行しているGPUについては言及していませんので、例としてGTX 980を使用します。これにはそれぞれ96Kbの共有メモリを持つ16個のSMがあるため、最大で16×3 = 48個のブロックを並行して実行できます。共有メモリを使用しなかった場合、常駐ワープの最大数によって、SMあたりのブロック数が4に制限され、64ブロックを並行して実行できます。

現在の既存のNvidia GPUでは、ブロック数を2倍にすると実行時間が約2倍になる理由を説明しているので、