记一次CUDA编程任务

shixiang

这个月6号开始，着手解决一个具有实际意义的计算任务，这里简单记录一下，供CUDA初学者参考，大神勿喷。

任务数据有9879896条，每条包含30个整数，任务是计算每两条数据之间的斯皮尔相关系数及其P值。原始数据只有500+MB，因此我并不认为这是个多么大的计算任务。随后稍加计算，我还是很惊呆的，要计算(9879896×9879895)÷2≈4.88亿亿组数据，但此时这还只是个数字概念，我也没意识到时间复杂度和空间复杂度的问题。

1. 计算规模初体验

数据格式：9879896行，30列，每列之间以空格符隔开，例如：

1. 0 2 0 2 0 0 0 0 0 0 0 40 0 0 35 0 0 53 0 44 0 0 0 0 0 0 0 0 0 0
   
2. 0 0 1 148 0 0 0 0 0 0 0 0 0 0 1133 0 1 0 0 1820 0 0 0 2 0 0 0 1 0 0
   
3. 0 0 0 33 1 0 0 0 0 0 0 0 0 0 231 0 0 0 0 402 0 0 0 0 0 0 0 0 0 0
   
4. 0 0 6 0 0 0 0 0 0 0 0 0 0 0 1 0 0 0 0 0 1 0 6 0 0 0 0 0 0 0
   
5. 0 0 0 0 0 0 0 0 5 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
   
6. ... ...
   
7. ... ...

复制代码

空间复杂度：单纯计算下结果大概有多大吧，每组计算结果包含相关系数和P值，若都以float（占4字节）精度存储，需要占用内存：4.88亿亿×8B≈400TB，当然，我们不具备这么大内存，因此无论以何种方式计算，都需要一批批地重复将数据载入内存、计算、存入硬盘这个过程，直到运算完成。那么，存入硬盘的结果会有400TB吗？不然，P值小于或等于0.05的结果才会需要输出，因此实际上会远远小于这个值，具体会小多少，先运行一批数据后才能做出估算。

时间复杂度：计算的组数规模是(n×(n-1))÷2，那么就看程序能跑多快了。我想先看看MATLAB多线程、Python多线程、Spark分布式计算能跑多快，是否能在最快时间内解决问题。

2. MATLAB多线程

MATLAB写起来最简单，计算相关系数和P值都不用操心，一行自带的函数调用就完成。打开MATLAB左下角的并行池，MATLAB将会自动寻找到机子上有的物理核心，并分配与物理核心数相同的worker。比如我的电脑是4核8线程，它只能开4个worker，不识别虚拟核心。

代码如下：

1. t1 = clock;
   
2. disp('>> loading ...');
   
3. A = importdata('D:/MASTER2016/5.CUDA/data-ID-top30-kv3.txt');
   
4. b = A'; %由于MATLAB只计算列与列之间的相关系数，因此需要转置操作
   
5. disp(etime(clock,t1));
   
6. 
   
7. num = size(b, 2);
   
8. disp('>> calculating ...');
   
9. fid = fopen('D:/MASTER2016/5.CUDA/result-matlab.txt', 'wt');
   
10. 
    
11. for i = 1 : num
    
12.     for j = i+1 : num
    
13.         [m, n] = corr(b(:, i), b(:, j), 'type', 'Spearman', 'tail', 'both');
    
14.         if isnan(n) || n>0.05
    
15.             continue;
    
16.         end
    
17.          
    
18.         fprintf(fid, 'X%d\tX%d\t%d\t%d\n', i, j, m, n);
    
19.     end
    
20. end
    
21. fclose(fid);
    
22. disp('>> OK!');

复制代码

这里我并没有考虑内存空间不够的问题，因为我只是想说明MATLAB的计算速度。开了多颗核心的情况下，MATLAB并没能完全压榨出所有的CPU性能，计算速度缓慢无比，更要命的是，它会越算越慢。据我估算，即使空间复杂度足够，MATLAB也要用超过20年的时间才能算完，这还是不考虑越算越慢的情况。

好了，此方案仅是打酱油。

3. Python多线程

Python语言由于本身的体质问题，Cython下不能调用多核，只能用多线程。理论上是这样，但还是有很多扩展包能够充分压榨出多核CPU性能，例如multiprocessing是其中的佼佼者。multiprocessing用起来也非常简单，考虑到CPU的多核运算下，每颗核心的算力还是很可观的，所有不能把每个计算组都拆成并行线程，那样内存的读写开销反而会使CPU一直在等待状态，不能一直满负载工作。鉴于此，我设计9879895组线程，每组代表某个特定行与剩下的各个数据行形成的数据组。这样每组线程下的运算量还是比较大的，能使CPU尽可能全在满负载状态。

代码如下：

1. # coding=utf-8
   
2. 
   
3. import math
   
4. import multiprocessing
   
