使用GPU加速 - 顯卡

請問有使用過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);
}


--