摩尔的预言唯有CUDA才是终极的CPU(三)

最新推荐文章于 2024-02-08 21:09:56 发布

cllm33533

最新推荐文章于 2024-02-08 21:09:56 发布

阅读量240

点赞数

文章标签： python 操作系统 c/c++

架构解析：从软件角度看CUDA

从软件的角度来看，CUDA是由C语言的扩展组成，这种架构当然要追溯到BrookGPU的时代，并且许多API的编写工作，都是从BrookGPU继承来的。CUDA扩展和延伸了许多BrookGPU的函数和变量。这里有一个非常典型的关键字__global__，它是一个前缀，在它后面的内容意味着是一个内核。这个功能将被CPU和执行单元所调用。__device__，标记有这个关键词的程序将会被GPU执行。但是仅仅在GPU中被调用。最后再介绍一个关键词__host__，这个前缀是可选的，它会指派CPU调用一个函数，并且由CPU执行它，换言之，它就像是一个非常传统的CPU执行函数。小熊在线www.beareyes.com.cn

对于__device__ 和 __global__ 函数来说，也有一些限制条件。他们不能被递归，也就是说他们不能自己调用自己。他们不能拥有可变数目的自变量。最后，要说一下__device__关键词，它必须驻留在GPU的显存空间内。理论上来说，绝对不可能获得它的物理地址。变量在控制存储区域的时候，也会有一些限制。__shared__ 是一个变量的前缀，当程序中出现这个关键词的时候，就表明变量要保存在流多重处理器中的共享存储区中。这与通过 __global__关键词调用会有些不同。这是因为执行配置中规定，在调用时，要指明grid栅格容器的大小，并指明它在哪个内核中，并且要为每一个block块指派所包含线程的尺寸。我们可以看看下面的这个例子：

__global__ void Func(float* parameter);

which will be called as follows:

Func<<< Dg, Db >>>(parameter);

其中Dg是grid栅格的尺寸，而Db是定义block块，这两个变量是CUDA的一种新的类型。小熊在线www.beareyes.com.cn

CUDA API其本质上来讲是由各种操作显存的函数组成的。cudaMalloc用来分配内存，cudaFree用来释放内存，cudaMemcpy用来互相拷贝内存和显存之间的数据。小熊在线www.beareyes.com.cn

在最后，我们将介绍一下CUDA程序的编译方式。这是非常有趣的，整个编译过程需要几个阶段。首先，所有的代码都要让CPU来处理，这些都会从文件中提取，并且他们都会通过标准的编译器。用于GPU处理的代码，首先要转换成中间媒介性语言——PTX。中间语言更像是一种汇编程序，并且能够中和潜在的无效代码。在最后的阶段，中间语言会转换成指令。这些指令会被GPU所认同，并且会以二进制的形式被执行。小熊在线www.beareyes.com.cn

PTX中间媒介语言

检验真理：实践CUDA

一旦你仔细阅读过NVIDIA的文档，你就会对CUDA的效能产生深刻的印象。有什么比自己亲自动手尝试写一个CUDA程序更让人兴奋的事儿呢？即使一切理论，一切技术，看起来都是那么的完美，但是实践起来，所有让人头疼的问题也就会立刻显现出来。如果你想自己动手尝试写一个CUDA程序，或者想更深入的了解它，那么阅读官方的帮助文档是最有效的手段。小熊在线www.beareyes.com.cn

事实上，着手写一个小项目，其实也挺简单的。网络上有许多高品质的免费工具可以参考使用。比如Visual C++ Express 2005，它具有非常完善的开发环境，几乎可以满足我们所有的需要。小熊在线www.beareyes.com.cn

最难的部分就是找一个比较简单的题材，让程序编写和实现起来并不那么十分的困难，同时又能充分发挥出CUDA的并行处理效能。小熊在线www.beareyes.com.cn

最终我们采用了一个代码块，它可以截取一个地图的高度值，并且通过相应的计算输出成普通的地图。由于编程技术有限，我们就不加入一些繁杂的细节功能了。我们只是安排了一些简单的循环：处理每一个地图上的像素，并且我们采用了矩阵式的算法，比较相邻的像素的颜色信息，使用一些稍微复杂一些的公式，来生成像素的颜色值。这类程序基本上可以充分发挥出CUDA的并行处理优势，下面我们就来做一个CUDA运算效能的简单测试。小熊在线www.beareyes.com.cn

