现在的位置: 首页 > 综合 > 正文

cuda 笔记

2014年02月10日 ⁄ 综合 ⁄ 共 7282字 ⁄ 字号 评论关闭

  CPU代码作用:在kernel启动前进行数据准备和设备初始化的工作,以及在kernel之间进行的一些串行运算。理想状况下 ,CPU串行代码的作用只是清理上一个kernel函数,并启动下一个kernel函数。

       CUDA并行计算函数kernel:它是整个CUDA程序中的一个可以被并行执行的步骤。

       kernel函数中存在两个层面的并行:Grid中的block间并行;block中的thread间并行。

       kernel:组织形式是grid;以block为单位执行。

Grid:表示一系列可以并行的block的集合;

         各个block之间无法通信,没有执行顺序;

         目前一个kernel中只有一个grid,未来DX11中将采用MIMD架构,允许一个kernel中存在多个不同的grid。

block:同一个block的线程需要共享数据,必须在同一个sm中发射;(在同一个时刻,一个sm中可有多个活动block)

          block中的每一个thread被发射到一个sp上;

          block的数量是处理核心的数量的几倍的时候,才能充分发挥GPU的运算能力:如果太少,无法体现其计算速度相较传统方式的优势。

Thread:有自己的私有寄存器和local memory;

             同一个block内的线程可以通过共享存储器和同步机制进行通信。

 

实际运行单元:warp(线程束),大小由硬件能力决定。tesla架构的gpu中为32。划分依据是block的ID,比如,0~31为一束。

                    32的warp:每发射一条warp指令,sm中的8个sp会将这条指令执行4遍。

CUDA编程key:

    程序以及warp中尽量避免分支:

                  warp中尽量避免使用分支:如果有分支,sm需要把每一个分支的指令发射到每一个sp上,再根据sp决定要不要执行,执行时间将

                  是所有分支之和。

     优化存储器访问:理想状态时所有存储器仅此那个传输的同时,GPU的各个核心也始终在进行计算。这需要合理划分程序。

CUDA软件体系:

        CUDA C:一种用c语言编写设备端代码的编程方式,包括对c的一些扩展和一个运行时库。

         nvcc编译器:分离源文件中的主机代码和设备代码;

                   主机代码由c文件形式输出,交流高性能编译器,如ICC、GCC或者其他合适的高性能编译器进行编译。或者在编译的最后阶段交给其他编译器生成.obj或者.o文件。

                   设备代码由nvcc编译成ptx代码或者二进制代码。

         ptx代码:类似于汇编余元,是为动态编译器jit设计的输入指令序列。

                      jit可以在使用不同的机器语言的显卡上运行同样的ptx,保证兼容性。

                      jit的输出受到硬件和其他一些因素影响,有不确定性。

                      对于需要确定代码的独立软件开发商,可以把代码编译成cuda二进制代码cubin,避免jit过程的不确定性。

CUDA运行时api:在驱动程序api基础上进行了封装,编程方便。

                        放在CUDArt包中;

                        函数以CUDA前缀;

CUDA驱动程序API:基于句柄的底层接口。

                          放在nvCUDA包里;

                          前缀为cu;

 

cuda自带的运算能力计算器:

C:/Documents and Settings/All Users/Application Data/NVIDIA Corporation/NVIDIA GPU Computing SDK/C/tools

cuda自带的实例程序:

包括矩阵相乘,粒子系统等示例程序,有些程序中带有cpu运算程序,方便进行比较;

C:/Documents and Settings/All Users/Application Data/NVIDIA Corporation/NVIDIA GPU Computing SDK/C/tools

VS2008中高亮显示dat文件:

C:/Documents and Settings/All Users/Application Data/NVIDIA Corporation/NVIDIA GPU Computing SDK/C/doc/syntax_highlighting/visual_studio_8

cuda说明文件,编程指南,cublas,referenceManual,CUFFT说明等:

C:/CUDA/doc

 

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存储器模型:

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

抱歉!评论已关闭.