用cuda优化一个双重for循环的问题

[加载中_无法显示]

我现在有如下一个程序段，因为n,m,很大的缘故，所以打算用cuda并行优化

for (i=0; i<n; i++)
	{
	    s = (y[i+1]-y[i])/m;
		for(j=1; j<= m; j++)
		{
			x[m*i+j] = (y[i]+j*s);
		}
	}

我想请问，这个双重循环应该怎么优化。
已知n=100 ，m=30000，
我想到先把S全部求出来，再传递给loop，

__global__ void loop(double *x, double *y, double *s, int n, int m)
{
	int tid = threadIdx.x + blockDim.x * blockIdx.x;
	while(tid <m*n)
		x[tid] = y[blockIdx.x] + s[blockIdx.x] * threadIdx.x;
}

启动核函数，
loop<<<n,m>>>(x,y,s,n,m);
我总感觉哪里不对，希望大家指正。

[加载中_无法显示]

此处n=100,m=28858

[yuanwcj]

楼主的想法没错,用cuda确实可以获得比较好的处理效果,但是简单看了遍,楼主的代码有2个问题 1. kernel行代码loop<<<n,m>>>(x,y,s,n,m)中,m=30000,这个有问题. n为grid中block的个数,m为block中线程的个数,不同版本的cuda中,m的个数大小不一样,但目前最大为2048,即每个block块中最大可能的线程个数为2048,超出这个值,kernel会出错直接返回. 2. int tid = threadIdx.x + blockDim.x * blockIdx.x;这一行有问题 blockDim.x为block中x维线程的个数,按照你的要求应该是和m对应的,但是实际上你的kernel代码是没有体现这一点的..,建议你再综合我说的第一条仔细想清楚..

[加载中_无法显示]

引用 2 楼 yuanwcj 的回复:

楼主的想法没错,用cuda确实可以获得比较好的处理效果,但是简单看了遍,楼主的代码有2个问题 1. kernel行代码loop<<<n,m>>>(x,y,s,n,m)中,m=30000,这个有问题. n为grid中block的个数,m为block中线程的个数,不同版本的cuda中,m的个数大小不一样,但目前最大为2048,即每个block块中最大可能的线程个数为2048,超出这个值,kernel会出错直接返回. 2. int tid = threadIdx.x + blockDim.x * blockIdx.x;这一行有问题 blockDim.x为block中x维线程的个数,按照你的要求应该是和m对应的,但是实际上你的kernel代码是没有体现这一点的..,建议你再综合我说的第一条仔细想清楚..

我按照你说的再去看看，感谢你的指导！

[yuanwcj]

引用 3 楼 asxqy714 的回复:

我按照你说的再去看看，感谢你的指导！

不客气,有问题可以继续询问 .刚开始接触cuda的时候,确实会有很多细节要注意 但是一通百通,知道原因之后就不会再犯类似的错误了

[加载中_无法显示]

引用 4 楼 yuanwcj 的回复:

[quote=引用 3 楼 asxqy714 的回复:] 我按照你说的再去看看，感谢你的指导！

不客气,有问题可以继续询问 .刚开始接触cuda的时候,确实会有很多细节要注意 但是一通百通,知道原因之后就不会再犯类似的错误了[/quote] 我将之前的代码进行了修改，先不考虑内存访问优化的问题。 我想着让每个线程计算n次，考虑到我的显卡1.2的计算能力，每个block分配128个线程，有6个SM，每个SM最多1024线程，所以分配48个block。

__global__ void loop(double *x, double *y, double *s, int n, int m)
{
	int tid = (threadIdx.x + 1) + blockDim.x * blockIdx.x;

	for (int i=0; i<n; i++ )
	{
		while (tid<=m)
		{
			x[m*i+tid] = y[i] + s[i] * tid;

			tid += blockDim.x*gridDim.x;

		}
	}
}
	// 启动核函数
	dim3 nblocks(48,1,1);
	dim3 nthreadsperblock(128,1,1);
	loop<<<nblocks , nthreadsperblock>>>(dev_x,dev_y,dev_s,n,m);	

关于一个线程计算n次这部分代码，我心里还有疑惑，不知道这样实现是否正确？

[yuanwcj]

引用 5 楼 asxqy714 的回复:

[quote=引用 4 楼 yuanwcj 的回复:] [quote=引用 3 楼 asxqy714 的回复:] 我按照你说的再去看看，感谢你的指导！

不客气,有问题可以继续询问 .刚开始接触cuda的时候,确实会有很多细节要注意 但是一通百通,知道原因之后就不会再犯类似的错误了[/quote] 我将之前的代码进行了修改，先不考虑内存访问优化的问题。 我想着让每个线程计算n次，考虑到我的显卡1.2的计算能力，每个block分配128个线程，有6个SM，每个SM最多1024线程，所以分配48个block。

__global__ void loop(double *x, double *y, double *s, int n, int m)
{
	int tid = (threadIdx.x + 1) + blockDim.x * blockIdx.x;

	for (int i=0; i<n; i++ )
	{
		while (tid<=m)
		{
			x[m*i+tid] = y[i] + s[i] * tid;

			tid += blockDim.x*gridDim.x;

		}
	}
}
	// 启动核函数
	dim3 nblocks(48,1,1);
	dim3 nthreadsperblock(128,1,1);
	loop<<<nblocks , nthreadsperblock>>>(dev_x,dev_y,dev_s,n,m);	

关于一个线程计算n次这部分代码，我心里还有疑惑，不知道这样实现是否正确？[/quote] 楼主 的意思可否理解成如下？通过48*128个线程，每个线程做28858/(48*128)个点的计算？ 这种方法本身是没有问题的，但是代码部分却是可以斟酌的 实际上，楼主的写法仍然是双重循环，这种写法的效率是相当低的。 可以直接写成如下形式： for(int i=0;i<m;i+=blockDim.x*gridDim.x) x[m*i+tid] = y[i] + s[i] * tid; 这和楼主的代码处理结果实际是等效的，但是在运算效率上却会提高很多 另外，有几个问题建议楼主注意： 1.需要处理的次数m=28858并不是线程个数P=48*128的整数倍，因此需要对部分数据做额外处理，这里有2种方法 一种是if(tid<m)的形式，通过对线程id来判断 另一种是数据补零的方式，在m后面补若干个0，保证每个线程处理的个数均是相同的,即每个线程处理(m+P-1)/P个点的数据 2.楼主的代码中用到了x=y+n*tid形式的数据，不建议楼主才用这种形式，如果是在share memory中，会造成严重的bank conflict 此外，楼主的代码创建核函数时并不需要使用3维网格形式，2维足以