我们的测试软件分为CPU和CUDA两个版本。两个程序算法之间相差并不大，我们并没有针对CPU架构重新开发新的算法。小熊在线www.beareyes.com.cn

测试图

最后咱们先别着急比较CPU版本和CUDA版本之间的测试结果，还是了解一下手头的编程工具CUDA SDK。由于我们第一次尝试编写CUDA的程序，因此我们并没有抱很高的期望。并且我们在编程方面也不够专业，在核心算法部分的代码也没有针对CPU做一些适当的优化。所以直接比较他们的结果，会让我们非常的震撼。

意料之中：惊人的测试结果

不过，我们并没有处理的措施，也没有处理的时间，只能任凭程序编写完成后，看看CUDA是否真的有很大的计算优势。程序的测试和开发，都是在一台笔记本电脑上完成的。它的配置为Core 2 Duo T5450处理器和一颗GeForce 8600M GT的显卡。操作系统为Vista。这台笔记本的配置并不算高端，远远没有现在的高端台式机速度快。但是测试结果非常有趣，因为这台笔记本的GPU处理器能力远远不能与NVIDIA的主打高端GPU产品相比，没有夸张的散热片，没有GTX 280那样的强悍配置，但是这颗仅集成有70M晶体管的CUDA GPU，却显示出了惊人的能力。这与我们预想的测试结果相吻合。小熊在线www.beareyes.com.cn

采用分辨率为2048x2048像素的图片，测试结果如下

CPU 1 thread: 1419 ms
CPU 2 threads: 749 ms
CPU 4 threads: 593 ms

GPU (8600M GT) blocks of 256 pixels: 109 ms
GPU (8600M GT) blocks of 128 pixels: 94 ms
GPU (8800 GTX) blocks of 128 pixels / 256 pixels: 31 ms

从测试结果来看，CUDA GPU的处理效能确实非常的强大。不过细心的读者也会发现，懒惰的程序员并没有针对CPU的线程进行优化，没有修改CPU内部的执行线程。最理想的测试环境，是处理器内部运行的线程数量同CUDA一样多。不过CPU内部仅有两个核心，我们只能运行2个独立的线程，他们的运行效能是非常强大的。最后，公平的说，四个线程的版本，确实出乎意料的快。而实际的测试成绩和我们脑海中所期待的，没有什么差别。小熊在线www.beareyes.com.cn

理论上来说，由于处理器创造了额外的线程，所以处理的效能会稍稍损失一些。但现实恰恰相反如何解释这种情况呢，也许是跟Windows的线程调度程序有关，多线程使得处理器始终处在忙碌的状态下，充分的利用了处理器的资源。小熊在线www.beareyes.com.cn

在任何情况下，结果都是线程的数量越多越好。当处理更小的材质规格时（512X512以下），处理效能提升越显著，并且在四个线程版本中，这种情况更加明显。不过GPU的处理效能仍然要比CPU快的多，8600GT要比双线程版本的运算效能快3倍以上。小熊在线www.beareyes.com.cn

CUDA工具

另一个让人吃惊的测试结果就是：即使是最慢的GPU运算结果，也是最快的CPU运算结果的6倍。对于一个未经人事的CUDA处女作程序来说，这个结果相当挑逗人。同时也请各位读者注意，如果想得到更好的测试结果，那么我们就要使用更小的blocks块，反之也是如此。小熊在线www.beareyes.com.cn

这里我来做一个简单的解释，我们的程序每个线程使用了14个寄存器，并且每个blocks块中封装了256个线程。由此，每个blocks块至少需要3584个寄存器，在一个多重处理器中，768个线程就已经饱和了。小熊在线www.beareyes.com.cn

在这种情况下，三个blocks块将会占用10572个寄存器。但是多重处理器仅有8192个寄存器，在整个CUDA的流水线中，也仅仅能跑动2个blocks块。小熊在线www.beareyes.com.cn

稍微改动一下程序，一个blocks块包含128个线程，每个线程处理1个像素。那么每个blocks块我们只需需要1792个寄存器。如果用总过8192个寄存器，除以1792的话，那么可以看出我们可以在系统中同时跑4个blocks块。小熊在线www.beareyes.com.cn

