Chinaunix

标题: Linus Torvalds ：忘掉那该死的并行吧！ [打印本页]
作者: wang290    时间: 2015-01-09 09:39
标题: Linus Torvalds ：忘掉那该死的并行吧！
在 Avoiding ping pong上，Linus Torvalds以其一贯高雅的调调抨击了“并行计算就是未来”的论调，并在原文和 Reddit上收获了数百条评论。虽然事情最终也没有一个结果，但是许多观点确实值得借鉴。

Linus论点如下：

    推崇并行只不过是浪费大家的时间，“并行更高效”这种理论纯属胡说八道。大容量缓存是高效的，如果缺少缓存，并行一些低等级微内核可以说是毫无意义，除下特定类型上大规模规则计算，比如图形处理。

    没有人会回到过去，那些复杂的乱序运行内核不会消失。扩展不可能无休止的进行，人们需求更多的移动性，那些叫嚣扩展到上千核心的论调纯属扯淡，无需理会。

    是有多么奇葩的思维才能幻想出这些神奇等等并行算法的用武之地？！

    对于并行来说，唯一的用武之地就是图形计算和服务器端，而并行计算在这些领域确实也得到了大量的应用。但是没有任何疑问，并行在其他领域毫无用武之地。

    所以，忘掉并行吧，它永远都不可能被大规模推广。对于终端用户来说，4核就差不多了，而在这个领域，如果不增加太多的能耗，你也无法塞入更多的内核。同时，也不会有智障去阉割内核，降低其大小和性能只为了多塞几个。通常情况下，阉割内核只是为了降低功耗，因此这里也不会有那么多阉割的内核让你使用。

    因此，讲究程序的并行性本质上就是错的，它基于了一个错误的前提，同时也只是一个早该过时的时髦术语。

    在图形计算和服务器端之外，并行并不是万金油。即使在未来全新的领域同样如此，因为你根本承担不起。如果你期望做低功耗计算机视觉，我敢肯定你一定不会在GP CPU上编码。你甚至不会去使用GPU，因为它的开销太大了。大部分情况下，你可能会选择一些特殊的硬件——可能会基于某些神经网络模型。

    放弃吧。“并行就是未来”的说法纯属胡说八道。

在看讨论之前，我们首先看一下Linus以reference counting为例说明了并行的复杂性（该部分转自CoolShell）

在Linus回复之前有人指出对象需要锁机制的情况下，引用计数的原子性问题：

    由于（对象）通过多线程方式及多种获取渠道，一般而言它需要自身维护一个互斥锁——否则引用计数就不要求是原子的，一个更高层次的对象锁足矣。

而Linus不那么认为：

    引用计数的问题在于你经常需要在对象数据上锁保护之前完成它。

问题有两种情况，它们锁机制是完全不一样的：

    object *reference* 对象引用
    object data 对象数据

对象数据保护一般是一个对象拥有一个锁，假设你没有海量扩展性问题，不然你需要一些外部大一点的锁（极端的例子，一个对象一个全局锁）。

但对象引用主要关于对象的寻找（移除或释放），它是否在哈希链，一棵树或者链表上。当对象引用计数降为零，你要保护的不是对象数据，因为对象没有在其它地方使用，你要保护的是对象的寻找操作。

而且查询操作的锁不可能在对象内部，因为根据定义，你还不知道这是什么对象，你在尝试寻找它。

因此一般你要对查询操作上锁，而且引用计数相对那个锁来说是原子的（译者注：查询锁不是引用计数所在的对象所有，不能保护对象引用计数，后面会解释为何引用计数变更时其所在对象不能上锁）。

当然这个锁是充分有效的，现在假设引用计数是非原子的，但你常常不仅仅使用一种方式来查询：你可能拥有其它对象的指针（这个指针又被其它对象的对象锁给保护起来），但同时还会有多个对象指向它（这就是为何你第一时间需要引用计数的理由）。

看看会发生什么？查询不止存在一个锁保护。你可以想象走过一张对象流程图，其中对象存在指向其它对象的指针，每个指针暗含了一次对象引用，但当你走过这个流程图，你必须释放源对象的锁，而你进入新对象时又必须增加一次引用。

而且为了避免死锁，你一般不能立即对新对象上锁——你必须释放源对象的锁，否则在一个复杂流程图里，你如何避免ABBA死锁（译者注：假设两个线程，一个是A->B，另一个B->;A，当线程一给A上锁，线程二给B上锁，此时两者谁也无法释放对方的锁）？

原子引用计数修正了这一点，当你从对象A到对象B，你会这样做：

    对象A增加一次引用计数，并上锁。
    对象A一旦上锁，A指向B的指针就是稳定的，于是你知道你引用了对象B。
    但你不能在对象A上锁期间给B上锁（ABBA死锁）。
    对象B增加一次原子引用计数。
    现在你可以扔掉对象A的锁（退出对象A）。
    对象B的原子引用计数意味着即使给A解锁期间，B也不会失联，现在你可以给B上锁。

