share memory的bank conflict分析

背景

在做高性能分析的时候，经常会出现一个什么bank conflict的名词，不仅是GPU的share memory会出现bank confict, 甚至连寄存器也会出现bank conflict, 那么这个是什么东西，下面进行一个系统的梳理。

硬件

从事计算机行业的同学肯定见过这个东西，没错这个玩意就是内存条，上面的小黑块就是内存颗粒，每一片称之为chip. 有兴趣的可以拆开自己的电脑看一下，这个chip大约有1mm左右的厚度，我们知道现在的半导体工艺一般都是纳米级别，所以这1mm后的内存，其实是由很多薄片累加一起的，每一个薄片叫做一个bank.图示如下：

接下来，每一个矩形bank可以继续拆成一个个二维的小方格，每次从一个bank取数和存数的话只有通过一根线（现在很多内存开了两个端口，分别做存数和取数），如果两个数据分布在不同的bank上（bank序列挨得近），那么一把就可以取出两个数据，如果两个数据在一个bank上，那么耗时就是2倍，因为只能串行拿出，也就是我们说的bank conflict，具体示意图如下，好了，硬件理解到这个程度已经可以了，下面来分析分析cuda中的share memory。

share_memory

share memory相比较device memory而言，share memory是在GPU芯片上的，而device memory是通过总线链接到GPU芯片上的，所以性能的高低一目了然，为此我们尽量要使用好share memory。share memory 有32个bank, 可以自己脑补上面图示的bank, 每一个小格子里面放4字节(也可以自己设置为8)大小的数据(int ,float),）被划分到32个bank中，每个bank的内存能同时读写(有两个port)，但是同一个bank的不同地址的数据则只能串行读写, 如果不容线程访问一个bank的一个地址也没问题，根据上面图示可以看到可以把一个数据放到不同的buffer里。例如__shared__ float data[32][32]，申请了1024个float数据，每个float正好是4字节，data按行存储，data[0][0]就位于第0个bank，data[0][1]位于第一个bank，以此类推.因此data[0][col]就被划分在了第col个bank中，即列数相同（data[0][1]和data[5][1]）的数据划分至了同一个bank中(自己脑补上面bank位置).如果一个warp的线程按列处理data那么就会造成bank conflict.

测试

#define WARPSIZE 32
//conflict
__global__ void kernel1(float* A) {__shared__ float data[32][32];int tid = threadIdx.x;int col = tid/WARPSIZE;int row = tid%WARPSIZE;//为了测试同一个warp的memory bank conflict, 这里连续的线程要处理不同行，但是同一列的数据。data[row][col] = 1.f;A[tid] = data[row][col];
}
//noconflict
__global__ void kernel2(float* A) {__shared__ float data[32][32];int tid = threadIdx.x;int row = tid/WARPSIZE;int col = tid%WARPSIZE;data[row][col] = 1.f;A[tid] = data[row][col];
}

整体测试代码：

#include<stdio.h>
#include<time.h>
#define WARPSIZE 32
__global__ void kernel1(float* A) {__shared__ float data[32][32];int tid = threadIdx.x;int col = tid/WARPSIZE;int row = tid%WARPSIZE;data[row][col] = 100.f;A[tid] = data[row][col];
}__global__ void kernel2(float* A) {__shared__ float data[32][32];int tid = threadIdx.x;int row = tid/WARPSIZE;int col = tid%WARPSIZE;data[row][col] = 100.f;A[tid] = data[row][col];
}int main() {int blocksize = 32*32;float* h_A = (float*)malloc(sizeof(float)*blocksize);float* d_A;cudaMalloc(&d_A, sizeof(float)*blocksize);kernel1<<<1, blocksize>>>(d_A);cudaDeviceSynchronize();cudaMemcpy(h_A, d_A, blocksize*sizeof(float), cudaMemcpyDeviceToHost);kernel2<<<1, blocksize>>>(d_A);cudaDeviceSynchronize();cudaMemcpy(h_A, d_A, blocksize*sizeof(float), cudaMemcpyDeviceToHost);    cudaFree(d_A);free(h_A);return 0;
}

