cuda的shared momery

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只访问不多于一个内存地址,那么只需要一次内存传输就可以了,否则需要多次传输,因此会降低内存带宽的使用。

 

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.mzph.cn/news/510132.shtml

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

图的表示方法和C++实现

图的表示最长用的两种方法是&#xff1a; 1&#xff09;、邻接矩阵表示法 2&#xff09;、邻接表表示 下面是两种构造图的方法 1&#xff09;邻接矩阵&#xff1a; [cpp] view plaincopy#include <iostream> #include <vector> using namespace std; //枚…

成为专业程序员的 6 个技巧

1.在你责怪别人之前&#xff0c;先检查自己的代码 先想一想自己的假设和其他人的假设。来自不同供应商的工具可能内置不同的假设&#xff0c;即便是相同的供应商对于不同的工具&#xff0c;其假设也可能不同。 当其他人正在报告一个你不能重复的问题的时候&#xff0c;去看看他…

HUE集成Hbase

目录 一、Hbase开启代理用户相关配置 二、代理用户授权认证 三、检查HUE在hue.ini文件中指定的HBASE的本地配置目录 一、Hbase开启代理用户相关配置 Cloudera Manager修改Hbase配置或Hbase配置文件-hbase-site.xml <property><name>hbase.thrift.support.proxyu…

OpenCV中cornerSubPixel()亚像素求精原理

采用的方法为最小二乘法&#xff1a; 首先我们要构建以下方程&#xff1a; 我们讨论角点的情况&#xff1a; q是我们要求的角点 p0和p1为q周围的点 &#xff08;q-pi&#xff09;为一个向量 Gi为pi处的梯度 所以满足一下公式 Gi*(q-pi)0 有以下两种情况&#xff1a; &a…

HBase 2.0 之修复工具 HBCK2 运维指南

HBase 2.0 之修复工具 HBCK2 运维指南 转载自&#xff1a;https://mp.weixin.qq.com/s/GVMWwB1WsKcdvZGfvX1lcA?spma2c4e.11153940.blogcont683107.11.49d762a815MegW 概述 目前社区已经发布了 HBase 的 2.0 版本&#xff0c;很多公司都希望去尝试新版本上的新功能&#xff0c…

html中article、section、aside的区别与联系

首先看看我做的图(PS:有点丑)&#xff0c;通俗易懂

优秀的程序员都避开了哪些坑?

程序员薪水有高有低&#xff0c;有的人一个月可能拿30K、50K&#xff0c;有的人可能只有2K、3K。同样有五年工作经验的程序员&#xff0c;可能一个人每月拿20K&#xff0c;一个拿5K。是什么因素导致了这种差异&#xff1f;我特意总结了容易导致薪水低的九大行为表现&#xff0c…

Springboot 集成 Swagger

1、问题描述 随着互联网技术的发展&#xff0c;现在的网站架构基本都由原来的后端渲染&#xff0c;变成了&#xff1a;前端渲染、先后端分离的形态&#xff0c;而且前端技术和后端技术在各自的道路上越走越远。 前端和后端的唯一联系&#xff0c;变成了API接口&#xff1b;API文…

email类型

<!DOCTYPE html> <html><head><meta charset"utf-8" /><title></title></head><body><form action"demo_form.php"method"get">请输入您的email地址: <input type"email" n…

优秀程序员的 18 大法则

经过多年的积累&#xff0c;我发现&#xff0c;下面这些基本的指导法则&#xff0c;可以帮助我成为一个更加高效的程序员。 程序设计法则&#xff0c;与设计和工程的原理密切相关。下面这些编程法则帮助我让我获益匪浅&#xff0c;所以我想分享给大家&#xff0c;希望也能帮助大…

url类型

<!DOCTYPE html> <html><head><meta charset"utf-8" /><title></title></head><body><form action"demo_form.php"method"get">请输入网址: <input type"url" name"us…

低效程序员的7个坏习惯

程序员总是想做到尽可能的高效&#xff0c;但很多人往往会觉得力不从心。这是因为他们在多年的编码过程中养成了一些不好的习惯。下面这7个坏习惯绝对是软件工程师需要改掉的。 1.缺乏激情 这已经是一个老生常谈的话题了&#xff0c;但却是真理。写了多年的代码后&#xff0c;程…

Storm消费Kafka异常 - topic其中两个分区达到某个值不进行消费,持续阻塞

Kafka消费storm&#xff0c;突然有两个分区无法消费数据(或重复消费无法提交offset) offset是我们自己进行管理&#xff0c;kafka日志也是正常没有报错&#xff0c;storm日志也是没有报错~ 就是卡住了 1.尝试将partition为0,1的offset记录删除&#xff0c;重新跑一遍&#xff…

为什么跳槽加薪会比内部调薪要高?

有网友在知乎提问&#xff1a; 最近在思考一个问题&#xff0c;为什么跳槽往往意味着加薪&#xff1f; 如果一个人确有价值&#xff0c;为什么在原来的公司没有在薪水上体现出来&#xff1f;如果没有价值&#xff0c;为什么跳槽以后就会加薪&#xff1f;还是可以单纯的解释为&a…

浏览器多代理配置 - SwitchyOmega

转自 https://www.switchyomega.com/settings/ 下载链接&#xff1a;https://proxy-switchyomega.com/download/ 情景模式 代理服务器 代理服务器可以支持 HTTP、HTTPS、SOCKS4、SOCKS5 代理协议。SOCKS 代理协议不支持验证。下图以配置 Shadowsocks 的 SOCKS5 代理协议为例。…

number类型

step&#xff1a;数字间隔 <!DOCTYPE html> <html><head><meta charset"utf-8" /><title></title></head><body><form action"demo_form.php"method"get">请输入数值: <input type&qu…

居然还能这样——程序员加薪的新方法

我的朋友A君是个典型的.NET开发人员&#xff0c;技术不错&#xff0c;人品也不错&#xff0c;在一家小公司&#xff08;姑且称为甲公司&#xff09;做项目开发&#xff0c;是技术骨干。 3个月前&#xff0c;他找到我说想跳槽&#xff0c;让我帮忙介绍工作。我说为什么想跳了&am…

range类型

输入包含一定范围内的数字 <!DOCTYPE html> <html><head><meta charset"utf-8" /><title></title></head><body><form action"demo_form.php"method"get">请输入数值: <input type&qu…

程序员真的很穷吗?

前几天一位做市场的同事跑过来问&#xff0c;池老师&#xff0c;我有一位朋友&#xff0c;快30了&#xff0c;想转行写程序&#xff0c;您觉得有戏吗&#xff1f;我看了看满目疮痍的他说&#xff0c;如果是你就没戏。 30多岁转行做程序员当然可行&#xff0c;毕竟历史上存在一些…

Hive 行转列,列传行 - Impala 暂不支持

注&#xff1a;Impala 不支持 lateral view explode 一、行转列 (对某列拆分&#xff0c;一列拆多行) 使用函数&#xff1a;lateral view explode(split(column, ,)) num eg: 如表&#xff1a;t_row_to_column_tmp 数据如下&#xff0c;对tag列进行拆分 SQL代码&#xff1a…