حل مشكلة Reduce الكلاسيكية مع OpenCl؟

Question

أحاول موازاة مشكلة Reduce الكلاسيكية (والتي يمكن أن تتوازى بشكل جيد مع MPI) مع OpenCl ، وهي تطبيق AMD. لكن النتيجة تزعجني.

اسمحوا لي أن أوجز المشكلة أولاً. هناك نوعان من البيانات التي تتدفق إلى النظام: مجموعة الميزات (30 معلمة لكل منهما) ومجموعة العينة (9000+ أبعاد لكل منها). إنها مشكلة كلاسيكية لتجديد الخريطة ، بمعنى أنني بحاجة إلى حساب درجة كل ميزة على كل عينة (MAP). وبعد ذلك ، يلخص النتيجة الإجمالية لكل ميزة (تقليل). هناك حوالي 10 آلاف الميزات و 30 ألف عينة.

حاولت طرقًا مختلفة لحل المشكلة. أولاً ، حاولت تحلل المشكلة بالميزات. المشكلة هي أن حساب النتيجة يتكون من الوصول إلى الذاكرة العشوائية (اختر بعض الأبعاد البالغ عددها 9000 وأحسابات الاستدلال/الطرح). بما أنني لا أستطيع تجميع الوصول إلى الذاكرة ، فإنه يكلف. ثم ، حاولت تحلل المشكلة من خلال العينات. المشكلة هي أنه لتلخيص النتيجة الإجمالية ، تتنافس جميع المواضيع على عدد قليل من متغيرات الدرجات. إنه يحافظ على الكتابة فوق النتيجة التي تبين أنها غير صحيحة. (لا يمكنني تنفيذ النتيجة الفردية أولاً وتلخيصها لاحقًا لأنها تتطلب 10K * 30K * 4 بايت).

الطريقة الأولى التي جربتها تعطيني نفس الأداء على وحدة المعالجة المركزية I7 860 مع 8 مؤشرات ترابط. ومع ذلك ، لا أعتقد أن المشكلة غير قابلة للحل: فهي تشبه بشكل ملحوظ مشكلة تتبع الشعاع (التي تنفذها من أجل حساب ملايين الأشعة ضد ملايين المثلثات). أيه أفكار؟

بالإضافة إلى ذلك ، أقوم بنشر بعض الكود لدي:

تتحلل حسب الميزة (الأعمال ، ولكن بطيئة):

__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(....);
}

Answer 1

أندرو كوك من HN هنا. من محاولتك الأولى ، أفهم الآن المشكلة بشكل أفضل ، وأرى أن اختيار العينة يعتمد على الميزة هو ما يقتلك هناك.

هل اختيار العينة من خلال الميزة عشوائية تمامًا ، أو هل يمكنك استغلال المنظمات في ذلك (ميزات الطلب بحيث تتم معالجة تلك العينات نفسها معًا)؟ هذا واضح ، لذلك أعتقد أنه غير ممكن.

لسوء الحظ ، أنا لا أفهم محاولتك الثانية.