CUDA学习笔记(十四) Constant Memory

转载至https://www.cnblogs.com/1024incn/tag/CUDA/

CONSTANT  MEMORY

constant Memory对于device来说只读但是对于host是可读可写。constant Memory和global Memory一样都位于DRAM,并且有一个独立的on-chip cache,比直接从constant Memory读取要快得多。每个SM上constant Memory cache大小限制为64KB。

constant Memory的获取方式不同于其它的GPU内存,对于constant Memory来说,最佳获取方式是warp中的32个thread获取constant Memory中的同一个地址。如果获取的地址不同的话,只能串行的服务这些获取请求了。

constant Memory使用__constant__限定符修饰变量。

constantMemory的生命周期伴随整个应用程序,并且可以被同一个grid中的thread和host中调用的API获取。因为constant Memory对device来说是可读的,所以只能在host初始化,使用下面的API:

cudaError_t cudaMemcpyToSymbol(const void *symbol, const void * src, size_t count, size_t offset, cudaMemcpyKind kind)

Implementing a 1D Stencil with Constant Memory

实现一个1维Stencil(数值分析领域的东,卷积神经网络处理图像的时候那个stencil),简单说就是计算一个多项式,系数放到constant Memory中,即y=f(x)这种东西,输入是九个点,如下:

{x − 4h, x − 3h, x − 2h, x − h, x, x + h, x + 2h, x + 3h, x + 4h}

在内存中的过程如下:

 

公式如下:

 

那么要放到constant Memory中的便是其中的c0、c1、c2 ……

因为每个thread使用九个点来计算一个点,所以可以使用shared memory来降低延迟。

__shared__ float smem[BDIM + 2 * RADIUS];

RADIUS定义了x两边点的个数,对于本例,RADIUS就是4。如下图所示,每个block需要RADIUS=4个halo(晕)左右边界:

 

#pragma unroll用来告诉编译器,自动展开循环。

 View Code

Comparing with the Read-only Cache

Kepler系列的GPU允许使用texture pipeline作为一个global Memory只读缓存。因为这是一个独立的使用单独带宽的只读缓存,所以对带宽限制的kernel性能有很大的提升。

Kepler的每个SM有48KB大小的只读缓存,一般来说,在读地址比较分散的情况下,这个只读缓存比L1表现要好,但是在读同一个地址的时候,一般不适用这个只读缓存,只读缓存的读取粒度为32比特。

有两种方式来使用只读缓存:

  • 使用__ldg限定
  • 指定特定global Memory称为只读缓存

下面代码片段对于第一种情况:

__global__ void kernel(float* output, float* input) {...output[idx] += __ldg(&input[idx]);...
}

下面代码对应第二种情况,使用__restrict__来指定该数据的要从只读缓存中获取:

void kernel(float* output, const float* __restrict__ input) {...output[idx] += input[idx];
}

一般使用__ldg是更好的选择。通过constant缓存存储的数据必须相对较小而且必须获取同一个地址以便获取最佳性能,相反,只读缓存则可以存放较大的数据,且不必地址一致。

下面的代码是之前stencil的翻版,使用过了只读缓存来存储系数,二者唯一的不同就是函数的声明:

 View Code

由于系数原本是存放在global Memory中的,然后读进缓存,所以在调用kernel之前,我们必须分配和初始化global Memory来存储系数,代码如下:

const float h_coef[] = {a0, a1, a2, a3, a4};
cudaMalloc((float**)&d_coef, (RADIUS + 1) * sizeof(float));
cudaMemcpy(d_coef, h_coef, (RADIUS + 1) * sizeof(float), cudaMemcpyHostToDevice);

下面是运行在TeslaK40上的结果,从中可知,使用只读缓存性能较差。

Tesla K40c array size: 16777216 (grid, block) 524288,32
3.4517ms stencil_1d(float*, float*)
3.6816ms stencil_1d_read_only(float*, float*, float const *)

总的来说,constant缓存和只读缓存对于device来说,都是只读的。二者都有大小限制,前者每个SM只能有64KB,后者则是48KB。对于读同一个地址,constant缓存表现好,只读缓存则对地址较分散的情况表现好。

The Warp Shuffle Instruction

之前我们有介绍shared Memory对于提高性能的好处,在CC3.0以上,支持了shuffle指令,允许thread直接读其他thread的寄存器值,只要两个thread在 同一个warp中,这种比通过shared Memory进行thread间的通讯效果更好,latency更低,同时也不消耗额外的内存资源来执行数据交换。

这里介绍warp中的一个概念lane,一个lane就是一个warp中的一个thread,每个lane在同一个warp中由lane索引唯一确定,因此其范围为[0,31]。在一个一维的block中,可以通过下面两个公式计算索引:

