我正在使用 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]);
}
有人可以帮助我吗?谢谢
最佳答案
通常会问一个比“有人能帮我吗?”更具体的问题显示的代码没有浮点运算(+、* 等),因此没有要计算的 CGMA(它为零)。
关于内存事务,您的代码有 4 个线程 block :
dim3 dimGrid(4,1,1);
每个线程 block 都可以在单独的多处理器上运行。每个 block 中有 10 个线程。下面一行代码:
d_Out[n]=d_In[n]+1;
将生成至少一个全局加载事务 (d_In
) 和一个全局存储事务 (d_Out
) 来为线程提供服务。第四个 block 将具有事件线程的全局索引 (n
) 为 30-35 的线程。当此 block 执行上述代码行时,它将生成两个 全局加载事务和两个 全局存储事务,因为线程需要两个缓存行来为其请求提供服务。所以这一行代码可能会生成 5 个全局加载事务和 5 个全局存储事务。
类似的原因,下一行代码:
d_rows[n]=crows;
可能会产生 5 个额外的全局商店交易。所以你的探查器输出:
1 gld_transactions Global Load Transactions 6 6 6
1 gst_transactions Global Store Transactions 11
我相信我已经解释了 6 个全局加载事务中的 5 个,以及 11 个全局存储事务中的 10 个。希望这足以让您了解这些数字的来源。
关于cuda - 使用 CUDA Profiler nvprof 进行内存访问,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/25945426/