看见了吗？原子引用计数使这种情况成为可能。是的，你想尽一切办法避免这种代价，比如，你也许把对象写成严格顺序的，这样你可以从A到B，绝不会从B到A，如此就不存在ABBA死锁了，你也就可以在A上锁期间给B上锁了。

但如果你无法做到这种强迫序列，如果你有多种方式接触一个对象（再一次强调，这是第一时间使用引用计数的理由），这样，原子引用计数就是简单又理智的答案。

如果你认为原子引用计数是不必要的，这就大大说明你实际上不了解锁机制的复杂性。

相信我，并发设计是困难的。所有关于“并行化如此容易”的理由都倾向于使用简单数组操作做例子，甚至不包含对象的分配和释放。

那些认为未来是高度并行化的人一成不变地完全没有意识到并发设计是多么困难。他们只见过Linpack，他们只见过并行技术中关于数组排序的一切精妙例子，他们只见过一切绝不算真正复杂的事物——对真正的用处经常是非常有限的。（译者注：当然，我无意借大神之口把技术宗教化。实际上Linus又在另一篇帖子中综合了对并行的评价。）

哦，我同意。我的例子还算简单，真正复杂的用例更糟糕。

我严重不相信未来是并行的。有人认为你可以通过编译器，编程语言或者更好的程序员来解决问题，他们目前都是神志不清，没意识到这一点都不有趣。

并行计算可以在简化的用例以及具备清晰的接口和模型上正常工作。你发现并行在服务器上独立查询里，在高性能计算(High-performance computing)里，在内核里，在数据库里。即使如此，人们还得花很大力气才能使它工作，并且还要明确限制他们的模型来尽更多义务（例如数据库要想做得更好，数据库管理员得确保数据得到合理安排来迎合局限性）。

当然，其它编程模型倒能派上用场，神经网络(neural networking)天生就是非常并行化的，你不需要更聪明的程序员为之写代码。

在未来，应用程序究竟会发展成什么样？与现在有着非常大的区别？还是基本上相同，这里我们不妨看一下讨论（更多讨论见Reddit和Avoiding ping pong）：

Martin Thompson：

一旦工作集的大小超过了缓存容量，更大的缓存毫无意义。在低延时领域，为了保证整个应用程序放到缓存，我们通常会不择手段，但是这绝对不是主流。使用更大的页，并让L2支持这些更大的页显然比实际缓存大小更有意义，当下我们已经可以看到很多大内存应用程序运行在Haswell上。

对比使用并行，通常情况下使用cache friendly 或者cache oblivious（实际上是cache friendly的升级）数据结构显然更具生产效率。时至今日，“如果把在Fork-Join和并行流上投入些许精力放到提供更好的通用数据结构上（比如Maps和Trees，cache friendly）是否会更划算”这样辩论已经不再是困扰。在所谓的“多核问题解决”上，对比FJ和并行流，主流应用程序显然可以获得更多的提升。在这里，并不是说FJ和并行流不是个好的解决方案，而是后者可以给投资带来更多的回报。

在并行和并发上也有很多实际的用例，其中Servlet模型就是一个很好的例子，甚至是PHP之类在服务器端上的扩展。当然，在这之上，管道的构建也是一个更为直观的模型。

当谈到数据结构上的并发存取时，数据结构可变需要被单独对待。如果数据结构是不可变的，或者支持无阻塞并发读取，那么在并行上将很容易扩展，也很容易被推断。并发修改任何有趣的远程数据结构（更不用说完整模型），管理起来都是非常复杂和困难的。抛开复杂性，任何从多个写入者到1个共享模型/状态的并发更新都会存在限制，这点已经被Universal Scalability Law证明。在核心越来越多的情况下，在需要扩展的情况下，我们经常和自己开玩笑——多个写入者到任何模型的更新是否是一个好主意。庆幸的是，在大多数开发的应用程序代码中，查询针对的模型通常都不会有变化。

基于产业并发存取共享状态的需求，一个严重的后果产生：我们通常都会同步的进行这个过程，并在一个分布式的环境中传播。在算法和方法设计时，我们需要拥抱异步方式以避免延时限制。通过异步方式，我们可以实现无阻塞访问，而基于强制隔离，我们可以让应用程序更好地执行，并拥有更好的弹性。带宽以高速提升，延时将趋于平稳。

新一年我对平台提供商的愿望清单是：基础设施将有更好的cache friendly和immutable，同时还具备让异步编程更容易的Append-Only 数据结构、更好的管道并发、无阻塞APIS（比如JDBC）、语言外延（比如支持state machines和continuations），以及可以做申明式查询的语言外延（比如 LINQ[3] for C#就可以提升）。同时也不要介意允许Java那种低等级访问，我们已经远超越了在浏览器沙箱中写程序的时代。

