1

NVidia GTX 680で実行しているカーネルがあり、グローバルメモリからローカルメモリに切り替えると実行時間が長くなりました。

有限要素レイトレーサーの一部である私のカーネルは、処理する前に各要素をローカルメモリにロードするようになりました。各要素のデータは、次の定義を持つstructfastTriangleに格納されます。

typedef struct fastTriangle {
    float cx, cy, cz, cw;
    float nx, ny, nz, nd;
    float ux, uy, uz, ud;
    float vx, vy, vz, vd;
} fastTriangle;

これらのオブジェクトの配列をカーネルに渡します。カーネルは次のように記述されています（簡潔にするために、無関係なコードを削除しました。

__kernel void testGPU(int n_samples, const int n_objects, global const fastTriangle *objects, __local int *x_res, __global int *hits) {
    // Get gid, lid, and lsize

    // Set up random number generator and thread variables

    // Local storage for the two triangles being processed
    __local fastTriangle triangles[2]; 

    for(int i = 0; i < n_objects; i++) {    // Fire ray from each object
        event_t evt = async_work_group_copy((local float*)&triangles[0], (global float*)&objects[i],sizeof(fastTriangle)/sizeof(float),0);

        //Initialise local memory x_res to 0's

        barrier(CLK_LOCAL_MEM_FENCE);
        wait_group_events(1, &evt);      


        Vector wsNormal = { triangles[0].cw*triangles[0].nx, triangles[0].cw*triangles[0].ny, triangles[0].cw*triangles[0].nz};

        for(int j = 0; j < n_samples; j+= 4) {
            // generate a float4 of random numbers here (rands

            for(int v = 0; v < 4; v++) {    // For each ray in ray packet
                //load the first object to be intesected
                evt = async_work_group_copy((local float*)&triangles[1], (global float*)&objects[0],sizeof(fastTriangle)/sizeof(float),0);

                // Some initialising code and calculate ray here
                // Should have ray fully specified at this point;


                for(int w = 0; w < n_objects; w++) {        // Check for intersection against each ray

                    wait_group_events(1, &evt);

                    // Check for intersection against object w


                    float det = wsDir.x*triangles[1].nx + wsDir.y*triangles[1].ny + wsDir.z*triangles[1].nz;
                    float dett = triangles[1].nd - (triangles[0].cx*triangles[1].nx + triangles[0].cy*triangles[1].ny + triangles[0].cz*triangles[1].nz);


                    float detpx = det*triangles[0].cx + dett*wsDir.x;
                    float detpy = det*triangles[0].cy + dett*wsDir.y;
                    float detpz = det*triangles[0].cz + dett*wsDir.z;


                    float detu = detpx*triangles[1].ux + detpy*triangles[1].uy + detpz*triangles[1].uz + det*triangles[1].ud;
                    float detv = detpx*triangles[1].vx + detpy*triangles[1].vy + detpz*triangles[1].vz + det*triangles[1].vd;


                    // Interleaving the copy of the next triangle
                    evt = async_work_group_copy((local float*)&triangles[1], (global float*)&objects[w+1],sizeof(fastTriangle)/sizeof(float),0);

                    // Complete intersection calculations

                } // end for each object intersected

                if(objectNo != -1) atomic_inc(&x_res[objectNo]);
            } // end for sub rays
        } // end for each ray
        barrier(CLK_LOCAL_MEM_FENCE);

        // Add all the local x_res to global array hits


        barrier(CLK_GLOBAL_MEM_FENCE);
    } // end for each object
}

このカーネルを最初に作成したとき、各オブジェクトをローカルメモリにバッファリングせず、代わりにグローバルメモリからアクセスしました。つまり、triangles[0].cxの代わりにobjects[i].cxを使用しました。

最適化に着手したとき、上記のようにローカルメモリを使用するように切り替えましたが、実行実行時間が約25％増加することがわかりました。

グローバルメモリ内のオブジェクトに直接アクセスするのではなく、ローカルメモリを使用してオブジェクトをバッファリングすると、パフォーマンスが低下するのはなぜですか？

4

1 に答える 1

2

ローカルメモリが高速化に役立つかどうかは、プログラムに大きく依存します。ローカル メモリを使用する場合は、次の 2 つの点を考慮する必要があります。

  1. データをグローバルからローカルにコピーし、ローカルからグローバルに再度コピーするときに、追加の計算が必要になります。

  2. 「barrier(...)」が 3 回ありますが、これらのバリアはパフォーマンス キラーです。すべての OpenCL タスクは、他のすべてのタスクをバリアで待機する必要があります。このようにして、並列処理が妨げられ、タスクが独立して実行されなくなります。

ローカル メモリは、計算でデータを何度も読み取る場合に最適です。ただし、高速な読み取りと書き込みでは、コピーと同期よりも多くのパフォーマンスを向上させる必要があります。

于 2012-09-24T06:55:00.257 に答える