Решите классическую проблему с уменьшением карты с OpenCl?

Question

Я пытаюсь параллельно классическую проблему с восстановлением карты (которая может хорошо параллельно с 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 здесь. С вашей первой попытки я теперь лучше понимаю проблему и вижу, что выбор выборки зависит от функции - это то, что вас там убивает.

Является ли выбор выборки по функции совершенно случайным, или вы можете использовать закономерности в этом (упорядочение функций, чтобы те, которые используют одни и те же образцы, обрабатывались вместе)? Это очевидно, поэтому я думаю, что это невозможно.

К сожалению, я не понимаю вашу вторую попытку.