AntiProtonBoy：

通常情况下，我对Linus是非常不感冒的，但是公平来讲，这次他说的确实很有道理。大量核心一般是用在大规模分布式应用系统中，比如说你想模仿一个神经网络。而在这个情况，你肯定也不会使用20万台个人电脑。他只是说在用户空间，30个小的核心并不会比4到8个高速核心快，因此并行化在这里并没有什么优势，也只有在遭遇瓶颈时才考虑到瓶颈。

Gabriele Svelto：

着眼当下移动领域，“足够快”很可能并非优化的终点。现在大部分使用电池的计算设备都在致力让用户能够获得一个更快的感知速度，从而在总体上节省电量。在这方面，某些并行算法完全处于劣势：在同等条件下，它们通过等价串行的方式，以增加计算（通常是通信）开销为代价来换取更快的速度。在实践中，并不是所有计算之外的开销都是等价的，因此，你还需要分摊一些固定的开销，不过整体更快的执行可能更加有效；但也正是这样，衡量是否要增加某个负载的速度将需要考虑更多变数。

Jeft：

Linus的说法可以说对，也可以说不对。事实上，人们期待可以更有效利用并行硬件的途径已经相当久了，所以不能把这个作为新事物来看。事实上我们需要的不仅是语言，如果你给它分配了太多工作，将从根本上挑战语言的基本结构，我们需求的语言是在需要的情况下可以最简单地并行，我们才刚刚开始。

Patrick Chase：

所有的一切都决定于容量和速度上的改变。当容量不足和（或者）算法集不稳定时，你使用的是商业硬件，而这十年的风格一直是GPU。当容量变高了，算法更稳定了，你开始考虑定制硬件（当下，一般复杂的ASIC定制大约是1000万美元或者更高；结果就是，你可以通过数学来发现哪个更有意义）。

如果只是容量变高了，算法还有一部分不稳定，那么定制一个包含了固定功能硬件和可编程硬件（DSPs、GPUs等等）的ASIC则非常有意义。这也是为什么高通公司为所有的Snapdragons都添加了“Hexagon”。

Maynard Handley：

当下，我们甚至没有开始程序员的再教育，让他们可以用更好的方式做事（更好的意思是抽象更匹配并行编程）。我们的语言、API以及工具仍然很糟糕，就像使用Fortran来做递归和指针一样。当下我们的工具并没有重构，这也让我们避免去关心某个函数调用链是否增加了一个新的参数等。

Patrick Chase：

针对Gabriele提出的“那些问题可以通过选择不同的语言解决”，你说的对，但是在现实世界中根本不可行。对比10年前，并行技术在难易度和开销上并没有什么根本上的突破。没有出现神奇的编译器，没有出现突破性的方法和语言，Amdahl法则并没有得到实质性的缓解。

序列化性能一定程度上受到了半导体工艺的限制，在这里我没看到任何微核心在equilibrium和optimum可以利用的因素。

因此，我觉得“anon”说的不错：并行方案只在必要的时候选用，提升单核性能则在任何可能的情况下。虽然这不是一个很好的愿景，但是却可以work。

Linus Torvalds：

我可以想象到人们在服务器领域已经使用上了60核心，但是我们不认为这是件值得推广的事情。我认为，在服务器端增加更多的缓存和集成更多的IO同样更具效率。

在客户端方面，仍然存在一些类似工作站的负载可以使用16核心，我认为借助它们，美术家确实可以更快地做Photoshop和视频编辑工作。但是从全局来看，这部分市场份额非常之小，从台式电脑市场萎缩就可见一斑。

因此，市场的趋势更应该是“4核心搭载大量的集成，既便宜又低功耗”。

但是，预测是困难的，特别是对未来，我们需要边走边看。

作者: yulihua49    时间: 2015-01-09 11:21
本帖最后由 yulihua49 于 2015-01-09 11:52 编辑
wang290 发表于 2015-01-09 09:39
在 Avoiding ping pong上，Linus Torvalds以其一贯高雅的调调抨击了“并行计算就是未来”的论调，并在原文和 ...

提升单核性能  ---  现在主频已经达到极限，还有什么办法能够提高性能吗？
对于终端用户来说，4核就差不多了 --  现在32核很普通了，128核也不稀奇了。还可以多CPU啊，还可以多服务器啊，在一个系统中，成千上万的核已经不是幻想。
在图形计算和服务器端之外  ---  把一半的计算机世界都除去了，这个论点还有什么意义吗？数十亿的客户端几乎都是由数百万的服务器提供支持，他们就像我们的左手和右手，各占半边天。云计算，大数据，哪个不靠服务器？哪个服务器不需要并行？
并行计算就靠优秀的软件，那些开发简单软件---把性能问题丢给硬件的论点，才是真正过时的论点。因为，性能，永远没有够的时候。


