2

私は、古典的なマップ削減の問題 (MPI とうまく並列できる) を OpenCL、つまり AMD 実装と並列化しようとしています。しかし、結果が気になります。

最初に問題について簡単に説明させてください。システムに流れ込むデータには、機能セット (それぞれに 30 個のパラメーター) とサンプル セット (それぞれに 9000 以上のディメンション) の 2 種類があります。これは、すべてのサンプル (マップ) のすべての機能のスコアを計算する必要があるという意味で、古典的なマップ削減問題です。そして、すべての機能の全体的なスコアを合計します (削減)。約 10,000 の機能と 30,000 のサンプルがあります。

問題を解決するためにさまざまな方法を試しました。まず、問題を特徴ごとに分解してみました。問題は、スコア計算がランダム メモリ アクセスで構成されていることです (9000 以上のディメンションのいくつかを選択し、プラス/マイナス計算を行います)。メモリアクセスを合体できないのでコストがかかります。次に、問題をサンプルごとに分解してみました。問題は、全体的なスコアを合計すると、すべてのスレッドがいくつかのスコア変数をめぐって競合することです。間違っていることが判明したスコアを上書きし続けます。（10k * 30k * 4バイトが必要なので、最初に個別のスコアを実行して後で合計することはできません）。

最初に試した方法では、8 スレッドの i7 860 CPU でも同じパフォーマンスが得られました。ただし、この問題が解決できないとは思いません。レイ トレーシングの問題 (数百万の三角形に対して数百万のレイを計算する) と非常によく似ています。何か案は？

さらに、私が持っているコードのいくつかを投稿しています：

機能ごとに分解します（機能しますが、遅いです）：

__kernel void __ccv_cl_pos_error_rate(__global unsigned int* err_rate,
__constant int* feature, __constant int* data, int num, __constant
unsigned int* w, int s, int isiz0, int isiz01, int step0, int step1)
{
    int igrid = get_global_id(0);
    __constant int* of = feature + igrid * 30;
    unsigned int e = 0;
    int k, i;
    int step[] = { step0, step1 };
    for (k = 0; k < num; k++)
    {
        __constant int* kd = data + k * isiz01;
        int pmin = kd[of[0] * isiz0 + of[1] + of[2] * step[of[0]]];
        int nmax = kd[of[3] * isiz0 + of[4] + of[5] * step[of[3]]];
        for (i = 0; i < 5; i++)
        {
            if (of[i * 6] >= 0)
                pmin = min(pmin, kd[of[i * 6] * isiz0 + of[i * 6 + 1] + of[i * 6 + 2] * step[of[i * 6]]]);
            if (of[i * 6 + 3] >= 0)
                nmax = max(nmax, kd[of[i * 6 + 3] * isiz0 + of[i * 6 + 4] + of[i * 6 + 5] * step[of[i * 6 + 3]]]);
        }
        if (pmin <= nmax)
            e += w[s + k];
    }
    err_rate[igrid] += e;
}

サンプルごとに分解しますが、機能しません:


__kernel void __ccv_cl_pos_error_rate(__global unsigned int* err_rate,
__constant int* feature, __constant int* data, int num, __constant
unsigned int* w, int s, int isiz0, int isiz01, int step0, int step1,
__local int* shared)
{
    int igrid = get_global_id(0);
    int lsize = get_local_size(0);
    int lid = get_local_id(0);
    unsigned int e = 0;
    int k, i;
    int ws = w[s + igrid];
    int step[] = { step0, step1 };
    for (k = 0; k < isiz01; k += lsize)
        if (k + lid < isiz01)
            shared[k + lid] = data[igrid * isiz01 + k + lid];
    barrier(....);
    for (k = 0; k < num; k++)
    {
        __constant int* of = feature + k * 30;
        int pmin = shared[of[0] * isiz0 + of[1] + of[2] * step[of[0]]];
        int nmax = shared[of[3] * isiz0 + of[4] + of[5] * step[of[3]]];
        for (i = 0; i < 5; i++)
        {
            if (of[i * 6] >= 0)
                pmin = min(pmin, shared[of[i * 6] * isiz0 + of[i * 6 + 1] + of[i * 6 + 2] * step[of[i * 6]]]);
            if (of[i * 6 + 3] >= 0)
                nmax = max(nmax, shared[of[i * 6 + 3] * isiz0 + of[i * 6 + 4] + of[i * 6 + 5] * step[of[i * 6 + 3]]]);
        }
        if (pmin <= nmax)
            err_rate[k] += ws; // here is wrong.
    }
    barrier(....);
}
4

1 に答える 1

1

hn のアンドリュー・クックです。あなたの最初の試みから、私は問題をよりよく理解するようになりました.

機能ごとのサンプルの選択は完全にランダムですか、それともその規則性を活用できますか (同じサンプルを使用する機能が一緒に処理されるように機能を順序付けます)? これは明らかなので、不可能だと思います。

残念ながら、私はあなたの 2 回目の試みを理解できません。

于 2010-04-01T11:01:51.833 に答える