Глобальная транзакция транзакции при доступе к объединенной памяти

Question

Я создал простое ядро ​​для проверки доступа к объединенной памяти, наблюдая количество транзакций в nvidia gtx980. Ядро,

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

    d_out[tid] = d_in[tid]; 
} 

При запуске этого со следующими конфигурациями ядра

#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); 

Поскольку доступа к данным в ядре полностью coalasced, а так как тип данных с плавающей точкой (4 байта) , количество загрузки/хранения сделок ожидается, можно найти следующим образом,

нагрузки транзакций Размер = 32 байта

Количество поплавков, которые могут быть загружены в транс действие = 32 байта/4 байт = 8

Количество операций, необходимых для загрузки данных 10240 = 10240/8 = 1280 сделок

Такое же количество сделок, как ожидается, для записи данных, а также.

Но при соблюдении nvprof показателей, были следующие результаты

gld_transactions 2560 
gst_transactions 1280 

gld_transactions_per_request 8.0 
gst_transactions_per_request 4.0 

Я не могу понять, почему это происходит в два раза больше операций, которые нужны для загрузки данных. Но когда дело доходит до эффективности загрузки/хранения, метрика выдаёт 100%

Что я упускаю здесь?

Answer 1

я воспроизвел свои результаты на 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 

Однако, на Windows, используя Nsight Visual Studio издание, я получаю значения, которые кажутся лучше:

Вы можете свяжитесь с NVIDIA, поскольку это может быть проблемой отображения в nvprof.