作者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