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

zhaohongjian000

回复 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

zhaohongjian000 发表于 2015-01-12 16:27
回复 16# windoze

我们的技术，几十核没啥问题。上千的核？还真是挑战，得继续研究。

linux_c_py_php

对于并行软件开发模式，锁最终会是瓶颈，随着cpu个数增多，性能无法线性。

不过这不是说没法提高了，只是不能再依靠mutex+shared_ptr搞并发了，必须靠副本，消息机制，就像erlang所倡导的那样。

wang290

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

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

大神是典型的工程师作风，往往不屑那些学院派，比如洪内核和微内核之争，大神还是很鄙视那些教授们的

windoze

回复 22# yulihua49

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

yulihua49

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

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

提示一下，怎么玩？

windoze

回复 26# yulihua49

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

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() );
     

复制代码

yulihua49

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

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

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

lost_templar

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

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