性能优化求助,shared memory 问题

[youleelove]

现在的情况是：每个block内16X16线程，每个线程取上下左右四个值来处理，所以shared memory要读18X18。

我打算让四个边框上的线程来读取shared memory里的边框上的数据，但是这样要用到四个条件判断,if(tx==0),if(ty==0). if(tx == blocksize-1), if(ty = blocksize-1).

但是代码跑起来后结果不对,不知道为何错误. 同时觉得应该有更好的办法解决shared memory的读入问题,求助版上的高手.
下面是未经优化的代码:

__global__ void Laplace_d (float *A, float *B, int *N_t){

	int tx = threadIdx.x; int ty = threadIdx.y;
	
	int j = blockIdx.x * blockDim.x + tx;
	int i = blockIdx.y * blockDim.y + ty;
	int index, left, right, top, bottom;
	int N=*N_t;

	index = i*N +j;
	left = i*N+ j-1;
	right = i*N+ j+1;
	top = (i-1)*N +j;
	bottom = (i+1)*N+j;
	if(i>0 && i<N-1 && j>0 && j<N-1){
		B[index]=0.25*( A[left]+A[right]+A[top]+A[bottom])*0.9+0.1*B[index];
	}
}

[youleelove]

能用一小段代码演示下怎么用18X18 来处理16X16的数据么？如果18X18的block size的话，那么就达不到full occupancy了吧？用16X16的block来处理14X14的数据呢？

[youleelove]

//I did a new version and this one works, so any comments are welcome. 
__global__ void Laplace_d (float *A, float *B, int *N_t){

	int tx = threadIdx.x; int ty = threadIdx.y;
	
	int j = blockIdx.x * blockDim.x + tx;
	int i = blockIdx.y * blockDim.y + ty;
	int index;

	int N=*N_t;

	index = i*N +j;

	//read data from global memory to shared memory
	__shared__ float sdata[BLOCK_SIZE][BLOCK_SIZE];
	sdata[ty][tx] = A[index];
	__syncthreads();

	if(i>0 && i<N-1 && j>0 && j<N-1){
		float left = 0, right =0, top = 0, bottom =0;
	
		//top
		if(ty > 0)
			top= sdata[ty-1][tx];
		else top = A[index-N];

		//left
		if(tx > 0)
			left = sdata[ty][tx-1];
		else left = A[index-1];

		//right
		if(tx < BLOCK_SIZE-1)
			right = sdata[ty][tx+1];
		else right = A[index+1];

		//bottom
		if(ty < BLOCK_SIZE-1)
			bottom = sdata[ty+1][tx];
		else bottom = A[index+N];
		
		B[index]=0.25*( left+right+top+bottom)*0.9+0.1*sdata[ty][tx];
	}
}

[无心人_过过小日子]

16x16计算14x14也行,但gmem读取时不对齐了,那个更好些,只有试验了.

[无心人_过过小日子]

[Quote=引用 5 楼 l7331014 的回复:]
16x16计算14x14也行,但gmem读取时不对齐了,那个更好些,只有试验了.
[/Quote]

有些没考虑仔细.请无视.呵呵.

[youleelove]

回Cyrosly：
我需要在主函数里这样进行调用：

while (k<ITERATIONS){
		Laplace_h(cpu_data2, cpu_data,N);
		Laplace_h(cpu_data, cpu_data2,N);
		k+=2;
	}

所以这里绑定成texture 貌似行不通，我对texture不太懂，正在看资料。

[youleelove]

[Quote=引用 7 楼 youleelove 的回复:]
回Cyrosly：
我需要在主函数里这样进行调用：C/C++ codewhile(k<ITERATIONS){
Laplace_h(cpu_data2, cpu_data,N);
Laplace_h(cpu_data, cpu_data2,N);
k+=2;
}
所以这里绑定成texture 貌似行不通，我对texture不太懂，正在看资料。
[/Quote]
代码错了，搞成CPU的了，下面是kernel的：

