并行程序设计---cuda memory

CUDA存储器模型:

GPU片内:register,shared memory;

host 内存: host memory, pinned memory.

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

register: 訪问延迟极低;

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

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

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

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

shared memory:訪问速度与寄存器相似。

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

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

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

声明keyword _shared_  int sdata_static[16];

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_ keyword。定义在全部函数之外。两种常数存储器的用法:直接在定义时初始化常数存储器。定义一个constant数组。然后使用函数进行赋值;

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

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

大型结构体或者数组。

无法确定大小的数组;

线程的输入和中间变量;

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

global memory:存在于显存中。也称为线性内存(显存能够被定义为线性存储器或者CUDA数组)。

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

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

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

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

时间: 2024-10-12 19:45:24

并行程序设计---cuda memory的相关文章

【CUDA并行程序设计系列(1)】GPU技术简介

http://www.cnblogs.com/5long/p/cuda-parallel-programming-1.html 本系列目录: [CUDA并行程序设计系列(1)]GPU技术简介 [CUDA并行程序设计系列(2)]CUDA简介及CUDA初步编程 [CUDA并行程序设计系列(3)]CUDA线程模型 [CUDA并行程序设计系列(4)]CUDA内存 [CUDA并行程序设计系列(5)]CUDA原子操作与同步 [CUDA并行程序设计系列(6)]CUDA流与多GPU 关于CUDA的一些学习资料

翻新并行程序设计的认知整理版(state of the art parallel)

近几年,业内对并行和并发积累了丰富的经验,有了较深刻的理解.但之前积累的大量教材,在当今的软硬件体系下,反而都成了负面教材.所以,有必要加强宣传,翻新大家的认知. 首先,天地倒悬,结论先行:当你需要并行时,优先考虑不需要线程间共享数据的设计,其次考虑共享Immutable的数据,最糟情况是共享Mutable数据.这个最糟选择,意味着最差的性能,最复杂啰嗦的代码逻辑,最容易出现难于重现的bug,以及不能测试预防的死锁可能性.在代码实现上,优先考虑高抽象级别的并行库(如C++11的future,PP

CUDA ---- Memory Model

Memory kernel性能高低是不能单纯的从warp的执行上来解释的.比如之前博文涉及到的,将block的维度设置为warp大小的一半会导致load efficiency降低,这个问题无法用warp的调度或者并行性来解释.根本原因是获取global memory的方式很差劲. 众所周知,memory的操作在讲求效率的语言中占有极重的地位.low-latency和high-bandwidth是高性能的理想情况.但是购买拥有大容量,高性能的memory是不现实的,或者不经济的.因此,我们就要尽量

OpenMP并行程序设计——for循环并行化详解

转载请声明出处http://blog.csdn.net/zhongkejingwang/article/details/40018735 在C/C++中使用OpenMP优化代码方便又简单,代码中需要并行处理的往往是一些比较耗时的for循环,所以重点介绍一下OpenMP中for循环的应用.个人感觉只要掌握了文中讲的这些就足够了,如果想要学习OpenMP可以到网上查查资料. 工欲善其事,必先利其器.如果还没有搭建好omp开发环境的可以看一下OpenMP并行程序设计--Eclipse开发环境的搭建 首

并行程序设计导论学习笔记——OpenMP(1)

使用OpenMP需要在编译器上打开OpenMP开关,并包含omp.h文件.我使用的是在Windows下的Visual Studio 2015,只需在工程选项中打开OpenMP支持就可以了.按照书上的说法,GCC增加参数-fopenmp就可以了. OpenMP有两个重要的函数: omp_get_thread_num() omp_get_num_threads() 他们的返回值都是无符号整数,第一个用来返回当前执行的线程编号,而第二个返回一共有的线程数量.和C语言的数组下标类似,当前执行的线程编号的

【Cuda并行编程之二】Cuda Memory Hierarchy_Cuda内存层次结构

要想编写高效的程序,那么一定要对内存结构有比较深刻的认识,就像C/C++里面的堆内存,栈内存,全局存储区,静态存储区,常量区等.Cuda是并行计算框架,而GPU的内存有限,那么如果想编写高效的Cuda程序,首先要对其内存结构有一个简单的认识. 首先我们先上一张图,然后通过解释一些名词和代码来进行解释. 各种存储器比较: 存储器  位置 拥有缓存 访问权限 变量生存周期 register GPU片内 N/A device可读/写 与thread相同 local memory 板载显存 无 devi

CUDA并行程序设计 开发环境搭建与远程调试

课题需要用到GPU加速.目前使用的台式电脑只有核心显卡,而实验室有一台服务器装有NVIDIA GTX980独显.因此,想搭建一个CUDA的开发环境,来实现在台式机上面开发cuda程序,程序在服务器而不必每次都跑去服务器上面. 目前找到的远程调试方案有三个: Windows平台下,通过Nsight Monitor连接 带有NVIDIA显卡的Windows笔记本电脑. Linux平台下,通过Nsight Eclipse Edition 连接Linux服务器 通过Xshell连接到服务器(操作系统为C

CUDA ---- Memory Access

Memory Access Patterns 大部分device一开始从global Memory获取数据,而且,大部分GPU应用表现会被带宽限制.因此最大化应用对global Memory带宽的使用时获取高性能的第一步.也就是说,global Memory的使用就没调节好,其它的优化方案也获取不到什么大效果. Aligned and Coalesced Access 如下图所示,global Memory的load/store要经由cache,所有的数据会初始化在DRAM,也就是物理的devi

CUDA memory

原文链接 CUDA存储器类型: 每个线程拥有自己的register and loacal memory; 每个线程块拥有一块shared memory; 所有线程都可以访问global memory; 还有,可以被所有线程访问的只读存储器:constant memory and texture memory 1.  寄存器Register 寄存器是GPU上的高速缓存器,其基本单元是寄存器文件,每个寄存器文件大小为32bit. Kernel中的局部(简单类型)变量第一选择是被分配到Register