第十节：CUDPP, 强大的数据平行CUDA库
Rob Farber 是西北太平洋国家实验室（Pacific Northwest National Laboratory）的高级科研人员。他在多个国家级的实验室进行大型并行运算的研究，并且是几个新创企业的合伙人。大家可以发邮件到[email protected]与他沟通和交流。
在关于CUDA（Compute Unified Device Architecture，即计算统一设备架构的简称）的系列文章的第9节，我探讨了如何使用CUDA拓展高等级语言（如Python）。在本小节中，我 要讨论CUDPP，也就是CUDA Data Parallel Primitives Library数据平行基元库。CUDPP是一个快速成熟起来的包，执行一些不是那么明显的算法以有效使用GPU进行基本的数据平行运算，如归类，流数据 结实化，和构建数据结构如树和求和面积表。我在这里探讨CUDPP，是因为它可能提供一些所需的功能以加速你的某个项目的发展。
我还介绍了创建“PLAN计划”的概念，用这个编程格式提供一个基于问题说明和目地硬件
的 优化执行配置。尽管不是一个优化编译器，对计划PLAN的使用可以极大增强程序员的能力，为CUDA启动的多个类型的GPU创建有效软件――此外，还可在 通用库框架范围之内，提供能力为特定的问题选定特定的优化代码。例如，NVIDIA cuFFT library在恰当时候可决定采用两个FFT算法中更为有效的一种。PLAN对CUDA或这一系列文章而言不是全新的概念了，它是一个常用的设计格式， 经受得住时间考验。
为什么使用CUDPP?
大部分人有一个库和方法的工具箱，我们用来进行一个操作。 简而言之，这些库提供了基本指令，我们可以用来快捷地执行一些计算任务。归类只是个样例，简易有效，可调用如qsort() 例行程序，以排序次序返回数据结构。NVIDIA cuBLAS和cuFFT库为一些不那么简单的任务提供相似的功能，如编程EFT和优化的BLAS功能。
CUDPP使用相同的理念提供优化的“同类中最好的“方法以执行基元运算，如平行前缀求和（扫描），平行排序（数字），平行减少和其它能够有效执行稀疏矩阵－向量相乘的方法，和其它运算。
平 行前缀扫描是一个基元，帮助有效解决平行问题，在这些平行问题里，每个输出都要求对输入有全局认识。例如，前缀求和（也就是扫描，前缀减少，或部分和）是 对列表进行的运算，每个结果列表里的元素都是，根据索引，把操作数列表的元素相加获得的。这看上去就是一个串行操作，因为每个结果都取决于所有之前的值， 如下所示：
定义：所有前缀求和运算取二进制关联运算符和n元素数组
given: [a0, a1, ..., an-1],returns:[a0, (a0 a1), ..., (a0 a1 ... ( an-1)].
样例: 如 是相加，那么在n元素数组上进行所有前缀求和运算
given [3, 1, 7, 0, 4, 1, 6, 3]returns [3, 4, 11, 11, 15, 16, 22, 25].
上 述所有前缀求和有许多用法，包括，但是不限于，归类，词法分析，字符串比较，多项式求值，流数据结实化和构建平行柱状图和数据结构（图形，树等）。不同类 型的调查文章提供了更多广泛的，详细的应用，如Guy Blelloch的 Prefix Sums and Their Applications。
显然，扫描的依序构建（可在CPU上单线程运行）是微不足道的。我们在输入数组里循环所有的元素，然后将数组的前元素值添加到输出数组的前元素的总和里，然后把和值写入输出数组的现有元素里。
1 void scan( float* output, float* input, int length)
2 {
3 output[0] = 0; // since this is a prescan, not a scan
4 for(int j = 1; j < length; ++j) {
5 output[j] = input[j-1] + output[j-1];
6 }
7 }
代码为长度为n的数组执行n次加法运算－n是所需加法的最小数，以生成扫描数组。如果扫描的平行版有效率，这就意味着相对于序列版来说，平行版没有 执行更多的加法运算（或操作）。换言之，两个执行都会有相同的复杂度，O(n)。CUDPP 声称要获得O(n)扫描运行时间，以澄清CUDPP的值，因为创建平行执行并非无关重要的。如需了解更多信息，请参与Shubhabrata Sengupta et al的Scan Primitives for GPU Computing。
CUDDP为1.0版提供：
分段扫描－平行执行多个可变长度扫描的算法。对算法如平行快速排序，平行稀疏矩阵－向量相乘和其它算法都有用；
平行稀疏矩阵－向量相乘（基于分段扫描）。稀疏矩阵操作很重要，因为这样GPU可以在空间内，以高效的计算方式，使用多个零（如稀疏矩阵）在矩阵上工作。因为大部分值都是零，大部分工作都可避免。类似地，也没必要浪费空间储存零。
改进的扫描算法，称为“warp扫描”，其性能提高，但代码更简单；
扫描和分段扫描现在支持加法，相乘，最大和最小运算子；
支持包容性扫描和分段扫描；
改进的，更为有用的cudppCompact()接口；
支持.反向协议；
支持CUDA2.0；
给予Mac OS X和Windows Vista更多支持
Nt 线程配置和GPU资源，再加上PLAN是存储和重新使用这些配置的便捷方法。此外，cuFFT优化也可以被采用，不过取决于所请求的FFT是不是二次方 幂。广泛应用的FFTW项目也利用了PLAN概念。FFTW广泛应用于各类平台。因为诸多原因，当开发需要在大量GPU构架上运行的通用解决方案 时，PLANS是一个有用的工具。
这个简单的CUDPP示例需要创建一个PLAN，以对目标GPU上的numElements元素进行正向专用浮 动求和扫描。通过填写CUDPPConfiguration struct，将之传递给规划，得以完成该任务。在本例中，告知规划算法(CUDPP_SCAN), 数据类型(CUDPP_FLOAT), 运算 (CUDPP_ADD), 和选项 (CUDPP_OPTION_FORWARD, CUDPP_OPTION_EXCLUSIVE)。方法cudppPlan和该配置被调用（配有最大数量的元素进行扫描numElements）。最后， 规划被告知我们仅希望通过传递1和0，为numRows和 rowPitch参数扫描单维数组。CUDPP文件向cudppPlan()提供了更多有关参数的细节。
CUDPPConfiguration config;
config.op = CUDPP_ADD;
config.datatype = CUDPP_FLOAT;
config.algorithm = CUDPP_SCAN;
config.options = CUDPP_OPTION_FORWARD | CUDPP_OPTION_EXCLUSIVE;
CUDPPHandle scanplan = 0;
CUDPPResult result = cudppPlan(&scanplan, config, numElements, 1, 0);
if (CUDPP_SUCCESS != result)
{
printf("Error creating CUDPPPlan\n");
exit(-1);
}
调用cudppPlan成功后，向计划对象inscanplan返回句柄（指针）。CUDPP然后开始运行，调用cudppScan(), 规划句柄，输出和输入设备数组，以及需扫描元素数都被传递给它。
1 // Run the scan 2 cudppScan (scanplan, d_odata, d_idata, numElements);
然后, 使用cudaMemcpy将扫描结果从d_odata复制回主机。通过计算CPU上的参考解决方案（通过computeSumScanGold())，验证GPU结果，然后对比CPU和GPU结果以确保正确性。
1 // allocate mem for the result on host side 2 float* h_odata = (float*) malloc( memSize); 3 // copy result from device to host 4 CUDA_SAFE_CALL( cudaMemcpy( h_odata, d_odata, memSize, 5 cudaMemcpyDeviceToHost) ); 6 // compute reference solution 7 float* reference = (float*) malloc( memSize); 8 computeSumScanGold( reference, h_idata, numElements, config); 9 // check result 10 CUTBoolean res = cutComparef( reference, h_odata, numElements); 11 printf( "Test %s\n", (1 == res) ? "PASSED" : "FAILED");
最后, 调用cudppDestroyPlan() 清理内存空间。主机然后使用using free() 和 cudaFree分别释放本地和设备数组，退出应用程序，因为simpleCUDPP已经完毕。
1 result = cudppDestroyPlan (scanplan);
2 if (CUDPP_SUCCESS != result)
3 {
4 printf("Error destroying CUDPPPlan\n");
5 exit(-1);
6 }
稀疏矩阵－向量相乘
CUDPP包含许多其它强大的功能，但是本文未做讨论。例如，使用CUDPP进行稀疏矩阵－向量相乘的简单测试代码是sptest.cu。从http://www.nada.kth.se/~tomaso/gpu08/sptest.cu.可进行下载。你可使用以下代码编译和运行：
# nvcc -I cudpp_1.0a/cudpp/include -o sptest sptest.cu -L cudpp_1.0a/lib -lcudpp # ./sptest