linus，老了，跟不上时代了。他那个linux，不就是做服务器的强项？做客户端，它拼得过WINDOWS吗？

linux，那个非抢占的优先级调度，一直是备受UNIX世界攻击的，对于做服务器也是很不利的。

并行计算中，的确存在很多难题，人们因此付出不懈的努力，并获得了有效的进展。
作者: hellioncu    时间: 2015-01-09 11:30
核太多，程序写得不好发挥不了作用。
作者: wang290    时间: 2015-01-09 18:47
yulihua49 发表于 2015-01-09 11:21
提升单核性能  ---  现在主频已经达到极限，还有什么办法能够提高性能吗？
对于终端用户来说，4核就差不 ...


大神你是理解不了的
作者: lost_templar    时间: 2015-01-10 08:00
正在折腾 cuda 的路过
作者: niao5929    时间: 2015-01-10 09:15
我觉得没有强大的本地化。说什么并行也是扯淡的。现在的所谓集群并行首先依托的是单个神经节点的本地化强大功能和更高效的网络。我们因该让神经节点具备一种模式设别功能。这样强大的本地化就可以很自然的处理各种复杂的情况。然后基于网络环境下的并行就自然形成了
作者: hwinlly    时间: 2015-01-10 22:55
大神和我一样，说话总是爱冲动。
作者: yulihua49    时间: 2015-01-11 19:27
niao5929 发表于 2015-01-10 09:15
我觉得没有强大的本地化。说什么并行也是扯淡的。现在的所谓集群并行首先依托的是单个神经节点的本地化强大 ...

？？并行有很多方法，你说的是－－－－－－－？
作者: zhaohongjian000    时间: 2015-01-12 10:13
回复 8# yulihua49

先说一下Linus，我觉得他有时候会矫枉过正。


大概几年前，工业界和学术界有这么一种提法：尽最大的可能增加核心的数目，不惜牺牲单个核心的性能。一个众核（many core）的概念开始流行。
很多人认为分支预测、高速缓存、乱序执行等技术非常耗电，采用众核的方法增加核心数目来增加性能，比提升单核性能在性能/功耗比上更合算。
所以产生了一些奇怪的硬件，动不动几百上千个核心，每个核心都非常挫。理论上，只要程序写的好，性能是没有问题的，好处是：
1.可以无限水平扩展，要更高性能只要增加核心就好了
2.省电

我觉得这套路很比较扯，毕竟写并行化程序是非常困难的，很多程序天生难以并行化，还有一些算法并行化后性能随着核心增加的提升并不理想。
我不清楚Linus是否主要针对这种技术吐槽，但感觉下面回复有一些是针对这种趋势的。
   
作者: yulihua49    时间: 2015-01-12 11:35
本帖最后由 yulihua49 于 2015-01-12 11:56 编辑
zhaohongjian000 发表于 2015-01-12 10:13
回复 8# yulihua49

先说一下Linus，我觉得他有时候会矫枉过正。

这不是什么观点问题，一切都是无奈之举，主频，受量子力学和光速限制，已经达到极限。
首先，单元已经是原子级，不能再小了，芯片尺寸，已经与主频的波长接近，主频也不能再提高。
分支预测、高速缓存、乱序执行也都没有放弃。
多核的芯片更大，所以主频略低。
有两种选择，较小的芯片和较高的主频，或者较大的芯片（多核）和较低的主频。
或者，用较小的芯片堆叠，并联，减少波的传播距离？散热怎么办？我也不是制造芯片的，说了不算。
不管怎么说，多核是（提高性能的）唯一途径。
那么，并行处理，交给软件了。（做硬件的，把球又抛给了软件）。
像12306那样的商业系统，除了并行，还能指望什么技术呢？
所以，面对几十个核（很普通的系统），往往是利用率很低，这就表现出软件的薄弱，就值得我们去研究，怎能忘掉那该死的并行呢？
作者: zhaohongjian000    时间: 2015-01-12 12:03
回复 10# yulihua49


    额，我说的重点在于，所谓众核经常大幅降低单核的性能，比如intel的一套众核芯片单个核心是十几年前的奔腾，性能放到现在不值一提。

    并行不是没有代价的，4个核心易于并行的算法也只能得到3.x的加速比，8个核心能有6倍加速就很牛逼了。在这种情况下，降低单核性能指望靠多核来弥补
是不现实的。
作者: yulihua49    时间: 2015-01-12 12:52
本帖最后由 yulihua49 于 2015-01-12 13:09 编辑
zhaohongjian000 发表于 2015-01-12 12:03
回复 10# yulihua49