5. import time
   
6. 
   
7. import scipy.stats as stats
   
8. 
   
9. 
   
10. def calculate2(i, X, all_glb, data_array_glb):
    
11.     all = all_glb.value
    
12.     result = []
    
13.     for j in range(i + 1, all):
    
14.         x = X
    
15.         y = data_array_glb[j]
    
16.         if math.fsum(x) == 0 or math.fsum(y) == 0:
    
17.             continue
    
18.         corr, p = stats.spearmanr(x, y)
    
19.         if p > 0.05:
    
20.             continue
    
21.         result.append([i + 1, j + 1, corr, p])
    
22.     return result
    
23. 
    
24. 
    
25. if __name__ == "__main__":
    
26. 
    
27.     multiprocessing.freeze_support()
    
28. 
    
29.     input_file = 'D:/MASTER2016/5.CUDA/data-ID-top30-kv3.txt'
    
30.     output_file = 'D:/MASTER2016/5.CUDA/result-python.txt'
    
31. 
    
32.     print '>> loading ...'
    
33.     start = time.clock()
    
34.     data = open(input_file)
    
35.     data_array = []
    
36.     for line in data:
    
37.         data_array.append(map(int, line.strip().split(' ')))
    
38.     data.close()
    
39.     print time.clock()-start, 's'
    
40. 
    
41.     print '>> calculating ...'
    
42.     results = []
    
43.     pool_size = 8
    
44.     pool = multiprocessing.Pool(processes=pool_size)
    
45.     all = len(data_array)
    
46.     manager = multiprocessing.Manager()
    
47.     all_share = manager.Value('i', int(all))
    
48.     data_array_share = manager.list(data_array)
    
49.     for i in range(all):
    
50.         data_X = data_array[i]
    
51.         results.append(pool.apply_async(calculate2, args=(i, data_X, all_share, data_array_share)))
    
52.     pool.close()
    
53.     pool.join()
    
54.     print time.clock() - start, 's'
    
55.     data_array = None
    
56. 
    
57.     print '>> saving ...'
    
58.     data2 = open(output_file, 'w')
    
59.     for res in results:
    
60.         temp_list = res.get()
    
61.         for temp in temp_list:
    
62.             data2.write('X'+str(temp[0])+'\t'+'X'+str(temp[1])+'\t'+str(temp[2])+'\t'+str(temp[3])+'\n')
    
63.     print time.clock()-start, 's'
    
64.     data2.close()

复制代码

这里，我依然没有考虑空间复杂度问题，因为要先看看计算能力是否能满足任务要求。注意需要安装multiprocessing包。Python的这个多线程下，确实能充分榨干CPU性能，风扇呼呼响，要命的是也存在越算越慢的问题。但是，即使CPU一直这么满负载运算，我粗略估算了下，也得要个14年+才能算完，也不算越算越慢的情况。

所以，此方案是打酱油2号。


4. Spark方案

Spark方案我并没有写完，因为写着写着就感觉到。。。肯定还是不行，CPU的算力也就那样了。就算调12台机器一起跑，也不适合用CPU下的线程模型解决问题了。

这种高并行的计算，要想取得最快计算速度，非GPU莫属。


5. CUDA方案

CUDA方案下，首先必须清晰地设计好线程模型，即：我需要用到几块GPU？我需要在每块GPU上设计多少个block？每个block设计多少个线程？每个线程分配多少运算量？这四个问题基本决定了CUDA程序的性能和复杂度。

CUDA是一种异构并行解决方案，即CPU用于控制，GPU用于主运算的方案。一个GPU有一个grid，每个grid里有大量block，每个block里有大量thread。在运算时，每个thread都是完全独立并行地运算，每个线程里的运算靠内核函数控制，这也是CUDA编程的核心，目前只能用CUDA C编写。因此JCUDA和PyCUDA做的只是内存分配这些CPU端控制的事情，还不能代替GPU端的CUDA C代码。



如上图，左边列是Host端，即CPU上执行的控制端，用于分配GPU内存空间，拷贝内存数据到GPU显存等等操作。右边列是Device端，即GPU上的并行模型，由grid，block，thread三者构成。不同型号GPU的最大block数和每个block中的最大thread不同，但是可以查询。在安装好CUDA Toolkit后，windows用户可以进入C:\ProgramData\NVIDIA Corporation\CUDA Samples\v8.0\1_Utilities\deviceQuery目录，打开相应版本的项目，执行运行查询。

比如我的机器：


基于此，我设计的线程模型是：比如数据是ROWS行，COLS列，那么有((ROWS-1)×ROWS)÷2组计算，每一行都要与从这行开始后面的每一行进行计算。开辟(ROWS-1)个block，编号0~(ROWS-1)对应着数据的行号。所以，对于第一行，行号是0，要与1~(ROWS-1)的每一行进行计算，一共有(ROWS-1)组，这些计算任务分配给第一块block的1024个线程上计算。依此类推。这样做并不是最佳的任务分配方案，因为不是公平分配，编号越靠后的block分配的任务越少。但是，这样做的好处是便于利用共享内存，加速每一个block内的计算。