在实践中，多重处理器的最佳的线程数为512个，而NVIDIA官方声称理论上最大为768。但是在GPU中执行的许多blocks块通过访问存储区进行数据通信就会更加的灵活，由此就不仅仅受限于寄存器的个数了。当操作更长的指令时会有延迟。系统能发布延迟指令，让另一个blocks块等待处理的结果，直到有可用的数据结果为止。四个blocks块的确会使系统具备更好的潜伏期，特别是当我们的程序造访若干个存储区的时候。

结果分析：CUDA还可以跑的更快

最后，尽管我们刚才说这个简单的CUDA程序测试并不是一次赛马。但是我们仍然无法抗拒8800GTX的诱惑力。我们还是忍不住在8800GTX上运行了这个程序，结果更是异常的惊人。8800GTX的通用计算结果，相当于移动版8600GT的三倍。而且测试的环境非常公平，并没有针对任何GPU做blocks块优化。小熊在线www.beareyes.com.cn

有贪心的读者也许并不满意这个结果，可能在期盼他们之间更大的性能差距，比如说4倍甚至更高。因为8800GTX具备128个ALU，要比移动版8600GT的32个ALU高出很多，并且他们的时钟频率也有较大的差异，8800GTX的核心频率为1.35GHz，而移动版8600GT仅有950MHz。但是在实际跑程序的时候，他们的性能差距并没有理论设想的那么大。小熊在线www.beareyes.com.cn

小编想，这里最有可能的一个假说就是：受限于显存的访问速度。更确切的说，对于最初的图片文件访问，像是在一个CUDA的多维数组阵列。这个名词听起来真是非常的深奥和复杂，其实只是一个纹理存取和运算的方法。它有着以下几个优势：

1、可以访问速度更快的纹理缓存

2、与CPU的数据操作方式不同，我们可以将数据封装在一个包里，不必用多边形处理图像文件。小熊在线www.beareyes.com.cn

由此，我们也可以利用免费的像素过滤机制给像素添加地址，例如在[0,1]的区间之内。而不必使用[0,宽度],[0,高度]这样冗长的数据格式。作为一个小熊在线忠实的读者，长期关注显卡频道的网友们都会知道，移动版8600具备16个材质处理单元，而8800GTX具备32个材质处理单元。因此他们两种体系架构之间有将近一倍的规格差距。并且由于两块显卡的核心频率不同，他们之间的性能差距也会拉大:

(32 x 0.575) / (16 x 0.475) = 2.4

这个比率接近3倍，这与我们实际的测试结果相吻合。这也就是为什么虽然他们的规格相差数倍，但是测试结果他们仅仅相差3倍。归根结底，都是由于G80材质处理单元中的ALU，成为了整个CUDA运算的瓶颈。与线程的数量和blocks块的设置并无太大关系。

性能分析工具

另外，还有更加令人鼓舞的成果。因为这是我们第一次编写CUDA程序，但是我们的运行环境并不是十分的理想。首先我们使用的是一台装有Vista操作系统的笔记本电脑，这就意味着我们只能使用还处在测试阶段的CUDA SDK 2.0版本。并且我们的NV显示驱动也仅仅是v174.55，这个测试驱动也是测试版。尽管如此，我们在编写、运行程序时，没有遇到任何的不愉快。只是在初次运行程序时，我们的程序有不少bug，程序企图访问已经分配的显存空间。显示器开始狂闪，接着就黑屏了。直到Vista自动重启了显示驱动服务，电脑的桌面才重新回到我们的眼帘。小熊在线www.beareyes.com.cn

但不得不承认，其实这也不足为奇，这仅仅是程序的一个小BUG，跟无数Windows应用程序一样，都是可以通过简单修复后补完的。小熊在线www.beareyes.com.cn

最后，我来谈谈对NVIDIA的一些小牢骚。NVIDIA CUDA所有文档，都在一本厚厚的百科全书大全中，我并没有找到一本稍微薄一些的快速入门指南。或者是一步步CUDA从入门到精通之类的文档，也没有针对Visual Studio的快速环境配置指南。在有不明白不理解的名词，登录NVIDIA的CUDA官方网站，居然没有一个关键词搜索框，我只能一篇一页的通读所有CUDA官网的文章。这里确实有失人性化。希望NVIDIA给广大的初学者，开更多方便之门，把CUDA的入门门槛降的再低一些就好了。小熊在线www.beareyes.com.cn

不过这些都是小问题，SDK中也有比较完成的程序例子可供各位参考探究。你可以从中找到如何构建一个迷你型CUDA应用程序的方法。对于初学者来说，这已经是一个不错的入门指南了。