你确认intel大幅降低单核性能？intel在2000年之前遇到了主频瓶颈，后来收买了digital公司的alpha技术，主频上去一大块，接近现在的水平。
同时采用了digital的乱序执行技术，这个当年我们是测过的，的确性能有显著提高。你确认他会退步？
多核CPU的主频有所降低，但不是大幅度的。

至于多核是3*还是多少，情况比较复杂，总之是软件的水平问题。
我们测试的，纯CPU的应用，在16核上，是16倍，一点没损失。但是到了系统级别，又有IO又有互斥，伸缩比就不一定了。但是一般说，会比单核强得多。
比如，我们的应用服务器系统，32核2.4G的机器，比4核3G的机器，要快至少6倍。那也很合算啊，整机价格才2倍。

为了解决互斥问题，人们发明了类似erlang语言，面向并行的，据说很有效。所以说球，又抛到了软件这边。
作者: zhaohongjian000    时间: 2015-01-12 14:14
回复 12# yulihua49


    是我没说清楚？你先了解一下众核的概念吧。当然这种东西没有推广，但很多厂商以及搞学术的认为这是趋势。

    为了防止你再次误解，我在强调一次，我说的这种东西不是你平时见到的服务器上的多核处理器，他们在核心之间的互联以及和内存的通讯方式上有很大区别。


    软件水平问题？你牛啊。。。很多算法多年都没有找到完美的并行算法，你说的这么轻巧。
作者: yulihua49    时间: 2015-01-12 15:46
本帖最后由 yulihua49 于 2015-01-12 15:59 编辑
zhaohongjian000 发表于 2015-01-12 14:14
回复 12# yulihua49

“众核”不是“多核”？那我就不清楚了。我说的普通的多核系统，在有限的条件下可以达到完美的并行度。
至于现在还未找到完美的并行方案，依旧是软件水平问题，包括最顶尖的学术研究，或者人类的智力都还不到位，还需要持续不断的努力与探索，怎能忘了那该死的并行？
尽管并行方案不完美，现在的成果，多核系统的性能，依然远远高于单核系统（在普通的商用服务器系统，而不仅是在研究机构）。
就拿ORACLE的RAC系统而言，无论在每个节点的多核服务器，还是在数据库引擎的总体性能，都显现了很好的性能的可伸缩性。

至于更多的核，更多的节点。如何达到满意的伸缩性，那就是要努力研究的了。

哦，请教下，intel的“众核”芯片是哪个型号？
作者: zhaohongjian000    时间: 2015-01-12 16:00
回复 14# yulihua49


    普通多核方案，是建立在尽量提高单核单核性能基础上的，依我之见工作的是不错的，尤其是业务逻辑完全可以水平分割的时候。
但有这么一群人认为不必提高单核性能，因为提高单核性能的技术（多发射、乱序、高速缓存）都非常费电。所以他们尝试制造了上千
个核的机器。我觉得Linus的批判对于这种方案完全合适（虽然我不知道他主要针对什么）。因为并行对软件开发来说是代价高昂的，
提升单核性能的技术还是非常重要的。
作者: windoze    时间: 2015-01-12 16:08
回复 14# yulihua49


他说的“众核”是指这个东西吧
作者: zhaohongjian000    时间: 2015-01-12 16:27
回复 16# windoze


    应该是这个，时间久了记不清了。还有这个：http://www.tilera.com/ 。

虽然有些亮点，但单核性能很挫是他们的共同特性。
作者: windoze    时间: 2015-01-12 16:36
回复 17# zhaohongjian000

这些Many Core架构的“CPU”本来就不是真正的“CPU”，你要是用通用CPU的标准去要求它们就有点过了……
那些“GPGPU”之类的你就当它们是一堆“协处理器”，这么一想就释然了。

话说你也只能这么想啊，Phi干脆连条件跳转指令都没有你说你不拿它当协处理器还能干神马啊～～～～～～
作者: zhaohongjian000    时间: 2015-01-12 16:44
回复 18# windoze


    我印象里这个东西是完整的CPU，也许记错了。而那个tilera更是一套完整的环境，上面可以跑linux，不是协处理器。
作者: yulihua49    时间: 2015-01-12 16:48
windoze 发表于 2015-01-12 16:08
回复 14# yulihua49


洋文，看着费劲，收藏，慢慢看。
作者: zhaohongjian000    时间: 2015-01-12 16:51
回复 18# windoze


    维基百科上的文字和我记忆里是相符的：
  1. The cores of Intel MIC are based on a modified version of P54C design, used in the original Pentium.[48] The basis of the Intel MIC architecture is to leverage x86 legacy by creating a x86-compatible multiprocessor architecture that can utilize existing parallelization software tools.[20] Programming tools include OpenMP, OpenCL,[49] Cilk/Cilk Plus and specialised versions of Intel's Fortran, C++[50] and math libraries.[51]
