Nsight Compute内存访问常用Metrics含义理解

比眉伴天荒 2022-12-24 06:57 183阅读 0赞

Nsight Compute 软件Source模块提供了精确到源代码行号的metrics参数，用于辅助性能调优，本篇基于访问共享内存的矩阵转置核函数的实现，记录一下对常用metrics含义的理解。

### Metrics含义 ###

**Memory L1 Transcations Global**:实际全局内存加载至L1缓存的内存交换次数，粒度128bytes  
**Memory L2 Transactions Global**:实际全局内存加载至L2缓存的内存交换次数，粒度32bytes,该参数的值应该是Memory L1 Transcations Global 的4倍  
**Memory Ideal L2 Transactions Global**:理论需要从全局内存加载至L2缓存的内存交换次数，当数值比Memory L2 Transactions Global小时，说明当前全局内存访问模式有效率问题  
**Memory L1 Transactions Shared**:L1缓存与共享内存的数据交换次数，粒度32bytes  
**Memory Ideal L1 Transactions Shared**:理论需要的L1缓存与共享内存的数据交换次数，当数值比Memory L1 Transactions Shared小时，说明存在Bank Conflict

### 代码实现 ###

核函数执行配置

dim3 block(32, 32, 1);
    dim3 grid(32, 32, 1);

无Bank Conflict的核函数，将共享内存设置为32 \* 33，避免从共享内存写入全局内存时发生冲突，实现如下：

__global__ void kSMMatrixT(float* d_src, float* d_dst, int ROWDIM, int COLDIM)
    { 
    	__shared__ float smTmp[SMDIM][SMDIM+1];
    	int srcxIdx = threadIdx.x + blockIdx.x * blockDim.x;
    	int srcyIdx = threadIdx.y + blockIdx.y * blockDim.y;
    	smTmp[threadIdx.y][threadIdx.x] = d_src[srcyIdx * COLDIM + srcxIdx];
    	__syncthreads();
    	int dstxIdx = blockIdx.y * blockDim.y + threadIdx.x;
    	int dstyIdx = blockIdx.x * blockDim.x + threadIdx.y;
    	d_dst[dstyIdx*ROWDIM + dstxIdx] = smTmp[threadIdx.x][threadIdx.y];
    
    }

存在Bank Conflict的核函数，共享内存设置为32 \*32 ，这种情况下从共享内存加载数据写入全局内存时，一个warp每个线程访问相同存储体的不同位置，发生冲突。代码实现如下：

__global__ void kBankCMatrixT(float* d_src, float* d_dst, int ROWDIM, int COLDIM)
    { 
    	__shared__ float smTmp[SMDIM][SMDIM];
    	int srcxIdx = threadIdx.x + blockIdx.x * blockDim.x;
    	int srcyIdx = threadIdx.y + blockIdx.y * blockDim.y;
    	smTmp[threadIdx.y][threadIdx.x] = d_src[srcyIdx * COLDIM + srcxIdx];
    	__syncthreads();
    	int dstxIdx = blockIdx.y * blockDim.y + threadIdx.x;
    	int dstyIdx = blockIdx.x * blockDim.x + threadIdx.y;
    	d_dst[dstyIdx * ROWDIM + dstxIdx] = smTmp[threadIdx.x][threadIdx.y];
    
    }

### 具体Metrics ###

BankConflict  
![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3lhbjMxNDE1_size_16_color_FFFFFF_t_70]  
无BankConflict  
![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3lhbjMxNDE1_size_16_color_FFFFFF_t_70 1]  
对比两个版本，BankConflict版本的Memory L1 Transactions Shared发生1048576次，Memory Ideal L1 Transactions Shared发生了32768，实际共享内存数据交换数是理论值的32倍，刚好符合代码的设计（共享内存写回内存时，一个warp的线程会访问共享内存同一个bank的32个位置）。而无Bank Conflict版本理论值与实际值一直，验证了没有Bank Conflict发生。

[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3lhbjMxNDE1_size_16_color_FFFFFF_t_70]: /images/20221120/cef62acdd2b243b48ccf50faebd34886.png
[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L3lhbjMxNDE1_size_16_color_FFFFFF_t_70 1]: /images/20221120/504026c32b0445fd89f76e15d650dd92.png