比如第一行，将数据第一行存入共享内存，那么它在与其他行分别计算的时候，直接从每个block内的共享内存读取数据，远远比从显存上的全局内存读取速度快得多。需要注意的是，每块block内的共享内存的大小也有硬件限制，上面截图中可以看到，GTX 950M的共享内存是49152B。

Talk is cheap. Show me the code:

1. #include <time.h>
   
2. #include <math.h>
   
3. #include <stdio.h>
   
4. #include <stdlib.h>
   
5. #include <assert.h>
   
6. #include "cuda_runtime.h"
   
7. #include "device_launch_parameters.h"
   
8. 
   
9. // 定义总数据矩阵的行数和列数
   
10. #define ROWS 15000
    
11. #define COLS 30
    
12. 
    
13. // 定义每一块内的线程个数，GT720最多是1024（必须大于总矩阵的列数：30）
    
14. #define NUM_THREADS 1024
    
15. 
    
16. 
    
17. bool InitCUDA()
    
18. {
    
19.     int count;
    
20.     cudaGetDeviceCount(&count);
    
21.     if (count == 0) {
    
22.         fprintf(stderr, "There is no device.\n");
    
23.         return false;
    
24.     }
    
25.     int i;
    
26.     for (i = 0; i < count; i++) {
    
27.         cudaDeviceProp prop;
    
28.         if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
    
29.             if (prop.major >= 1) {
    
30.                 break;
    
31.             }
    
32.         }
    
33.     }
    
34.     if (i == count) {
    
35.         fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
    
36.         return false;
    
37.     }
    
38.     cudaSetDevice(i);
    
39.     return true;
    
40. }
    
41. 
    
42. __device__ float meanForRankCUDA(int num)
    
43. {
    
44.     float sum = 0;
    
45.     for (int i = 0; i <= num; i++) {
    
46.         sum += i;
    
47.     }
    
48.     return sum / (num + 1);
    
49. }
    
50. 
    
51. 
    
52. __device__ float meanForArrayCUDA(float array[], int len)
    
53. {
    
54.     float sum = 0;
    
55.     for (int i = 0; i < len; i++) {
    
56.         sum += array[i];
    
57.     }
    
58.     return sum / len;
    
59. }
    
60. 
    
61. 
    
62. __device__ float spearmanKernel(int Xarray[], int Yarray[])
    
63. {
    
64.     //1，对原先的数据进行排序，相同的值取平均值
    
65.     float Xrank[30];
    
66.     float Yrank[30];
    
67.     int col = 30;
    
68. 
    
69.     for (int i = 0; i < col; i++) {
    
70.         int bigger = 1;
    
71.         int equaer = -1;
    
72.         for (int j = 0; j < col; j++) {
    
73.             if (Xarray[i] < Xarray[j]) {
    
74.                 bigger = bigger + 1;
    
75.             }
    
76.             else if (Xarray[i] == Xarray[j]) {
    
77.                 equaer = equaer + 1;
    
78.             }
    
79.         }
    
80.         Xrank[i] = bigger + meanForRankCUDA(equaer);
    
81.     }
    
82.     for (int i = 0; i < col; i++) {
    
83.         int bigger = 1;
    
84.         int equaer = -1;
    
85.         for (int j = 0; j < col; j++) {
    
86.             if (Yarray[i] < Yarray[j]) {
    
87.                 bigger = bigger + 1;
    
88.             }
    
89.             else if (Yarray[i] == Yarray[j]) {
    
90.                 equaer = equaer + 1;
    
91.             }
    
92.         }
    
93.         Yrank[i] = bigger + meanForRankCUDA(equaer);
    
94.     }
    
95. 
    
96.     //2，计算斯皮尔曼相关性系数
    
97.     float numerator = 0;
    
98.     float denominatorLeft = 0;
    
99.     float denominatorRight = 0;
    
100.     float meanXrank = meanForArrayCUDA(Xrank, col);
     
101.     float meanYrank = meanForArrayCUDA(Yrank, col);
     
102.     for (int i = 0; i < col; i++) {
     
103.         numerator += (Xrank[i] - meanXrank) * (Yrank[i] - meanYrank);
     
104.         denominatorLeft += powf(Xrank[i] - meanXrank, 2);
     
105.         denominatorRight += powf(Yrank[i] - meanYrank, 2);
     
106.     }
     
107.     float corr = 0;
     
108.     if ((denominatorLeft != 0) && (denominatorRight != 0)) {
     
109.         corr = numerator / sqrtf(denominatorLeft * denominatorRight);
     
110.     }
     
111.     return corr;
     
112. }
     
