Confus au sujet du profilage des résultats de mon noyau OpenCL (DSP fenêtre Fonction)

Question

J'ai terminé un noyau de fonction de fenêtre dans OpenCL. Fondamentalement, une fonction de fenêtre applique seulement un ensemble de coefficients sur un autre ensemble de nombres pièce par pièce (Wikipedia explique mieux). J'ai pu bourrer le tableau de flotteur coefficient de fenêtre dans le cache constant pour la plupart des cas.

J'attendais mes résultats de Compute Prof pour montrer que l'hôte à l'appareil et le dispositif d'accueil des transferts de mémoire prendrait plus de 95% du temps de traitement. Pour presque tous mes cas, il est seulement 80% du temps de traitement. Je suis en train d'écrire et de lire un tableau de 4,2 millions de flotteur et du conseil d'administration et d'écrire un autre réseau de flotteurs qui reste généralement bien au-dessous d'un million.

quoi que ce soit fait dans le poisson regard du noyau? Tous les avis sur si elle est un problème qui devrait courir plus vite sur un GPU qu'une unité centrale de traitement en premier lieu (je suis pas encore à 100% à ce sujet). Je suis un peu étonné que les raisons pour lesquelles mon gld_efficiency et vol stationnaire gst_efficiency entre 0,1 et 0,2. J'ai fait ce noyau avec la mémoire globale G80 coalescent à l'esprit. Ma mémoire globale semble globale à un débit bien 40gbs. Le noyau est assez simple et est affiché ci-dessous.

__kernel void window(__global float* inputArray, // first frame to ingest starts at 0.  Sized to nFramesToIngest*framesize samples
    __constant float* windowArray, // may already be partly filled
    int windowSize, // size of window frame, in floats
    int primitivesPerDataFrame, //amount of primitives in each frame of inputArray parameter
    int nInFramesThisCall, //each thread solves a frame, so this integer represent how many threads this kernel launches
    int isRealNumbers //0 for complex, non-zero for real 
)
{
int gid = get_global_id(0) + get_global_size(0) * get_global_id(1);

if(gid < nInFramesThisCall) //make sure we don't execute unnecessary threads
{
    if(isRealNumbers)
    {
        for(int i = 0; i < primitivesPerDataFrame; i++)
        {
            int inputArrayIndex = (gid*primitivesPerDataFrame)+i;
            inputArray[inputArrayIndex] = inputArray[inputArrayIndex] * windowArray[i%windowSize];
        }
    }
    else //complex
    {
        for(int i = 0; i < primitivesPerDataFrame; i++)
        {
            int inputArrayIndex = (gid*primitivesPerDataFrame)+i;
            inputArray[inputArrayIndex] = inputArray[inputArrayIndex] * windowArray[i%windowSize/2];
        }
    }
}

}

Answer 1

Combien de threads (le terme OpenCL est travail-éléments, en passant) utilisez-vous? Vous avez besoin d'au moins quelque chose dans les centaines de charger un gros GPU efficacement.

Vous dites que vous voulez utiliser l'accès mémoire coalisée, mais une charge avec un décalage comme

int inputArrayIndex = (gid*primitivesPerDataFrame)+i;

ne sera pas que cela soit possible dans la plupart des cas. G80 NVidia a des restrictions assez sévères en matière de coalescent, voir pour plus d'informations « Guide OpenCL meilleures pratiques ». Fondamentalement, le travail à des éléments d'une chaîne ont des éléments d'accès d'un bloc aligné 64 ou 128 octets d'une certaine façon en même temps à faire des charges et les magasins se produisent coalescées.

Ou pour vous donner un exemple: si primitivesPerDataFrame est 16, les charges et les magasins d'une chaîne sont faites à des décalages espacés 16 éléments en dehors, ce qui rend toute coalescent efficace impossible

.