référence OpenCL - Avis sur les paramètres pour varier

Question

je voudrais faire référence à l'exécution de la réduction de la somme en deux étapes avec OpenCL (de cette AMD link) sur une radeon HD 7970 XT Tahiti. Initialement, j'ai pris une première version de code où je n'ai pas utilisé la première boucle qui effectue une réduction d'un tableau d'entrée de taille N à un tableau de sortie de taille NworkItems. Voici cette première boucle dans le code du noyau:

int global_index = get_global_id(0); 
    float accumulator = 0; 
    // Loop sequentially over chunks of input vector 
    while (global_index < length) { 
    float element = buffer[global_index]; 
    accumulator += element; 
    global_index += get_global_size(0); 
    } 

Donc, avec cette première version, j'ai mesuré le temps d'exécution en fonction de la taille du tableau d'entrée (qui est égal au nombre total de threads) et pour différentes tailles de travail groupe. Voici les résultats:

Maintenant, je voudrais faire une référence où j'utilise cette boucle initiale ci-dessus. Mais je ne sais pas quels paramètres je dois varier.

on dit que AMD recommande De this link, un multiple de 64 pour la taille d'un groupe de travail (32 pour NVIDIA).

De plus, à partir du dernier commentaire sur this other link, il est recommandé de définir la taille du groupe de travail comme: WorkGroup size = (Number of total threads)/(Compute Units). Sur ma carte GPU, j'ai 32 unités de calcul.

Je voudrais obtenir des conseils pour savoir quels paramètres serait intéressant de faire varier afin de comparer runtimes dans cette deuxième version (avec la première boucle de réduction). Par exemple, je peux prendre des valeurs différentes pour le rapport (N size of input array)/(total NworkItems) et une valeur fixe pour WorkGroup size (voir expression ci-dessus),

ou faire au contraire, i.e. dois-je modifier la valeur pour WorkGroup size et fixer le (N size of input array)/(total NworkItems) ratio?

Toute idée est la bienvenue, Merci

Answer 1

Vous devez la somme des données locales au lieu de propager des données, pour faciliter le transfert de mémoire (coalescées accès aux données). Donc, utilisez-le à la place:

int chunk_size = length/get_global_size(0)+(length%get_global_size(0) > 0); //Will give how many items each work item needs to process 
    int global_index = get_group_id(0)*get_local_size(0)*chunk_size + get_local_id(0); //Start at this address for this work item 
    float accumulator = 0; 

    for(int i=0; i<chunk_size; i++) 
    // Loop sequentially over chunks of input vector 
    if (global_index < length) { 
     float element = buffer[global_index]; 
     accumulator += element; 
     global_index += get_local_size(0); 
    } 
    } 

Vous devriez également utiliser des tailles de deux puissances pour faciliter la mise en cache.