113. 
     
114. 
     
115. __global__ static void spearCUDAShared(const int* a, size_t lda, float* c, size_t ldc, float* d, size_t ldd)
     
116. {
     
117.     extern __shared__ int data[];
     
118.     const int tid = threadIdx.x;
     
119.     const int row = blockIdx.x;
     
120.     int i, j;
     
121.     // 同步第1行~倒数第二行到共享内存，行数由block个数（总数据矩阵的行数-1）控制，每个block共享一行数据
     
122.     if (tid < 30) {
     
123.         data[tid] = a[row * lda + tid];
     
124.     }
     
125.     __syncthreads();
     
126. 
     
127.     int cal_per_block = gridDim.x - row; // 每个块分担的计算量
     
128.     int cal_per_thread = cal_per_block / blockDim.x + 1; // 每个线程分担的计算量
     
129.     // 分配各线程计算任务，通过for循环控制在一个线程需要计算的组数
     
130.     for (i = row + cal_per_thread * tid; i < (row + cal_per_thread * (tid + 1)) && i < gridDim.x; i++) {
     
131.         int j_row[30]; // 存放总数据矩阵的第j行
     
132.         for (j = 0; j < 30; j++) {
     
133.             j_row[j] = a[(i + 1)*lda + j];
     
134.         }
     
135.         float corr = spearmanKernel(data, j_row);
     
136.         c[row * ldc + (i + 1)] = corr;
     
137.         float t_test = 0;
     
138.         if (corr != 0) t_test = corr*(sqrtf((30 - 2) / (1 - powf(corr, 2))));
     
139.         d[row * ldd + (i + 1)] = t_test;
     
140.         //printf("block号：%d, 线程号：%d, 计算组：%d-%d, id号：%d, block个数：%d, 每块线程个数：%d, 该块总计算量：%d, 该块中每个线程计算量：%d, corr: %lf, %d, %d, %d - %d, %d, %d\n", row, tid, row, i + 1, (row*blockDim.x + tid), gridDim.x, blockDim.x, cal_per_block, cal_per_thread, corr, data[0], data[1], data[29], j_row[0], j_row[1], j_row[29]);
     
141.     }
     
142. }
     
143. 
     
144. 
     
145. clock_t matmultCUDA(const int* a, float* c, float* d)
     
146. {
     
147.     int *ac;
     
148.     float *cc, *dc;
     
149.     clock_t start, end;
     
150.     start = clock();
     
151. 
     
152.     size_t pitch_a, pitch_c, pitch_d;
     
153.     // 开辟a、c、d在GPU中的内存
     
154.     cudaMallocPitch((void**)&ac, &pitch_a, sizeof(int)* COLS, ROWS);
     
155.     cudaMallocPitch((void**)&cc, &pitch_c, sizeof(float)* ROWS, ROWS);
     
156.     cudaMallocPitch((void**)&dc, &pitch_d, sizeof(float)* ROWS, ROWS);
     
157.     // 复制a从CPU内存到GPU内存
     
158.     cudaMemcpy2D(ac, pitch_a, a, sizeof(int)* COLS, sizeof(int)* COLS, ROWS, cudaMemcpyHostToDevice);
     
159. 
     
160.     spearCUDAShared << <ROWS - 1, NUM_THREADS, sizeof(int)* COLS >> > (ac, pitch_a / sizeof(int), cc, pitch_c / sizeof(float), dc, pitch_d / sizeof(float));
     
161. 
     
162.     cudaMemcpy2D(c, sizeof(float)* ROWS, cc, pitch_c, sizeof(float)* ROWS, ROWS, cudaMemcpyDeviceToHost);
     
163.     cudaMemcpy2D(d, sizeof(float)* ROWS, dc, pitch_d, sizeof(float)* ROWS, ROWS, cudaMemcpyDeviceToHost);
     
164.     cudaFree(ac);
     
165.     cudaFree(cc);
     
166. 
     
167.     end = clock();
     
168.     return end - start;
     
169. }
     
170. 
     
171. 
     
172. void print_int_matrix(int* a, int row, int col) {
     
173.     for (int i = 0; i < row; i++) {
     
174.         for (int j = 0; j < col; j++) {
     
175.             printf("%d\t", a[i * col + j]);
     
176.         }
     
177.         printf("\n");
     
178.     }
     
179. }
     
180. 
     
181. 
     
182. void print_float_matrix(float* c, int row, int col) {
     
183.     for (int i = 0; i < row; i++) {
     
184.         for (int j = 0; j < col; j++) {
     
185.             printf("%f\t", c[i * col + j]);
     
186.         }
     
187.         printf("\n");
     
188.     }
     
189. }
     
190. 
     