laneID = threadIdx.x % 32

warpID = threadIdx.x / 32

例如,在同一个block中的thread1和33拥有相同的lane索引1。

Variants of the Warp Shuffle Instruction

有两种设置shuffle的指令:一种针对整型变量,另一种针对浮点型变量。每种设置都包含四种shuffle指令变量。为了交换整型变量,使用过如下函数:

int __shfl(int var, int srcLane, int width=warpSize);

该函数的作用是将var的值返回给同一个warp中lane索引为srcLane的thread。可选参数width可以设置为2的n次幂,n属于[1,5]。

eg:如果shuffle指令如下:

int y = shfl(x, 3, 16);

则,thread0到thread15会获取thread3的数据x,thread16到thread31会从thread19获取数据x。

当传送到shfl的lane索引相同时,该指令会执行一次广播操作,如下所示:

 

另一种使用shuffle的形式如下:

int __shfl_up(int var, unsigned int delta, int width=warpSize)

该函数通过使用调用方的thread的lane索引减去delta来计算源thread的lane索引。这样源thread的相应数据就会返回给调用方,这样,warp中最开始delta个的thread不会改变,如下所示:

 

第三种shuffle指令形式如下:

int __shfl_down(int var, unsigned int delta, int width=warpSize)

该格式是相对__shfl_down来说的,具体形式如下图所示:

 

最后一种shuffle指令格式如下:

int __shfl_xor(int var, int laneMask, int width=warpSize)

这次不是加减操作,而是同laneMask做抑或操作,具体形式如下图所示:

 

所有这些提及的shuffle函数也都支持单精度浮点值,只需要将int换成float就行,除此外,和整型的使用方法完全一样。

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

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

相关文章

【C++笔记】C++继承

【C笔记】C继承 一、继承的概念二、继承的语法和权限三、父类和子类成员之间的关系3.1、子类赋值给父类(切片)3.2、同名成员 四、子类中的默认成员函数4.1、构造函数4.2、拷贝构造4.3、析构函数 五、C继承大坑之“菱形继承”5.1、什么是“菱形继承”5.2、解决方法 一、继承的概…

数据结构-- 并查集

0. 引入 并查集是来解决等价问题的数据结构。 离散数学中的二元关系。 等价关系需满足自反性、对称性、传递性。 a ∈ S , a R a a R b & b R a a R b ∩ b R c > a R c a \in S, aRa \\ aRb \& bRa \\ aRb \cap bRc >aRc a∈S,aRaaRb&bRaaRb∩bRc>a…

【Opencv】OpenCV使用CMake和MinGW的编译安装出错解决

编译时出现的错误: mingw32-make[1]: *** [modules/core/CMakeFiles/opencv_core.dir/all] Error 2 Makefile:161: recipe for target ‘all’ failed mingw32-make: *** [all] Error 2解决方法: 根据贴吧老哥的解答,发现是mingw版本有问题导…

【JAVA学习笔记】43 - 枚举类

项目代码 https://github.com/yinhai1114/Java_Learning_Code/tree/main/IDEA_Chapter11/src/com/yinhai/enum_ 〇、创建时自动填入版权 作者等信息 如何在每个文件创建的时候打入自己的信息以及版权呢 菜单栏-File-setting-Editor-File and Code Templaters -Includes-输入信…

自动化测试07Selenium01

目录 什么是自动化测试 Selenium介绍 Selenium是什么 Selenium特点 工作原理 SeleniumJava环境搭建 Selenium常用的API使用 定位元素findElement CSS选择语法 id选择器:#id 类选择 .class 标签选择器 标签名 后代选择器 父级选择器 自己选择器 xpath …

2.1 向量与线性方程组