最终结论：CUDA才是终极的CPU（上）

最终结论：CUDA才是终极的CPU（上）

NVIDIA自从发布CUDA之后，就推出了GeForce 8800系列显示卡。在当时，NV承诺的性能是极端的诱人，但是我们要时刻保持警惕，别被商家给忽悠了。毕竟这是个先河，还从来没有人搞过GPGPU。而且目前可以用的资源还太少，对于GPGPU的研究也仅仅是在初期阶段。不过我要说的是，GPGPU的潜力是无限的。小熊在线www.beareyes.com.cn

早在2007年初，还仅仅只有一个试用版本的SDK，但是CNUDA的更新也十分的频繁，这证明NVIDIA正在专注通用计算事业。今天的CUDA发展的很好，SDK已经有2.0的测试版本可以使用了，而且它支持目前主流的操作系统，如Windows Vista、XP、Linux和Mac OS X。并且NVIDIA也推出了专门的CUDA官方网站。支持多语言，全世界的开发人员都可以参与到其中来。小熊在线www.beareyes.com.cn

从个人的层面上讲，小编我对CUDA的印象是给予非常的肯定和鼓励。如果你非常熟悉GPU的架构和特性，那么你的智商也足以学会简单编程，通过API去实现自己的CUDA应用程序。并且，你不必费心的去考虑硬件的细节，只需要善用好这几千个并行线程就可以实现伟大的程序。我们一开始编写这个程序的时候，还不敢确定CUDA有多强的性能，通过这个小程序，我们发现它确实要比CPU的性能强大许多。小熊在线www.beareyes.com.cn

因此，CUDA并不是研究员勾引消费者去购买GeForce显卡的一种骗人的把戏。CUDA的实用性是任何C语言程序员都可以读懂的。即使是普通的电脑爱好者，只要在这方面投入一些时间和精力，也会步入这个并行计算的殿堂。不过我觉得NVIDIA应该提供更多，更好品质的开发帮助文档，来帮助用户入门。小熊在线www.beareyes.com.cn

最终结论：CUDA才是终极的CPU（下）

最终结论：CUDA才是终极的CPU（下）

NVIDIA的CUDA会成为未来的一个主流的API，主要就是因为一个字：便捷。我们都知道，未来的IT领域是并行计算的天下。无论是硬件、软件都会朝着这个方向发展，这是业内技术的大趋势。目前的这些开发范例，也还只是创建一些线程，不过开发人员一定要小心的处理共享资源，因为现在的处理器也朝着多核心的方向发展着。小熊在线www.beareyes.com.cn

但是在数年以后，处理器的核心数，也不可能达到几百个。不过目前NVIDIA是第一个做到了这些的厂商。不过一枝独秀的狭义性也显现出来，目前也仅仅有NVIDIA的GF8、9系列和高端的Quadro/Tesla才支持CUDA程序。小熊在线www.beareyes.com.cn

Tesla高端计算解决方案

虽然可以运行CUDA的NVIDIA显示卡遍及世界各个角落，不过NVIDIA也并不满足现状。毕竟它的竞争对手们也不会袖手旁观。AMD也提供了自己的SDK，名叫Stream Computing。尽管至今没有一个实际的东西，Intel也发布了自己的解决方案——Ct。这三个竞争者，就像三国一样，三足鼎立，互相促进又互相牵制。另外在暗中伺机窥探的还有一个大佬，微软。他们也推出了自己的API，印有“微软出品”的字样相信这更容易让众多的开发者接受。小熊在线www.beareyes.com.cn

AMD Stream Computing

NVIDIA仍然有很多机遇与挑战。虽然在技术方面CUDA有许多先天的优势，但是想要说服系统平台开发商，和崇尚兼容性，稳定性为首要素的软件开发人员来支持CUDA也不是一件容易的事儿。不过随着CUDA SDK 2.0的发布，相信会有更多的应用程序开发商愿意支持这项技术。同时小编我也深信，在未来的计算机系统中，GPU的性能大幅超越了CPU，而处理器仅仅起到一个调度协调的作用，真正具有强大运算性能的部件已经远非CPU可比，而CUDA才是终极形态的CPU。

来自 “ ITPUB博客 ” ，链接：http://blog.itpub.net/14741601/viewspace-410824/，如需转载，请注明出处，否则将追究法律责任。

转载于:http://blog.itpub.net/14741601/viewspace-410824/