第七章 来玩泡泡吧 今天我们来玩泡泡排序法吧! ◆ 演算法 为了能够平行化, 使用交错的方式对配对的资料进行排序, 如下图所示, 假设 有 10 笔资料, 我们先对位址 (0 1),(2 3),... 的 0 based 资料对进行排序, 然後再对位址 (1 2),(3 4),... 的 1 based 资料对进行排序, 如此一来, 每次进程可让最小的往前浮上两个单位, 总共要执行 N/2 次才会浮到最前面, 每次比对数量为 N-1, 所以是 O(N^2/2) 的泡泡法. (0 1)(2 3)(4 5)(6 7)(8 9) 0-based 0 (1 2)(3 4)(5 6)(7 8) 9 1-based 若资料量为奇数时, 例如 N=11, 则执行的方式如下, (0 1)(2 3)(4 5)(6 7)(8 9) 10 0-based 0 (1 2)(3 4)(5 6)(7 8)(9 10) 1-based 所以端点处和偶数时略有不同, 玩泡泡时要注意一下. ◆ 单一区块范例 以下示范用共享记忆体来做泡泡排序, 每个执行绪负责比较一对相邻的资料, 因为单一区块最大的执行绪个数是 512, 所以这个版本最多的输入资料数为 N = 512*2+1 = 1025 每次进程先比较 0-based 配对, 再比较 1-based 配对, 总共要执行 N/2 次, 也就是 blockDim 次, 程式码放在附录 bubble1.cu, 在 GTX 8800 Ultra 上 执行的结果如下 (host 为Intel Q6600) time[gpu]: 0.72 ms time[host]: 1.36 ms ratio: 1.88889 x ------------------------ test[gpu]: pass test[host]: pass 效能只有 1.8 倍, 因为它只用到一个区块, 也就是只用到 G80 上面 16 个 多处理器 (MP, multiprocessors) 的其中之一, 如果能够使用很多个区块 不仅效能能提昇数倍, 而且也能加大排序资料量. ◆ 练习 (1) 将范例改成用全域记忆体的版本, 并和使用共享记忆体的范例比较效能. (2) 将泡泡排序改为多区块版本, 首先将资料分段, 然後使用多区块对每段 进行排序, 其分段法仿照范例, 例如要排序 10*512 笔资料, 便可设定 10 个区块, 每个区块使用 256 个执行绪做区块的泡泡排序, 整个网格 先做 0-based, 再做 256-based, 每次发出的区块如下 0-based (0~511), (512~1023), (1024~1535), ... 256-based (256~767), (768~1279), (1280~1791), ... 因为每次最小的群组都会往前步进一个区块, 所以对 10*512 而言, 只要 让它 loop 个 10 次, 就可以把资料全部排序好了. (Note: 网格发出可在 host 上控制, 结果於下礼拜公布) (3) 试估计多区块版本的效能, 并与结果比较. (需考虑 MP 个数) ======================================================================== bubble1.cu ======================================================================== #include <cuda.h> #include <time.h> #include <stdio.h> #include <stdlib.h> #include <math.h> //---------------------------------------------- // 排序 N 个 float 元素 (N=1~1025) // 使用到的记忆体大小为 SIZE 个 bytes // 只使用单一区块 // 区块大小为 N/2 //---------------------------------------------- #define N 1024 #define SIZE (N*sizeof(float)) #define GRID 1 #define BLOCK (N/2) #define testLoop 1000 //测试效能时的 loop 数 //---------------------------------------------- // 交换函式 (host 和 kernel 都可以使用) // 因为加了 __host__ 和 __device__ 两个标签 //---------------------------------------------- inline __host__ __device__ void swap(float& a, float& b){ float c=a; a=b; b=c; } //---------------------------------------------- // 泡泡的 kernel (由小到大排列 N 个元素 a->r) //---------------------------------------------- __global__ void bubble(float *r, float *a){ //*** blockDim=N/2 *** int j=threadIdx.x; //j=0,1,2,...blockDim-1 int k=2*threadIdx.x; //k=0,2,4,...2*(blockDim-1) 配对的基底索引 //配置共享记忆体 __shared__ float s[N+20]; //载入资料到共享记忆体 __syncthreads(); //同步化执行绪, 加速载入速度 (合并读取 coalesced) s[j]=a[j]; //使用全部执行绪一起载入前半段 (0~N/2-1) s[j+N/2]=a[j+N/2]; //使用全部执行绪一起载入後半段 (N/2~N-1) if(j==0){ //若 N 为奇数时, 还要多载入一个尾巴, 只使用第 0 个执行绪 s[N-1]=a[N-1]; } //开始泡泡排序 for(int loop=0; loop<=N/2; loop++){ //排列 0 based 配对资料 (0,1) (2,3) (4,5) .... __syncthreads(); //同步化确保共享记忆体已写入 if(s[k]>s[k+1]){ swap(s[k],s[k+1]); } //排列 1 based 配对资料 (1,2) (3,4) (5,6) .... __syncthreads(); //同步化确保共享记忆体已写入 if(s[k+1]>s[k+2]){ if(k<N-2) //若 N 为偶数时, 最後一个执行绪不作用 swap(s[k+1],s[k+2]); } } //转出资料到全域记忆体 __syncthreads(); r[j]=s[j]; r[j+N/2]=s[j+N/2]; if(j==0){ r[N-1]=s[N-1]; } } //---------------------------------------------- // 泡泡的 host 函数 //---------------------------------------------- void bubble_host(float *r, float *a){ //载入资料 for(int k=0; k<N; k++){ r[k]=a[k]; } for(int loop=0; loop<=N/2; loop++){ //排列 0 based 配对资料 for(int k=0; k<N-1; k+=2){ if(r[k]>r[k+1]){ swap(r[k],r[k+1]); } } //排列 1 based 配对资料 for(int k=1; k<N-1; k+=2){ if(r[k]>r[k+1]){ swap(r[k],r[k+1]); } } } } //---------------------------------------------- // 主程式 //---------------------------------------------- int main(){ //配置 host 记忆体 float *a=(float*)malloc(SIZE); float *b=(float*)malloc(SIZE); float *c=(float*)malloc(SIZE); //初始化 for(int k=0; k<N; k++){ a[k]=k; c[k]=0; } //对阵列 a 洗牌 srand(time(0)); for(int k=0; k<2*N; k++){ int i=rand()%N; int j=rand()%N; swap(a[i],a[j]); } //配置 device 记忆体 float *ga, *gc; cudaMalloc((void**)&ga, SIZE); cudaMalloc((void**)&gc, SIZE); //载入 (顺便载入 c 来清空装置记忆体内容) cudaMemcpy(ga, a, SIZE, cudaMemcpyHostToDevice); cudaMemcpy(gc, c, SIZE, cudaMemcpyHostToDevice); //测试 kernel 效能 double t0=(double)clock()/CLOCKS_PER_SEC; for(int k=0; k<testLoop; k++){ //呼叫 kernel (此为单一 block 的版本) bubble<<<1,BLOCK>>>(gc,ga); //同步化执行绪, 避免还没做完, 量到不正确的时间 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; //显示计算时间, 并比较 printf("time[gpu]: %g ms\n",t0*1000); printf("time[host]: %g ms\n",t1*1000); printf("ratio: %g x\n",t1/t0); //读出 device 资料 cudaMemcpy(c, gc, SIZE, cudaMemcpyDeviceToHost); //测试 device 结果的正确性 printf("------------------------\n"); bool flag=true; for(int k=0; k<N; k++){ if(c[k]!=k){ flag=false; break; } } printf("test[gpu]: %s\n",flag?"pass":"fail"); //测试 host 结果的正确性 flag=true; for(int k=0; k<N; k++){ if(b[k]!=k){ flag=false; break; } } printf("test[host]: %s\n",flag?"pass":"fail"); //释放记忆体 cudaFree(ga); cudaFree(gc); free(a); free(b); free(c); return 0; } ======================================================================== --

※ 发信站: 批踢踢实业坊(ptt.cc) ◆ From: 114.45.213.247 ※ a5000ml:转录至看板 C_and_CPP 10/29 21:05 ※ uf2000uf:转录至某隐形看板 10/30 09:58