一、行图像与列图像 线性代数的中心问题是求解线性方程组。线性的意思是这些方程的未知数是一次的,即每个未知数只会乘数字,而不会出现 x x x 与 y y y 相乘的项。下面是一个由两个未知数组成的方程组: 两个方程 两个未知数 { x − 2 y 1…

关于QTableWidget的it所占内存的释放问题

说起QTableWidget,我们一般都是用QTableWidget进行数据显示或者使用QTableWidget其他控件组合进行设置。 这里主要说说数据展示,我们在进行数据展示时,数据量分为两种有限和无限。不管是QTableWidget中插入的数据有限还是数据无限。在数据多时…

【从0到1设计一个网关】什么是网关?以及为什么需要自研网关?

文章目录 什么是网关?网关类型网关的优缺点目前的网关解决方案有哪些?为什么要自研Gateway网关?自研网关需要注意什么? 注: 这篇文章作为我的网关的第一篇文章,并不涉及任何代码,只是提及了网关…

LeetCode 155. 掷骰子等于目标和的方法数:动态规划

【LetMeFly】1155.掷骰子等于目标和的方法数:动态规划 力扣题目链接:https://leetcode.cn/problems/number-of-dice-rolls-with-target-sum/ 这里有 n 个一样的骰子,每个骰子上都有 k 个面,分别标号为 1 到 k 。 给定三个整数 …

如何在国际腾讯云服务器上设置IP节点的密码?

跟着云计算技术的发展,越来越多的企业和个人开始运用云服务器来存储和处理数据。腾讯云服务器是一款非常受欢迎的云服务器产品,它提供了高效稳定的计算能力。在运用腾讯云服务器的过程中,咱们可能需求设置IP节点的暗码,以便维护咱…

杂谈:DC对Verilog和SystemVerilog语言的支持

DC对Verilog和SystemVerilog语言的支持 设计语言用哪种?Design Compiler对二者的支持简单的fsm电路测试测试结果对比写在最后 设计语言用哪种? 直接抛出结论:先有电路,后为描述。设计端而言,没有语言的高低好坏&#…

Gin vs Beego: Golang的Web框架之争

前言 Golang作为一门高效且简洁的语言,已经在Web开发领域得到了广泛的应用。Gin和Beego是Golang中两个著名的Web框架,它们都提供了一系列强大的功能,帮助开发者构建高性能的Web应用。本文将对Gin和Beego进行全面的对比,帮助开发者…

SQL SERVER连接oracle数据库几种方法

--1 方式 --查询oracle数据库中的表 SELECT * FROM OPENDATASOURCE( MSDAORA, Data SourceGE160;User IDDAIMIN;PasswordDAIMIN )..DAIMIN.JOBS 举一反三:在查询分析器中输入: SELECT * FROM OPENDATASOURCE( MSDAORA, Data SourceORCL;User…

vite.config.js文件配置代理设置VITE_APP_BASE_API

.env.development文件 ENV development # base api VITE_APP_BASE_API /dev-api.env.production文件 ENV production # base api VITE_APP_BASE_API /apidefine: {process.env: {VITE_APP_BASE_API: https://xxx.com}},server: {hmr: true, // vue3 vite配置热更新不用手动…

妙手ERP功能更新丨Shopee全球产品支持使用定价模板修改价格、Ozon新增SKU模板 、Temu采集箱支持添加货源链接......

为了给卖家朋友带来更好的使用体验,更高效地运营跨境店铺,妙手ERP在近两周优化了以下多项功能。 01、平台授权模块 TikTok - 支持授权美国跨境店及ACCU店铺,TikTok平台各功能均可使用 02、产品模块优化 全平台 - 支持产品重量单位切换 Sho…

Java架构师内功数据库

目录 1 导学2 数据库基本概念2.1 数据库系统2.2 三级模式-两级映像2.3 数据库设计2.4 数据模型2.4.1 E-R模型2.4.2 关系模型2.5 关系代数3 规范化和并发控制3.1 函数依赖3.2 键与约束3.3 范式3.3.1 第一范式1NF3.3.2 第二范式3.3.3 第三范式3.4 模式分解3.5 并发控制3.6 封锁协…

centos搭建elastic集群

1、环境可以在同一台集群上搭建elastic,也可以在三台机器上搭建,这次演示的是在同一台机器搭建机器。 2、下载elastic :https://www.elastic.co/cn/downloads/past-releases#elasticsearch 2、​​​​​​ tar -zxvf elasticsearch-xxx-版…

KNN(K近邻)水仙花的分类(含答案)

题目 以下采用K-NN算法来解决水仙花的分类问题,每个样本有两个特征,第一个为水仙花的花萼长度,第二个为水仙花 的花萼宽度,具体数据见表, 1)设置k3, 采用欧式距离,分析分类精度为多少…

vue如何使用冻结对象提升代码效率及其原理解析

先给大家伙整个实际工作中一定会碰到的问题 如下vue dome ,它的代码非常简单功能也1非常简单,就是一个按钮,点击后会显示有多少条数据 来看看源码, html部分就是一个按钮绑定了一个loadData事件,然后在p标签内展示了这个myData这个数据的长度 <template><div id&quo…

详解如何快速查询Node.js版本信息(六种方法)

Node.js是一款基于Chrome V8引擎的快速、轻量级的JavaScript运行时。随着应用程序规模越来越庞大&#xff0c;Node.js版本的更新也日益频繁。这篇文章旨在帮助开发者们快速查询Node.js版本信息。 一、使用node命令查询版本信息 node命令提供了一个参数-v&#xff0c;可以在控…