GPU(CUDA)学习日记(九)------ CUDA存储器模型

标签: cuda存储bindingcache编程api
2012-09-27 10:53 1677人阅读 评论(1) 收藏 举报
分类:
GPU(16) CUDA(16) 高性能计算(7) 动态加载(12) 并行运算(7)

CUDA存储器模型:

GPU片内:register,shared memory;

板载显存:local memory,constant memory, texture memory, texture memory,global memory;

host 内存: host memory, pinned memory.

register: 访问延迟极低;

基本单元:register file (32bit/each)

计算能力1.0/1.1版本硬件:8192/SM;

计算能力1.2/1.3版本硬件: 16384/SM;

每个线程占有的register有限,编程时不要为其分配过多私有变量;

local memory:寄存器被使用完毕,数据将被存储在局部存储器中;

大型结构体或者数组;

无法确定大小的数组;

线程的输入和中间变量;

定义线程私有数组的同时进行初始化的数组被分配在寄存器中;

shared memory:访问速度与寄存器相似;

实现线程间通信的延迟最小;

保存公用的计数器或者block的公用结果;

硬件1.0~1.3中,16KByte/SM,被组织为16个bank;

声明关键字 _shared_  int sdata_static[16];

global memory:存在于显存中,也称为线性内存(显存可以被定义为线性存储器或者CUDA数组);

cudaMalloc()函数分配,cudaFree()函数释放,cudaMemcpy()进行主机端与设备端的数据传输;

初始化共享存储器需要调用cudaMemset();

二维三维数组:cudaMallocPitch()和cudaMalloc3D()分配线性存储空间,可以确保分配满足对齐要求;

cudaMemcpy2D(),cudaMemcpy3D()与设备端存储器进行拷贝;

host内存:分为pageable memory 和 pinned memory

pageable memory: 通过操作系统API(malloc(),new())分配的存储器空间;、

pinned memory:始终存在于物理内存中,不会被分配到低速的虚拟内存中,能够通过DMA加速与设备端进行通信;

cudaHostAlloc(), cudaFreeHost()来分配和释放pinned memory;

使用pinned memory优点:主机端-设备端的数据传输带宽高;

某些设备上可以通过zero-copy功能映射到设备地址空间,从GPU直接访问,省掉主存与显存间进行数据拷贝的工作;

pinned memory 不可以分配过多:导致操作系统用于分页的物理内存变, 导致系统整体性能下降;通常由哪个cpu线程分配,就只有这个线程才有访问权限;

cuda2.3版本中,pinned memory功能扩充:

portable memory:让控制不同GPU的主机端线程操作同一块portable memory,实现cpu线程间通信;使用cudaHostAlloc()分配页锁定内存时,加上cudaHostAllocPortable标志;

write-combined Memory:提高从cpu向GPU单向传输数据的速度;不使用cpu的L1,L2 cache对一块pinned memory中的数据进行缓冲,将cache资源留给其他程序使用;在pci-e总线传输期间不会被来自cpu的监视打断;在调用cudaHostAlloc()时加上cudaHostAllocWriteCombined标志;cpu从这种存储器上读取的速度很低;

mapped memory:两个地址:主机端地址(内存地址),设备端地址(显存地址)。  可以在kernnel程序中直接访问mapped memory中的数据,不必在内存和显存之间进行数据拷贝,即zero-copy功能;在主机端可以由cudaHostAlloc()函数获得,在设备端指针可以通过cudaHostGetDevicePointer()获得;通过cudaGetDeviceProperties()函数返回的canMapHostMemory属性知道设备是否支持mapped memory;在调用cudaHostAlloc()时加上cudaHostMapped标志,将pinned memory映射到设备地址空间;必须使用同步来保证cpu和GPu对同一块存储器操作的顺序一致性;显存中的一部分可以既是portable memory又是mapped memory;在执行CUDA操作前,先调用cudaSetDeviceFlags()(加cudaDeviceMapHost标志)进行页锁定内存映射。

constant memory:只读地址空间;位于显存,有缓存加速;64Kb;用于存储需要频繁访问的只读参数 ;只读;使用_constant_ 关键字,定义在所有函数之外;两种常数存储器的使用方法:直接在定义时初始化常数存储器;定义一个constant数组,然后使用函数进行赋值;

texture memory:只读;不是一块专门的存储器,而是牵涉到显存、两级纹理缓存、纹理拾取单元的纹理流水线;数据常以一维、二维或者三维数组的形式存储在显存中;缓存加速;可以声明大小比常数存储器大得多;适合实现图像树立和查找表;对大量数据的随机访问或非对齐访问有良好的加速效果;在kernel中访问纹理存储器的操作成为纹理拾取(texture fetching);纹理拾取使用的坐标与数据在显存中的位置可以不同,通过纹理参照系约定二者的映射方式;将显存中的数据与纹理参照系关联的操作,称为将数据与纹理绑定(texture binding);显存中可以绑定到纹理的数据有:普通线性存储器和cuda数组;存在缓存机制;可以设定滤波模式,寻址模式等;

