cuda - 全局内存加载/存储效率和全局内存合并

Question

我有以下简单的代码:

#include<stdio.h>

#define BLOCKSIZE_X 32
#define BLOCKSIZE_Y 1

int iDivUp(int a, int b) { return ((a % b) != 0) ? (a / b + 1) : (a / b); }

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
    if (code != cudaSuccess) 
    {
        fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}

__global__ void kernel0(float *d_a, float *d_b, const unsigned int M, const unsigned int N)
{
    const int tidx = threadIdx.x + blockIdx.x * blockDim.x;
    const int tidy = threadIdx.y + blockIdx.y * blockDim.y;

    if ((tidx < M)&&(tidy < N)) {

        d_b[tidy * M + tidx] = d_a[tidy * M + tidx];

    }

}

void main()
{
    const unsigned int M = 32;
    const unsigned int N = 1;

    float *d_a; cudaMalloc((void**)&d_a, M*N*sizeof(float));
    float *d_b; cudaMalloc((void**)&d_b, M*N*sizeof(float));

    dim3 dimGrid(iDivUp(M, BLOCKSIZE_X), iDivUp(N, BLOCKSIZE_Y));
    dim3 dimBlock(BLOCKSIZE_X, BLOCKSIZE_Y);

    kernel0<<<dimGrid, dimBlock>>>(d_a, d_b, M, N);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    cudaDeviceReset();

}

执行 32 float 的两个数组之间的赋值s。我试图了解全局内存合并访问和全局内存加载/存储效率以及其他指标/事件之间的关系。

Visual Profiler 显示以下指标:

Global Memory Load Efficiency = 50%
Global Memory Store Efficiency = 100%

Global Memory Load Efficiency 的值(value)让我感到惊讶。我本来以为 100%两种情况下的效率，因为我相信我正在执行完美合并的内存访问。所以我的问题是:

为什么在执行合并内存访问时，全局内存加载效率为 50%，而全局内存存储效率为 100%？

我还调查了其他可能对报告有用的指标/事件:

gld_inst_32bit = 32 (Number of 32-bit global memory load transactions)
gst_inst_32bit = 32 (Number of 32-bit global memory store transactions)

事实上，我正在请求加载/写入 32 float s。

uncached global load transaction = 0 (Number of uncached global load transactions)
l1 global load miss = 2 (Number of global load misses in L1 cache)

根据我的(可能是错误的)理解，上述两个事件似乎是矛盾的。在 l1 的情况下缓存未命中，我原以为第一个事件与 0 不同.

gld_request = 1 (Number of executed global load instructions per warp in a SM)
gst_request = 1 (Number of executed global store instructions per warp in a SM)

这似乎与我正在执行完美合并的内存访问的事实一致。

反汇编后的代码如下:

/*0000*/         MOV R1, c[0x1][0x100];                          /* 0x2800440400005de4 */
/*0008*/         S2R R3, SR_CTAID.Y;                             /* 0x2c0000009800dc04 */
/*0010*/         S2R R4, SR_TID.Y;                               /* 0x2c00000088011c04 */
/*0018*/         IMAD R4, R3, c[0x0][0xc], R4;                   /* 0x2008400030311ca3 */
/*0020*/         S2R R0, SR_CTAID.X;                             /* 0x2c00000094001c04 */
/*0028*/         ISETP.LT.U32.AND P0, PT, R4, c[0x0][0x2c], PT;  /* 0x188e4000b041dc03 */
/*0030*/         S2R R2, SR_TID.X;                               /* 0x2c00000084009c04 */
/*0038*/         IMAD R0, R0, c[0x0][0x8], R2;                   /* 0x2004400020001ca3 */
/*0040*/         ISETP.LT.U32.AND P0, PT, R0, c[0x0][0x28], P0;  /* 0x18804000a001dc03 */
/*0048*/    @!P0 BRA.U 0x78;                                     /* 0x40000000a000a1e7 */
/*0050*/     @P0 IMAD R2, R4, c[0x0][0x28], R0;                  /* 0x20004000a04080a3 */
/*0058*/     @P0 ISCADD R0, R2, c[0x0][0x20], 0x2;               /* 0x4000400080200043 */
/*0060*/     @P0 ISCADD R2, R2, c[0x0][0x24], 0x2;               /* 0x4000400090208043 */
/*0068*/     @P0 LD R0, [R0];                                    /* 0x8000000000000085 */
/*0070*/     @P0 ST [R2], R0;                                    /* 0x9000000000200085 */
/*0078*/         EXIT;                                           /* 0x8000000000001de7 */