while (k<ITERATIONS){
		Laplace_d<<<dimGrid, dimBlock>>>(device1, device2,n);
		Laplace_d<<<dimGrid, dimBlock>>>(device2, device1,n);
		k+=2;
	}

[无心人_过过小日子]

试试下面,没调试过哦.
启动时block数是((N-2)/BLOCKSIZE+1,(N-2)/BLOCKSIZE+1),线程数是(BLOCKSIZE+2,BLOCKSIZE+2).

__global__ void Laplace_d (float *A, float *B, int *N_t){

int tx = threadIdx.x; int ty = threadIdx.y;
int x = tx - 1; int y = ty - 1;

int j = blockIdx.x * BLOCK_SIZE + tx;
int i = blockIdx.y * BLOCK_SIZE + ty;
int index;

int N=*N_t;

index = i*N +j;

__shared__ float sdata[BLOCK_SIZE+2][BLOCK_SIZE+2];
if((i<=N-1)&&(j<=N-1)){
sdata[ty][tx] = A[index];
}
__syncthreads();

if((i<N-1)&&(j<N-1)){
if((x>=0)&&(y>=0)&(x<BLOCK_SIZE)&&(y<BLOCK_SIZE)){
B[index]=0.25*(sdata[ty-1][tx]+sdata[ty+1][tx]+sdata[ty][tx-1]+sdata[ty][tx+1])*0.9+0.1*sdata[ty][tx];
}
}
}

[无心人_过过小日子]

补充说明:
A的坐标变量[j][i],范围是0<->N-1.
sdata的坐标变量[ty][tx],范围是0<->B+1.
B的坐标变量[j][i],范围是1<->N-2.
要计算的点的坐标变量[y][x],范围是0<->B-1.

启动时的数据分块按要计算的点的范围分成BxB的块.

[Cyrosly]

[Quote=引用 7 楼 youleelove 的回复:]
回Cyrosly：
我需要在主函数里这样进行调用：
C/C++ codewhile (k<ITERATIONS){
Laplace_h(cpu_data2, cpu_data,N);
Laplace_h(cpu_data, cpu_data2,N);
k+=2;
}


所以这里绑定成texture 貌似行不通，我对texture不太懂，正在看资料。
[/Quote]
texture可以的

[无心人_过过小日子]

没看到你的错误代码,估计问题如下:
1)每次读取块要和上行读取块/左边读取块没重合(该有一个数据重叠的)?
2)在读取地方if?应该是在计算地方用if判断的.
3)地址关系的对应不正确?

建议线程改为18x18线程,但计算结果还是按16x16分割.这样更方便些.

[Cyrosly]

#define DIMX 16
#define DIMY 16

//用纹理避免边界元素的非合并访问
//B  :D
//A  :绑定到纹理coalesed
//N_t:F

__global__ 
void laplace(float* D,const float* F,const float weight)
{
    int tidx=__mul24(blockDim.x,blockIdx.x)+threadIdx.x;
    int tidy=__mul24(blockDim.y,blockIdx.y)+threadIdx.y;
    const int pitch=__mul24(gridDim.x,blockDim.x);
    const int gloc=__mul24(pitch,tidy)+tidx;

    __shared__ float smem[DIMY+2][DIMX+2];

    tidx=threadIdx.x+1;
    tidy=threadIdx.y+1;
    smem[tidy][tidx]=tex1Dfetch(coalesed,gloc);
    if(threadIdx.x==0){
        smem[tidy][0     ]=tex1Dfetch(coalesed,gloc-1  ); //不需要担心索引越界，因为线型纹理在索引超出范围后会返回0，对结果没有影响
         smem[tidy][DIMX+1]=tex1Dfetch(coelesed,gloc+DIM);
    }
    if(threadIdx.y==0){
        smem[0     ][tidx]=tex1Dfetch(coelesed,gloc-pitch);
        smem[DIMY+1][tidx]=tex1Dfetch(coalesed,gloc+__mul24(pitch,DIMY));
    } __syncthreads();

    D[gloc]=0.25f*weight*(smem[tidy][tidx-1]+smem[tidy][tidx+1]+smem[tidy-1][tidx]+smem[tidy+1][tidx])+(1.f-weight)*F[gloc];
}
    

