Linus Torvalds ：忘掉那该死的并行吧！

wang290

在 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核心搭载大量的集成，既便宜又低功耗”。

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

liuyeid

感觉不知道在说什么，如果并行有临界资源，一般都不会比单进程号多少，大多数环境下都是更挫

yulihua49

linux_c_py_php 发表于 2015-01-12 20:10
对于并行软件开发模式，锁最终会是瓶颈，随着cpu个数增多，性能无法线性。

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

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

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

所以说，并行的方法有很多种。怎能忘掉它？

VIP_fuck

回复 28# lost_templar

高大上


   

lost_templar

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

量子力学的东西，没指望你能看懂

yulihua49

lost_templar 发表于 2015-01-13 22:33
回复 26# yulihua49

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

没看懂，cuda是啥，你要介绍一下，你那个程序的功能是啥。

lost_templar

回复 26# yulihua49

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

1. 
   
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;
   
6. 
   
7.     if ( tilt_index >= tilt_size ) return;
   
8. 
   
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;
    
16. 
    
17.         unsigned long dimdim = dim*dim;
    
18. 
    
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;
    
26. 
    
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_;
    
31. 
    
32.         //1)
    
33.         kernel_assert( (compose_a<<<1, dim>>>( ug, ar, diag, thickness, a_, dim )) );
    
34.     cuda_assert( cudaDeviceSynchronize() );
    
35. 
    
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() );
    
42. 
    
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() );
    
50. 
    
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() );
    
58. 
    
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() );
    
93. 
    
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() );
     
103. 
     
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() );
     
113. 
     
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() );
     
120. 
     
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.                 }
     
130. 
     
131. 
     
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() );
     

复制代码

windoze

回复 26# yulihua49

妥妥的MPI+并行算法，目前没有什么更好的办法。

yulihua49

windoze 发表于 2015-01-13 14:24
回复 22# yulihua49

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

提示一下，怎么玩？

windoze

回复 22# yulihua49

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