本页使用了标题或全文手工转换

CUDA

维基百科,自由的百科全书
跳到导航 跳到搜索
CUDA
开发者 Nvidia
初始版本 2007年6月23日,​11年前​(2007-06-23
稳定版本
稳定版本
9.2
(2018年5月16日,​4个月前​(2018-05-16
操作系统 Windows XP以上,
macOS, Linux
类型 GPGPU
许可协议 免费软件
网站 developer.nvidia.com/cuda-zone

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

概要[编辑]

CUDA处理流例子
1.将主存的处理数据复制入显存中
2. CPU指令驱动GPU
3. GPU每一核心并行处理
4. GPU将显存结果传回主存

GeForce 8800 GTX为例,其核心拥有128个内处理器。利用CUDA技术,就可以将那些内处理器串通起来,成为线程处理器去解决数据密集的计算。而各个内处理器能够交换、同步和共享数据。利用NVIDIA的C-编译器,通过驱动程序,就能利用这些功能。亦能成为流处理器,让应用程序利用进行运算。

GeForce 8800 GTX显示卡的运算能力可达到520GFlops,如果建设SLI系统,就可以达到1TFlops。[4]

但程序员在利用CUDA技术时,须分开三种不同的存储器,要面对繁复的线程层次,编译器亦无法自动完成多数任务,以上问题就提高了开发难度。而将来的G100会采用第二代的CUDA技术,提高效率,降低开发难度。

目前,已有软件厂商利用CUDA技术,研发了一个Adobe Premiere Pro的插件。通过插件,用户就可以利用显示核心去加速H.264/MPEG-4 AVC的编码速度。速度是单纯利用CPU作软件加速的7倍左右。

在NVIDIA收购AGEIA后,NVIDIA获取相关的物理加速技术,即是PhysX物理引擎。配合CUDA技术,显卡可以模拟成一颗PhysX物理加速芯片[5]。目前,全系列的GeForce 8显示核心都支持CUDA。而NVIDIA亦不会再推出任何的物理加速卡,显卡将会取代相关产品。

为了将CUDA推向民用,NVIDIA会举行一系列的编程比赛,要求参赛者开发程序,充分利用CUDA的计算潜能。但是,要将GPGPU普及化,还要看微软能否在Windows操作系统中,提供相关的编程接口。[6]

2008年8月,NVIDIA推出CUDA 2.0[7]。2010年3月22日,NVIDIA推出CUDA 3.0,仅支持Fermi及之后的架构[8]

CUDA是一种由NVIDIA提出的并由其制造的图形处理单元(GPUs)实现的一种并行计算平台及程序设计模型。CUDA给程序开发人员提供了直接访问CUDA GPUs中的虚拟指令集和并行计算组件的存储器。

使用CUDA技术,GPUs可以用来进行通用处理(不仅仅是图形);这种方法被称为GPGPU。与CPUs不同的是,GPUs有着侧重以较慢速度运行大量并发线程的并发流架构,而非快速运行单一线程。

软件发展者可以通过CUDA加速库,编译器指令(如OpenACC)以及匹配工业标准的程序设计语言(如C,C++和Fortran)扩展对CUDA平台进行操作。C/C++程序师可以使用“CUDA C/C++”,使用“NVCC”——NVIDIA基于LLVM的C/C++编译器进行编译;Fortran程序师可以使用“CUDA Fortran”,使用PGI公司的PGI CUDA Fortran编译器进行编译。除了库、编译器指令、CUDA C/C++和CUDA Fortran,CUDA平台还支持其它计算接口,如Khronos Group的OpenCL,Microsoft的DirectCompute,以及C++AMP。其协力厂商封装也可用于Python,Perl,Fortran,Java,Ruby,Lua,Haskell,MATLAB,IDL及Mathematica的原生支持。

在电脑游戏行业中,GPUs不仅用于进行图形渲染,而且用于游戏物理运算(物理效果如碎片、烟、火、流体),比如PhysX和Bullet。在计算生物学与密码学等领域的非图形应用上,CUDA的加速效果达到了可以用数量级来表示的程度。

CUDA同时提供底层API与高阶API。最初的CUDA软件发展包(SDK)于2007年2月15日公布,支持Microsoft Windows和Linux。而后在第二版中加入了对Mac OS X的支持,取代了2008年2月14日发布的测试版。所有G8x系列及以后的NVIDIA GPUs皆支持CUDA技术,包括GeForce,Quadro和Tesla系列。CUDA与大多数标准操作系统兼容。Nvidia声明:根据二进制兼容性,基于G8x系列开发的程序无需修改即可在未来所有的Nvidia显卡上运行。

优点[编辑]

在GPUs(GPGPU)上使用图形APIs进行传统通用计算,CUDA技术有下列几个优点:

  • 分散读取——代码可以从存储器的任意地址读取
  • 统一虚拟内存(CUDA 4)
  • 共用存储器——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 C程序设计指南3.1》-附录D.6]
  • 双精度浮点(CUDA计算能力1.3及以上)与IEEE754标准有所差异:倒数、除法、平方根仅支持舍入到最近的偶数。单精确度中不支持反常值(denormal)及sNaN(signaling NaN);只支持两种IEEE舍入模式(舍位与舍入到最近的偶数),这些在每条指令的基础上指定,而非控制字码;除法/平方根的精度比单精确度略低。

应用[编辑]

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

支持的产品[编辑]

所有基于G80及之后架构的民用与专业显卡或运算模块皆支持CUDA技术[12]

示例[编辑]

下列的示例是如何用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;
}

下列的例子是用Python改写. Python相关的消息可取自PyCUDA.

import pycuda.driver as drv
import numpy
import pycuda.autoinit

mod = drv.SourceModule("""
__global__ void multiply_them(float *dest, float *a, float *b)
{
  const int i = threadIdx.x;
  dest[i] = a[i] * b[i];
}
""")

multiply_them = mod.get_function("multiply_them")

a = numpy.random.randn (400).astype(numpy.float32)
b = numpy.random.randn (400).astype(numpy.float32)

dest = numpy.zeros_like (a)
multiply_them(
        drv.Out(dest), drv.In(a), drv.In(b),
        block=(400,1,1))

print dest-a*b

更多的Python的矩阵相乘问题可取自pycublas.

import numpy
from pycublas import CUBLASMatrix
A = CUBLASMatrix(numpy.mat([[1,2,3],[4,5,6]],numpy.float32))
B = CUBLASMatrix(numpy.mat([[2,3],[4,5],[6,7]],numpy.float32))
C = A*B
print C.np_mat()

参考文献[编辑]

相关条目[编辑]

外部链接[编辑]