2

基本的に、ランダム/予測不可能な方法でグローバル メモリにアクセスする OpenCL カーネルを作成しています (現在、レイ トレーサーのパス トレーシング コンポーネントは最適化されていません)。これは、CPU に対する GPU の並列化されたパフォーマンスの利点をほぼ完全に無効にしています (参考までに、私は i7-2630QM CPU、GTX 560m GPU で実行しています - 以下のパフォーマンス数値)。微調整/テストを容易にするために、このメモリ アクセス パターンをシミュレートする「テスト」カーネルを作成しました。基本的に、処理する三角形座標の大きな配列とインデックスのリストを GPU に提供します。インデックスごとに、その三角形と 63 の後に ray-triangle 交差を実行し、octree 内のオブジェクトの反復を模倣します。

「グローバル」の代わりに読み取り専用テクスチャ メモリを使用する合体メモリ アクセス、ループ展開、ワーク グループ サイズとスレッド分散の微調整、ローカル メモリとバリア、および関数の手動インライン化など、さまざまな最適化を試しました。これらはすべて、せいぜい段階的なパフォーマンスの向上をもたらしました。カーネルを実行する前にインデックスを並べ替えると大幅に高速化されますが、オクツリー トラバーサルの場合、これには反復ごとに GPU での再並べ替えが必要であり、他の要因と組み合わせると、それが非常に役立つかどうか疑問に思います。

修正できる大きな穴があるかどうかを把握しようとしています-データ型の誤用、目に見えない最適化、ドライバーが古すぎる（1Dテクスチャを許可しないOpenCL 1.0を使用）など-または期待している場合私が使用しているハードウェアを考えると、パフォーマンスが大幅に向上しています (レイ トレーシング側のさまざまな最適化はまだ行われていませんが、それを掘り下げる前に、このより一般的な問題を解決したいと思います)。事前に洞察や提案をいただければ幸いです。

64 個の三角形の 409,600 ブロック (409,600 スレッドとして実行) のパフォーマンス数値 (秒):

CPU (Single Thread):

Unsorted:   2.21
Sorted:     1.48 

GPU:

        Sorted  Unsorted
Texture     0.07    0.15
Global      0.02    0.25

コード：

#define IMG_WIDTH_MINUS_ONE 32767
#define IMG_HEIGHT_LOG_2 15
#define SUB(dest,v1,v2) \
      dest[0]=v1[0]-v2[0]; \
      dest[1]=v1[1]-v2[1]; \
      dest[2]=v1[2]-v2[2];

#define EPSILON 0.00001

#define CROSS(dest,v1,v2) \
      dest[0]=v1[1]*v2[2]-v1[2]*v2[1]; \
      dest[1]=v1[2]*v2[0]-v1[0]*v2[2]; \
      dest[2]=v1[0]*v2[1]-v1[1]*v2[0];

#define DOT(v1,v2) (v1[0]*v2[0]+v1[1]*v2[1]+v1[2]*v2[2])

__kernel void square(
   __global int4 *inputIndeces,
   __read_only image2d_t image,
   __global float* output,
   const unsigned int count)
{

    int global_id = get_global_id(0);
    float r_orig[3];
    float r_dir[3];
    float4 trianglePoints[3];
    int cpuStartIndex = inputIndeces[global_id].x;
    int outputIndex = inputIndeces[global_id].w;
    output[outputIndex] = 0.0;
    r_orig[0] = 0.0;
    r_orig[1] = 0.0;
    r_orig[2] = 500.0;
    r_dir[0] = 0.0;
    r_dir[1] = 0.0;
    local int counter;
    counter = 0;
    r_dir[2]= -1.0;
    float tvec[3], pvec[3], qvec[3], edgeA[3], edgeB[3];
    float det, inv_det, t, u, v;
    #pragma unroll 64
    for (int ind=cpuStartIndex;ind<cpuStartIndex+64;++ind) {



        int tIndex = ind<<2;

        int2 coords[3];


        coords[0] = (int2)(tIndex & IMG_WIDTH_MINUS_ONE,tIndex >> IMG_HEIGHT_LOG_2);
        coords[1] = (int2)((tIndex + 1) & IMG_WIDTH_MINUS_ONE,(tIndex + 1) >> IMG_HEIGHT_LOG_2);
        coords[2] = (int2)((tIndex + 2) & IMG_WIDTH_MINUS_ONE,(tIndex + 2) >> IMG_HEIGHT_LOG_2);

        trianglePoints[0] = read_imagef(image, sampler, coords[0]);
        trianglePoints[1] = read_imagef(image, sampler, coords[1]);
        trianglePoints[2] = read_imagef(image, sampler, coords[2]);

        edgeA[0] = (trianglePoints[0].w - trianglePoints[0].x);
        edgeA[1] = (trianglePoints[1].x - trianglePoints[0].y);
        edgeA[2] = (trianglePoints[1].y - trianglePoints[0].z);

        edgeB[0] = (trianglePoints[1].z - trianglePoints[0].x);
        edgeB[1] = (trianglePoints[1].w - trianglePoints[0].y);
        edgeB[2] = (trianglePoints[2].x - trianglePoints[0].z);

        CROSS(pvec,r_dir,edgeB);
        det = DOT(edgeA, pvec);

        if (det > -EPSILON && det < EPSILON) {
             continue;
        }

        inv_det = 1.0 / det;

        tvec[0] = r_orig[0] - trianglePoints[0].x;
        tvec[1] = r_orig[1] - trianglePoints[0].y;
        tvec[2] = r_orig[2] - trianglePoints[0].z;

        u = DOT(tvec, pvec) * inv_det;
        if (u < 0.0 || u > 1.0) {
                continue;
            }

            CROSS(qvec,tvec,edgeA);
            v = DOT(r_dir, qvec) * inv_det;
            if (v < 0.0 || u + v > 1.0) {
                continue;
            }
            t = DOT(edgeB, qvec) * inv_det;
            if (t > 0.001) {
                ++counter;
            }
        else {
            continue;
        }

    }
    output[outputIndex] = (float)counter;
}
4

0 に答える 0