Я пытаюсь распараллелить классическую задачу сокращения карт (которая может хорошо сочетаться с MPI) с OpenCL, а именно реализацию AMD. Но результат меня беспокоит.
Позвольте мне сначала кратко рассказать о проблеме. Существует два типа данных, которые поступают в систему: набор функций (30 параметров для каждого) и набор образцов (9000+ измерений для каждого). Это классическая проблема уменьшения карты в том смысле, что мне нужно вычислить оценку каждого объекта в каждом образце (карта). А затем суммируйте общий балл по каждой функции (уменьшить). Есть около 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(....);
}