编译：

nvcc -O3 bankconflict.cu -o bankconflict

分析

利用下面命令分析耗时：

nvprof ./bankconflict

利用下面代码分析冲突事件：

nvprof --events shared_ld_bank_conflict,shared_st_bank_conflict ./bankconflict

可以得到下面的结果：

可以看到存数据(shared_st_bank_conflict=store bank conflict)的时候，会有992=31*32个冲突，因为每一个warp中的线程第一个线程是没有冲突的，所以是31*32，和预期的一模一样。不过这里有一个问题，kernel1最后执行了A[tid] = data[row][col]，按道理来说应该也存在load bank conflict.但是为什么使用nvprof显示的结果却没有呢？
原因是我们编译的时候使用了-O3编译优化，编译器优化了我们的程序，减少了bank conflict的次数.可以通过禁止编译优化来观察结果，重新编译. //-g是debugcpu,-G是debug GPU并且关闭所有优化
nvcc -g -G bankconflict.cu -o bankconflict
测试：
nvprof --events shared_ld_bank_conflict,shared_st_bank_conflict ./bankconflict

参考：https://www.cnblogs.com/deepllz/p/11490544.html

share memory的bank conflict分析相关推荐

Share memory中bank conflict问题
Share memory是片上资源,生命周期是整个block中,它的数据读写十分快,有1个cycle latency.在Share memory中,经常存在bank conflict问题,如果没有ba ...
bank conflict
存储体冲突(bank conflict):当被访问的存储体没有恢复时又出现对该存储体新访问的现象. 简介目前 CUDA 装置中,每个 multiprocessor 有 16KB 的 shared me ...
CUDA bank 及bank conflict
bank 是CUDA中一个重要概念,是内存的访问时一种划分方式,在CPU中,访问某个地址的内存时,为了减少读写内次次数,访问地址并不是随机的,而是一次性访问bank内的内存地址,类似于内存对齐一样,一 ...
难理解的bank conflict
之前看Nvidia-OpenCL-SDK里有一个例子讲到过bank conflict,但没怎么明白,它选择的是用奇数来避免. #define BLOCK_DIM 16// This kernel is ...
CUDA:矩阵乘法的实现(Share Memory)
本文参加2022CUDA on Platform线上训练营学习笔记矩阵乘法的GPU端实现一.矩阵乘法(Matrix Multiply)基础二.矩阵乘法的CPU端实现三.矩阵乘法的GPU端实现( ...
使用Memory Analyzer tool(MAT)分析内存泄漏（一）
使用Memory Analyzer tool(MAT)分析内存泄漏(一) (2010年05月21日) 发表于 Java博客前言的前言 :本文是自 2005 年 8 月以来,首次在一个月之内发布三篇文 ...
【Android 内存优化】使用 Memory Analyzer ( MAT ) 工具分析内存 ( MAT 工具使用 | 最大对象 | 类实例个数 | 引用与被引用 | GC Roots 最短链 )
文章目录一. 内存中最大的对象二. 查看每个类的对象实例的个数三. 查看对象的引用与被引用四. 查看对象到 GC Roots 的最短距离 1. 选择 Merge Shortest Paths ...
CA/TA通信的share memory设计思想解读
引流关键词: WSM, 共享内存,sharememory,share memory, optee.ATF.TF-A.Trustzone.optee3.14.MMU.VMSA.cache.TLB.arm ...
使用Memory Analyzer tool(MAT)分析内存泄漏
http://www.blogjava.net/rosen/archive/2010/05/21/321575.html http://www.blogjava.net/rosen/archive/2 ...

share memory的bank conflict分析

背景

硬件

share_memory

测试

分析

share memory的bank conflict分析相关推荐

最新文章

热门文章