用户
 找回密码
 立即注册
pk82418 该用户已被删除
发表于 2014-12-1 21:00:34
68710
最近换了新电脑,发现个郁闷的问题。同样一个kenel函数,原来vs2008+cuda5.5 在gtx 580上只要运行0.3秒不到,现在vs2012+cuda6.5 在 gtx titan black上却需要3秒以上。按道理说后者的计算能力应该要比前者强大得多啊,为何计算速度反而不到原来的十分之一?

代码如下

__global__ void meshmap1_device(int* mapindex_gpu,float* mapvalue_gpu,float* node1_gpu,int* innode1_gpu,float* node2_gpu,int* element2_gpu,int* inelement2_gpu)
{
    const int tid = threadIdx.x;
const int bid = blockIdx.x;

int temp1=gridDim.x/8;
    int temp2=floorf(bid/temp1);
temp1=bid%temp1;
    const int loopnum=(temp2==7?loopnum2_gpu[0]:loopnum1_gpu[0]);
const int elementoffset=temp2*loopnum1_gpu[0];
const int nodeoffset=temp1*blockDim.x;
int i,j,k,nodeindex[4];
float x[4],y[4],z[4],Ve;
__shared__ float minvalue[256];
__shared__ int minindex[256];
if (nodeoffset+tid<innodecount1_gpu[0])
{
  minvalue[tid]=10;
  j=innode1_gpu[nodeoffset+tid];
  const float xx=node1_gpu[j];
  const float yy=node1_gpu[j+nodecount1_gpu[0]];
  const float zz=node1_gpu[j+2*nodecount1_gpu[0]];

  for (i=0;i<loopnum;i++)
  {
   k=inelement2_gpu[elementoffset+i];
   nodeindex[0]=element2_gpu[k];
   nodeindex[1]=element2_gpu[k+elementcount2_gpu[0]];
   nodeindex[2]=element2_gpu[k+2*elementcount2_gpu[0]];
   nodeindex[3]=element2_gpu[k+3*elementcount2_gpu[0]];

   
   x[0]=node2_gpu[nodeindex[0]];y[0]=node2_gpu[nodeindex[0]+nodecount2_gpu[0]];z[0]=node2_gpu[nodeindex[0]+2*nodecount2_gpu[0]];
   x[1]=node2_gpu[nodeindex[1]];y[1]=node2_gpu[nodeindex[1]+nodecount2_gpu[0]];z[1]=node2_gpu[nodeindex[1]+2*nodecount2_gpu[0]];
   x[2]=node2_gpu[nodeindex[2]];y[2]=node2_gpu[nodeindex[2]+nodecount2_gpu[0]];z[2]=node2_gpu[nodeindex[2]+2*nodecount2_gpu[0]];
   x[3]=node2_gpu[nodeindex[3]];y[3]=node2_gpu[nodeindex[3]+nodecount2_gpu[0]];z[3]=node2_gpu[nodeindex[3]+2*nodecount2_gpu[0]];
   
   Ve=abs(x[1]*y[2]*z[3]+y[1]*z[2]*x[3]+x[2]*y[3]*z[1]-z[1]*y[2]*x[3]-y[1]*x[2]*z[3]-z[2]*y[3]*x[1]-x[0]*y[2]*z[3]-y[0]*z[2]*x[3]-x[2]*y[3]*z[0]+z[0]*y[2]*x[3]+y[0]*x[2]*z[3]+x[0]*y[3]*z[2]+x[0]*y[1]*z[3]+x[1]*y[3]*z[0]+x[3]*y[0]*z[1]-z[0]*y[1]*x[3]-x[1]*y[0]*z[3]-y[3]*z[1]*x[0]-x[0]*y[1]*z[2]-x[1]*y[2]*z[0]-x[2]*y[0]*z[1]+z[0]*y[1]*x[2]+x[1]*y[0]*z[2]+y[2]*z[1]*x[0]);
   
   x[0]=xx;y[0]=yy;z[0]=zz;
   Ve=Ve-abs(x[1]*y[2]*z[3]+y[1]*z[2]*x[3]+x[2]*y[3]*z[1]-z[1]*y[2]*x[3]-y[1]*x[2]*z[3]-z[2]*y[3]*x[1]-x[0]*y[2]*z[3]-y[0]*z[2]*x[3]-x[2]*y[3]*z[0]+z[0]*y[2]*x[3]+y[0]*x[2]*z[3]+x[0]*y[3]*z[2]+x[0]*y[1]*z[3]+x[1]*y[3]*z[0]+x[3]*y[0]*z[1]-z[0]*y[1]*x[3]-x[1]*y[0]*z[3]-y[3]*z[1]*x[0]-x[0]*y[1]*z[2]-x[1]*y[2]*z[0]-x[2]*y[0]*z[1]+z[0]*y[1]*x[2]+x[1]*y[0]*z[2]+y[2]*z[1]*x[0]);
   x[1]=xx;y[1]=yy;z[1]=zz;x[0]=node2_gpu[nodeindex[0]];y[0]=node2_gpu[nodeindex[0]+nodecount2_gpu[0]];z[0]=node2_gpu[nodeindex[0]+2*nodecount2_gpu[0]];
   Ve=Ve-abs(x[1]*y[2]*z[3]+y[1]*z[2]*x[3]+x[2]*y[3]*z[1]-z[1]*y[2]*x[3]-y[1]*x[2]*z[3]-z[2]*y[3]*x[1]-x[0]*y[2]*z[3]-y[0]*z[2]*x[3]-x[2]*y[3]*z[0]+z[0]*y[2]*x[3]+y[0]*x[2]*z[3]+x[0]*y[3]*z[2]+x[0]*y[1]*z[3]+x[1]*y[3]*z[0]+x[3]*y[0]*z[1]-z[0]*y[1]*x[3]-x[1]*y[0]*z[3]-y[3]*z[1]*x[0]-x[0]*y[1]*z[2]-x[1]*y[2]*z[0]-x[2]*y[0]*z[1]+z[0]*y[1]*x[2]+x[1]*y[0]*z[2]+y[2]*z[1]*x[0]);
   x[2]=xx;y[2]=yy;z[2]=zz;x[1]=node2_gpu[nodeindex[1]];y[1]=node2_gpu[nodeindex[1]+nodecount2_gpu[0]];z[1]=node2_gpu[nodeindex[1]+2*nodecount2_gpu[0]];
   Ve=Ve-abs(x[1]*y[2]*z[3]+y[1]*z[2]*x[3]+x[2]*y[3]*z[1]-z[1]*y[2]*x[3]-y[1]*x[2]*z[3]-z[2]*y[3]*x[1]-x[0]*y[2]*z[3]-y[0]*z[2]*x[3]-x[2]*y[3]*z[0]+z[0]*y[2]*x[3]+y[0]*x[2]*z[3]+x[0]*y[3]*z[2]+x[0]*y[1]*z[3]+x[1]*y[3]*z[0]+x[3]*y[0]*z[1]-z[0]*y[1]*x[3]-x[1]*y[0]*z[3]-y[3]*z[1]*x[0]-x[0]*y[1]*z[2]-x[1]*y[2]*z[0]-x[2]*y[0]*z[1]+z[0]*y[1]*x[2]+x[1]*y[0]*z[2]+y[2]*z[1]*x[0]);
   x[3]=xx;y[3]=yy;z[3]=zz;x[2]=node2_gpu[nodeindex[2]];y[2]=node2_gpu[nodeindex[2]+nodecount2_gpu[0]];z[2]=node2_gpu[nodeindex[2]+2*nodecount2_gpu[0]];
   Ve=Ve-abs(x[1]*y[2]*z[3]+y[1]*z[2]*x[3]+x[2]*y[3]*z[1]-z[1]*y[2]*x[3]-y[1]*x[2]*z[3]-z[2]*y[3]*x[1]-x[0]*y[2]*z[3]-y[0]*z[2]*x[3]-x[2]*y[3]*z[0]+z[0]*y[2]*x[3]+y[0]*x[2]*z[3]+x[0]*y[3]*z[2]+x[0]*y[1]*z[3]+x[1]*y[3]*z[0]+x[3]*y[0]*z[1]-z[0]*y[1]*x[3]-x[1]*y[0]*z[3]-y[3]*z[1]*x[0]-x[0]*y[1]*z[2]-x[1]*y[2]*z[0]-x[2]*y[0]*z[1]+z[0]*y[1]*x[2]+x[1]*y[0]*z[2]+y[2]*z[1]*x[0]);
   Ve=abs(Ve);
   
   if (minvalue[tid]>Ve)
   {
    minvalue[tid]=Ve;
    minindex[tid]=k;
   }

  __syncthreads();
     mapvalue_gpu[nodeoffset+tid+temp2*innodecount1_gpu[0]]=minvalue[tid];
     mapindex_gpu[nodeoffset+tid+temp2*innodecount1_gpu[0]]=minindex[tid];
}
}
使用道具 举报 回复
发新帖
您需要登录后才可以回帖 登录 | 立即注册