Nombre total de transactions de chargement en accès mémoire coalescé

Question

J'ai créé un noyau simple pour tester l'accès mémoire coalescé en observant le nombre de transactions, dans la carte nvidia gtx980. Le noyau est,

__global__ 
void copy_coalesced(float * d_in, float * d_out) 
{ 
    int tid = threadIdx.x + blockIdx.x*blockDim.x; 

    d_out[tid] = d_in[tid]; 
} 

Quand je lance cela avec les configurations du noyau suivantes

#define BLOCKSIZE 32 

int data_size = 10240;     //always a multiply of the BLOCKSIZE 
int gridSize = data_size/BLOCKSIZE; 

copy_coalesced<<<gridSize, BLOCKSIZE>>>(d_in, d_out); 

Depuis l'accès aux données dans le noyau est entièrement coalasced, et puisque le type de données est flottante (4 octets) , Le nombre de charge/Transactions de magasin attendus peuvent être trouvés comme suit,

charge transaction Taille = 32 octets

nombre de flotteurs qui peut être chargé par trans Action = 32 octets/4 octets = 8

Nombre de transactions nécessaires pour charger des données 10240 = 10240/8 = 1280 transactions

Le même montant des transactions sont attendues pour l'écriture des données aussi bien.

Mais en observant les mesures nvprof, les résultats suivants ont été

gld_transactions 2560 
gst_transactions 1280 

gld_transactions_per_request 8.0 
gst_transactions_per_request 4.0 

Je ne peux pas comprendre pourquoi il faut deux fois les opérations dont il a besoin pour le chargement des données. Mais quand il s'agit de l'efficacité de chargement/stockage à la fois les métriques donne 100%

Qu'est-ce qui me manque ici?

Answer 1

Je reproduit vos résultats sur linux,

1 gld_transactions    Global Load Transactions    2560 
1 gst_transactions    Global Store Transactions    1280 
1 l2_tex_read_transactions  L2 Transactions (Texture Reads)  1280 
1 l2_tex_write_transactions L2 Transactions (Texture Writes)  1280 

Toutefois, sur Windows à l'aide nSight édition Visual Studio, je reçois des valeurs qui semblent être mieux:

Vous pouvez contactez NVIDIA car il pourrait s'agir simplement d'un problème d'affichage dans nvprof.