zoukankan      html  css  js  c++  java
  • 并行程序设计---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()与设备端存储器进行拷贝;

     

  • 相关阅读:
    angular4升级angular5问题记录之this.location.back()
    angular4升级angular5问题记录之No NgModule metadata found for 'AppModule'
    Angular4图片上传预览路径不安全问题
    Angular4.0引入laydate.js日期插件方法
    Angular4.0用命令行创建组件服务出错
    IE10以下的img标签问题
    关于for循环删除数组内容出现的问题
    关于onmouseover和onmouseout的bug
    纯lua实现Base64加密与解密
    SciTE如何修改背景色
  • 原文地址:https://www.cnblogs.com/gccbuaa/p/7231364.html
Copyright © 2011-2022 走看看