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

标签: cuda存储bindingcache编程api
2010-12-14 01:33 1223人阅读 评论(0) 收藏 举报
分类:
CUDA(26)
GPU片内:register,shared memory; 板载显存:local memory,constant memory, texture memory, texture memory,global memory; host 内存: host memory, pinned memory.

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数组;存在缓存机制;可以设定滤波模式,寻址模式等;

CUDA学习笔记之 CUDA存储器模型相关推荐

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

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

  2. CUDA学习笔记之程序优化

    CUDA学习笔记之程序优化 标签: cuda优化conflict存储算法数学计算 2010-01-05 17:18 5035人阅读 评论(4) 收藏 举报 分类: CUDA(6) 版权声明:本文为博主 ...

  3. 深度学习(三十六)异构计算CUDA学习笔记(1)

    异构计算CUDA学习笔记(1) 原文地址:http://blog.csdn.net/hjimce/article/details/51506207 作者:hjimce 近日因为感觉自己在深度学习工程化 ...

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

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

  5. CUDA学习笔记(持续更新——蜗速)

    CUDA学习笔记(持续更新--蜗速) 1.CUDA 程序实现流程如下 2.内存管理 3.核函数 4.全局数据访问唯一索引 5.设备管理 附录代码 1.CUDA 程序实现流程如下 将数据从CPU内存拷贝 ...

  6. WinSock学习笔记3:Select模型

    WinSock学习笔记3:Select模型 unit Unit1; interface uses   Windows, Messages, SysUtils, Variants, Classes, G ...

  7. Python数据挖掘学习笔记】九.回归模型LinearRegression简单分析氧化物数据

    #2018-03-23 16:26:20 March Friday the 12 week, the 082 day SZ SSMR [Python数据挖掘学习笔记]九.回归模型LinearRegre ...

  8. 2020-4-20 深度学习笔记20 - 深度生成模型 3 (实值数据上的玻尔兹曼机)

    第二十章 深度生成模型 Deep Generative Models 中文 英文 2020-4-17 深度学习笔记20 - 深度生成模型 1 (玻尔兹曼机,受限玻尔兹曼机RBM) 2020-4-18 ...

  9. 2020-4-22 深度学习笔记20 - 深度生成模型 5 (有向生成网络--sigmoid信念网络/可微生成器网络/变分自编码器VAE/生产对抗网络GAN/生成矩匹配网络)

    第二十章 深度生成模型 Deep Generative Models 中文 英文 2020-4-17 深度学习笔记20 - 深度生成模型 1 (玻尔兹曼机,受限玻尔兹曼机RBM) 2020-4-18 ...

最新文章

  1. ssm配置socket_ssm框架中集成websocket实现服务端主动向客户端发送消息
  2. MySQL探秘(七):InnoDB行锁算法
  3. 三、垃圾收集之判断对象是否存活
  4. m_Orchestrate learning system---二十、如何写代码不容易犯错
  5. flex布局学习笔记
  6. 卷积神经网络(CNN:Convolutional Neural Network)
  7. Java计基础---常用类之StringBuilder类--String、StringBuilder和StringBuffer 的区别
  8. Android WebView:这是一份全面 详细的WebView学习指南
  9. elf section类型_探索ELF可执行文件的“干货”:段头表和段的基本介绍
  10. 全国计算机等级考试3月份报名时间,2021年3月全国计算机等级考试报名时间公布...
  11. 光学计算机的工作原理,使用光学计算机的人工智能超分辨率
  12. 计算机视觉大佬--何凯明
  13. php生成五星红旗,php基于GD库画五星红旗的方法_PHP
  14. 英伟达服务器系统显卡驱动,NVIDIA显卡驱动
  15. macbook 鼠标光标乱跳解决
  16. 贪吃蛇c语言存档读档,刚学C语言,想写一个贪吃蛇的代码
  17. 移动web适配的方式
  18. 钉钉直播 — 为什么直播屏幕共享显示黑屏?(仅针对mac端)
  19. 非诚勿扰php男嘉宾,非诚勿扰 php
  20. python 图片锐化_Python图像处理介绍--图像模糊与锐化

热门文章

  1. what is the thinking routine of the open source?
  2. 如何积累自己的技术认知
  3. Creator Upload NFT sequencial diagram
  4. 真实感水面绘制-资料收集
  5. 安卓上为什么不能用system.io.file读取streammingAssets目录下的文件
  6. 大数相乘(大数阶乘模板)
  7. 邀您共赴数据库学术顶会ICDE 2019——阿里云专场 零距离接触达摩院数据库“最强大脑”...
  8. SLF4j、log4j管理系统日志(Maven)
  9. 在表空间有足够free space的情况下出现ORA-1652
  10. NYOJ 625 笨蛋的难题(二)