[請益] 使用GPU加速

看板VideoCard (顯卡板)作者holyyoung (嘿嘿羊)時間16年前 (2009/07/29 18:10)推噓2(2推 0噓 2→)

留言4則, 4人參與討論串1/1

請問有使用過CUDA的大大最近剛在使用把程式改成4個thread的平行化但是程式卻慢百倍!!! 按照道理說如果改4個thread 若扣掉搬資料的時間理應還是會比原來程式快一些這個不知道是為何會慢那麼多?? 感謝指教!!! -- =====================原來程式碼改的地方========================= for (j = 0; j < BLOCK_SIZE; j++) { memcpy(&m5[0],&img->cof[i0][j0][j][0], BLOCK_SIZE * sizeof(int)); m7 = &(img->m7[j][0]); m6[0] = m5[0] + m5[2]; m6[1] = m5[0] - m5[2]; m6[2] = (m5[1] >> 1) - m5[3]; m6[3] = m5[1] + (m5[3] >> 1); m7[0] = m6[0] + m6[3]; m7[1] = m6[1] + m6[2]; m7[2] = m6[1] - m6[2]; m7[3] = m6[0] - m6[3]; } for (i = 0; i < BLOCK_SIZE; i++) { ipos = i + ioff; m5[0]=img->m7[0][i]; m5[1]=img->m7[1][i]; m5[2]=img->m7[2][i]; m5[3]=img->m7[3][i]; m6[0] = m5[0] + m5[2]; m6[1] = m5[0] - m5[2]; m6[2] = (m5[1] >> 1) - m5[3]; m6[3] = m5[1] + (m5[3] >> 1); } ===============================改成kernal=================== __global__ void idctcuda_h(int *dev_A,int *dev_B) { int id=threadIdx.x; dev_B[id*4+0]=dev_A[id*4+0] +dev_A[id*4+2]; dev_B[id*4+1]=dev_A[id*4+0] -dev_A[id*4+2]; dev_B[id*4+2]=(dev_A[id*4+1])>>1-dev_A[id*4+3]; dev_B[id*4+3]=dev_A[id*4+1] +(dev_A[id*4+3])>>1; dev_A[0*4+id]=dev_B[id*4+0]+dev_B[id*4+3]; dev_A[1*4+id]=dev_B[id*4+1]+dev_B[id*4+2]; dev_A[2*4+id]=dev_B[id*4+1]-dev_B[id*4+2]; dev_A[3*4+id]=dev_B[id*4+0]-dev_B[id*4+3]; __syncthreads(); } __global__ void idctcuda_v(int *dev_A,int *dev_B) { int id=threadIdx.x; dev_B[id*4+0]=dev_A[id*4+0] +dev_A[id*4+2]; dev_B[id*4+1]=dev_A[id*4+0] -dev_A[id*4+2]; dev_B[id*4+2]=(dev_A[id*4+1])>>1-dev_A[id*4+3]; dev_B[id*4+3]=dev_A[id*4+1] +(dev_A[id*4+3])>>1; dev_A[0*4+id]=dev_B[id*4+0]+dev_B[id*4+3]; dev_A[1*4+id]=dev_B[id*4+1]+dev_B[id*4+2]; dev_A[2*4+id]=dev_B[id*4+1]-dev_B[id*4+2]; dev_A[3*4+id]=dev_B[id*4+0]-dev_B[id*4+3]; __syncthreads(); } __global__ void idctcuda_h(int *dev_A,int *dev_B) { int id=threadIdx.x; dev_B[id*4+0]=dev_A[id*4+0] +dev_A[id*4+2]; dev_B[id*4+1]=dev_A[id*4+0] -dev_A[id*4+2]; dev_B[id*4+2]=(dev_A[id*4+1])>>1-dev_A[id*4+3]; dev_B[id*4+3]=dev_A[id*4+1] +(dev_A[id*4+3])>>1; dev_A[0*4+id]=dev_B[id*4+0]+dev_B[id*4+3]; dev_A[1*4+id]=dev_B[id*4+1]+dev_B[id*4+2]; dev_A[2*4+id]=dev_B[id*4+1]-dev_B[id*4+2]; dev_A[3*4+id]=dev_B[id*4+0]-dev_B[id*4+3]; __syncthreads(); } __global__ void idctcuda_v(int *dev_A,int *dev_B) { int id=threadIdx.x; dev_B[id*4+0]=dev_A[id*4+0] +dev_A[id*4+2]; dev_B[id*4+1]=dev_A[id*4+0] -dev_A[id*4+2]; dev_B[id*4+2]=(dev_A[id*4+1])>>1-dev_A[id*4+3]; dev_B[id*4+3]=dev_A[id*4+1] +(dev_A[id*4+3])>>1; dev_A[0*4+id]=dev_B[id*4+0]+dev_B[id*4+3]; dev_A[1*4+id]=dev_B[id*4+1]+dev_B[id*4+2]; dev_A[2*4+id]=dev_B[id*4+1]-dev_B[id*4+2]; dev_A[3*4+id]=dev_B[id*4+0]-dev_B[id*4+3]; __syncthreads(); } void itranscuda(int *regist) { int data_size = 16*sizeof(int); int *dev_A ,*dev_B; int i; cudaMalloc( (void**)&dev_A, data_size ); cudaMalloc( (void**)&dev_B, data_size ); cudaMemcpy( dev_A, img, data_size, cudaMemcpyHostToDevice ); //dim3 dimblock(2,2); //dim3 dimgrid(w/dimblock.x,h/dimblock.y); idctcuda_h<<<1,4 >>>( dev_A,dev_B); idctcuda_v<<<1,4 >>>( dev_A,dev_B); cudaMemcpy( regist, dev_A, data_size, cudaMemcpyDeviceToHost ); cudaFree(dev_A); } -- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 140.123.112.5