Résolvez un problème classique de Map-Reduce avec OpenCL?

Question

J'essaie de parallèle à un problème de Map-Reduce classique (qui peut bien parallèle avec MPI) avec OpenCL, à savoir l'implémentation AMD. Mais le résultat me dérange.

Permettez-moi d'abord du problème. Il existe deux types de données qui circulent dans le système: l'ensemble de fonctionnalités (30 paramètres pour chacun) et l'ensemble d'échantillons (9000+ dimensions pour chacun). Il s'agit d'un problème de réduction de carte classique dans le sens où je dois calculer le score de chaque fonctionnalité sur chaque échantillon (carte). Et puis, résumez le score global pour chaque fonctionnalité (réduction). Il y a environ 10 000 caractéristiques et 30k échantillons.

J'ai essayé différentes façons de résoudre le problème. Tout d'abord, j'ai essayé de décomposer le problème par des fonctionnalités. Le problème est que le calcul du score se compose d'accès à la mémoire aléatoire (choisissez certaines des dimensions de 9000+ et des calculs DO plus / soustraction). Comme je ne peux pas fusionner l'accès à la mémoire, cela coûte. Ensuite, j'ai essayé de décomposer le problème par des échantillons. Le problème est que pour résumer le score global, tous les fils sont en concurrence pour quelques variables de score. Il continue d'écraser le score qui se révèle incorrect. (Je ne peux pas réaliser un score individuel en premier et résumer plus tard car il nécessite 10k * 30k * 4 octets).

La première méthode que j'ai essayée me donne la même performance sur le processeur i7 860 avec 8 threads. Cependant, je ne pense pas que le problème soit insoluble: il est remarquablement similaire au problème de traçage des rayons (pour lequel vous effectuez un calcul qui des millions de rayons contre des millions de triangles). Des idées?

De plus, je publie une partie du code que j'ai:

Décomposer par fonction (fonctionne, mais lent):

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

Décomposer par échantillon, pas travailler:

__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 de HN ici. Dès votre première tentative, je comprends mieux le problème et je vois que le choix d'échantillon dépend de la fonctionnalité est ce qui vous tue là-bas.

La sélection d'échantillon par fonctionnalité est-elle complètement aléatoire, ou pouvez-vous exploiter les régularités en cela (fonctionnalités de commande pour que ceux qui utilisent les mêmes échantillons soient traités ensemble)? C'est évident, donc je suppose que ce n'est pas possible.

Malheureusement, je ne comprends pas votre deuxième tentative.