本文的原文出处为:http://blog.csdn.net/ouczoe/article/details/5125621 帮助很大!谢谢原创作者!

GPU(CUDA)学习日记(九)------ CUDA存储器模型相关推荐

  1. CUDA学习笔记之 CUDA存储器模型

    CUDA学习笔记之 CUDA存储器模型 标签: cuda存储bindingcache编程api 2010-12-14 01:33 1223人阅读 评论(0) 收藏 举报 分类: CUDA(26) GP ...

  2. cuda学习笔记5——CUDA实现图像形态学腐蚀、膨胀

    cuda学习笔记5--CUDA实现图像形态学腐蚀.膨胀 代码 linux如何编译cuda和opencv代码 耗时情况 代码 #include "cuda_runtime.h" #i ...

  3. CUDA学习(九十七)

    显式同步和逻辑GPU活动: 请注意,即使内核在上述示例中快速运行并在CPU触及y之前完成,也需要显式同步. Unified Memory使用逻辑活动来确定GPU是否空闲. 这与CUDA编程模型保持一致 ...

  4. CUDA学习(九十一)

    CUDA环境变量: 有关多进程服务的环境变量记录在GPU部署和管理指南的MultiProcess Service部分.

  5. cuda学习(5):使用cuda核函数加速warpaffine

    1. warpaffine 介绍 放射变换(warpaffine),主要解决图像的缩放和平移来处理目标检测中常见的预处理行为. 比如有一张猫的图片,对图片进行letterbox变换,将图片进行缩放,并 ...

  6. Cuda 学习教程:Cuda 程序初始化

    Cuda程序初始化 目前,cuda里面没有对设备的初始化函数InitDevice(),只能每次调用的api函数的时候,加载设备的上下文,自动进行初始化,这将带来问题: First函数调用的时候,需要自 ...

  7. 推荐系统与深度学习(九)——NCF模型原理

    公众号后台回复"图书",了解更多号主新书内容 作者:livan 来源:数据python与算法 NCF原理 随着深度学习的推广,一些近期的研究将深度学习应用到了推荐中,然而很多工作都 ...

  8. C++ Primer Plus学习(九)——内存模型和名称空间

    内存模型和名称空间 单独编译 存储持续性.作用域和链接性 名称空间 单独编译 程序一般可以分为三部分: 头文件:包含结构声明和使用这些结构的函数的原型: 源代码文件:包含与结构有关的函数的代码: 源代 ...

  9. Activity学习日记(九)

    目录 Handler的基本使用方法示例 listview基本使用 简单聊天程序客户端(java)[服务器可以找epoll服务器日志] SendBd_有序广播和无需广播_获取电池电量 serviceTe ...

  10. Cuda学习笔记(一)——sm流处理器簇对blocks的调度策略

    由于GPU目前在各行各业的广泛应用,无论是深度学习.大数据.云计算等都离不开GPU的并行加速,前阵子自学了Cuda-c编程,希望将来的研究工作能够用得上. Cuda系列总共有4篇,这里主要用于记录本人 ...

最新文章

  1. Centos安装GD库
  2. python subprocess库 终端命令行命令
  3. letswave7中文教程3:脑电数据预处理-ICA去除伪影
  4. 如何通过Google查找专业文献资料 [转]^_^!!
  5. linux shell 缺少 ps 命令
  6. [云炬创业管理笔记]第一章测试4
  7. C/C++入门易错点及常用小技巧
  8. 没人性!让我们这么早来上班
  9. (27)StyleLint—CSS代码格式校验
  10. PreparedStatement JDBC域处理/SQl攻击
  11. 学计算机的能看出批图吗,高手教你:如何看出一张图有没有被 P 过?
  12. 关于 TCP window size
  13. hadoop相关问题-stop-all.sh
  14. 海马苹果助手ipad版_资源分享苹果越狱工具
  15. 智慧城市大数据分析系统解决方案
  16. 2019年中国大学生计算机设计大赛国赛答辩
  17. 利用com.googlecode.libphonenumber解析手机号归属地
  18. iPhone设备上安装beta版本系统,在浏览器中搜索网址 beta.apple.com
  19. 高速PCB设计系列基础知识67 | 尺寸与公差标注内容与通用知识
  20. 基于android的共享车位app

热门文章

  1. 屏蔽朋友圈的第一天的感悟
  2. Final Cut Pro模版网站
  3. DM 源码阅读系列文章(二)整体架构介绍
  4. 3.6 mkpasswd命令
  5. SQL 2005启用组件Ad Hoc Distributed Queries
  6. 【深入Cocos2d-x】探索Cocos2d-x中的内存管理-引用计数和自动释放池
  7. HTTP referer/HTTP referrer
  8. C++ 面试知识总结
  9. CodeForces 416B
  10. 把 分数化为循环小数 和 把循环小数化为分数 的方法