Использование CUDA Profiler nvprof для доступа к памяти

Question

Я использую nvprof для получения количества обращений к глобальной памяти для следующего кода CUDA. Количество загрузок в ядре равно 36 (доступ к массиву d_In), а количество ячеек в ядре 36 + 36 (для доступа к массиву d_Out и массиву d_rows). Таким образом, общее количество загрузок глобальной памяти составляет 36, а количество глобальных хранилищ памяти - 72. Однако, когда я просматриваю код с помощью профилировщика nvprof CUDA, он сообщает следующее: (В основном я хочу вычислить вычисление на глобальный доступ к памяти (CGMA))

 1    gld_transactions  Global Load Transactions   6   6   6 
     1    gst_transactions  Global Store Transactions   11   11   11 
     1   l2_read_transactions   L2 Read Transactions   133   133   133 
     1   l2_write_transactions   L2 Write Transactions   24   24   24 


#include <stdio.h> 
#include "cuda_profiler_api.h" 
__constant__ int crows; 

__global__ void kernel(double *d_In, double *d_Out, int *d_rows){ 
     int tx=threadIdx.x; 
     int bx=blockIdx.x; 
     int n=bx*blockDim.x+tx; 
     if(n < 36){ 
       d_Out[n]=d_In[n]+1; 
       d_rows[n]=crows; 
     } 
     return; 
} 

int main(int argc,char **argv){ 

    double I[36]={1,5,9,2,6,10,3,7,11,4,8,12,13,17,21,14,18,22,15,19,23,16,20,24,25,29,33,26,30,34,27,31,35,28,32,36}; 

    double *d_In; 
    double *d_Out; 
    int *d_rows; 

    double Iout[36]; 
    int rows=5; 
    int h_rows[36]; 

    cudaMemcpyToSymbol(crows,&rows,sizeof(int)); 
    cudaMalloc(&d_In,sizeof(double)*36); 
    cudaMalloc(&d_Out,sizeof(double)*36); 
    cudaMalloc(&d_rows,sizeof(int)*36); 

    cudaMemcpy(d_In,I,sizeof(double)*36,cudaMemcpyHostToDevice); 

    dim3 dimGrid(4,1,1); 
    dim3 dimBlock(10,1,1); 

    cudaProfilerStart(); 
    kernel<<<dimGrid,dimBlock>>>(d_In,d_Out,d_rows); 
    cudaProfilerStop(); 

    cudaMemcpy(Iout,d_Out,sizeof(double)*36,cudaMemcpyDeviceToHost); 
     cudaMemcpy(h_rows,d_rows,sizeof(int)*36,cudaMemcpyDeviceToHost); 


    int i; 
    for(i=0;i<36;i++) 
     printf("%f %d\n",Iout[i],h_rows[i]); 


} 

Может кто-нибудь мне помочь? Спасибо

Answer 1

Обычно принято задавать вопрос, что-то более конкретное, чем «Может кто-нибудь мне помочь?» Ваш код, как показано, не имеет операций с плавающей запятой (+, * и т. Д.), Поэтому для вычисления нет CGMA (оно равно нулю).

Что касается операций памяти, ваш код имеет 4 threadblocks:

dim3 dimGrid(4,1,1); 

Каждый threadblock может выполняться на отдельном многопроцессорных. В каждом блоке есть 10 потоков. Следующая строка кода:

  d_Out[n]=d_In[n]+1; 

будет генерировать по меньшей мере одну глобальную транзакцию нагрузки (d_In) и один глобальный магазин сделки (d_Out) для обслуживания нити. Четвертый блок будет иметь потоки, глобальные индексы (n) для активных потоков будут 30-35. Когда этот блок выполняет указанную выше строку кода, он будет генерировать две глобальные нагрузки и две транзакции глобального хранилища, потому что потоки требуют двух каналов для обслуживания своих запросов. Таким образом, эта одна строка кода может генерировать 5 глобальных транзакций нагрузки и 5 глобальных транзакций хранилища.

По тем же причинам, следующая строка кода:

  d_rows[n]=crows; 

может генерировать 5 дополнительных глобальных транзакций магазина. Так ваша продукция профилировщика:

1    gld_transactions  Global Load Transactions   6   6   6 
    1    gst_transactions  Global Store Transactions   11 

Я считаю, что я объяснил 5 из 6 глобальных транзакций нагрузки, и 10 из 11 глобальных магазинов сделок. Надеюсь, этого достаточно, чтобы дать вам представление о происхождении этих чисел.