bande passante mémoire effective d'un accès mémoire strided

Question

Supposons que j'ai un noyau qui effectue un accès mémoire strided comme suit:

__global__ void strideExample (float *outputData, float *inputData, int stride=2) 
{ 
     int index = (blockIdx.x * blockDim.x + threadIdx.x) * stride; 
     outputData[index] = inputData[index]; 
} 

Je comprends que les accès avec une taille foulée de 2 entraînera une efficacité de charge/stockage de 50% , puisque la moitié des éléments impliqués dans la transaction ne sont pas utilisés (perte de bande passante). Comment procédons-nous pour calculer l'efficacité de chargement/stockage pour les grandes foulées? Merci d'avance!

Answer 1

En général:

load efficiency = requested loads/effective loads 

Où requested loads est le nombre d'octets que le logiciel demandé de lire et effective loads est le nombre d'octets que le matériel effectivement devait lire. La même formule s'applique aux magasins.

accès parfaitement coalisées ont une efficacité de 1.

Vos demandes de code exactement (blockIdx.x * blockDim.x + threadIdx.x) * sizeof(float) octets. En supposant que outputData est correctement aligné (comme le sont les pointeurs retournés par cudaMalloc), le matériel devra lire (blockIdx.x * blockDim.x + threadIdx.x) * sizeof(float) * stride octets, arrondi à la taille de la transaction (128 octets pour SM/L1, 32 octets pour L1/L2). En supposant que votre taille de bloc est assez grande, l'arrondi à la taille de la transaction devient négligeable et vous pouvez simplifier l'équation à 1/stride, ce qui donne dans ce cas une efficacité de charge d'environ ~ 16,7%.