1

並列に実行できるスレッド/ブロックの数は限られていると思います。つまり、スレッド/ブロックが多すぎると、それらの一部が一部の処理ユニットで順次実行されます。次の例を作成する必要があります。としましょう、私はいくつか持っていkernel<<<B, N>>>()ます。の実行時間kernel<<<1,1>>>()は t0 です。

kernel<<<B, N>>>()最初のタスクは、実行時間t ~ t0である B と N の最大値を見つけることです。そして、実行時間をkernel<<<B, 2*N>>>()(またはkernel<<<2*B, N>>>()) t1 ~ 2*t にしたいと考えています。

私は 448 個の CUDA コア (14 個の SM) を備えた Tesla C2075 を所有しており、占有率 1 の例を構築したいと考えています。

これは可能ですか? はいの場合、カーネル関数はどのように見えるべきですか?

4

1 に答える 1

0

CUDA では、スレッドは個々の SM でワープとしてスケジュールされます。各ワープには最大 32 のスレッドを含めることができます。スケジューラは、SM 内でワープを並行して実行しようとします。特定のワープのデータの準備が整っていない場合、データは利用可能になるまでスケジューラによって保持されます。あなたの質問は関係がcudaEvent_tあります。(カーネルの実行時間を測定するために)を使用して達成しようとしていることを実行できると思います。

起動構成はkernel<<<B,Tnum>>>(arg1...argn);、アルゴリズムでどの程度の並列処理を利用できるかによって完全に異なります。また、スレッドの数は、カーネルを起動して得られる最適な実行時間に基づいて決定する必要があります。

多くの場合、128/256スレッドを使用して複数のブロックを起動するだけで、最適な速度アップを実現できます。例を挙げると、サイズの 2 つの配列の個々の要素を10243 番目の配列に追加するとします。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 が提供する占有率計算ツールがあることをお伝えしたいと思います。

于 2012-10-07T03:30:45.430 に答える