第一节
CUDA 让你可以一边使用熟悉的编程概念，一边开发可在GPU上运行的软件。
Rob Farber 是西北太平洋国家实验室（Pacific Northwest National Laboratory）的高级科研人员。他在多个国家级的实验室进行大型并行运算的研究，并且是几个新创企业的合伙人。大家可以发邮件到rmfarber@gmail.com与他沟通和交流。
您是否有兴趣在使用高级语言（比如C编程语言）编程时，通过标准多核处理器将性能提升几个数量级？您是否期待拥有跨多个设备的伸缩能力？
很多人（包括我自己）都通过使用NVIDIA的CUDA（Compute Unified DeviceArchitecture，即计算统一设备架构的简称）获得了这种高性能和可伸缩性，以编写廉价的多线程GPU程序。我特别强调“编程”是因为CUDA是为您的工作服务的架构，它不会“强迫”您的工作适应有限的一组性能库。使用CUDA，您可以发挥您的才能，设计软件以便在多线程硬件上获得最佳性能——并从中获得乐趣，因为计算正确的映射是很有意思的，而且软件开发环境十分合理和直观。
本文是这一系列文章的第一节，介绍了CUDA的功能（通过使用代码）和思维过程，帮助您将应用程序映射到多线程硬件（比如GPU）以获得较大的性能提升。当然，并不是所有问题都可以有效映射到多线程硬件，因此我会介绍哪些可以进行有效映射，哪些不能，而且让您对哪些映射可以运行良好有个常识性的了解。
“CUDA编程”和“GPGPU编程”并不相同（尽管CUDA运行在GPU上）。以前，为GPU编写软件意味着使用GPU语言编程。我的一个朋友曾将这一过程描述为将数据从您的肘部拉到眼前。CUDA允许使用熟悉的编程概念开发可以在GPU上运行的软件。通过将软件直接编译到硬件（例如，GPU汇编语言），可避免图形层API的性能开销，这样可以提供更出色的性能。
您可以任选一种CUDA设备。图1和图2分别显示了运行在一个笔记本和一个台式机的离散GPU上的CUDA 多体模拟（N-body simulation）程序。
图 1: 运行在使用了Quadro FX 570M的笔记本上的多体天体模拟程序。
图 2: 运行在使用了GeForce 8800 GTS 512MB的台式机上的多体天体模拟程序。
CUDA真的可以将应用程序性能提高一到两个数量级——或者这只是一种夸张，而非现实呢？
CUDA是一种相当新的技术，但是在一些书中和网络上，已经有很多样例突出介绍了这种技术，在使用当前商用GPU硬件时对性能的极大提升。图表1和表2总结了NVIDIA和BeckmanInstitute网站上相关内容。CUDA的核心是，让程序员能够使数千线程保持忙碌的工作状态。目前这一代NVIDIAGPU能够有效地支持大量线程，因此它们可以将应用程序性能提高一到两个数量级。这些图形处理器的价位区别很大，几乎所有的人都使用得起。较新的主板将通过提供更大的内存带宽、异步数据传输、原子操作和双精度浮点计算等多项硬件技术改进，扩展CUDA的功能。随着技术的不断进步，CUDA软件环境将不断扩展，最终GPU和“多核”处理器之间的区别也会逐渐消失。作为程序开发人员，我们可以预计，具有成千上万活动线程的应用程序将变得很常见而且CUDA将会运行在多个平台上，包括一般用途的处理器。
|
应用程序样例 |
URL |
应用程序加速 |
|
地震数据库（Seismic Database） |
66x到100x |
|
|
移动电话天线仿真（Mobile Phone Antenna Simulation） |
45x |
|
|
分子动态学（Molecular Dynamics） |
21x 到100x |
|
|
神经元仿真（Neuron Simulation） |
100x |
|
|
MRI 处理（MRI processing） |
245x 到 415x |
|
|
大气干扰运仿真（Atmospheric Cloud Simulation） |
50x |
表 1: NVIDIA总结，www.nvidia.com/object/IO_43499.html
|
GPU 性能结果, 2008年3月 GeForce8800GTX w/ CUDA 1.1, Driver 169.09 |
||
|
运算/算法 |
算法类 |
加速 vs. Intel QX6700 CPU |
|
微荧光光解法（Fluorescence microphotolysis） |
迭代矩阵（Iterative matrix） / 模板（stencil） |
12x |
|
对列表计算（Pairlist calculation） |
粒子对距离测试（Particle pair distance test） |
10x到11x |
|
对列表更新（Pairlist update） |
粒子对距离测试（Particle pair distance test） |
5x 到15x |
|
分子动态学非健合力运算（ Molecular dynamics nonbonded force calculation） |
多体断电力运算（N-body cutoff force calculations） |
10x 到20x |
|
断电电子密度量（Cutoff electron density sum）
|
粒子－网W/断电（Particle-grid w/ cutoff） |
15x 到23x |
|
断电潜能总结（Cutoff potential summation） |
粒子－网W/断电（Particle-grid w/ cutoff） |
12x 到21x |
|
直接库仑总结（Direct Coulomb summation） |
粒子－网（Particle-grid w/ cutoff） |
44x |
表 2: Beckman Institute表格，自 www.ks.uiuc.edu/Research/vmd/publications/siam2008vmdcuda.pdf
在20世纪80年代，我还是Los Alamos NationalLaboratory的科研人员，当时我有幸使用了拥有多达65,536个平行处理器的ThinkingMachines超级计算机。CUDA被证明是天生用于现代大量平行（即高线程）环境的框架。它的性能优势显而易见。我的一段生产代码，现在用CUDA编写并且运行在NVIDIA GPU上，与2.6-Ghz四核Opteron系统相比，具有明显的线形伸缩和几乎两个数量级的速度提升。
启用CUDA的图形处理器作为主计算机内的联合处理器运行。这意味着每个GPU都被认为有其自己的内存和处理元素，它们是与主计算机分开的。要进行有效的工作，数据必须在主计算机的内存空间和CUDA设备之间传输。因此，性能结果必须包括IO时间才更有意义。同事们喜欢将其称为“诚实的数据”，因为它们会更准确地反映将要交付生产的性能应用程序。
我坚持与现有技术相比，一到两个数量级的性能提升是一个巨变，可以在很大程度上改变运算的某些方面。例如，以前可能需要花费一年时间的运算任务现在只要几天就可以完成，几个小时的运算突然变得可交互了，因为使用新技术它们可以在几秒钟内完成，过去不易处理的实时处理任务现在变得极易处理。最后，它为具有正确技能集和能力的顾问和工程师们提供了良好的机会，使他们可以编写高线程（或大量平行）软件。对于您来说，这种计算能力又能给您的职业、应用程序或实时处理需求带来哪些好处呢？
开始不需要任何成本，您只需要从CUDA Zone主页下载CUDA（查找”获取CUDA”）。然后，按照您的特定的操作系统安装指导操作。您甚至不需要图形处理器，因为你可以直接使用软件模拟器，在你的笔记本或者工作站上运行，开始工作。当然，使用启动CUDA的GPU，可以获得更好的性能。或许你的计算机应该有一个这样的GPU了。在CUDAZone 主页上查看支持CUDA的GPU链接（支持CUDA的GPU包括共享的片上内存和线程管理）。
如果要购买一个新的图形处理器卡，我建议您依次阅读以下文章，因为我将探讨不同的硬件特性（如内存带宽、注册数量、原子操作等）将如何影响应用程序的性能。这样有助于您为应用程序选择恰当的硬件。另外，CUDA Zone论坛提供了关于CUDA各个方面的大量信息，包括购买哪些硬件。
安装完毕后，CUDA Toolkit将提供一个合理的C语言程序开发工具集，它包括：
- nvcc C编译器；
- GPU 的CUDA FFT和BLAS库
- 性能分析器
- alpha 版本（截至2008年3月）的GPU的gdb调试器
- CUDA运行时驱动程序（现在还可以在标准的NVIDIA GPU驱动程序中得到)
- CUDA编程手册
nvccC编译器完成了将C代码转换成将运行在GPU或模拟器上的可执行程序的大部分工作。幸好，汇编语言编程不要求达到很高的性能。下面的文章将介绍从其它高级语言，包括C++、FORTRAN和Python使用CUDA的内容。我假设您熟悉C/C++。不需要有平行编程或CUDA经验。这与现有CUDA文档是一致的。
创建和运行CUDA C语言程序与创建和运行其它C编程环境的工作流是一样的。面向Windows和Linux环境的明确构建和运行说明在CUDA文档中。简言之，这一工作流就是：
- 使用最喜欢的编辑器创建或编辑CUDA程序。注意：CUDA C 语言程序的后缀为.cu。
- 使用nvcc编译程序创建可执行程序（NVIDIA提供了带有示例的完整makefiles。通常用于CUDA设备时您只需键入make，用于模拟器时只需键入make emu=1）。
- 运行可执行程序。
表1是一个带您入门的简单CUDA程序。它只是一个简单的程序，调用CUDA API将数据移入和移出CUDA设备。并没有添加新内容，以免在学习如何使用工具构建和运行CUDA程序时发生混淆。在下一篇文章中，我将介绍如何开始使用CUDA设备执行一些工作。
1 // moveArrays.cu
2 //
3 // demonstrates CUDA interface to data allocation on device (GPU)
4 // and data movement between host (CPU) and device.
5
6
7 #include <stdio.h>
8 #include <assert.h>
9 #include <cuda.h>
10
11 int main(void)
12 {
13 float *a_h, *b_h; // pointers to host memory
14
15 float *a_d, *b_d; // pointers to device memory
16
17 int N = 14;
18 int i;
19 // allocate arrays on host
20
21 a_h = (float *)malloc(sizeof(float)*N);
22 b_h = (float *)malloc(sizeof(float)*N);
23 // allocate arrays on device
24
25 cudaMalloc((void **) &a_d, sizeof(float)*N);
26 cudaMalloc((void **) &b_d, sizeof(float)*N);
27 // initialize host data
28
29 for (i=0; i<N; i++) {
30 a_h = 10.f+i;
31 b_h = 0.f;
32 }
33 // send data from host to device: a_h to a_d
34
35 cudaMemcpy(a_d, a_h, sizeof(float)*N, cudaMemcpyHostToDevice);
36 // copy data within device: a_d to b_d
37
38 cudaMemcpy(b_d, a_d, sizeof(float)*N, cudaMemcpyDeviceToDevice);
39 // retrieve data from device: b_d to b_h
40
41 cudaMemcpy(b_h, b_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
42 // check result
43
44 for (i=0; i<N; i++)
45 assert(a_h == b_h);
46 // cleanup
47
48 free(a_h); free(b_h);
49 cudaFree(a_d); cudaFree(b_d);
试下这些开发工具吧。对初学者的一些建议：可以使用printf语句看看在模拟器下运行时（使用make emu=1构建可执行程序）GPU上会发生什么。还可以随意试验调试器的alpha版本。