CUDA（Compute Unified Device Architecture，统一计算架构[1]）是由NVIDIA所推出的一种集成技术，是该公司对于GPGPU的正式名称。通过这个技术，用户可利用NVIDIA的GeForce 8以后的GPU和较新的Quadro GPU进行计算。亦是首次可以利用GPU作为C-编译器的开发环境。NVIDIA营销的时候[2]，往往将编译器与架构混合推广，造成混乱。实际上，CUDA可以兼容OpenCL或者自家的C-编译器。无论是CUDA C-语言或是OpenCL，指令最终都会被驱动程序转换成PTX代码，交由显示核心计算。

优点

在GPUs（GPGPU）上使用图形APIs进行传统通用计算，CUDA技术有下列几个优点：
分散读取——代码可以从内存的任意地址读取
统一虚拟内存（CUDA 6）
共用内存——CUDA公开一个快速的共用存储区域（每个处理器48K），使之在多个进程之间共用。其作为一个用户管理的高速缓存，比使用纹理查找可以得到更大的有效带宽。
与GPU之间更快的下载与回读全面支持整型与位操作，包括整型纹理查找

限制

  • CUDA不支持完整的C语言标准。它在C++编译器上运行主机代码时，会使一些在C中合法（但在C++中不合法）的代码无法编译。
  • 不支持纹理渲染（CUDA 3.2及以后版本通过在CUDA数组中引入“表面写操作”——底层的不透明数据结构——来进行处理）
  • 受系统主线的带宽和延迟的影响，主机与设备内存之间数据复制可能会导致性能下降（通过过GPU的DMA引擎处理，异步内存传输可在一定范围内缓解此现象）
  • 当线程总数为数千时，线程应按至少32个一组来运行才能获得最佳效果。如果每组中的32个进程使用相同的运行路径，则程序分支不会显著影响效果；在处理本质上不同的任务时，SIMD运行模型将成为一个瓶颈（如在光线追踪算法中遍历一个空间分区的数据结构）
  • 与OpenCL不同，只有NVIDIA的GPUs支持CUDA技术
  • 由于编译器需要使用优化技术来利用有限的资源，即使合法的C/C++有时候也会被标记并中止编译
  • CUDA（计算能力1.x）使用一个不包含递归、函数指针的C语言子集，外加一些简单的扩展。而单个进程必须运行在多个不相交的内存空间上，这与其它C语言运行环境不同。
    CUDA（计算能力2.x）允许C++类功能的子集，如成员函数可以不是虚拟的（这个限制将在以后的某个版本中移除）
  • 双精度浮点（CUDA计算能力1.3及以上）与IEEE754标准有所差异：倒数、除法、平方根仅支持舍入到最近的偶数。单精确度中不支持反常值（denormal）及sNaN（signaling NaN）；只支持两种IEEE舍入模式（舍位与舍入到最近的偶数），这些在每条指令的基础上指定，而非控制字码；除法/平方根的精度比单精确度略低。[11]

应用

利用CUDA技术，配合适当的软件（例如MediaCoder、Freemake Video Converter），就可以利用显示核心进行高清视频编码加速。视频解码方面，同样可以利用CUDA技术实现。此前，NVIDIA的显示核心本身已集成PureVideo单元。可是，实现相关加速功能的一个微软API－DXVA，偶尔会有加速失效问题。所以利用CoreAVC配合CUDA，变相在显示核心上实现软件解码，解决兼容性问题。另外，配合适当的引擎，显示核心就可以计算光线跟踪。NVIDIA就放出了自家的Optix实时光线跟踪引擎，通过CUDA技术利用GPU计算光线跟踪。

CUDA C/C++语言迚行扩展后形成的变种。

主要包括以下几个方面：
• 函数类型限定符
• 变量类型限定符
• 执行配置运算符<<< , , , >>>
• 五个内建变量，
• 其它的还有数学函数，原子函数，纹理读取、绑定
函数，内建栅栏，存储器fence函数等。

函数限定符

用来确定函数是在CPU还是在GPU上执行，以及这个函数是从CPU调用还是从GPU调用。
• device表示从GPU上调用，在GPU上执行。
• global表示在CPU上调用，在GPU上执行，也就是所谓的内核(kernel)函数；在计算能力3.0及以后的设备上， global可调用global。
• host表明在CPU上调用，在CPU上执行，这是默认时的情况，也是传统的C/C++函数。

CUDA程序的一般模式

分配主机存储器并初始化
分配设备存储器
将已初始化的主机存储器内容复制到已分配的设备存储器上GPU进行计算
将计算完的结果从设备复制到主机上处理复制到主机上的数据

CUDA线程模式

网格（Grid）
• 一Kernel映射一网格
• 网格在设备上执行
• 划分为线程块
线程块（Block）
• 发射到SM上执行
• 利用共享存储器通信
• 划分为线程
线程（Thread）
• 映射到SP上执行

五个内建变量

运行时获得网格和块的尺寸及线程索引等信息
• gridDim：包含三个元素x,y,z的结构体，表示网格在
x,y,z方向上的尺寸，对应于执行配置中的第一个参数。
在计算能力2.0的设备上只支持其前两维
• blockDim：包含三个元素x,y,z的结构体，表示块在
x,y,z方向上的尺寸，对应于执行配置的第二个参数
• blockIdx：包含三个元素x,y,z的结构体，分别表示当前
线程所在块在网格中x,y,z方向上的索引
• threadIdx：包含三个元素x,y,z的结构体，分别表示当
前线程在其所在块中x,y,z方向上的索引
• warpSize：表明warp的尺寸，在计算能力1.0的设备中，
这个值是24，在1.0以上的设备中，这个值是32。

示例

用C++自GPU的image数组中获取纹理（texture）

cudaArray* cu_array;
texture<float, 2> tex;

// Allocate array
cudaChannelFormatDesc description = cudaCreateChannelDesc<float>();
cudaMallocArray(&cu_array, &description, width, height);

// Copy image data to array
cudaMemcpy(cu_array, image, width*height*sizeof(float), cudaMemcpyHostToDevice);

// Bind the array to the texture
cudaBindTextureToArray(tex, cu_array);

// Run kernel
dim3 blockDim(16, 16, 1);
dim3 gridDim(width / blockDim.x, height / blockDim.y, 1);
kernel<<< gridDim, blockDim, 0 >>>(d_odata, height, width);
cudaUnbindTexture (tex);

__global__ void kernel(float* odata, int height, int width)
{
   unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
   unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
   float c = tex2D(tex, x, y);
   odata[y*width+x] = c;
}