191. void read_ints(int* a) {
     
192.     FILE* file = fopen("D:\\MASTER2016\\5.CUDA\\data-ID-top30-kv.txt", "r");
     
193.     int i = 0;
     
194.     int count = 0;
     
195. 
     
196.     fscanf(file, "%d", &i);
     
197.     while (!feof(file))
     
198.     {
     
199.         a[count] = i;
     
200.         count++;
     
201.         if (count == ROWS*COLS) break;
     
202.         fscanf(file, "%d", &i);
     
203.     }
     
204.     fclose(file);
     
205. }
     
206. 
     
207. 
     
208. int main()
     
209. {
     
210.     int *a; // CPU内存中的总数据矩阵，ROWS行，COLS列
     
211.     float *c; // CPU内存中的相关系数结果矩阵，ROWS行，ROWS列
     
212.     float *d; // CPU内存中的T值结果矩阵，ROWS行，ROWS列
     
213.     a = (int*)malloc(sizeof(int)* COLS * ROWS);
     
214.     c = (float*)malloc(sizeof(float)* ROWS * ROWS);
     
215.     d = (float*)malloc(sizeof(float)* ROWS * ROWS);
     
216. 
     
217.     clock_t start = clock();
     
218.     printf(">> loading ... rows: %d, cols: %d", ROWS, COLS);
     
219.     read_ints(a);
     
220.     clock_t end = clock() - start;
     
221.     printf("\nTime used: %.2f s\n", (double)(end) / CLOCKS_PER_SEC);
     
222. 
     
223.     //print_int_matrix(a, ROWS, COLS);
     
224.     //printf("\n");
     
225. 
     
226.     printf(">> calculating ... ");
     
227.     printf("\n---------------------------------------");
     
228.     printf("\ntotal groups: %lld", (long long)ROWS*(ROWS - 1) / 2);
     
229.     printf("\ntotal threads: %d (blocks) * 1024 = %d", (ROWS - 1), (ROWS - 1) * 1024);
     
230.     printf("\ntotal space complexity: %lld MB", (long long)((ROWS / 1024) * (ROWS / 1024) * 8));
     
231.     printf("\n---------------------------------------");
     
232.     if (!InitCUDA()) return 0;
     
233.     clock_t time = matmultCUDA(a, c, d);
     
234.     double sec = (double)(time + end) / CLOCKS_PER_SEC;
     
235.     printf("\nTime used: %.2f s\n", sec);
     
236. 
     
237.     printf(">> saving ... ");
     
238.     FILE *f = fopen("D:\\MASTER2016\\5.CUDA\\result-c-2.txt", "w");
     
239.     for (int i = 0; i < ROWS; i++) {
     
240.         for (int j = i + 1; j < ROWS; j++) {
     
241.             float t_test = d[i * ROWS + j];
     
242.             if (t_test >= 2.042) {
     
243.                 fprintf(f, "X%d\tX%d\t%f\t%lf\n", i + 1, j + 1, c[i * ROWS + j], t_test);
     
244.             }
     
245.         }
     
246.     }
     
247.     fclose(f);
     
248.     end = clock() - start;
     
249.     printf("OK\nTime used: %.2f s\n", (double)(end) / CLOCKS_PER_SEC);
     
250. 
     
251.     //printf(">> 相关系数结果矩阵: \n");
     
252.     //print_float_matrix(c, ROWS, ROWS);
     
253.     //printf(">> T值结果矩阵: \n");
     
254.     //print_float_matrix(d, ROWS, ROWS);
     
255. 
     
256.     getchar();
     
257.     return 0;
     
258. }

复制代码

需要指出的是，上面程序保存为filename.cu文件，执行nvcc -o filename filename.cu编译，执行filename即可运行。其中ROWS是从总数据文件中读取的行数，用于控制数据规模调试程序，如果ROWS大于或等于总数据行数，那么就是读取整个文件了。

由于空间复杂度太高，也就是最开始提到的，那么下面做些调整，加个控制参数，每次只计算一定的行数，使显存满载但不超出即可。相应地，内核函数中的索引号，保存文件的函数都需要做些微调，代码如下：

1. #include <time.h>
   
2. #include <math.h>
   
3. #include <stdio.h>
   
4. #include <stdlib.h>
   
5. #include <assert.h>
   
6. #include "cuda_runtime.h"  
   
7. #include "device_launch_parameters.h"
   
8. 
   
9. // 定义总数据矩阵的行数和列数
   
10. #define ROWS 1000
    
11. #define COLS 30
    
12. 
    
13. // 控制一次计算占用显存的大小：CONTROL_ROWS*ROWS*8（字节）< 显存
    
14. #define CONTROL_ROWS 45
    
15. 
    
16. // 定义每一块内的线程个数，GT720最多是1024
    
17. #define NUM_THREADS 1024
    
18. 
    
19. 
    
20. bool InitCUDA()
    
