===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,占有率和寄存器的压力:
*通过多核多线程的使用潜伏期被隐藏
*限制同时运行线程的一些因素:
**寄存器: