CUDA编程

函数申明：

• 　　__global__ void KernelFunc()　　Executed:device　　Callable:host
• 　　__device__ float DeviceFunc()　　......:device　　　　......:device
• 　　__host__ float HostFunc()　　　　......:host　　　　　......:host

　　__global__:　

　　　　　　返回值必须是void

　　__device__:

　　　　　　曾今默认内联，现在有一些变化。

　　Global和Device函数：

1. 1. 尽量少用递归
   2. 不要使用静态变量
   3. 少用malloc
   4. 小心通过指针实现的函数调用

• 向量数据类型：
  • 　　char[1-4],uchar[1-4]
  • 　　short[1-4],ushort[1-4]
  • 　　int[1-4],uint[1-4]
  • 　　long[1-4],ulong[1-4]
  • 　　longlong[1-4],ulonglong[1-4]
  • 　　floa[1-4]
  • 　　double1,double2
• 向量数据类型
  • 　　同时适用于host和device代码，通过函数make_<type name>构造

int2 i2 = make_int2(1, 2);

float4 f4 = make_float4(1.0f, 2.0f, 3.0f, 4.0f);

• • 通过.x, .y, .z,and .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,......
• 线程同步
  • 快内线程可以同步
  • 调用__syncthreads创建一个barrier栅栏
  • 每个线程在调用点等待块内所有线程执行到这个地方，然后所有线程继续执行后续指令

Mds[i] = Md[j]

__syncthreads()

func(Mds[i], Mds[i+1])

• 线程调度
  • Wrap一块内的一组线程　　
    • G80/GT200 - 32个线程
    • 运行于同一个SM
    • 线程调度的基本单元
    • threadIdx值连续
    • 一个实现细节 - 理论上
      • WrapSize
• 内存模型
  • 寄存器Registers
    • 每个线程专用
    • 快速，片上，可读写
    • 增加Kernal的寄存器用量，会导致什么结果？
  • 寄存器Register
    • 每个SM
      • 多达768threads
      • 8K个寄存器
  • 局部存储器Local Memory
    • 存储于global memory
      • 作用域每个thread
    • 用于存储自动变量数组
      • 通过常量索引访问
  • 共享存储器Shared Memory
    • 每个块
    • 快速，片上，可读写
    • 全速随机访问
    • 每个SM包括8个block，16KB共享存储器
  • 全局存储器Global Memory
    • 长延时（100个周期)
    • 片外，可读写
    • 随机访问影响性能
    • Host主机端课读写
    • GT200
      • 带宽：150GB/s
      • 容量：4GB
    • G80 - 86.4GB/s
  • 常量存储器Constant Memory
    • 短延时，高带宽，当所有线程访问同一位置是只读
    • 存储于gloabl memory但是有缓存
    • Host主机可读写
    • 容量:64KB

1. Gloabl and constant变量
   • Host 可以通过以下函数访问：
     • cudaGetSymbolAddress()
     • cudaGetSymbolSize()
     • cudaMemcpyToSymbol()
     • cudaMemcpyFromSymbol()
   • constants变量必须在函数外声明