1

私はいくつかの CUDA プログラムに取り組んでおり、定数メモリを使用して計算を高速化したかったのですが、定数メモリを使用するとコードが約 30% 遅くなることがわかりました。

コンスタント メモリはワープ全体に読み取りをブロードキャストするのに適していることを知っており、プログラムでそれを利用できると考えました。

定数メモリ コードは次のとおりです。

__constant__ float4 constPlanes[MAX_PLANES_COUNT];

__global__ void faultsKernelConstantMem(const float3* vertices, unsigned int vertsCount, int* displacements, unsigned int planesCount) {

    unsigned int blockId = __mul24(blockIdx.y, gridDim.x) + blockIdx.x;
    unsigned int vertexIndex = __mul24(blockId, blockDim.x) + threadIdx.x;

    if (vertexIndex >= vertsCount) {
        return;
    }

    float3 v = vertices[vertexIndex];
    int displacementSteps = displacements[vertexIndex];

    //__syncthreads();

    for (unsigned int planeIndex = 0; planeIndex < planesCount; ++planeIndex) {
        float4 plane = constPlanes[planeIndex];
        if (v.x * plane.x + v.y * plane.y + v.z * plane.z + plane.w > 0) {
            ++displacementSteps;
        }
        else {
            --displacementSteps;
        }
    }

    displacements[vertexIndex] = displacementSteps;
}

グローバル メモリ コードは同じですが、1 つ多くのパラメーター (プレーンの配列へのポインター) があり、グローバル配列の代わりにそれを使用します。

私はそれらの最初のグローバルメモリが読み取ると思った

float3 v = vertices[vertexIndex];
int displacementSteps = displacements[vertexIndex];

スレッドの「非同期化」を引き起こす可能性があり、定数メモリ読み取りのブロードキャストを利用しないため、__syncthreads(); を呼び出してみました。定数メモリを読み取る前ですが、何も変更されませんでした。

なにが問題ですか？前もって感謝します！

システム：

  • CUDA ドライバーのバージョン: 5.0
  • CUDA 機能: 2.0

パラメーター：

  • 頂点数: ~250 万
  • 飛行機の数: 1024

結果：

  • 定数メモリ バージョン: 46 ミリ秒
  • グローバル メモリ バージョン: 35 ミリ秒

編集：

そのため、定数メモリを高速化する方法として、次のような多くのことを試しました。

1) 2 つのグローバル メモリ読み取りをコメント アウトして、影響があるかどうかを確認します。グローバル メモリはさらに高速でした。

2) CM キャッシュを利用するために、スレッドごとにより多くの頂点 (8 から 64) を処理します。これは、スレッドごとに 1 つの頂点よりもさらに遅くなりました。

2b) 共有メモリを使用してディスプレイスメントと頂点を保存します。最初にそれらすべてをロードし、すべてのディスプレイスメントを処理して保存します。ここでも、示されている CM の例よりも低速です。

この経験の後、私は CM 読み取りブロードキャストがどのように機能し、コードでどのように正しく「使用」できるかを本当に理解していません。このコードは、おそらく CM では最適化できません。

EDIT2：

微調整の別の日、私は試しました：

3) スレッドごとにより多くの頂点 (8 から 64) をメモリ合体で処理します (すべてのスレッドは、システム内のスレッドの総数に等しいインクリメントで進みます) -- これにより、インクリメントが 1 の場合よりも良い結果が得られますが、速度は向上しません。

4) この if ステートメントを置き換えます

if (v.x * plane.x + v.y * plane.y + v.z * plane.z + plane.w > 0) {
    ++displacementSteps;
}
else {
    --displacementSteps;
}

このコードを使用して分岐を避けるために、少しの数学で「予測できない」結果が得られます。

float dist = v.x * plane.x + v.y * plane.y + v.z * plane.z + plane.w;
int distInt = (int)(dist * (1 << 29));  // distance is in range (0 - 2), stretch it to int range
int sign = 1 | (distInt >> (sizeof(int) * CHAR_BIT - 1));  // compute sign without using ifs
displacementSteps += sign;

残念ながら、これは if so if を使用するよりもかなり遅い (~30%) ので、思ったほど大きな悪ではありません。

EDIT3：

この問題はおそらく定数メモリを使用しても改善できないと結論付けています。これらは私の結果です*：

グローバルおよびコンスタント メモリ パフォーマンスのグラフ

*時間は、15 回の独立した測定値の中央値として報告されています。定数メモリがすべてのプレーン (4096 および 8192) を保存するのに十分な大きさでない場合、カーネルが複数回呼び出されました。

4

1 に答える 1

2

計算機能 2.0 チップには 64k の固定メモリがありますが、各マルチプロセッサには 8k の固定メモリ キャッシュしかありません。コードには、16k の定数メモリすべてへのアクセスを必要とする各スレッドがあるため、キャッシュ ミスによってパフォーマンスが低下しています。プレーン データの定数メモリを効果的に使用するには、実装を再構築する必要があります。

于 2013-03-06T07:33:13.933 に答える