编辑

我的配置:CUDA 6.5、GeForce GT540M、Windows 7。

如果我增加 M来自 32至 64启动两个块并使我的卡的两个可用流式多处理器忙碌，然后全局内存负载效率变为 100%这些是新的指标/事件:

gld_inst_32bit = 64 
gst_inst_32bit = 64 

uncached global load transaction = 0 
l1 global load miss = 2 

gld_request = 2
gst_request = 2

gld_inst_32bit的增加, gst_inst_32bit , gld_request和 gst_request是预期的和一致的，因为现在我正在加载 7 存储 64 float s 和 2全局内存加载/存储合并请求。但我还是不明白 uncached global load transaction和 l1 global load miss可以保持不变，同时全局内存负载吞吐量变化提供 100%效率。

编辑

开普勒 K20c 的结果 M=32 :

Global Memory Load Efficiency = 100%
Global Memory Store Efficiency = 100%
gld_inst_32bit = 64 
gst_inst_32bit = 64 
gld_request = 1
gst_request = 1
uncached global load transaction = 1
l1 global load miss = 0
l1 global load hit = 0

现在，Visual Profiler 报告了一个未缓存的全局加载事务，但没有 l1全局负载未命中。

编辑

我对这个问题进行了更多调查，增加了 M 的值。并保留 BLOCKSIZE_X固定的。

当块数为奇数时，即我的GT540M卡的两个Streaming Multiprocessor负载不平衡，则全局内存负载效率小于 100% , 否则为 100%在偶数情况下。全局内存加载效率缓慢趋于 100%只要在奇数情况下增加块数。

如果我禁用 L1使用 -Xptxas -dlcm=cg 编译缓存，正如@Jez 所建议的那样，全局内​​存负载效率始终等于 100% ，因为它是全局内存存储效率。我知道全局内存存储不使用 L1缓存，但只有 L2。

一些图片显示，对于 M 的不同值，全局内存加载效率的行为
M=32



M=64



M=96



M=128

M=160

M=192


请注意 M是 32 的整数倍允许通过单个扭曲加载整个缓存行。

通过禁用 L1 ， 我有:
M=32

M=64

M=96


编辑 - 特斯拉 C2050 的结果

M = 32    33.3%
M = 64    28.6%
M = 96    42.9%
M = 128   57.1%
M = 160   71.4%
M = 192   85.7%
M = 224  100%
M = 256  114%
M = 288   90%

同样，如果我禁用 L1缓存，我有 100%所有情况下的全局内存负载效率。

Answer 1

计数器的准确性

NVIDIA 分析器可以收集原始计数器和指标。许多指标要求内核多次执行。理想情况下，分析器能够在一次通过中收集指标的所有原始计数器，但鉴于性能监控系统的局限性，这根本不可能。

在 Fermi 架构上，用于收集全局内存负载效率和全局内存存储效率的 L1 统计数据的 HWPM 系统只能观察到每个 GPC 的 1 个 L1 单元。对于 GF100 (C2050)，这相当于 25% 的观察。

如果工作负载没有完全填满机器并且每次通过的单位工作量相同，则分析器将无法提供准确的结果。

在 Kepler 架构上，HWPM 系统可以从每个 L1 收集 L1 统计信息，但仍然对 L2 有一些限制，这可能会导致小的差异。

在 Maxwell 架构上，内存系统明显不同，因为全局、本地和表面请求现在都通过统一的 L1/TEX 缓存。

缓存 VS。未缓存

在 Fermi 架构中，所有全局加载/存储都通过 L1 缓存。未缓存的全局加载/存储仍然通过 L1，使用 LSU 事务，并需要进行标记查找以使缓存行无效。原子是通过 L1 的唯一一种不会使 L1 缓存失效的全局访问形式。

Kepler 架构有一些小的变化。默认情况下，大多数芯片不缓存全局内存访问，因此所有全局加载都未缓存。在 GK110 和 GK208 芯片上，可以使用新的 LDG 指令通过 TEX 缓存加载全局数据。

从 SM 到 L1 的缓存和未缓存全局加载事务为 128 字节。

从 L1 到 L2 的缓存全局加载事务作为 4 个 32B 请求完成。

从 L1 到 L2 的未缓存全局加载事务以 32B 请求的最小数量完成。