GPU架构
GPU特别适用于 密集计算，高度可并行计算，图形学
晶体管主要被用于 执行计算，而不是缓存数据，控制指令流
GPU计算的历史
2001/2002 —— 研究人员把GPU当做数据并行协处理器
GPGPU领域从此诞生
2007—— NVIDIA 发布 CUDA
CUDA 统一计算设备架构
GPGPU发展成 GPU Computing
2008—— Khronos 发布OpenCL 规范
CUDA术语
Host——主机端，通常指CPU，采用标准C语言编程，C++，Python
Device——设备端，通常指GPU(数据可并行)
采用标准C的扩展语言编程
Host和Device拥有各自的存储器
CUDA编程包括主机端和设备端两部分代码
Kernel——数据并行处理函数(核函数)，Kernel是指执行在GPU上的完整代码
通过调用Kernel函数在设备端创建轻量级线程，线程由硬件负责创建并调度
函数声明
__global__ void KernelFunc() 只能从主机端调用，在设备上执行，入口函数
__device__ float DeviceFunc() 设备 设备
__host__ float HostFunc() 主机 主机
__global__ 返回值类型必须是 void
__device__ 曾经默认内联，现在有些变化
__device__ 和 __host__ 可以同时使用
CUDA/GPU编程模型
CPU/GPU互动模型
GPU线程组织模型
GPU存储模型
基本的编程问题
CPU-GPU交互
各自的物理内存
通过PCIE总线互连
交互开销较大
并行编程模型
共享存储模型
线程模型
消息传递模型
数据并行模型
具体实例
OpenMp
MPI
SPMD：单指令多数据相似，每一个代码段处理多个数据
MPMD：多指令多数据
线程层次
Grid——一维或多维线程块(block)
一维或二维
Block——一组线程
一维，二维或三维
一个 Grid 里的每个 Block 的线程数是一样的
CUDA存储器类型(访存速度由大到小)
Register 寄存器
shared memory 共享内存
local memory 局部存储器
global memory 全局存储器
constant memory 常量存储器
texture memory 纹理存储器
内存模型
变量声明 存储器 作用域 生命期
必须是单独的自动变量不能是数组 register thread kernel
自动变量数组 local thread kernel
__shared__ int sharedVar shared block kernel
__device__ int globalVar global grid application
__constant__ intconstantVar constant grid application
global 和 constant 变量
Host 可以通过以下函数访问
cudaGetSymbolAddress()
cudaGetSymbolSize()
cudaMemcpyToSymbol()
cudaMemcpyFromSymbol()
Constants 变量必须在函数外声明
向量数据类型
char [1-4], uchar[1-4]
short [1-4],
int [1-4],
long [1-4],
longlong [1-4],
float [1-4]
double1, double2
同时适用 host 和 device 代码
通过make_<type name> 构造
int2 i2 = make_int2(1,2);
float4 f4 = make_float(1.0f,2.0f,3.0f,4.0f);
通过 .x .y .z 和 .w 访问
int2 i2 = make_int2(1,2);
int x = i2.x;
int y = i2.y;
数学函数
部分函数列表
sqrt,rsqrt
exp,log
sin,cos,tan,sincos
asin,acos,atan2
trunc,ceil,floor
Intrinsic function 内建函数
仅面向 device 设备端
更快,但精度降低
以__为前缀
__exp,__log,__sin,__pow.....