resolver um problema clássico de redução de mapa com opencl?

Question

Estou tentando paralelo a um problema clássico de redução de mapa (que pode ser bem paralelo ao MPI) com OpenCL, ou seja, a implementação AMD.Mas o resultado me incomoda.

Deixe-me resumir o problema primeiro.Existem dois tipos de dados que fluem para o sistema:o conjunto de recursos (30 parâmetros para cada) e o conjunto de amostras (mais de 9.000 dimensões para cada).É um problema clássico de redução de mapa, no sentido de que preciso calcular a pontuação de cada recurso em cada amostra (Mapa).E então, some a pontuação geral de cada recurso (Reduzir).Existem cerca de 10 mil recursos e 30 mil amostras.

Tentei diferentes maneiras de resolver o problema.Primeiro, tentei decompor o problema por características.O problema é que o cálculo da pontuação consiste em acesso aleatório à memória (escolha algumas das mais de 9.000 dimensões e faça cálculos de adição/subtração).Como não consigo unir o acesso à memória, isso custa.Então, tentei decompor o problema por amostras.O problema é que, para somar a pontuação geral, todos os threads competem por poucas variáveis ​​de pontuação.Ele continua substituindo a pontuação que acaba sendo incorreta.(Não posso realizar a pontuação individual primeiro e resumir depois porque requer 10k * 30k * 4 bytes).

O primeiro método que tentei me deu o mesmo desempenho na CPU i7 860 com 8 threads.No entanto, não acho que o problema seja insolúvel:é notavelmente semelhante ao problema de rastreamento de raios (para o qual você calcula milhões de raios contra milhões de triângulos).Alguma ideia?

Além disso, estou postando alguns dos códigos que tenho:

decompor por recurso (funciona, mas é lento):

__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;
}

decompor por amostra, não funcionar:

__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

Andrew Cooke, daqui.desde a sua primeira tentativa, agora entendo melhor o problema e vejo que ter a escolha da amostra dependendo do recurso é o que está matando você.

a seleção da amostra por recurso é completamente aleatória ou você pode explorar regularidades nisso (ordenar recursos para que aqueles que usam as mesmas amostras sejam processados ​​juntos)?isso é óbvio, então acho que não é possível.

infelizmente, não entendo sua segunda tentativa.