最近换了新电脑，发现个郁闷的问题。同样一个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=temp2loopnum1_gpu[0];
const int nodeoffset=temp1blockDim.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+2elementcount2_gpu[0]];
nodeindex[3]=element2_gpu[k+3elementcount2_gpu[0]];
x[0]=node2_gpu[nodeindex[0]];y[0]=node2_gpu[nodeindex[0]+nodecount2_gpu[0]];z[0]=node2_gpu[nodeindex[0]+2nodecount2_gpu[0]];
x[1]=node2_gpu[nodeindex[1]];y[1]=node2_gpu[nodeindex[1]+nodecount2_gpu[0]];z[1]=node2_gpu[nodeindex[1]+2nodecount2_gpu[0]];
x[2]=node2_gpu[nodeindex[2]];y[2]=node2_gpu[nodeindex[2]+nodecount2_gpu[0]];z[2]=node2_gpu[nodeindex[2]+2nodecount2_gpu[0]];
x[3]=node2_gpu[nodeindex[3]];y[3]=node2_gpu[nodeindex[3]+nodecount2_gpu[0]];z[3]=node2_gpu[nodeindex[3]+2nodecount2_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]+2nodecount2_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]+2nodecount2_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]+2nodecount2_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+temp2innodecount1_gpu[0]]=minvalue[tid];
mapindex_gpu[nodeoffset+tid+temp2innodecount1_gpu[0]]=minindex[tid];
}
}