[Cyrosly]

哎，还不明白那我就写出来

while(k<ITERATIONS){
    bind_texture(tex,data[SRC]);
    laplace_solver(data[DST],N);
    ++k;
}

[Cyrosly]

还要加上最后一句SWAP()交换SRC和DST

[dtx08]

敢问Cryosly是如何在devise中实现SWAP()的？

[dtx08]

[Quote=引用 13 楼 Cyrosly 的回复:]
还要加上最后一句SWAP()交换SRC和DST
[/Quote]

敢问Cryosly是如何在devise中实现SWAP()的？

[Cyrosly]

[Quote=引用 15 楼 dtx08 的回复:]
引用 13 楼 Cyrosly 的回复:
还要加上最后一句SWAP()交换SRC和DST



敢问Cryosly是如何在devise中实现SWAP()的？
[/Quote]

外循环不是在GPU上的

[dtx08]

[Quote=引用 16 楼 Cyrosly 的回复:]
引用 15 楼 dtx08 的回复:

引用 13 楼 Cyrosly 的回复:
还要加上最后一句SWAP()交换SRC和DST


敢问Cryosly是如何在devise中实现SWAP()的？


外循环不是在GPU上的
[/Quote]
这个我自然知道。你是用cudaMemcpy吗？我用的是cudaMemcpy(fromdevicetodevice)测试了下，这样挪移的成本完全抵消了texture的速度。你有快的吗？能否把代码写出来了？

[youleelove]

程序搞定，结果没有错误，Cyrosly谢谢谢谢谢谢，哈哈。

texture < float, 1, cudaReadModeElementType> coalesed;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();

__global__ 
void Laplace_d(float* D,int* N)
{

    int tidx=blockDim.x*blockIdx.x+threadIdx.x;
    int tidy=blockDim.y*blockIdx.y+threadIdx.y;
    const int pitch=*N;
    const int gloc=pitch*tidy+tidx;

    __shared__ float smem[BLOCK_SIZE+2][BLOCK_SIZE+2];

    int sidx=threadIdx.x+1;
    int sidy=threadIdx.y+1;
    smem[sidy][sidx]=tex1Dfetch(coalesed,gloc);
    if(threadIdx.x==0){
        smem[sidy][0     ]=tex1Dfetch(coalesed,gloc-1  ); //不需要担心索引越界，因为线型纹理在索引超出范围后会返回0，对结果没有影响
         smem[sidy][BLOCK_SIZE+1]=tex1Dfetch(coalesed,gloc+BLOCK_SIZE);
    }
    if(threadIdx.y==0){
        smem[0][sidx]=tex1Dfetch(coalesed,gloc-pitch);
        smem[BLOCK_SIZE+1][sidx]=tex1Dfetch(coalesed,gloc+__mul24(pitch,BLOCK_SIZE));
    } 
	__syncthreads();


	if(tidx > 0 && tidx < pitch-1 && tidy > 0 && tidy < pitch -1){//之前的代码忘记了边界情况，后来改正了
		D[gloc]=0.25*0.9*(smem[sidy][sidx-1]+smem[sidy][sidx+1]+smem[sidy-1][sidx]+smem[sidy+1][sidx])+0.1*smem[sidy][sidx];
	}
    
}

最后附上调用函数的代码

while (k<ITERATIONS){

		cudaBindTexture(0,&coalesed, device2,&channelDesc,size);
		Laplace_d<<<dimGrid, dimBlock>>>(device1,n);
		cudaUnbindTexture(&coalesed);

		cudaBindTexture(0,&coalesed, device1,&channelDesc,size);
		Laplace_d<<<dimGrid, dimBlock>>>(device2,n);
		cudaUnbindTexture(&coalesed);
		k+=2;
	}

[youleelove]

sigh~~~,结果反而慢了