some questions about local memory misses and texture use

[amanda_sheng]

these days i have some headach questions on the global memory bandwidth. My device is Tesla C2050 os:win7 64bit

when i use texture memory, the Computer visual profiler report

       <b>Achieved global memory throughput:  109.46 ( Peak global memory throughput(GB/s):  144.00 )</b>

but the time cost is higher than in the case without texture, more important is that the compute visual profiler report that:

             <b> The replays due to local memory cache misses are high</b>

if i use global memory directly, the time cost is lower and there is no local memory cache misses but the global memory bandwidth is very low comparing with peak value.

Is there something i can do to improve the global memory bandwith,causing i think i did use coalescing type to access global memory.

some code below:

global void BackProj(float* _Img_Memory,unsigned int _startAngle,unsigned int _filtLen)

{

//__shared__ float sdata[256];

//__shared__ float backShare[256*8];

unsigned int betaID = blockIdx.y/(Img_size>>8);

//unsigned int betaID = blockIdx.y;

float beta = (fmodf((betaID+_startAngle),FullAng))/180.0f*PI;

unsigned int X_id;

UINT Y_id;

unsigned int _R_ID,ID_in_Memory;

float _R;

float X_beta,Y_beta;

float Lbeta;

float sbeta,cbeta;

// X_id = threadIdx.x;

// sdata[X_id]=_weightProj[betaID*_filtLen+X_id];

__sincosf(beta,&sbeta,&cbeta);

// __syncthreads();

X_id = threadIdx.x + blockIdx.x * blockDim.x;

cbeta =cbeta*0.2f;

sbeta = sbeta *0.2f;

//sbeta=__sinf(beta);

//cbeta=__cosf(beta);

#pragma unroll

//==========Y_id = i;

for (unsigned int i=0;i < 256;i++){

	Y_id = ((blockIdx.y<<8)+i);	

	ID_in_Memory = X_id+ Y_id*Img_size;

	Y_id = ((blockIdx.y<<8)+i)&(Img_size-1);	

	

	Y_beta = D-(sbeta*(256.5f-X_id)-cbeta*(256.5f-Y_id));

	X_beta = (256.5f-X_id)*cbeta+(256.5f-Y_id)*sbeta;

	Lbeta = 1.0f/Y_beta;

	//ID_in_Memory = X_id+ ((blockIdx.y<<9)+i)*Img_size;

	X_beta = atan(X_beta*Lbeta);

	// when 2^126< y < 2^128, the results of __fdividef(x,y) is wrong

	_R = __fdividef(X_beta,deltaR) + det_center;

	_R_ID = int(_R);

	_R = _R - _R_ID;

	

//	_Img_Memory[ID_in_Memory] = Lbeta * Lbeta *((1.0f-_R)*_weightProj[betaID*_filtLen + _R_ID] 

//																		+ _R*_weightProj[betaID*_filtLen + _R_ID+(UINT)1]);

	

	_Img_Memory[ID_in_Memory] = Lbeta * Lbeta *((1.0f-_R)*tex2D(texRef,_R_ID,betaID) + _R*tex2D(texRef,_R_ID+(UINT)1,betaID));

	// BACKSHARE(threadIdx.x+i*blockDim.x)= Lbeta * Lbeta *((1.0f-_R)*SDATA(_R_ID) + _R*SDATA(_R_ID+1));

//	_Img_Memory[ID_in_Memory]=Lbeta * Lbeta *((1.0f-_R)*sdata[_R_ID] + _R*sdata[_R_ID+1]);

	

}

}