随着多核CPU和众核GPU的到来，并行编程已经得到了业界越来越多的重视，CPU-GPU异构程序能够极大提高现有计算机系统的运算性能，对于科学计算等运算密集型程序有着非常重要的意义。
Jetson tk1 GK20a GPU中拥有192个CUDA核（单独的ALU），因此非常适合并行计算，而且数值计算的速度远远优于CPU。
CUDA是一个完整的GPGPU解决方案，提供了硬件的直接访问接口，而不必像传统方式一样必须依赖图形API接口（OpenGL和Direct 3D）来实现GPU的访问。
在架构上采用了一种全新的计算体系结构来使用GPU提供的硬件资源，从而给大规模的数据计算应用提供了一种比CPU更加强大的计算能力。
CUDA采用C语言作为编程语言提供大量的高性能计算指令开发能力，使开发者能够在GPU的强大计算能力的基础上建立起一种效率更高的密集数据计算解决方案。
从CUDA体系结构的组成来说，包含了三个部分：开发库、运行期环境和驱动。
开发库是基于CUDA技术所提供的应用开发库。
目前CUDA提供了两个标准的数学运算库——CUFFT（离散快速傅立叶变换）和CUBLAS（离散基本线性计算）的实现。
这两个数学运算库所解决的是典型的大规模的并行计算问题，也是在密集数据计算中非常常见的计算类型。
开发人员在开发库的基础上可以快速、方便的建立起自己的计算应用。
运行期环境提供了应用开发接口和运行期组件，包括基本数据类型的定义和各类计算、类型转换、内存管理、设备访问和执行调度等函数。
基于CUDA开发的程序代码在实际执行中分为两种，一种是运行在CPU上的宿主代码（Host Code），一种是运行在GPU上的设备代码（Device Code）。
不同类型的代码由于其运行的物理位置不同，能够访问到的资源不同，因此对应的运行期组件也分为公共组件、宿主组件和设备组件三个部分，基本上囊括了所有在GPGPU开发中所需要的功能和能够使用到的资源接口，开发人员可以通过运行期环境的编程接口实现各种类型的计算。
由于目前存在着多种GPU版本的NVidia显卡，不同版本的GPU之间都有不同的差异，因此驱动部分基本上可以理解为是CUDA-enable的GPU的设备抽象层，提供硬件设备的抽象访问接口。
CUDA提供运行期环境也是通过这一层来实现各种功能的。
目前基于CUDA 开发的应用必须有NVIDIA CUDA-enable的硬件支持。
一个.cu文件内既包含CPU程序（称为主机程序），也包含GPU程序（称为设备程序）。
如何区分主机程序和设备程序？根据声明，凡是挂有“__global__”或者“__device__”前缀的函数，都是在GPU上运行的设备程序，不同的是__global__设备程序可被主机程序调用，而__device__设备程序则只能被设备程序调用。
没有挂任何前缀的函数，都是主机程序。
主机程序显示声明可以用__host__前缀。
设备程序需要由NVCC进行编译，而主机程序只需要由主机编译器（如VS2008中的cl.exe，Linux上的GCC）。
主机程序主要完成设备环境初始化，数据传输等必备过程，设备程序只负责计算。
主机程序中，有一些“cuda”打头的函数，这些都是CUDA Runtime API，即运行时函数，主要负责完成设备的初始化、内存分配、内存拷贝等任务。
我们前面第三节用到的函数cudaGetDeviceCount()，cudaGetDeviceProperties()，cudaSetDevice()都是运行时API。
线程并行将线程的概念引申到CUDA程序设计中，我们可以认为线程就是执行CUDA 程序的最小单元，在GPU上每个线程都会运行一次该核函数。
但GPU上的线程调度方式与CPU有很大不同。
CPU上会有优先级分配，从高到低，同样优先级的可以采用时间片轮转法实现线程调度。
GPU上线程没有优先级概念，所有线程机会均等，线程状态只有等待资源和执行两种状态，如果资源未就绪，那么就等待；一旦就绪，立即执行。
当GPU资源很充裕时，所有线程都是并发执行的，这样加速效果很接近理论加速比；而GPU资源少于总线程个数时，有一部分线程就会等待前面执行的线程释放资源，从而变为串行化执行。
块并行块并行相当于操作系统中多进程的情况，CUDA有线程组（线程块）的概念，将一组线程组织到一起，共同分配一部分资源，然后内部调度执行。
线程块与线程块之间，毫无瓜葛。
这有利于做更粗粒度的并行。
我们的任务有时可以采用分治法，将一个大问题分解为几个小规模问题，将这些小规模问题分别用一个线程块实现，线程块内可以采用细粒度的线程并行，而块之间为粗粒度并行，这样可以充分利用硬件资源，降低线程并行的计算复杂度。
线程块、线程组织图如下所示。
多个线程块组织成了一个Grid，称为线程格（经历了从一位线程，二维线程块到三维线程格的过程)
流并行流可以实现在一个设备上运行多个核函数。
前面的块并行也好，线程并行也好，运行的核函数都是相同的（代码一样，传递参数也一样）。
而流并行，可以执行不同的核函数，也可以实现对同一个核函数传递不同的参数，实现任务级别的并行。
CUDA中的流用cudaStream_t类型实现，用到的API有以下几个：cudaStreamCreate(cudaStream_t * s)用于创建流，cudaStreamDestroy(cudaStream_t s)用于销毁流，cudaStreamSynchronize()用于单个流同步，cudaDeviceSynchronize()用于整
个设备上的所有流同步，cudaStreamQuery()用于查询一个流的任务是否已经完成。
前面介绍了三种利用GPU实现并行处理的方式：线程并行，块并行和流并行。
在这些方法中，各个线程所进行的处理是互不相关的，即两个线程不回产生交集，每个线程都只关注自己的一亩三分地，对其他线程毫无兴趣，就当不存在。
当然，实际应用中，这样的例子太少了，也就是遇到向量相加、向量对应点乘这类才会有如此高的并行度，而其他一些应用，如一组数求和，求最大（小）值，各个线程不再是相互独立的，而是产生一定关联，线程2可能会用到线程1的结果，这时就需要利用线程通信技术了。
线程通信在CUDA中有三种实现方式：
1. 共享存储器；
2. 线程同步；
3. 原子操作；
注意的是，位于同一个Block中的线程才能实现通信，不同Block中的线程不能通过共享内存、同步进行通信，而应采用原子操作或主机介入。
CUDA的架构：在CUDA架构下，线程的最小单元是thread，多个thread组成一个block，多个block再组成一个grid，不同block之间的thread不能读写同一shared memory共享内存，因此，block里面的thread之间的通信和同步所带来的开销是比较大的。
SM以32个Thread 为一组的Warp 来执行Thread。
Warp内的线程是静态的，即在属于同一个warp 内的thread之间进行通信，不需要进行栅栏同步(barrier)。
Fermi的设计根据G80和GT200的架构作出的很多缺陷来改变。
在Fermi中，每个SM中的数量不再是GT200的8个SP，而是变成了32个SP，NVIDIA现在又称之为CUDA Core，总共具有16个SM，所以总共有512个SP。
而在GT200中，是30个SM，240个SP。