Compute Unified Device Architecture
--------------------------------------------------------
GPU更适用于并行计算:
更多的晶体管用于数据计算,而不是用于数据缓存和流程控制;
大量的数据,高强度的运算,内存读写延迟影响不大,不需要复杂的数据缓存机制;
数据的并行处理;

主要的问题:
通过API控制GPU;
GPU能够任意读取显存,却不能任意写入;
内存带宽的限制;
--------------------------------------------------------
CUDA仍然使用在driver之后的API;
DRAM memory addressing 可以分为两种形式:Gather,同一块内存数据可以送往不同的计算单元(ALU);Scatter,一个计算单元能够将数据写入不同的内存区域;CUDA可以让GPU做到这一点;
DRAM解决内存带宽的方法仍然是cache,线程之间通过它们来共享数据。避免overfetch和round-trips to DRAM。
(The bridge device may transmit several read request packets 402 before receiving the first read return packet 404. This behavior is known as prefetch and tends to improve the performance of the device by maintaining signal bandwidth on the bus at levels close to the maximum levels which the bus is capable of sustaining (e.g. improving bus bandwidth efficiency).

One disadvantage of prefetch is that a number or read return packets may arrive for the component after it has indicated that it will no longer accept returned data. Due to prefetch, the bus bridge may have transmitted additional read request packets 402 to the target component which result in returned data which the requesting component does not accept. The additional read return packets are known as overfetch. Overfetch is a side effect of prefetch and tends to waste bandwidth on the shared bus.

Efficient operation of the shared bus may be achieved by utilizing prefetch while attempting to minimize overfetch. Traditional approaches have adjusted the number of prefetched packets according to buffer capacity in the bus bridge component. However, this approach may not adequately compensate for the negative impact of overfetch on bus efficiency. )
-------------------------------------------------------------------------------------------
可以并行处理数据的function叫做kernel。它被编译为设备可执行代码,下载到设备上执行。
通过API利用设备DMA(Direct Memory Access)实现host内存和device内存的互相拷贝。
一些线程能够高度共享数据,在kernel中可以为它们设置同步点。这些线程组成了一个1,2或3维的block,block中的每个线程都有一个ID。
执行相同kernel的相同大小的block能够组成一个grid。不同block中的线程不能直接通信和同步。grid是二维的,其中的每个block都有一个ID。
---------------------------------------------------------------------------------------------
每个thread有register,local mem;
每个block有shared mem;
每个grid有global mem,constant mem,texture mem(后两者是只读的);kernel执行期间他们不可以被改变。
----------------------------------------------------------------------------------------------
硬件上,每一个multiprocessor有一组register,和共享的一个memory;还有一个只读的constant cache和只读的texture cache;
一个multiprocessor能够充当一个或多个block;他能批次处理多少个block取决于一个线程要多少个register和一个block需要多少shared mem;
一个multprocessor上的SIMD fashion线程组成一个warp,线程数目数目叫做warp size。一个block应该由多个warp执行。
一个multprocessor上的所有active block(能被一次批次处理的blocks)的所有warp之间被调度为轮询执行。
----------------------------------------------------------------------------------------------
给定一个grid中的thread个数,如何划分block的个数呢?
首先,它至少和multiprocessor个数相同;但是block太少,会因为线程同步使multiprocessor处于空闲;
另一方面,如果block太多,即block中的线程数目太少,设备内存的读取会成为瓶颈;
----------------------------------------------------------------------------------------------
显存中,有一部分叫做primary surface,用作输出到显示设备。当分辨率改变时候,需要的primary surface也会相应改变;于是影响了CUDA显存的分配形式,可能会导致CUDA程序crash。


-------------------------------- 
VARIABLE TYPE QUALIFIERS
-------------------------------
 device code中,一般的变量声明,都存在与register中,或者local mem中;指针的使用需要分辨出是shared mem还是global mem;不要试图在hostdevice之间共享内存指针,会出现内存段错误。
__device__
变量存在于device,未使用其它qualifier情况下:
global mem中;应用程序生命期;能够由每个线程,以及host读写;
__constant__:存在于constant mem中;应用程序生命期;能够由每个线程,以及host读写;
__shared__:存在于block shared mem中;block生命周期;由block中的线程读写;
这些变量都不能用于结构;不能用于host上的参数,临时变量。
---------------
TEXTURE:
---------------
首先声明一个Texture Reference,类型,维数和读数据方式(8bit16bit是否作为normalize float读入);
运行时fetching:默认通过[0,N)作为索引取样,或可以normalize[0,1)addressing mode; filtering;
Texture也可以binddevice linear Mem 而不是 CUDA Array,但失去了运行时fecthing的多种选择;
Texture Binddevice mem需要对齐,由API返回偏移量。