复制代码
从奔腾的核心修改的，可以直接跑之前的程序。
作者: yulihua49    时间: 2015-01-12 16:54
zhaohongjian000 发表于 2015-01-12 16:27
回复 16# windoze

我们的技术，几十核没啥问题。上千的核？还真是挑战，得继续研究。
作者: linux_c_py_php    时间: 2015-01-12 20:10
对于并行软件开发模式，锁最终会是瓶颈，随着cpu个数增多，性能无法线性。

不过这不是说没法提高了，只是不能再依靠mutex+shared_ptr搞并发了，必须靠副本，消息机制，就像erlang所倡导的那样。
作者: wang290    时间: 2015-01-13 13:41
linux_c_py_php 发表于 2015-01-12 20:10
对于并行软件开发模式，锁最终会是瓶颈，随着cpu个数增多，性能无法线性。

不过这不是说没法提高了，只是 ...


大神是典型的工程师作风，往往不屑那些学院派，比如洪内核和微内核之争，大神还是很鄙视那些教授们的
作者: windoze    时间: 2015-01-13 14:24
回复 22# yulihua49

上千核的程序就不是这么写了……
作者: yulihua49    时间: 2015-01-13 19:35
windoze 发表于 2015-01-13 14:24
回复 22# yulihua49

上千核的程序就不是这么写了……

提示一下，怎么玩？
作者: windoze    时间: 2015-01-13 20:31
回复 26# yulihua49

妥妥的MPI+并行算法，目前没有什么更好的办法。
作者: lost_templar    时间: 2015-01-13 22:33
回复 26# yulihua49