21. {
    
22.     int count;
    
23.     cudaGetDeviceCount(&count);
    
24.     if (count == 0) {
    
25.         fprintf(stderr, "There is no device.\n");
    
26.         return false;
    
27.     }
    
28.     int i;
    
29.     for (i = 0; i < count; i++) {
    
30.         cudaDeviceProp prop;
    
31.         if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
    
32.             if (prop.major >= 1) {
    
33.                 break;
    
34.             }
    
35.         }
    
36.     }
    
37.     if (i == count) {
    
38.         fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
    
39.         return false;
    
40.     }
    
41.     cudaSetDevice(i);
    
42.     return true;
    
43. }
    
44. 
    
45. __device__ float meanForRankCUDA(int num)
    
46. {
    
47.     float sum = 0;
    
48.     for (int i = 0; i <= num; i++) {
    
49.         sum += i;
    
50.     }
    
51.     return sum / (num + 1);
    
52. }
    
53. 
    
54. 
    
55. __device__ float meanForArrayCUDA(float array[], int len)
    
56. {
    
57.     float sum = 0;
    
58.     for (int i = 0; i < len; i++) {
    
59.         sum += array[i];
    
60.     }
    
61.     return sum / len;
    
62. }
    
63. 
    
64. 
    
65. __device__ float spearmanKernel(int Xarray[], int Yarray[])
    
66. {
    
67.     //1，对原先的数据进行排序，相同的值取平均值
    
68.     float Xrank[30];
    
69.     float Yrank[30];
    
70.     int col = 30;
    
71. 
    
72.     for (int i = 0; i < col; i++) {
    
73.         int bigger = 1;
    
74.         int equaer = -1;
    
75.         for (int j = 0; j < col; j++) {
    
76.             if (Xarray[i] < Xarray[j]) {
    
77.                 bigger = bigger + 1;
    
78.             }
    
79.             else if (Xarray[i] == Xarray[j]) {
    
80.                 equaer = equaer + 1;
    
81.             }
    
82.         }
    
83.         Xrank[i] = bigger + meanForRankCUDA(equaer);
    
84.     }
    
85.     for (int i = 0; i < col; i++) {
    
86.         int bigger = 1;
    
87.         int equaer = -1;
    
88.         for (int j = 0; j < col; j++) {
    
89.             if (Yarray[i] < Yarray[j]) {
    
90.                 bigger = bigger + 1;
    
91.             }
    
92.             else if (Yarray[i] == Yarray[j]) {
    
93.                 equaer = equaer + 1;
    
94.             }
    
95.         }
    
96.         Yrank[i] = bigger + meanForRankCUDA(equaer);
    
97.     }
    
98. 
    
99.     //2，计算斯皮尔曼相关性系数
    
100.     float numerator = 0;
     
101.     float denominatorLeft = 0;
     
102.     float denominatorRight = 0;
     
103.     float meanXrank = meanForArrayCUDA(Xrank, col);
     
104.     float meanYrank = meanForArrayCUDA(Yrank, col);
     
105.     for (int i = 0; i < col; i++) {
     
106.         numerator += (Xrank[i] - meanXrank) * (Yrank[i] - meanYrank);
     
107.         denominatorLeft += powf(Xrank[i] - meanXrank, 2);
     
108.         denominatorRight += powf(Yrank[i] - meanYrank, 2);
     
109.     }
     
110.     float corr = 0;
     
111.     if ((denominatorLeft != 0) && (denominatorRight != 0)) {
     
112.         corr = numerator / sqrtf(denominatorLeft * denominatorRight);
     
113.     }
     
114.     return corr;
     
115. }
     
116. 
     
117. 
     
118. __global__ static void spearCUDAShared(const int* a, size_t lda, float* c, size_t ldc, float* d, size_t ldd, int cols, int start)
     
119. {
     
120.     extern __shared__ int data[];
     
121.     const int tid = threadIdx.x;
     
122.     const int row = blockIdx.x;
     
123. 
     
124.     int i, j;
     
125.     // 同步第1行~倒数第二行到共享内存，行数由block个数控制，每个block共享一行数据
     
126.     if (tid < 30) {
     
127.         data[tid] = a[(start + row) * lda + tid];
     
128.     }
     
129.     __syncthreads();
     
130. 
     
131.     int cal_per_block = cols - (start + row); // 每个块分担的计算量
     
132.     int cal_per_thread = cal_per_block / blockDim.x + 1; // 每个线程分担的计算量
     
133.     // 分配各线程计算任务，通过for循环控制在一个线程需要计算的组数
     
134.     for (i = row + cal_per_thread * tid; i < (row + cal_per_thread * (tid + 1)) && i < cols; i++) {
     
135.         int j_row[30]; // 存放总数据矩阵的第j行
     
136.         for (j = 0; j < 30; j++) {
     
137.             j_row[j] = a[(start + i + 1)*lda + j];
     
138.         }
     
139.         float corr = spearmanKernel(data, j_row);
     
140.         c[row * ldc + (start + i + 1)] = corr;
     
141.         float t_test = 0;
     
142.         if (corr != 0) t_test = corr*(sqrtf((30 - 2) / (1 - powf(corr, 2))));
     
143.         d[row * ldd + (start + i + 1)] = t_test;
     
144.         //printf("block号：%d, 线程号：%d, 计算组：%d-%d, id号：%d, block个数：%d, 每块线程个数：%d, 该块总计算量：%d, 该块中每个线程计算量：%d, corr: %lf, %d, %d, %d - %d, %d, %d\n", row, tid, row, i + 1, (row*blockDim.x + tid), gridDim.x, blockDim.x, cal_per_block, cal_per_thread, corr, data[0], data[1], data[29], j_row[0], j_row[1], j_row[29]);
     
145.     }
     
146. }
     
