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,一经查实,立即删除!

相关文章

串口通信参数

串口是一种非常通用的设备通信的协议&#xff08;不要与通用串行总线Universal Serial Bus(USB)混淆&#xff09;。大多数计算机包含两个基于RS232的串口。串口同时也是仪器仪表设备通用的通信协议&#xff1b;很多GPIB兼容的设备也带有RS-232口。同时&#xff0c;串口通信协议…

马虎的算式

小明是个急性子&#xff0c;上小学的时候经常把老师写在黑板上的题目抄错了。 有一次&#xff0c;老师出的题目是&#xff1a;36 x 495 ? 他却给抄成了&#xff1a;396 x 45 ? 但结果却很戏剧性&#xff0c;他的答案竟然是对的&#xff01;&#xff01; 因为 36 * 495 …

identifier __ldg is undefined

__ ldg&#xff08;&#xff09; 内置仅适用于计算能力3.5&#xff08;或更高版本&#xff09;架构。 这意味着&#xff1a; 必须在计算3.5&#xff08;或更新版本&#xff09;GPU上运行必须针对计算3.5&#xff08;或更新版本&#xff09;的GPU进行编译不能为旧体系结构编译。…

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…

html中contentEditable属性

contentEditable的主要功能是允许用户在线编辑元素中的内容&#xff0c;contentEditable是一个布尔型属性&#xff0c;因此在使用时候赋予初值&#xff0c;默认由inherit状态确定(看其父元素状态) 示例代码&#xff1a; <!doctype html> <html> <head> <…

Storm tick 功能

1. tick的功能 Apache Storm中内置了一种定时机制——tick&#xff0c;它能够让任何bolt的所有task每隔一段时间&#xff08;精确到秒级&#xff0c;用户可以自定义&#xff09;收到一个来自__systemd的__tick stream的tick tuple&#xff0c;bolt收到这样的tuple后可以根据业…

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; //枚…

html中contextmenu属性

contexmenu属性用于定义<div>元素的上下文菜单&#xff0c;所谓上下文菜单就是用书右键点击元素时候出现。 示例代码&#xff1a; <!doctype html> <html> <head> <meta charset"utf-8"> <title> </title> </head>…

Hive 分区表操作 创建、删除

删除某个分区指定数据 ALTER TABLE tableName DROP PARTITION (times 201851); ALTER TABLE user_portrait_task DROP PARTITION (times > 201801);

成为专业程序员的 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;通俗易懂

ros学习网站

1.turtlebot ROS平台介绍和学习资源 https://www.ncnynl.com/turtlebot2.html 2.古月居 https://www.guyuehome.com/ http://www.fetorobot.com/NewsDetail/1133238.html# 官方&#xff1a; http://wiki.ros.org/ sudo sh -c . /etc/lsb-release && echo "d…

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

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