正在吐血调试的一段，并行在四块显卡上，有 2880 + 2496 + 2496 + 2496 个核

  2. __global__ void
  3. make_individual_pattern_intensity_diff( double* cuda_ug, unsigned long* cuda_ar, double* cuda_diag, double thickness, unsigned long* cuda_dim, double* cuda_I_exp, double* cuda_I_diff, unsigned long column_index, double2* cuda_cache, unsigned long max_dim, unsigned long tilt_size )
  4. {
  5.     unsigned long const tilt_index = blockDim.x * blockIdx.x + threadIdx.x;

  7.     if ( tilt_index >= tilt_size ) return;

  9.     unsigned long const dim = *(cuda_dim + tilt_index);
  10.     double* ug = cuda_ug;
  11.     unsigned long* ar = cuda_ar + tilt_index * max_dim * max_dim;
  12.     double* diag = cuda_diag + tilt_index * max_dim;
  13.     double* I_exp = cuda_I_exp + tilt_index * max_dim;
  14.     double* I_diff = cuda_I_diff + tilt_index * max_dim;
  15.     double2* cache = cuda_cache + 6 * tilt_index * max_dim * max_dim;

  17.         unsigned long dimdim = dim*dim;

  19.         //cache should be of size 6*N^2
  20.         double2* a_ = cache;
  21.         double2* aa_ = a_ + dimdim;
  22.         double2* aaa_ = aa_ + dimdim;
  23.         double2* p1 = aaa_ + dimdim;
  24.         double2* p2 = p1 + dimdim;
  25.         double2* p3 = p2 + dimdim;

  27.         //reuse memory in latter steps, when a_, aa_ and aaa_ are idle
  28.         double2* p2p3 = a_;
  29.         double2* s = aa_;
  30.         double2* s_ = aaa_;

  32.         //1)
  33.         kernel_assert( (compose_a<<<1, dim>>>( ug, ar, diag, thickness, a_, dim )) );
  34.     cuda_assert( cudaDeviceSynchronize() );

  36.         //2)
  37.     //TODO
  38.     double* the_norm = (double*)aa_;
  39.     kernel_assert( (Dznrm2<<<1,128>>>( dimdim, a_, the_norm )) );
  40.     //kernel_assert( (Dasum<<<1,128>>>( dimdim, a_, the_norm )) );
  41.     cuda_assert( cudaDeviceSynchronize() );

  43.         //double const ratio = (*the_norm) * 53.71920351148152;
  44.         double const ratio = (*the_norm) / 5.371920351148152;
  45.         unsigned long const scaler = ratio < 1.0 ? 0 : ceil(log2(ratio));
  46.         unsigned long const scaling_factor =  1 << scaler;
  47.         double const scale = scaling_factor;
  48.         kernel_assert( (Zscal<<<1, 128>>>( dimdim, 1.0/scale, a_ )) );    //a_ /= scale
  49.     cuda_assert( cudaDeviceSynchronize() );

  51.         //3)
  52.     dim3 const mm_grids( (dim+15)/16, (dim+15)/16 );
  53.     dim3 const mm_threads( 16, 16 );
  54.     kernel_assert( (Zgemm<<<mm_grids, mm_threads>>>( aa_, a_, a_, dim, 1.0 )) );
  55.     cuda_assert( cudaDeviceSynchronize() );
  56.     kernel_assert( (Zgemm<<<mm_grids, mm_threads>>>( aaa_, aa_, a_, dim, 1.0 )) );
  57.     cuda_assert( cudaDeviceSynchronize() );

  59.         //4)
  60.         /*solve(_Z^9+9*_Z^8+72*_Z^7+504*_Z^6+3024*_Z^5+15120*_Z^4+60480*_Z^3+181440*_Z^2+362880*_Z+362880 = 0)
  61.          * Returns:
  62.          *  2.697333461536989227389605+5.184162062649414177834087*I,     //c1
  63.          *  -.3810698456631129990312942+4.384644533145397950369203*I,    //c2
  64.          *  -2.110839800302654737498705+3.089910928725500922777702*I,    //c3
  65.          *  -3.038648072936697089212469+1.586801195758838328803868*I,    //c4
  66.          *  -3.333551485269048803294274,                                 //c5
  67.          *  -3.038648072936697089212469-1.586801195758838328803868*I,    //c6
  68.          *  -2.110839800302654737498705-3.089910928725500922777702*I,    //c7
  69.          *  -.3810698456631129990312942-4.384644533145397950369203*I,    //c8
  70.          *  2.697333461536989227389605-5.184162062649414177834087*I             //c9
  71.          *
  72.          *  expand((x-c1)*(x-c2)*(x-c3))  >> p1                                                                                                   (                     p1_c                             )
  73.          *          x^3-.205423815571221490859606*x^2-(12.65871752452031305098099*I)*x^2-58.21460179641193947200471*x-(3.189848964212376356715960*I)*x-19.71085376106750328141397+94.20645646169128946503649*I
  74.          *
  75.          *  expand((x-c4)*(x-c5)*(x-c6))  >> p2   (         p2_c            )
  76.          *          x^3+9.410847631142442981719212*x^2+39.17363072664900708597702-6.123261017392618755198919*10^(-24)*I+32.01029973951970099352671*x+(4.*10^(-24)*I)*x
  77.          *
  78.          *  expand((x-c7)*(x-c8)*(x-c9))  >> p3                                                                                                  (                         p3_c                         )
  79.          *          x^3-.205423815571221490859601*x^2+(12.65871752452031305098099*I)*x^2-58.21460179641193947200470*x+(3.18984896421237635671600*I)*x-19.71085376106750328141404-94.20645646169128946503646*I
  80.          *
  81.          *  expand((x-c1)*(x-c2)*(x-c3)*(x-c4)*(x-c5)*(x-c6)*(x-c7)*(x-c8)*(x-c9))
  82.          *          3.628800000000000000000003*10^5-1.365022562699469279472268*10^(-19)*I+3.628800000000000000000003*10^5*x+x^9+9.00000000000000000000000*x^8+72.00000000000000000000006*x^7+503.9999999999999999999995*x^6+3024.000000000000000000002*x^5+15120.00000000000000000000*x^4+60479.99999999999999999995*x^3+1.814400000000000000000001*10^5*x^2-(5.*10^(-22)*I)*x^6-(1.*10^(-20)*I)*x^4-(1.0*10^(-19)*I)*x^3+(2.*10^(-24)*I)*x^8-(3.0*10^(-19)*I)*x^2-(7.*10^(-21)*I)*x^5-(4.*10^(-19)*I)*x+(2.*10^(-23)*I)*x^7
  83.          */
  84.         //4 - p1)
  85.     kernel_assert( (Zcopy<<<1,128>>>( dimdim, aaa_, p1 )) );
  86.     cuda_assert( cudaDeviceSynchronize() );
  87.     kernel_assert( (Zaxpy<<<1,128>>>( dimdim, -0.205423815571221490859606, -12.65871752452031305098099, p1, aa_ )) );
  88.     cuda_assert( cudaDeviceSynchronize() );
  89.     kernel_assert( (Zaxpy<<<1,128>>>( dimdim, -58.21460179641193947200471, -3.189848964212376356715960, p1, a_ )) );
  90.     cuda_assert( cudaDeviceSynchronize() );
  91.         kernel_assert( (sum_diag<<<1,dim>>>( p1, dim, -19.71085376106750328141397, 94.20645646169128946503649 )) );
  92.     cuda_assert( cudaDeviceSynchronize() );

  94.         //4 - p2)
  95.     kernel_assert( (Zcopy<<<1,128>>>( dimdim, aaa_, p2 )) );
  96.     cuda_assert( cudaDeviceSynchronize() );
  97.     kernel_assert( (Zaxpy<<<1,128>>>( dimdim, 9.410847631142442981719212, 0.0, p2, aa_ )) );
  98.     cuda_assert( cudaDeviceSynchronize() );
  99.     kernel_assert( (Zaxpy<<<1,128>>>( dimdim, 32.01029973951970099352671, 0.0, p2, a_ )) );
  100.     cuda_assert( cudaDeviceSynchronize() );
  101.         kernel_assert( (sum_diag<<<1,dim>>>( p2, dim, 39.17363072664900708597702, 0.0  )) );
  102.     cuda_assert( cudaDeviceSynchronize() );

  104.         //4 - p3)
  105.     kernel_assert( (Zcopy<<<1,128>>>( dimdim, aaa_, p3 )) );
  106.     cuda_assert( cudaDeviceSynchronize() );
  107.     kernel_assert( (Zaxpy<<<1,128>>>( dimdim, -0.205423815571221490859601, 12.65871752452031305098099, p3, aa_ )) );
  108.     cuda_assert( cudaDeviceSynchronize() );
  109.     kernel_assert( (Zaxpy<<<1,128>>>( dimdim, -58.21460179641193947200470, 3.18984896421237635671600, p3, a_ )) );
  110.     cuda_assert( cudaDeviceSynchronize() );
  111.         kernel_assert( (sum_diag<<<1,dim>>>( p3, dim, -19.71085376106750328141404, -94.20645646169128946503646 )) );
  112.     cuda_assert( cudaDeviceSynchronize() );

  114.         //4 - s)
  115.         // s = 1/602.39521910453439454428( p1 * ( 1/602.39521910453439454428 * p2 * p3 ) ) = (p1 p2 p3)/362880
  116.     kernel_assert( (Zgemm<<<mm_grids, mm_threads>>>( p2p3, p2, p3, dim, 0.0016600397351866578333 )) );
  117.     cuda_assert( cudaDeviceSynchronize() );
  118.     kernel_assert( (Zgemm<<<mm_grids, mm_threads>>>( s, p1, p2p3, dim, 0.0016600397351866578333 )) );
  119.     cuda_assert( cudaDeviceSynchronize() );

  121.         //5)
  122.         for ( unsigned long index = 0; index != scaler; ++index )
  123.                 {
  124.             kernel_assert( (Zgemm<<<mm_grids, mm_threads>>>( s_, s, s, dim, 1.0 )) );
  125.                 cuda_assert( cudaDeviceSynchronize() );
  126.                         double2* tmp = s_;
  127.                         s_ = s;
  128.                         s = tmp;
  129.                 }


  132.         //6)
  133.         //kernel_assert( (extract_intensity_diff<<<1,dim>>>( s, I_exp, I_diff, dim, column_index )) );
  134.     double const ac_offset = cuda_ug[0];
  135.     double const dc_offset = cuda_ug[1];
  136.         kernel_assert( (extract_intensity_diff_with_offset<<<1,dim>>>( s, I_exp, I_diff, dim, column_index, ac_offset, dc_offset )) );
  137.         cuda_assert( cudaDeviceSynchronize() );
复制代码
作者: yulihua49    时间: 2015-01-14 13:18
lost_templar 发表于 2015-01-13 22:33
回复 26# yulihua49

正在吐血调试的一段，并行在四块显卡上，有 2880 + 2496 + 2496 + 2496 个核

没看懂，cuda是啥，你要介绍一下，你那个程序的功能是啥。
作者: lost_templar    时间: 2015-01-14 17:41
yulihua49 发表于 2015-01-14 13:18
没看懂，cuda是啥，你要介绍一下，你那个程序的功能是啥。


量子力学的东西，没指望你能看懂
作者: VIP_fuck    时间: 2015-01-14 18:00
回复 28# lost_templar

高大上


   
作者: yulihua49    时间: 2015-01-16 21:23
本帖最后由 yulihua49 于 2015-01-16 21:32 编辑
linux_c_py_php 发表于 2015-01-12 20:10
对于并行软件开发模式，锁最终会是瓶颈，随着cpu个数增多，性能无法线性。

不过这不是说没法提高了，只是 ...

这种“众核”的东西，简单的做法是构建很多的虚拟机，云平台用。

假如4096核，我可以这么用，建100个进程，每进程41线程，10000个客户端。这样可以支持100W用户的接入。
为了容错，至少2台服务器，可以200W接入。前端设负载均衡器。

所以说，并行的方法有很多种。怎能忘掉它？
作者: liuyeid    时间: 2015-02-03 23:02
感觉不知道在说什么，如果并行有临界资源，一般都不会比单进程号多少，大多数环境下都是更挫



欢迎光临 Chinaunix (http://bbs.chinaunix.net/) Powered by Discuz! X3.2