作者holyyoung (嘿嘿羊)
看板VideoCard
標題[請益] 使用GPU加速
時間Wed Jul 29 18:10:46 2009
請問有使用過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
1F：推 vixen:程式怎麼改得? 這要po原始碼才能看出吧? 07/29 19:35
※ 編輯: holyyoung 來自: 140.123.112.5 (07/30 11:53)
2F：→ serlontw:資料不夠大光程式預備的時間,用單執行緒程式就跑完了.. 07/30 17:31
3F：推 sdk:4個thread太少了...好歹也來個4000 or 40000.... 08/01 01:11
4F：→ holyyoung:4000耶 不過程式很難可以改成100條以上 08/02 21:17