我正在尝试将一个经典的 map-reduce 问题(可以与 MPI 很好地并行)与 OpenCL,即 AMD 实现并行。但结果让我很困扰。
我先简单介绍一下问题。有两种类型的数据流入系统:特征集(每个参数 30 个)和样本集(每个参数 9000 多个维度)。这是一个经典的 map-reduce 问题,因为我需要计算每个样本(Map)上每个特征的得分。然后,总结每个特征的总体得分(减少)。大约有 10k 个特征和 30k 个样本。
我尝试了不同的方法来解决这个问题。首先,我尝试通过特征来分解问题。问题是分数计算由随机内存访问组成(选择 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(....);
}