cuda的global memory介绍

CUDA Memory Model

对于程序员来说,memory可以分为下面两类:

  • Programmable:我们可以灵活操作的部分。
  • Non-programmable:不能操作,由一套自动机制来达到很好的性能。

在CPU的存储结构中,L1和L2 cache都是non-programmable的。对于CUDA来说,programmable的类型很丰富:

  • Registers
  • Shared memory
  • Local memory
  • Constant memory
  • Texture memory
  • Global memory

下图展示了memory的结构,他们各自都有不用的空间、生命期和cache。

Global Memory

global Memory是空间最大,latency最高,GPU最基础的memory。“global”指明了其生命周期。任意SM都可以在整个程序的生命期中获取其状态。global中的变量既可以是静态也可以是动态声明。可以使用__device__修饰符来限定其属性。global memory的分配就是之前频繁使用的cudaMalloc,释放使用cudaFree。global memory驻留在devicememory,可以通过32-byte、64-byte或者128-byte三种格式传输。这些memory transaction必须是对齐的,也就是说首地址必须是32、64或者128的倍数。优化memory transaction对于性能提升至关重要。当warp执行memory load/store时,需要的transaction数量依赖于下面两个因素:

  1. Distribution of memory address across the thread of that warp 就是前文的连续
  2. Alignment of memory address per transaction 对齐

一般来说,所需求的transaction越多,潜在的不必要数据传输就越多,从而导致throughput efficiency降低。

对于一个既定的warp memory请求,transaction的数量和throughput efficiency是由CC版本决定的。对于CC1.0和1.1来说,对于global memory的获取是非常严格的。而1.1以上,由于cache的存在,获取要轻松的多。

下面代码是通过可以通过32-byte、64-byte或者128-byte三种格式传输,优化传输效率:

template <typename T>
__device__ inline uint32_t pack_uint8x4(T x, T y, T z, T w){uchar4 uint8x4;uint8x4.x = static_cast<uint8_t>(x);uint8x4.y = static_cast<uint8_t>(y);uint8x4.z = static_cast<uint8_t>(z);uint8x4.w = static_cast<uint8_t>(w);return load_as<uint32_t>(&uint8x4);
}template <unsigned int N>
__device__ inline void store_uint8_vector(uint8_t *dest, const uint32_t *ptr);template <>
__device__ inline void store_uint8_vector<1u>(uint8_t *dest, const uint32_t *ptr){dest[0] = static_cast<uint8_t>(ptr[0]);
}template <>
__device__ inline void store_uint8_vector<2u>(uint8_t *dest, const uint32_t *ptr){uchar2 uint8x2;uint8x2.x = static_cast<uint8_t>(ptr[0]);uint8x2.y = static_cast<uint8_t>(ptr[0]);store_as<uchar2>(dest, uint8x2);
}template <>
__device__ inline void store_uint8_vector<4u>(uint8_t *dest, const uint32_t *ptr){store_as<uint32_t>(dest, pack_uint8x4(ptr[0], ptr[1], ptr[2], ptr[3]));
}template <>
__device__ inline void store_uint8_vector<8u>(uint8_t *dest, const uint32_t *ptr){uint2 uint32x2;uint32x2.x = pack_uint8x4(ptr[0], ptr[1], ptr[2], ptr[3]);uint32x2.y = pack_uint8x4(ptr[4], ptr[5], ptr[6], ptr[7]);store_as<uint2>(dest, uint32x2);
}template <>
__device__ inline void store_uint8_vector<16u>(uint8_t *dest, const uint32_t *ptr){uint4 uint32x4;uint32x4.x = pack_uint8x4(ptr[ 0], ptr[ 1], ptr[ 2], ptr[ 3]);uint32x4.y = pack_uint8x4(ptr[ 4], ptr[ 5], ptr[ 6], ptr[ 7]);uint32x4.z = pack_uint8x4(ptr[ 8], ptr[ 9], ptr[10], ptr[11]);uint32x4.w = pack_uint8x4(ptr[12], ptr[13], ptr[14], ptr[15]);store_as<uint4>(dest, uint32x4);
}

 例子代码见:

https://github.com/Alexjqw/cuda-learning

 

参考:https://www.cnblogs.com/1024incn/p/4564726.html

 

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

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

相关文章

Mybase到期 破解

1.找到mybase安装目录 C:\Users\Zsh\AppData\Local\wjjsoft\nyfedit6 2.打开nyfedit.ini 3. 内容中查找 App.UserLic.FirstUseOn 删除 4.保存打开mybase即可

Android中给按钮同时设置背景和圆角示例代码

前言 最近在做按钮的时候遇到在给按钮设置一张图片作为背景的同时还要自己定义圆角&#xff0c;最简单的做法就是直接切张圆角图作为按钮就可以了&#xff0c;但是如果不这样该怎么办呢&#xff0c;看代码&#xff1a; 下面来看效果图 一、先建一个圆角的shape文件&#xff1a;…

awk使用方法

awk是行处理器: 相比较屏幕处理的优点&#xff0c;在处理庞大文件时不会出现内存溢出或是处理缓慢的问题&#xff0c;通常用来格式化文本信息 awk处理过程: 依次对每一行进行处理&#xff0c;然后输出 awk命令形式: awk [-F|-f|-v] ‘BEGIN{} //{command1; command2} END{}’…

cude的__ldg使用

一般使用__ldg是更好的选择。通过constant缓存存储的数据必须相对较小而且必须获取同一个地址以便获取最佳性能&#xff0c;相反&#xff0c;只读缓存则可以存放较大的数据&#xff0c;且不必地址一致。 __global__ void transpose2(const real* A, real* B, const int N) {co…

cuda的shared momery

CUDA SHARED MEMORY 在global Memory部分&#xff0c;数据对齐和连续是很重要的话题&#xff0c;当使用L1的时候&#xff0c;对齐问题可以忽略&#xff0c;但是非连续的获取内存依然会降低性能。依赖于算法本质&#xff0c;某些情况下&#xff0c;非连续访问是不可避免的。使用…

图的表示方法和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 代理协议为例。…