※ [本文转录自 VideoCard 看板] 作者: a5000ml (咖啡里的海洋蓝) 看板: VideoCard 标题: [分享] CUDA 程式设计(9) -- 泡泡龙 part 2 时间: Thu Nov 6 03:08:19 2008 第 9 章 多区块的泡泡排序 运动完回来了，继续来玩上次的泡泡吧，这次要做多区块的版本 ================================================================= ◆ SIMT 对传统演算法带来的挑战 ================================================================= 上礼拜的演算法，大家可能会觉得不太习惯，所以开始之前，要解释一下： //传统的泡泡排序 (单链连续) ----------------------- for(int m=N; m>0; m--){ for(int k=0; k<m-1; k++){ if(a[k]>a[k+1]){ swap(a[k],a[k+1]); } } } //交错的泡泡排序 (双链错开) ------------------------ for(int m=N/2+1; m>0; m--){ __syncthreads(); for(int k=0; k<N-1; k+=2){ //0-based if(r[k]>r[k+1]){ swap(r[k],r[k+1]); } } __syncthreads(); for(int k=1; k<N-1; k+=2){ //1-based if(r[k]>r[k+1]){ swap(r[k],r[k+1]); } } } //------------------------------------------------- 这样的更动的理由，就是在上一篇 (8) SIMT vs OpenMP 中提及的： (1)程式码正交性 (2)序列化问题 SIMT 要避免执行绪搞在一起，不然竞争之下，就不知道写入的是什麽哩， 例如排列 (5,4,3)，两个 SP 同时作用，平行读取到的是同一个时段的值， 传统的泡泡法，(5,4) 要交换，(4,3)也要交换，所以就会变成 (4,?,4)， 中间的就看那个 SP 比较慢写入，可能是 5 也可能是 3，不过不管是谁， 结果都是错的，所以才需要错开，让执行绪处理的资料没有重叠 (正交化)， 然後在中间塞入 __syncthreads()，好让两个 base 循序的部份得以正确。 很多传统的好演算法都是基於循序过程，就拿排序来说好了，像 HeapSort、 QuickSort，循序的味道都很重，是否都能用 SIMT 来做，是值得研究的问题， 但软体和硬体总是交互成长，我想 SIMT 应该会刺激演算法的革新，而演算法 发展所遇到的困难，也会带动 SIMT 架构上的改进。 ================================================================= ◆ 多区块的泡泡排序 ================================================================= 多区块的泡泡，要分两层来做 (1) BLOCK 层的泡泡 -- 就是上礼拜的单一区块版本 (2) GRID 层的泡泡 -- 需要用 HOST 来派遣 例如下面要排序 16 个数字 9,5,8,4, 7,6,3,2, 6,8,7,4, 3,2,1,0 我们先把它分成 4 个区块，每个区块就用上次的单区块版本来排序 (1a)，这时可以设定 gridDim=4， 让这些区块使用多个 MP 同时排序，完成後再把每个区块的结果切成两半，跟另一个 相邻的半区块结合，形成一个完整的区块，设定 gridDim=3 对这三块进行排序 (1b)， 我们留意到每经过一个 ab 流程，最小的半区块 (尾巴的01) 就会往前步进一个区块， 令 G=N/BLOCK，只需经过 G-1 个 ab 周期，便可将它推到最前面的区块，然後再 过一个周期，就可完成区块内排序，整个工作就完成了，需要跑 G 次。 --------------------- 9584 7632 6874 3210 ------ (1a) 排列 4 个区块 4589 2367 4678 0123 45 8923 6746 7801 23 ------ (1b) 排列 3 个区块 45 2389 4667 0178 23 --------------------- 4523 8946 6701 7823 ------ (2a) 2345 4689 0167 2378 23 4546 8901 6723 78 ------ (2b) 23 4456 0189 2367 78 --------------------- 2344 5601 8923 6778 ------ (3a) 2344 0156 2389 6778 23 4401 5623 8967 78 ------ (3b) 23 0144 2356 6789 78 --------------------- 2301 4423 5667 8978 ------ (4a) 0123 2344 5667 7889 01 2323 4456 6778 89 ------ (4b) 01 2233 4456 6778 89 --------------------- 程式码列於最後面的附录中，值得注意的是不同区块间无法在 device 层级同步化， 而这种交错的方式在各进程之间有循序性，需要前面的资料正确写入，所以才把它 拆成多个 host 控制的 kernel 呼叫，中间塞入 cudaThreadSynchronize() 这个 API 来同步。 ================================================================= ◆ 演算法效能评估 ================================================================= (1) 单区块排序 B 个元素，需经过 B/2 个进程，每个进程比较 B 个元素， 所以时间约为 T1 = O(B^2/2) = alpha * B^2/2 (令系数为 alpha) (2) 排序 N=G*B 笔资料, 要排 G 个区块, 每个区块排序 B 个, 每次的 ab进程要发出 2G-1 个区块, 回圈 G 次, 若全部用单区块版本来排序， 所需的时间约为 TG = (2G-1)*G*T1 = O((G*B)^2) = O(N^2) 看起来比 Host 要差, 因为由(1)可知，Host 所需的时间约为 TH = O(N^2/2) = beta * N^2/2 (令系数为 beta) 多区块的版本引入 factor 2 (3) 考虑多处理器 MP 带来的增速 (工作量分担), 所需的时间变为 TM = TG/MP = O(N^2/MP) = alpha * N^2/MP (4) 考虑单区块和 HOST 的效能比 R1 = beta/alpha，估计多区块和 HOST 的 效能比约为 RG = TH/TM = R1*MP/2 (5) 上次执行的结果 R1 约为 1.9 倍，所以多区块版本和 HOST 的效能比约为 RG ~ MP 可以用来测试显示卡有几个多处理器 (MP, multiprocessors), 当然这个 1.9 倍是以 Intel Q6600 做准，其它 CPU 还是用 TM 来看会比较单纯 Note: 泡泡排序的核心每载入 & 输出 N 个元素，就要对它进行 N^2/2 个操作, 平均每传送一个元素，所搭配的运算量为 N/4，当 N 是 512 时 N/4=128 所以它的瓶颈是在运算(或者说是 branch)，不是记忆体传输。 ================================================================= ◆ 测试结果 ================================================================= 以下是一些机器上的测试结果，大家可以看到跟我们上面的估计差不多， 和主机的效能比约等於 multiprocessor 的个数，而且因为它的瓶颈在计算， 所以透过 PCI-E 传输资料的时间跟本不重要 (详见 code) 8800 GTX Ultra vs. Intel(R) Core(TM)2 Quad CPU Q6600 @ 2.40GHz ------------------------------------- NUM=131072 GRID=128 BLOCK=512 ------------------------------------- test[gpu]: pass test[host]: pass time[gpu]: 1330 ms time[host]: 22940 ms ratio: 17.2481 x ------------------------------------- NUM=262144 GRID=256 BLOCK=512 ------------------------------------- test[gpu]: pass test[host]: pass time[gpu]: 5330 ms time[host]: 91850 ms ratio: 17.2326 x GTX 280 vs. Intel(R) Core(TM)2 Quad CPU Q9400 @ 2.66GHz ------------------------------------- NUM = 262144 GRID=256 BLOCK=512 ------------------------------------- test[gpu]: pass test[host]: pass time[gpu]: 2920 ms time[host]: 84070 ms ratio: 28.7911 x ------------------------------------- NUM = 524288 GRID=512 BLOCK=512 ------------------------------------- test[gpu]: pass test[host]: pass time[gpu]: 11470 ms time[host]: 348880 ms ratio: 30.4167 x C1060 vs. Intel(R) Core(TM)2 Quad CPU Q9300 @ 2.50GHz ------------------------------------- NUM=131072 GRID=128 BLOCK=512 ------------------------------------- test[gpu]: pass test[host]: pass time[gpu]: 830 ms time[host]: 22450 ms ratio: 27.0482 x ------------------------------------- NUM=262144 GRID=256 BLOCK=512 ------------------------------------- test[gpu]: pass test[host]: pass time[gpu]: 2900 ms time[host]: 90260 ms ratio: 31.1241 x Note: 网格数若设太小，效能会比较差，因为 MP 无法吃尽网格，会闲置， 而且这些额外的时间比较没有机会被其它成功填满的网格行程分担掉， 还有无法整除的情况，如 GTX 280 系列的 MP 有 30 个，网格数若设 128， 则 12830=4…8 余下的 8 个区块在执行时其它的 22 个 MP 闲置着， 所以平均效能就会比较弱一点 (这边测试的 8800 GTX 网格数都可以整除， 所以效能会比预计的稍微高一点，这是其中一个原因) ================================================================= ◆ 附录: 多区块的泡泡程式码 ================================================================= #include <cuda.h> #include <time.h> #include <stdio.h> #include <stdlib.h> #include <math.h> //---------------------------------------------- // 请用 BLOCK 和 GRID 设定 NUM 的大小 //---------------------------------------------- #define BLOCK 512 //区块大小 #define GRID 32 //网格大小 #define NUM (2*BLOCK*GRID) //需要排列的元素个数 #define SIZE (NUM*sizeof(float)) #define testLoop 1 //---------------------------------------------- inline __host__ __device__ void swap(float& a, float& b){ float c=a; a=b; b=c; } //---------------------------------------------- // 单一区块的泡泡排序 (偶数版, 只排 BLOCK*2 个元素, 要对 blockIdx 定位) //---------------------------------------------- __global__ void bubble(float *r, float *a){ const int j=threadIdx.x; const int k=2*threadIdx.x; a=a+blockIdx.x*BLOCK*2; r=r+blockIdx.x*BLOCK*2; __shared__ float s[BLOCK*2+10]; __syncthreads(); s[j]=a[j]; s[j+BLOCK]=a[j+BLOCK]; for(int loop=0; loop<=BLOCK; loop++){ __syncthreads(); if(s[k]>s[k+1]){ swap(s[k],s[k+1]); } __syncthreads(); if(j<BLOCK-1) if(s[k+1]>s[k+2]){ swap(s[k+1],s[k+2]); } } __syncthreads(); r[j]=s[j]; r[j+BLOCK]=s[j+BLOCK]; } //---------------------------------------------- // Host 的泡泡排序 //---------------------------------------------- void bubble_host(float *r, float *a){ for(int k=0; k<NUM; k++){ r[k]=a[k]; } for(int loop=0; loop<=NUM/2; loop++){ for(int k=0; k<NUM-1; k+=2){ if(r[k]>r[k+1]){ swap(r[k],r[k+1]); } } for(int k=1; k<NUM-1; k+=2){ if(r[k]>r[k+1]){ swap(r[k],r[k+1]); } } } } //---------------------------------------------- int main(){ srand(time(0)); float *a=(float*)malloc(SIZE); float *b=(float*)malloc(SIZE); float *c=(float*)malloc(SIZE); printf("------------------------\n"); printf("NUM = %d\n",NUM); printf("------------------------\n"); for(int k=0; k<NUM; k++){ a[k]=k; c[k]=0; } for(int k=0; k<2*NUM; k++){ int i=rand()%NUM; int j=rand()%NUM; swap(a[i],a[j]); } float *gc; cudaMalloc((void**)&gc, SIZE); cudaMemcpy(gc, a, SIZE, cudaMemcpyHostToDevice); //测试 GPU 多区块泡泡效能 double t0=(double)clock()/CLOCKS_PER_SEC; for(int k=0; k<testLoop; k++){ //连 PCI-E 载入的时间算进去都划算 cudaMemcpy(gc, a, SIZE, cudaMemcpyHostToDevice); cudaThreadSynchronize(); for(int g=0; g<GRID; g++){ //进程 a bubble<<<GRID,BLOCK>>>(gc,gc); cudaThreadSynchronize(); //进程 b: 区块数少1, 工作点平移半个区块 bubble<<<GRID-1,BLOCK>>>(gc+BLOCK,gc+BLOCK); cudaThreadSynchronize(); } cudaMemcpy(c, gc, SIZE, cudaMemcpyDeviceToHost); cudaThreadSynchronize(); } t0=((double)clock()/CLOCKS_PER_SEC-t0)/testLoop; //测试 Host 泡泡效能 double t1=(double)clock()/CLOCKS_PER_SEC; for(int k=0; k<testLoop; k++){ bubble_host(b,a); } t1=((double)clock()/CLOCKS_PER_SEC-t1)/testLoop; //测量准确性 bool flag=true; for(int k=0; k<NUM; k++){ if(c[k]!=k){ flag=false; break; } } printf("test[gpu]: %s\n",flag?"pass":"fail"); flag=true; for(int k=0; k<NUM; k++){ if(b[k]!=k){ flag=false; break; } } printf("test[host]: %s\n",flag?"pass":"fail"); printf("time[gpu]: %g ms\n",t0*1000); printf("time[host]: %g ms\n",t1*1000); printf("ratio: %g x\n",t1/t0); printf("\n"); cudaFree(gc); free(a); free(b); free(c); return 0; } --

※ 发信站: 批踢踢实业坊(ptt.cc) ◆ From: 114.45.208.148 --

※ 发信站: 批踢踢实业坊(ptt.cc) ◆ From: 114.45.208.148