147. 
     
148. 
     
149. clock_t matmultCUDA(const int* a, float* c, float* d, int start_index, int control_rows)
     
150. {
     
151.     int *ac;
     
152.     float *cc, *dc;
     
153.     clock_t start, end;
     
154.     start = clock();
     
155. 
     
156.     size_t pitch_a, pitch_c, pitch_d;
     
157.     // 开辟a、c、d在GPU中的内存
     
158.     cudaMallocPitch((void**)&ac, &pitch_a, sizeof(int)* COLS, ROWS);
     
159.     cudaMallocPitch((void**)&cc, &pitch_c, sizeof(float)* ROWS, control_rows);
     
160.     cudaMallocPitch((void**)&dc, &pitch_d, sizeof(float)* ROWS, control_rows);
     
161.     // 复制a从CPU内存到GPU内存
     
162.     cudaMemcpy2D(ac, pitch_a, a, sizeof(int)* COLS, sizeof(int)* COLS, ROWS, cudaMemcpyHostToDevice);
     
163. 
     
164.     spearCUDAShared << <control_rows, NUM_THREADS, sizeof(int)* COLS >> > (ac, pitch_a / sizeof(int), cc, pitch_c / sizeof(float), dc, pitch_d / sizeof(float), ROWS - 1, start_index);
     
165.    
     
166.     cudaMemcpy2D(c, sizeof(float)* ROWS, cc, pitch_c, sizeof(float)* ROWS, control_rows, cudaMemcpyDeviceToHost);
     
167.     cudaMemcpy2D(d, sizeof(float)* ROWS, dc, pitch_d, sizeof(float)* ROWS, control_rows, cudaMemcpyDeviceToHost);
     
168.     cudaFree(ac);
     
169.     cudaFree(cc);
     
170.     cudaFree(dc);
     
171. 
     
172.     end = clock();
     
173.     return end - start;
     
174. }
     
175. 
     
176. 
     
177. void print_int_matrix(int* a, int row, int col) {
     
178.     for (int i = 0; i < row; i++) {
     
179.         for (int j = 0; j < col; j++) {
     
180.             printf("%d\t", a[i * col + j]);
     
181.         }
     
182.         printf("\n");
     
183.     }
     
184. }
     
185. 
     
186. 
     
187. void print_float_matrix(float* c, int row, int col) {
     
188.     for (int i = 0; i < row; i++) {
     
189.         for (int j = 0; j < col; j++) {
     
190.             printf("%f\t", c[i * col + j]);
     
191.         }
     
192.         printf("\n");
     
193.     }
     
194. }
     
195. 
     
196. void read_ints(int* a, char *input_file) {
     
197.     FILE* file = fopen(input_file, "r");
     
198.     int i = 0;
     
199.     int count = 0;
     
200. 
     
201.     fscanf(file, "%d", &i);
     
202.     while (!feof(file))
     
203.     {
     
204.         a[count] = i;
     
205.         count++;
     
206.         if (count == ROWS*COLS) break;
     
207.         fscanf(file, "%d", &i);
     
208.     }
     
209.     fclose(file);
     
210. }
     
211. 
     
212. void clear_ints(char * out_file) {
     
213.     FILE *f = fopen(out_file, "w");
     
214.     fclose(f);
     
215. }
     
216. 
     
