Shared Memory,Warp和Bank Conflict

1. 概念

        Shared memory是片上存储器,因此与local memory或global memory相比更高的带宽和更低的延迟。前提是线程之间没有 bank conflicts

        为了实现高带宽,共享内存被划分为大小相等的内存模块,称为Banks,可以同时访问。因此,任何由 n 个地址组成的内存读取或写入请求都可以同时提供服务,从而产生比单个模块带宽高 n 倍的总带宽。

        但是,如果warp内多个线程的内存请求的两个地址位于同一Bank中,则存在bank conflict,并且必须序列化访问。硬件将具有bank conflict的内存请求拆分为根据需要尽可能多的单独的无冲突请求,从而将吞吐量降低一个系数,该系数等于单独的内存请求数。如果单独内存请求的数量为 n,则初始内存请求称为导致 n 路存储体冲突。此处的一个例外是,当 warp 中的多个线程对同一共享内存位置进行寻址时,从而导致广播。在这种情况下,来自不同存储区的多个广播被合并到从请求的共享内存位置到线程的单个多播中。(注意区分:同一Bank(冲突) 和 同一位置(广播))

        因此,为了获得最佳性能,了解内存地址如何映射到内存存储体非常重要,以便安排内存请求,从而最大限度地减少存储体冲突。

        在计算能力为 5.x 或更高版本的设备上,每个Bank的每个时钟周期的带宽为 32 位,并且连续的 32 位字被分配给连续的Bank。Warp大小为 32 个线程,Bank的数量也为 32 个,因此Warp中的任何线程之间都可能发生Bank冲突。

2. NVIDIA Tesla V100的Bank组织形式

        计算能力7.5,Maximum amount of shared memory per SM为64KB。Bank数量为32,Bank宽度为32-bit(4B)。故每个Bank可保存:

\frac{64\times 1024}{32\times 4}=512

个整型或单精度浮点型数。或者说:Bank组织成了512行32列的矩阵。

3. Bank Conflict示例

        图 21 显示了共享内存跨步访问(strided access)的一些示例。适用于计算能力为 3.x(在 32 位模式下)或计算能力为 5.x 、 6.x、7.x 的设备。

 左:线性寻址,步幅为一个 32 -bit(无bank conflict)。

中间:线性寻址,步幅为两个 32 -bit(双向bank conflict)。

右:线性寻址,步幅为三个 32 -bit(无bank conflict)。

        图 22 显示了涉及广播机制的一些内存读取访问示例。计算能力为 3.x、5.x 、 6.x、7.x 的设备的示例。

 左:通过随机排列实现无冲突访问。

中间:无冲突访问,因为线程 3、4、6、7 和 9 访问Bank5 中的同一位置。属于广播

右:无冲突广播访问(线程访问Bank内的相同位置)。

(更多可参考:https://blog.csdn.net/endlch/article/details/47043069

  • 0
    点赞
  • 3
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
CUDA cooperative group中,可以使用以下两种方式获取warp id和lane id: 1. 使用`thread_group_tile`的`thread_rank()`和`thread_lane_id()`方法获取当前线程在warp中的id和lane id。例如: ```cuda #include <cooperative_groups.h> using namespace cooperative_groups; __global__ void kernel() { thread_block_tile<32> tile = tiled_partition<32>(this_thread_block()); int warp_id = tile.meta_group_id(); int lane_id = tile.thread_rank(); // ... } ``` 在上述代码中,`tiled_partition<32>(this_thread_block())`表示将当前线程所在的block分成32个线程一组的warp,然后使用`tile.meta_group_id()`获取当前线程所在的warp id,使用`tile.thread_rank()`获取当前线程在warp中的lane id。 2. 使用`thread_group`的`this_thread()`方法获取当前线程所在的warp,并使用`thread_group`的`thread_rank()`和`thread_lane_id()`方法获取当前线程在warp中的id和lane id。例如: ```cuda #include <cooperative_groups.h> using namespace cooperative_groups; __global__ void kernel() { thread_group g = this_thread_block(); int warp_id = g.thread_rank() / 32; int lane_id = g.thread_rank() % 32; // ... } ``` 在上述代码中,`this_thread_block()`表示获取当前线程所在的block,然后使用`g.thread_rank() / 32`获取当前线程所在的warp id,使用`g.thread_rank() % 32`获取当前线程在warp中的lane id。 需要注意的是,以上两种方式都是在CUDA cooperative group中使用的,如果在普通的kernel函数中使用可能会导致错误。
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值