CUDA SHARED MEMORY

在global Memory部分，数据对齐和连续是很重要的话题，当使用L1的时候，对齐问题可以忽略，但是非连续的获取内存依然会降低性能。依赖于算法本质，某些情况下，非连续访问是不可避免的。使用shared memory是另一种提高性能的方式。

GPU上的memory有两种：

· On-board memory

· On-chip memory

global memory就是一块很大的on-board memory，并且有很高的latency。而shared memory正好相反，是一块很小，低延迟的on-chip memory，比global memory拥有高得多的带宽。我们可以把他当做可编程的cache，其主要作用有：

· An intra-block thread communication channel 线程间交流通道

· A program-managed cache for global memory data可编程cache

· Scratch pad memory for transforming data to improve global memory access patterns

本文主要涉及两个例子作解释：reduction kernel，matrix transpose kernel。

shared memory（SMEM）是GPU的重要组成之一。物理上，每个SM包含一个当前正在执行的block中所有thread共享的低延迟的内存池。SMEM使得同一个block中的thread能够相互合作，重用on-chip数据，并且能够显著减少kernel需要的global memory带宽。由于APP可以直接显式的操作SMEM的内容，所以又被称为可编程缓存。

由于shared memory和L1要比L2和global memory更接近SM，shared memory的延迟比global memory低20到30倍，带宽大约高10倍。

当一个block开始执行时，GPU会分配其一定数量的shared memory，这个shared memory的地址空间会由block中的所有thread 共享。shared memory是划分给SM中驻留的所有block的，也是GPU的稀缺资源。所以，使用越多的shared memory，能够并行的active就越少。

关于Program-Managed Cache：在C语言编程里，循环（loop transformation）一般都使用cache来优化。在循环遍历的时候使用重新排列的迭代顺序可以很好利用cache局部性。在算法层面上，我们需要手动调节循环来达到令人满意的空间局部性，同时还要考虑cache size。cache对于程序员来说是透明的，编译器会处理所有的数据移动，我们没有能力控制cache的行为。shared memory则是一个可编程可操作的cache，程序员可以完全控制其行为。

Shared Memory Allocation

我们可以动态或者静态的分配shared Memory，其声明即可以在kernel内部也可以作为全局变量。

其标识符为：__shared__。

下面这句话静态的声明了一个2D的浮点型数组：

__shared__ float tile[size_y][size_x];

如果在kernel中声明的话，其作用域就是kernel内，否则是对所有kernel有效。如果shared Memory的大小在编译器未知的话，可以使用extern关键字修饰，例如下面声明一个未知大小的1D数组：

extern __shared__ int tile[];

由于其大小在编译器未知，我们需要在每个kernel调用时，动态的分配其shared memory，也就是最开始提及的第三个参数：

kernel<<<grid, block, isize * sizeof(int)>>>(...)

应该注意到，只有1D数组才能这样动态使用。

Memory Banks

为了获得高带宽，shared Memory被分成32（对应warp中的thread）个相等大小的内存块，他们可以被同时访问。不同的CC版本，shared memory以不同的模式映射到不同的块（稍后详解）。如果warp访问shared Memory，对于每个bank只访问不多于一个内存地址，那么只需要一次内存传输就可以了，否则需要多次传输，因此会降低内存带宽的使用。

cuda的shared momery相关推荐

详解PyTorch编译并调用自定义CUDA算子的三种方式
点击上方"3D视觉工坊",选择"星标" 干货第一时间送达在上一篇教程中,我们实现了一个自定义的CUDA算子add2,用来实现两个Tensor的相加.然后用Py ...
CUDA ---- GPU架构（Fermi、Kepler）
GPU架构 SM(Streaming Multiprocessors)是GPU架构中非常重要的部分,GPU硬件的并行性就是由SM决定的. 以Fermi架构为例,其包含以下主要组成部分: CUDA co ...
CUDA编程——GPU架构，由sp，sm，thread，block，grid，warp说起
目录 1.从硬件看 2.从软件看 3.对应关系 4.SIMT和SIMD 掌握部分硬件知识,有助于程序员编写更好的CUDA程序,提升CUDA程序性能,本文目的是理清sp,sm,thread,block, ...
6.CUDA编程手册中文版---附录AB
附录A 支持GPU设备列表更多精彩内容,请扫描下方二维码或者访问https://developer.nvidia.com/zh-cn/developer-program 来加入NVIDIA开发者计划 ...
通过设置PYTORCH_CUDA_ALLOC_CONF中的max_split_size_mb解决Pytorch的显存碎片化导致的CUDA:Out Of Memory问题
问题的出现最近在基友的带动下开始投身ai绘画的大潮,于是本地部署了stable diffusion web ui,利用手上的24G显存开始了愉快的跑高分辨率图片之旅.然而某天在用inpaint功能修 ...
CUDA10.1安装 +VS2015开发环境搭建
基本信息 operator system: win10 系统中已经安装了CUDA8,可以同时安装两个版本.再安装CUDA10和仅安装CUDA10一样.在使用的时候选择CUDA10即可. 官网下载软件 ...
【网络】UCX（Unified Communication X )|统一抽象通信接口
目录 UCX 的意义 UCX 通信接口简介支持的传输(协议) UCX社区 UCX 编程模型简介建立连接内存注册异步任务处理(重点) 使用UCX 编译debug版本构建RPM包构建DEB 包 ...
IOS音视频（四十六）离线在线语音识别方案
IOS音视频(四十六)离线在线语音识别方案 IOS音视频(四十六)离线在线语音识别方案方案一:Siri语音识别 Siri语音识别简介 Siri语音识别功能类介绍 Siri语音识别功能集成方案二:百 ...
【Python-GPU加速】基于Numba的GPU计算加速（一）基本
Numba是一个可以利用GPU/CPU和CUDA 对python函数进行动态编译,大幅提高执行速度的加速工具包. 利用修饰器@jit,@cuda.jit,@vectorize等对函数进行编译 JIT: ...

cuda的shared momery

CUDA SHARED MEMORY

Shared Memory Allocation

Memory Banks

cuda的shared momery相关推荐

最新文章

热门文章