217. void cal_and_save(int i, int *a, char *out_file, int control_rows) {
     
218.     float *c; // CPU内存中的相关系数结果矩阵，ROWS行，ROWS列
     
219.     float *d; // CPU内存中的T值结果矩阵，ROWS行，ROWS列
     
220.     c = (float*)malloc(sizeof(float)* control_rows * ROWS);
     
221.     d = (float*)malloc(sizeof(float)* control_rows * ROWS);
     
222. 
     
223.     clock_t time = matmultCUDA(a, c, d, i, control_rows);
     
224. 
     
225.     FILE *f = fopen(out_file, "a");
     
226.     for (int m = 0; m < control_rows; m++) {
     
227.         for (int n = i + m + 1; n < ROWS; n++) {
     
228.             float t_test = d[m * ROWS + n];
     
229.             if (t_test >= 2.042) {
     
230.                 fprintf(f, "X%d\tX%d\t%f\t%lf\n", i + m + 1, n + 1, c[m * ROWS + n], t_test);
     
231.             }
     
232.         }
     
233.     }
     
234.     fclose(f);
     
235. 
     
236.     //printf(">> 相关系数结果矩阵: \n");
     
237.     //print_float_matrix(c, CONTROL_ROWS, ROWS);
     
238.     //printf(">> T值结果矩阵: \n");
     
239.     //print_float_matrix(d, CONTROL_ROWS, ROWS);
     
240. 
     
241.     free(c);
     
242.     free(d);
     
243. }
     
244. 
     
245. int main()
     
246. {
     
247.     int *a; // CPU内存中的总数据矩阵，ROWS行，COLS列
     
248.     a = (int*)malloc(sizeof(int)* COLS * ROWS);
     
249. 
     
250.     char *input_file = "D:\\MASTER2016\\5.CUDA\\data-ID-top30-kv.txt";
     
251.     char *out_file = "D:\\MASTER2016\\5.CUDA\\result-c.txt";
     
252. 
     
253.     clock_t start = clock();
     
254.     printf(">> loading ... rows: %d, cols: %d", ROWS, COLS);
     
255.     read_ints(a, input_file);
     
256.     clear_ints(out_file);
     
257.     clock_t end = clock() - start;
     
258.     printf("\nTime used: %.2f s\n", (double)(end) / CLOCKS_PER_SEC);
     
259. 
     
260.     //print_int_matrix(a, ROWS, COLS);
     
261.     //printf("\n");
     
262. 
     
263.     printf(">> calculating ... ");
     
264.     printf("\n---------------------------------------");
     
265.     printf("\ntotal groups: %lld", (long long)ROWS*(ROWS - 1) / 2);
     
266.     printf("\ntotal threads: %d (blocks) * 1024 = %d", (ROWS - 1), (ROWS - 1) * 1024);
     
267.     printf("\ntotal space complexity: %lld MB", (long long)((CONTROL_ROWS / 1024) * (ROWS / 1024) * 8));
     
268.     printf("\n---------------------------------------");
     
269. 
     
270.     if (!InitCUDA()) return 0;
     
271. 
     
272.     int i;
     
273.     for (i = 0; i < ROWS - 1; i += CONTROL_ROWS) {
     
274.         printf("\n>> calculating and saving ... id: %d ... ", i);
     
275.         cal_and_save(i, a, out_file, CONTROL_ROWS);
     
276.         end = clock() - start;
     
277.         printf("Time used: %.2f s", (double)(end) / CLOCKS_PER_SEC);
     
278.     }
     
279. 
     
280.     // 不能整除的非整数部分需要计算
     
281.     //i -= CONTROL_ROWS;
     
282.     //int control_rows = ROWS - 1 - i;
     
283.     //printf("\n%d", control_rows);
     
284.     //if (control_rows > 0) {
     
285.     //    printf("\n>> calculating and saving ... id: %d ... ", i);
     
286.     //    cal_and_save(i, a, out_file, control_rows);
     
287.     //    end = clock() - start;
     
288.     //    printf("Time used: %.2f s", (double)(end) / CLOCKS_PER_SEC);
     
289.     //}
     
290. 
     
291.     printf("\nFinished.\n");
     
292. 
     
293.     getchar();
     
294.     return 0;
     
295. }

复制代码

到现在，由于空间复杂度过高而显存不够的问题通过增加时间复杂度的方法基本解决了。当然在显存足够的情况下，还是一次性算完是最快的，实测CUDA提速100+倍，数据量越大提速越明显。原因一是你必须被逼着按照CUDA的并行模型来写程序，二是GPU的架构设计确实更适合超大并行程序的加速。游戏画面渲染就是这样，你可以想成一个block控制一块屏幕的渲染，每块block的每个线程控制几个像素格的渲染，而这些图像渲染完全可以是独立并行的，GPU的设计初衷，即是增加核心数，不玩命升频率，增加显存带宽，使成百上千的核心数的并行计算能力得到充分释放。

但是目前的程序当然也不是完美的，我没有考虑如何隐藏内存与显存之间数据的传输延迟，没有考虑多块GPU如何联动运算。后面我会思考这些。

无疑，并行计算是计算的未来。而异构并行计算，也将是所有架构师必须增加的学习库。


附件：《CUDA专家手册  GPU编程权威指南》、《CUDA 编程指南4.0中文版》
链接：http://pan.baidu.com/s/1qYMyQfU 密码：75ya