0

一意のブロックインデックスを必要とし、3Dグリッド全体でスケーリングできるように、可能な限りスケーリングする必要があるカーネルがあるとします。

計算はかなり複雑に見えます。1つのスレッドだけで計算を実行し、共有メモリに保存します。それは良い考えですか？
すべての文献で、常にレジスタに格納されていますが、共有メモリの欠点は何ですか？

よくわかりませんが、レジスタが1サイクルであるのに対し、共有メモリの読み取り/書き込みアクセスは4サイクルですか？

それ以外の：

__global__ kernel()
{
    //get unique 3D block index
    const unsigned long long int blockId = blockIdx.x //1D
        + blockIdx.y * gridDim.x //2D
        + gridDim.x * gridDim.y * blockIdx.z; //3D
}

多分使用:(ブロックのx次元のみが使用されると仮定）

__global__ kernel()
{
    __shared__ unsigned long long int blockId_s;

    if(threadIdx.x == 0)
        blockId_s = blockIdx.x //1D
            + blockIdx.y * gridDim.x //2D
            + gridDim.x * gridDim.y * blockIdx.z; //3D
    __syncthreads();
}

これにより、スレッドごとに1つのレジスタが節約され、コンピューティング機能1.xではコストがかかります。

私にはテストがなく、パフォーマンスに良いか悪いかわかりません。__syncthreads()cc 1.xで使用可能なもう1つのレジスタは引数ですが、ステートメントを使用するとパフォーマンスが少し遅くなるはずです。

4

2 に答える 2

2

この場合、答えは「はい」でしたが、主な理由は、結果のコードが使用するレジスターが少なく、これにより全体的な占有率が高くなり、速度がいくらか向上したためです。これは、最初のワープにいくつかの分岐の相違と同期プリミティブがあるコードにも関わらずでした。

ただし、これを普遍的なルールと見なすべきではありません。確実にする唯一の方法は、コードを記述し、ターゲットGPUでベンチマークすることです。

この回答は、この質問を未回答のキューから削除するために、コミュニティwikiエントリとしてコメントから追加されました。

于 2014-05-26T05:55:42.743 に答える
-2

共有男性の場合、blockId_sにアクセスするたびにn-wayバンク競合（n = WarpSize）が発生します。[CUDA Cプログラミングガイド、5.3.2.3]

そして、+共有メモリは登録よりも遅いです。

于 2012-04-28T19:43:33.977 に答える