合并内存访问时的全局加载事务计数答案

【问题标题】：Global load transaction count when in coalesced memory access合并内存访问时的全局加载事务计数
【发布时间】：2016-10-16 02:57:52
【问题描述】：

我创建了一个简单的内核，通过观察 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);

由于内核中的数据访问是完全合并的，并且由于数据类型是浮点数（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%

我在这里错过了什么？

【问题讨论】：

你是如何分配d_in和d_out的？
cudaMalloc(&d_in, sizeof(float)*data_size); 和 cudaMalloc(&d_out, sizeof(float)*data_size);

标签： cuda nvprof

【解决方案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

但是，在使用 NSIGHT Visual Studio 版本的 Windows 上，我得到的值似乎更好：

您可能需要联系 NVIDIA，因为这可能只是 nvprof 中的显示问题。

【讨论】：