CUDA では、スレッドは個々の SM でワープとしてスケジュールされます。各ワープには最大 32 のスレッドを含めることができます。スケジューラは、SM 内でワープを並行して実行しようとします。特定のワープのデータの準備が整っていない場合、データは利用可能になるまでスケジューラによって保持されます。あなたの質問は関係がcudaEvent_t
あります。(カーネルの実行時間を測定するために)を使用して達成しようとしていることを実行できると思います。
起動構成はkernel<<<B,Tnum>>>(arg1...argn);
、アルゴリズムでどの程度の並列処理を利用できるかによって完全に異なります。また、スレッドの数は、カーネルを起動して得られる最適な実行時間に基づいて決定する必要があります。
多くの場合、128/256
スレッドを使用して複数のブロックを起動するだけで、最適な速度アップを実現できます。例を挙げると、サイズの 2 つの配列の個々の要素を1024
3 番目の配列に追加するとします。1 ブロックのカーネル関数は次のようになります。
__global__ void kadd(int *c,int *a,int *b)
{
unsigned int tid = threadIdx.x;//Since only one block of 1024 threads suffices
if(tid < MAXNUM) //MAXNUM = 1024
c[tid] = a[tid]+ b[tid];
}
そして、起動構成は次のようになります
kadd<<<1,1024>>>(c,a,b);
ただし、これは GPU の SM の 1 つでブロックを実行するだけであり、GPU リソースを十分に活用していないことを意味します。GPU をさらに活用するには、複数のブロックとスレッドを使用できます。カーネルは次のようになります。
__global__ void kadd(int *c,int *a,int *b)
{
unsigned int tid = blockIDx.x * blockDim.x + threadIdx.x;//Since multiple blocks are used
if(tid < MAXNUM) //MAXNUM = 1024
c[tid] = a[tid]+ b[tid];
}
対応する起動構成は次のようになります
kadd<<<8,128>>>(c,a,b);
これにより、スレッド8
のブロックがそれぞれ起動128
されます。アルゴリズムの要件に基づいて、この起動構成をいじることができます。GPU を最大限に活用するために、 launching2D
またはgridによってこれらの起動構成をさらに調べることができます。3D
そのため、カーネルのタイミングを調整すると、要件に最適な構成が得られます。これは、共有メモリの使用、グローバルメモリの合体アクセスなどによっても変化します。最後に、より高い占有率を実現するために、ブロックとスレッドの最適な組み合わせを見つけるために使用できる、NVIDIA が提供する占有率計算ツールがあることをお伝えしたいと思います。