写的一个加速程序，可以正常运行，表现也可以，只是用profiler测试时，global S/L效率始终很低（25%/6%）, 但程序中仅有的几次对global的访问可以确定是合并的，不得其解。
程序比较长，我把关键的部分贴一下吧，重点关注标注出来的有关全局内存访问部分就可以。
其中 #define DOUBLE_OR_FLOAT double , 未见定义的都是放在constant mem里的全局变量。
线程和线程块都是一维的。线程块的大小设为32。
global void Flash_dev(DOUBLE_OR_FLOAT* mPres, DOUBLE_OR_FLOAT* mHC_Zc,
int* cellFlashStatus, int* mCellStatus, DOUBLE_OR_FLOAT* mSat, int* mTrivalK,
DOUBLE_OR_FLOAT* mHC_K, DOUBLE_OR_FLOAT* xcp)
{
shared DOUBLE_OR_FLOAT x[tpbNCOMPS], y[tpbNCOMPS], zshare[tpbNCOMPS], Kshare[tpbNCOMPS];
int ID = blockIdx.x * blockDim.x + threadIdx.x;
int ib = ID;
if( ib>=nCells[0] ) return;
//copy global 2 shared，全局变量保证了合并访问
for (int i=0;i<nHCComps[0];i++) {
y[i*tpb+threadIdx.x] = xcp[nCells[0]*i+ib];
x[i*tpb+threadIdx.x] = xcp[(nHCComps[0]+i)*nCells[0]+ib];
zshare[i*tpb+threadIdx.x] = mHC_Zc[nCells[0]*i+ib];
Kshare[i*tpb+threadIdx.x] = mHC_K[nCells[0]*i+ib];
}
__syncthreads();
bool inputK=false,trivalK;
DOUBLE_OR_FLOAT a[2],b[2],c[2],A[2],B[2],Z[2],Phase_MW[2],Den[2],l;
DOUBLE_OR_FLOAT* da_dxi = (DOUBLE_OR_FLOAT*)malloc(sizeof(DOUBLE_OR_FLOAT)*nHCComps[0]*2);
[b] DOUBLE_OR_FLOAT* f = (DOUBLE_OR_FLOAT*)malloc(sizeof(DOUBLE_OR_FLOAT)*nHCComps[0]*2);
DOUBLE_OR_FLOAT pres = mPres[ib];
int oldCellStatus = mCellStatus[ib], fluidType;
if(mTrivalK[ib] != 1) inputK = true; [/b]
[b]
//以下调用的每个device函数都不会访问global mem
if(oldCellStatus == 1) //0 gas, 1 oil, 2 gasoil, 3 singleP
{
fluidType = calcStab(pres,zshare,a,b,c,A,B,Z,da_dxi,
Kshare,x,y,f,Phase_MW,Den,1,inputK,trivalK,ib);
if (fluidType == 3)
fluidType = calcStab(pres,zshare,a,b,c,A,B,Z,da_dxi,
Kshare,x,y,f,Phase_MW,Den,0,false,trivalK,ib);
}
else
{
fluidType = calcStab(pres,zshare,a,b,c,A,B,Z,da_dxi,
Kshare,x,y,f,Phase_MW,Den,0,inputK,trivalK,ib);
}
if (fluidType == 3)
{
trivalK = true;
if(oldCellStatus==1)
fluidType = singlePStatus(A,B,Z);
else
fluidType = oldCellStatus;
}
else if (fluidType == 2)
{
bool flashFlag = flash_final(pres,zshare,a,b,c,A,
B,Z,da_dxi,Kshare,x,y,f,Phase_MW,Den,l,ib);
if(!flashFlag)
{
trivalK = true;
if(oldCellStatus == 1)
fluidType = singlePStatus(A,B,Z);
else fluidType = oldCellStatus;
}
}
free(da_dxi);
free(f);
}[/b]
以上程序测试得到（有效线程数32384，也就是nCell[0]=32384,nHCComps[0]=9,kernal运行时间800ms）
Global S/L efficiency 24.9%/6.3%
Global S/L request through output 610/855 Mb/s
问题就是，不关是效率还是请求吞吐量看起来都不合理。当前版本的程序里甚至没有对global mem的store操作，而所有load也都是连续线程访问连续位置。request load through output = （849+8+4+4）*32384/1024/1024/0.8 = 11.7Mb/s。
这些该怎么解释呢？希望请教请教大家。