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；不要试图在host和device之间共享内存指针，会出现内存段错误。
__device__ 变量存在于 device ，未使用其它 qualifier 情况下：
在 global mem 中；应用程序生命期；能够由每个线程，以及 host 读写；
__constant__ ：存在于 constant mem 中；应用程序生命期；能够由每个线程，以及 host 读写；
__shared__ ：存在于 block shared mem 中； block 生命周期；由 block 中的线程读写 ;
这些变量都不能用于结构；不能用于 host 上的参数，临时变量。
---------------------------
FUNCTION TYPE QUALIFIERS
---------------------------
__device__ ： device 调用， device 执行；
__global__ ：作为 kernel ；由 host 调用，由 device 执行；
__host__ ： host 调用， host 执行；（默认值）
其中 __device__ 和 __host__ 可以同时使用，将产生两套编译结果。
__device__ 其实是 inline 函数，不能使用函数指针。
在 device 上执行的函数，不支持递归，不支持静态，不能使用可变数目参数；
__global__ 是最主要的类型，它需要指定执行配置信息；它通过 shared memory 传递参数； device 在执行完毕之前，函数就已经返回；它不能有返回值。
__global__ void Func(float* parameter);
must be called like this:
Func<<< Dg, Db, Ns >>>(parameter);
分别代表 Dimension Grid ； Dimension Block ； Number 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 Mem：tex1Dfetch；
Texturing from CUDA Arrays：tex1D, tex2D；
Atomic Functions：read-modify-write on 32bit integers in global mem;
-------------------------
HOST RUNTIME COMPONENT：
-------------------------
两套API：DRIVER（cu）或者RUNTIME（cuda），只可以使用其中某一套。
RUNTIME做了初始化工作，context management，module management；
COMMON
两套API都可以列举设备及其性能；一般应该一个线程对应一个device；
Mem可以分配为linear mem或者CUDA arrays：前者是32bit地址空间，可以自由使用指针；CUDA arrays的内存分配不透明，device只能够读取。
__global__和大多数runtime函数都是异步的，并不是执行完毕才返回。CudaThreadsSynchronize（）和cuCtxSynchronize（）强制runtime等待直到所有任务完毕。只有在host和device之间的内存拷贝，和OPENGL/DIRECT3D的协作，map/unmap，free mem这些操作不是异步的。
-----------
RUNTIME API
-----------
初始化：当第一次调用某条runtime函数时候，执行初始化工作；
Device Management：cudaGetDeviceCount; cudaGetDevicePorperties; cudaSetDevice用于为host线程选择一个对应的设备，它应当在所有__global__函数之前调用。
Memory Management：cudaMalloc; cudaMallocPitch用于二维情况; cudaFree;
CudaMallocArray(需要channel desc); cudaFreeArray;
CudaGetSymbolAddress；cudaGetSymbolSize；用于取得global mem的变量指针；
cudaMemCpy系列函数可以从host或device到device；
cudaMallocHost；cudaFreeHost用于page-locked的host内存，这样的内存在host和device之间交换数据时候会快些（比起malloc分配的分页内存）；
Texture Reference Management：cudaBindTexture；cudaUnbindTexture；
调试：emulation模式为block中的每一个thread建立一个host线程；但是它不是simulating，一些问题仍难以发现。比如host可能会保存float的计算中间量为更高精度；（可以试着声明变量为volatile强制令其单精度存储；编译选项；_controlfp）
----------
DRIVER API
----------
略
--------------------------------------------------------
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；不要试图在host和device之间共享内存指针，会出现内存段错误。
__device__ 变量存在于 device ，未使用其它 qualifier 情况下：
在 global mem 中；应用程序生命期；能够由每个线程，以及 host 读写；
__constant__ ：存在于 constant mem 中；应用程序生命期；能够由每个线程，以及 host 读写；
__shared__ ：存在于 block shared mem 中； block 生命周期；由 block 中的线程读写 ;
这些变量都不能用于结构；不能用于 host 上的参数，临时变量。
---------------
TEXTURE:
---------------
首先声明一个Texture Reference，类型，维数和读数据方式（8bit，16bit是否作为normalize float读入）;
运行时fetching：默认通过[0,N)作为索引取样，或可以normalize到[0,1)；addressing mode; filtering;
Texture也可以bind到device linear Mem 而不是 CUDA Array，但失去了运行时fecthing的多种选择；
Texture Bind到device 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
-------------------------------
__global__ void Func(float* parameter);
must be called like this:
Func<<< Dg, Db, Ns >>>(parameter);
分别代表 Dimension Grid ； Dimension Block ； Number 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 Mem：tex1Dfetch；
Texturing from CUDA Arrays：tex1D, tex2D；
Atomic Functions：read-modify-write on 32bit integers in global mem;
-------------------------
HOST RUNTIME COMPONENT：
-------------------------
两套API：DRIVER（cu）或者RUNTIME（cuda），只可以使用其中某一套。
RUNTIME做了初始化工作，context management，module management；
COMMON
两套API都可以列举设备及其性能；一般应该一个线程对应一个device；
Mem可以分配为linear mem或者CUDA arrays：前者是32bit地址空间，可以自由使用指针；CUDA arrays的内存分配不透明，device只能够读取。
__global__和大多数runtime函数都是异步的，并不是执行完毕才返回。CudaThreadsSynchronize（）和cuCtxSynchronize（）强制runtime等待直到所有任务完毕。只有在host和device之间的内存拷贝，和OPENGL/DIRECT3D的协作，map/unmap，free mem这些操作不是异步的。
-----------
RUNTIME API
-----------
初始化：当第一次调用某条runtime函数时候，执行初始化工作；
Device Management：cudaGetDeviceCount; cudaGetDevicePorperties; cudaSetDevice用于为host线程选择一个对应的设备，它应当在所有__global__函数之前调用。
Memory Management：cudaMalloc; cudaMallocPitch用于二维情况; cudaFree;
CudaMallocArray(需要channel desc); cudaFreeArray;
CudaGetSymbolAddress；cudaGetSymbolSize；用于取得global mem的变量指针；
cudaMemCpy系列函数可以从host或device到device；
cudaMallocHost；cudaFreeHost用于page-locked的host内存，这样的内存在host和device之间交换数据时候会快些（比起malloc分配的分页内存）；
Texture Reference Management：cudaBindTexture；cudaUnbindTexture；
调试：emulation模式为block中的每一个thread建立一个host线程；但是它不是simulating，一些问题仍难以发现。比如host可能会保存float的计算中间量为更高精度；（可以试着声明变量为volatile强制令其单精度存储；编译选项；_controlfp）
----------
DRIVER API
----------
略
转载于:https://blog.51cto.com/general/241857
1万+
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