---------------------------
 FUNCTION TYPE QUALIFIERS
---------------------------
__device__device调用,device执行;
__global__:作为kernel;由host调用,由device执行;
__host__host调用,host执行;(默认值)
其中__device____host__可以同时使用,将产生两套编译结果。

__device__
其实是inline函数,不能使用函数指针。
device上执行的函数,不支持递归,不支持静态,不能使用可变数目参数;
__global__是最主要的类型,它需要指定执行配置信息;它通过shared memory传递参数;device在执行完毕之前,函数就已经返回;它不能有返回值。
-------------------------------- 
EXECUTION CONFIGURATION for global function
-------------------------------
As an example, a function declared as
__global__ void Func(float* parameter);
must be called like this:
Func<<< Dg, Db, Ns >>>(parameter);
分别代表Dimension GridDimension BlockNumber of Bytes in Shared Mem包括静态分配和动态分配;其中动态内存分配方法:
extern __shared__ float shared[];
所有以这种方式声明的变量都初始指向同一内存区域,需要我们自己划分内存区域
__device__ void func() // __
{ short* array0 = (short*)array;
  float* array1 = (float*)&array0[1];
  int* array2 = (int*)&array1[64];

----------------------------------
COMMON RUNTIME COMPONENT

----------------------------------
BUILT-IN VARIABLES
不可写,不可使用指针
gridDim : dim3
blockIdx : uint3
blockDim : dim3
threadIdx : uint3
支持的向量类型:char,uchar,short,ushort,int,uint,long,ulong,float
数学函数;
clock_t clock()时钟周期计数;使用计时时候注意,多线程是通过时分来实现的。
------------------------
DEVICE RUNTIME FUNCTION

------------------------
数学函数:使用
__前缀;编译选项 –use_fast_math可以实现一个速度快但精度不高的版本;
Void __syncthreads():同步block中的所有thread;使用在条件判断中要特别注意!
Type Conversion__float2int[rn,rz,ru,rd]round to nearest even; round towards zero; round up; round down
Type Casting__int_as_float__float_as_int
Texturing from Device Memtex1Dfetch
Texturing from CUDA Arraystex1D, tex2D
Atomic Functionsread-modify-write on 32bit integers in global mem;
-------------------------
HOST RUNTIME COMPONENT

-------------------------
两套APIDRIVERcu)或者RUNTIMEcuda),只可以使用其中某一套。
RUNTIME做了初始化工作,context managementmodule management
COMMON
两套API都可以列举设备及其性能;一般应该一个线程对应一个device
Mem可以分配为linear mem或者CUDA arrays:前者是32bit地址空间,可以自由使用指针;CUDA arrays的内存分配不透明,device只能够读取。
__global__和大多数runtime函数都是异步的,并不是执行完毕才返回。CudaThreadsSynchronize()和cuCtxSynchronize()强制runtime等待直到所有任务完毕。只有在hostdevice之间的内存拷贝,和OPENGL/DIRECT3D的协作,map/unmapfree mem这些操作不是异步的。

-----------
RUNTIME API
-----------
初始化:当第一次调用某条runtime函数时候,执行初始化工作;
Device ManagementcudaGetDeviceCount; cudaGetDevicePorperties; cudaSetDevice用于为host线程选择一个对应的设备,它应当在所有__global__函数之前调用。
Memory ManagementcudaMalloc; cudaMallocPitch用于二维情况; cudaFree; 
       CudaMallocArray(需要channel desc); cudaFreeArray; 
       CudaGetSymbolAddresscudaGetSymbolSize;用于取得global mem的变量指针;
       cudaMemCpy系列函数可以从hostdevicedevice
       cudaMallocHostcudaFreeHost用于page-lockedhost内存,这样的内存在hostdevice之间交换数据时候会快些(比起malloc分配的分页内存);
Texture Reference ManagementcudaBindTexturecudaUnbindTexture
调试:emulation模式为block中的每一个thread建立一个host线程;但是它不是simulating,一些问题仍难以发现。比如host可能会保存float的计算中间量为更高精度;(可以试着声明变量为volatile强制令其单精度存储;编译选项;_controlfp
----------
DRIVER API
----------