CUDA基础学习笔记

===CUDA介绍===
CUDA是一种可扩展的并行编程模型，使得程序不用通过重新编译就可以在任何数量处理器上运行。
支持异步系统（如CPU+GPU）,CPU和GPU是各自拥有DRAM的独立器件
以100多个核心，1000多条并行线程来衡量
主要让程序专门用于并行算法，而不是一种并行编程语言的结构

===理解CUDA的并行运算===
CUDA中的并行抽象的关键：

1，同时进行线程的等级制度
**有多个线程组成的平行kernel
所有线程都在执行相关的程序，用平行线程来代替顺序循环
**线程组成线程块
同一个块中的线程可以合作并且共享内存
**线程快组成Grid（栅格）
每一个线程，线程块都具有唯一的ID号，分别是threadIdx和blockIdx

Kernel含有的一些内容:
*Device code:
**如__global__关键字表示代码运行在device端；
**还有一系列的浮点数参量 float* A...
**threadIdx.x和blockDim.x，blockIdx.x组合表示的id号
**基于带id号i参量的一些运算：C[i]=A[i]+B[i]
*Host code：
如<<<N/256,256>>>: 用N/256个线程块运行，每个block有256个线程。

2，轻重量级别的同步单元

*同步和协调：
线程块里面的线程可以通过barries来进行同步。这可以通过一个叫做__syncthreads()的内部函数来实现，强制一个线程块中所有的线程，必须等到每一个线程都允许执行的状态
如：
   //将Asub和Bsub从device内存加载到共享内存（数据准备）
   As[row][col] = GetElement(Asub, row, col);
   Bs[row][col] = GetElement(Bsub, row, col);
  
   //等带所有的数据都加载完毕才开始执行计算（一致）
   __syncthreads();
  
   //线程全部等待集合完毕，开始运算（运行）
   for (int e = 0; e < BLOCK_SIZE; ++e)
   Cvalue += As[row][e] * Bs[e][col];
  
   //等待前一次所有计算都完成才开始下一轮新的运算（等待计算结束）
   __syncthreads();

*线程块通过原子内存运算符来调整，实现线程的调整
如atomicInc()用于增加共享队列的指针

*Kernel之间存在隐含的barrier

3，联合线程的存储共享模型

*内存的模型：
**对于单个线程：单线程存取的本地内存--DDR
   仅对某些自动变量的访问：
   a,不能确定是用常数量索引的数组;b,可能消耗太多寄存器空间的大结构或数组,c,超过可用的寄存器的任何变量
**对于每个线程块：线程块中所有线程都可以访问的共享内存：__shared__ 开头--GPU内部
   a,变量通过block共享：__shared__ int *begin, *end;
   b,中间结果存储:__shared__ int a; a= b;....b=a;最后结果又返回
   c,各个线程之间的通讯值。scratch[threadIdx.x]=... scratch[threadIdx.x - 1];
**对于每个device:随着连续的kernel而可以和对应kernel存取的全局内存cudaMalloc()用于分配；cudaFree()释放:DDR
**对于多个device::cuDeviceGetCount()和cuDeviceGet()：DDR cudaMemcpy()用于CPU和GPU之间的数据转移

4，对于C语言的最小扩充：
*特殊语法和函数
**头声明关键字，说明存储类型：
   __global__ void KernelFunc(...); // kernel callable from host
   __device__ void DeviceFunc(...); // function callable on device
   __device__ int GlobalVar; // variable in device memory
   __shared__ int SharedVar; // in per-block shared memory

**扩充函数语法符号，声明平行kernel启动：
   KernelFunc<<<500, 128>>>(...); // 500 blocks, 128 threads each
**用特殊的变量为kernel内部的线程进行鉴别：
   dim3 threadIdx; dim3 blockIdx; dim3 blockDim;
**kernel代码内部固有的扩展：
   __syncthreads(); // barrier使同步

*GPU代码中的特征：
**标准的数学函数：sin{f}, pow{f}, atan{f}, ceil{f}, min, sqrt{f}。。。
**原子内存操作符：atomicAdd, atomicMin, atomicAnd, atomicCAS。。。
**纹理使用：
   texture<float,2> my_texture;//声明纹理
   float4 texel = texfetch(my_texture, u, v);
*运行时间的支持：
**明确内存分配并返回指向GPU内存的指针
   cudaMalloc(), cudaFree()
**明确内存的拷贝：包括host和device之间，device之间的拷贝：
   cudaMemcpy(), cudaMemcpy2D(), ...
**纹理的管理：
   cudaBindTexture(), cudaBindTextureToArray(), ...
**OpenGL和DirectX互相操作性：
   cudaGLMapBufferObject(), cudaD3D9MapVertexBuffer(),
5，CUDA编译：
*CUDA的编译过程：
**通用的操作：
  C/C++ CUDA程序（*.cu）-->NVCC-->CPU code
**专业处理：                |--->PTX Code-->PTX to 目标转换器-->目标device(GPU1...GPUn)

 

通过目标转换器，PTX code能够适用于不同版本的显卡

*CUDA的库：
**CUBLAS:基本线性代数子程序库，不需要CUDA驱动的直接操作。
**CUFFT:快速傅立叶变换库，无需CUDA驱动的操作，直接实现。

*CUDPP：数据并行单元
减小，扫描，归类。。。

6，优化
*Kernel优化：
**全局内存吞吐量的最大化
**有效利用共享内存
**发散扭曲最小化
**固有指令
*GPU和CPU的优化:
*PCIE吞吐量最大化
*内存异步拷贝

7，占有率和寄存器的压力：
*通过多核多线程的使用潜伏期被隐藏
*限制同时运行线程的一些因素：
**寄存器：