怎样实现多个kernel之间的运算?

我想用4个 KERNEL,第一个kernel用来存一个一维数组叫database 大小是N,,第二个kernel也用来存一个一维数组叫sample 大小是M,这里的N 远远大于M,第三个kernel用来产生一个二维矩阵s0的大小是 【M N】M是行,N是列, S0 里面的每个元素是根据以下规则确定的,将database看成行向量,sample看为列向量,S0的每一个元素是比较sample里面每一个元素是否和database里面每一个元素是否相等后的结果,相等就是一个正数值,不相等就是一个负数值。第四个kernel用来产生另外一个矩阵S,里面的每一个元素值是由以下步骤确定的,例如S(i, i)的值等于max{S0(i, j), S0(i-1, j), S0(i, j-1)}。如何利用kernel之间的同步完成上述功能?或者有什么别的更好的方法。因为block之间的同步会很复杂。十分感谢。

调用顺序
kernel1
kernel2
kernel3
kernel4

如果你的线程比较少,那么可以考虑kernel1和kernel2并发,也就是concurrent kernel
而kernel3和kernel4也是可以并发的,所以只需要并发kernel1和kernel2,再并发kernel3和kernel4。

如果你的kernel能够占用所有的SM资源,那么直接按顺序调用就好了!

1:生成矩阵一般做好用cudaMemcpy复制到GPU显存就可以了,不用写成kernel,除非是你需要并行生成数据才需要写kernel在显存里面生成数据。我估计你前面两个矩阵是某种现成的测量数据,所以我认为kernel1,2是不需要的。

2:比较生成S0,这个需要写一个kernel。这个计算过程大致考虑了一下,是比较容易实现的,访存效率也比较容易保证,访存数据很容易被合并访问或广播。比如一个block给M个线程(如果M有一两百个数据规模),每个线程先读入M中对应位置的数据,然后和N中的数据对比,写入值。每个block比如说计算N中1000个数据。

3:生成S,直接上M行N列的线程(划分为若干个block),直接读取比较,合并访问问题不大,cache应该有很大帮助。

4:如有需要,也可以考虑生成S0和S合并到一起实现,不过似乎较为麻烦,请LZ自行根据算法思考决定。

另外,我认为LZ叙述中的kernel3,4是不能并发的。因为kernel4依赖于kernel3的数据,而并发的时候kernel3,4分别在两个流里面,并不保证两者之间的相互顺序。

另外,觉得LZ有点文不对题。kernel计算完成的时候是自动同步的,block内部可以用syncthreads同步,而如果需要block同步,那么可以通过kernel结束来同步。

LZ的问题,前两个kernel可能是不必要的,如果确需kernel实现,是可以采用kernel并发的,不过这种只调用一次的kenrel,规模也不是很大,并发意义也不大。后两个kernel必须是顺序的,因为有结果依赖,自己kernel计算结束的时候,就保证的结果已经完成。如果问题有一定规模,那么GPU占用率应该还是可以的。

嗯,没仔细看。。。kernel4必须在kernel3之后调用

嗯KERNEL 4 的S是要根据KERNEL3 计算的,我是初学,探索阶段,头疼,最近NY飓风,晚点附上CODE

global void myKernel1( char** gpu_sample, char** gpu_data, float **gpu_s0) // 888888
{
dim3 dimGrid1;
dim3 dimBlock1;
dimBlock1.x = dimBlock1.y = BLOCK_SIZE;
dimGrid1.x = dimGrid1.y = GRID_SIZE;

int i1 = threadIdx.x + blockIdx.x * dimBlock1.x;

int j1 = threadIdx.y + blockIdx.y * dimBlock1.y;

if( i1 > N || j1 > M ) return;

while ( j1 < (M+1) && i1 < (N+1) )
{
if(gpu_sample[0][j1] == gpu_data[i1][0])
{

		gpu_s0[i1+1][j1+1] = 5;
	}
	else 
								
		gpu_s0[i1+1][j1+1] = -3;

}

}

global void myKernel2( float **gpu_s, float **gpu_s0)
{
dim3 dimGrid2;
dim3 dimBlock2;
dimBlock2.x = dimBlock2.y = BLOCK_SIZE;
dimGrid2.x = dimGrid2.y = GRID_SIZE;

/* shared float sTemp0[ threadsPerBlock ][threadsPerBlock];//??? */

int i2 = threadIdx.x + blockIdx.x * dimBlock2.x;

int j2 = threadIdx.y + blockIdx.y * dimBlock2.y;
float w = -4;
float zero = 0;

gpu_s[i2][0] = gpu_s[0][j2] = 0;  // 初始化矩阵S的第一行第一列为 0。


/*if ( j2 < (M+1) && i2 < (N+1) )
sTemp0[threadIdx.x][threadIdx.y] = gpu_s0[i2][j2]; //????????
__syncthreads();*/

   	
if( i2 > N || j2 > M ) return;
while ( j2 < (M+1) && i2 < (N+1) )
{
					
	gpu_s[i2][j2] = max(gpu_s[i2-1][(j2-1)] + gpu_s0[i1+1][j1+1], //?????????
					  gpu_s[i2][(j2-1)] + w, 
					  gpu_s[(i2-1)][j2] + w, 
					  zero);
}

}

这是我写的两个kernel的代码。 矩阵S0在kernel 1 中产生,怎样在kernel 2 中调用S0矩阵。在kernel 2 中我想算出一个矩阵S, 这个矩阵是由S0得到的。我写的代码不对,用问号标注出了不解的地方。求指导。十分感谢。

1:先不说实现效率和逻辑正确性如何,你两个kernel里面的while循环都死循环了,线程索引使while条件为真的线程会一直执行该循环,由于循环内部也没有跳出循环的判断,所以会运行到地老天荒。

2:第一个kernel里面你的gpu_sample可以只开辟为一维数组即可,当然这个倒不算错误。

3:第二个kernel里面,你注释掉了shared memory的声明和使用部分,并注释为疑问。这属于算法实现问题,无法直接回答,按照顶楼的简要叙述,也可以不使用shared memory。

4:第二个kernel中,你对max函数注释为疑问,由于文中没有max函数的定义以及我们无从知晓你算法细节,所以此问题也无法回答。

5:此外,你两个kernel函数的参数都使用了二级指针,这应该是不正确的。

最后,建议LZ提供更多的代码信息,算法描述信息和调试信息等,以便更好地解决问题。

嗯…死循环了…

我在程序的前面定义了M等于40, N等于400.你所说的进入了死循环是指没有给出M和N?这里我不太理解,为什么会死循环?另外,我写了一个deviced_kernel函数来计算max. 这部分应该没有错误。我想知道在kernel 1中产生的s0怎么用在kernel 2当中来产生s矩阵。就是说怎么在kernel 2 中调用kernel 1 中产生的s0?现在程序编译的时候没有错误,只有warning。

int i2 = threadIdx.x + blockIdx.x * dimBlock2.x;
int j2 = threadIdx.y + blockIdx.y * dimBlock2.y;

while ( j2 < (M+1) && i2 < (N+1) )
{

gpu_s[i2][j2] = max(gpu_s[i2-1][(j2-1)] + gpu_s0[i1+1][j1+1], //???
gpu_s[i2][(j2-1)] + w,
gpu_s[(i2-1)][j2] + w,
zero);
}

比如说这里,你的线程号(i,j)总有落在m+1和n+1范围内的吧,比如线程号为(0,0),那这个循环要怎么跳出?

还有,CUDA对于二维数组的支持并不是很好,你在内核里用char **的方式,要使用正确的cudaMalloc才能实现。所以我建议还是用一维数组吧,然后进行一维、二维互相映射的方式解决二维的访问形式。

这个问题我马上就去FIGURE OUT。我想知道在kernel 1中产生的s0怎么用在kernel 2当中来产生s矩阵。就是说怎么在kernel 2 中调用kernel 1 中产生的s0?我用的二维数组,我确定是对的,用的CUDAMALLOC分配空间,但关键问题是怎么在kernel 2 中调用kernel 1 中产生的s0? 跪求!

此外,我补充一下ICE版主的帖子:

楼主您真的确定是要使用float **gpu_s0吗?

如果真的如此,您的gpu_s0 [ i ] [ j ]真的是您想要的顺序吗?(一般都是行优先的,不过这个无所谓,先假设你的i,j无实际行、列意义)。

如果还真的如此,您的float **gpu_s0需要(1)先指向行、列指针的数组,(2)然后里面的每个指针指向每行、列,然后您才能访问元素。

如果只是简单的分配了行sizeof(元素)大小的数组,建议使用float *,而不是float **, 后者需要继续额外分配一下指针的数组,并初始化这些指针们。

祝您调试愉快。真是对ICE的补充。[/i]

求问,什么叫做一维数组二维数组的映射,能否给个例子啊。我不是很清楚这个概念是什么意思。
多谢。还有重点问题是怎么在kernel 2 中调用kernel 1 中的S0。这是最关键的问题。多谢多谢。

一维数组的二维映射简单说就是申请一维的数组或者malloc一段一维的连续存储空间,自己实现二维索引的寻址和访问。

关于在kernel2中调用kernel1生成的s0,大致流程如下:

step1:启动kernel1,生成s0,结束kernel1。
step2:启动kernel2,此时将s0的指针通过参数传递给kernel2,在kernel2中使用s0,结束kernel2。

就是一般的通过指针访问数组。

祝您编程愉快~

这是我的两个kernel的代码。编译后,有warning,但是没有error。另外,我的老板说。kernel 1 里面的s0是可以直接用在 kernel 2 里面。我把warning 贴出来,大家帮忙看看。这要怎么改。多谢多谢。warning的地方我在代码中标注出来了。求指点。

global void myKernel1( char** gpu_sample, char** gpu_data, float **gpu_s0)
{
dim3 dimGrid1;
dim3 dimBlock1;
dimBlock1.x = dimBlock1.y = BLOCK_SIZE;
dimGrid1.x = dimGrid1.y = GRID_SIZE;

int i1 = threadIdx.x + blockIdx.x * dimBlock1.x;

int j1 = threadIdx.y + blockIdx.y * dimBlock1.y;

if( i1 > N || j1 > M ) return;

while ( j1 < (M+1) && i1 < (N+1) )
{
if(gpu_sample[0][j1] == gpu_data[i1][0]) //warning
{

		gpu_s0[i1+1][j1+1] = 5;
	}
	else 
								
		gpu_s0[i1+1][j1+1] = -3;

i1 += blockDim.x * gridDim.x;
j1 += blockDim.y * gridDim.y;
}

}

global void myKernel2( float **gpu_s0, float **gpu_s )
{
dim3 dimGrid2;
dim3 dimBlock2;
dimBlock2.x = dimBlock2.y = BLOCK_SIZE;
dimGrid2.x = dimGrid2.y = GRID_SIZE;

float w = -4;									
float zero = 0;

shared float shared[ threadsPerBlock ][threadsPerBlock];
int i2 = threadIdx.x + blockIdx.x * dimBlock2.x;
int j2 = threadIdx.y + blockIdx.y * dimBlock2.y;

while( j2 < (M+1) && i2 < (N+1) )
{
  shared[threadIdx.x][threadIdx.y] = gpu_s0[i2][j2];  	// warning
  
  i2 += blockDim.x * gridDim.x;
  j2 += blockDim.y * gridDim.y;
}

__syncthreads();
	

if( j2 < (M+1) && i2 < (N+1) )
gpu_s[i2][0] = gpu_s[0][j2] = 0;


/*if ( j2 < (M+1) && i2 < (N+1) )
sTemp0[threadIdx.x][threadIdx.y] = gpu_s0[i2][j2]; //????????
__syncthreads();*/

   	
if( i2 > N || j2 > M ) return;
while ( j2 < (M+1) && i2 < (N+1) )
{
					
	gpu_s[i2][j2] = max(gpu_s[i2-1][(j2-1)] + shared[threadIdx.x][threadIdx.y], //?????????
					  gpu_s[i2][(j2-1)] + w, 
					  gpu_s[(i2-1)][j2] + w, 
					  zero);		//warning
    i2 += blockDim.x * gridDim.x;
	j2 += blockDim.y * gridDim.y;
}

}

warning 提示:

./test_10_15_2012.cu(155): Warning: Cannot tell what pointer points to, assuming global memory space
./test_10_15_2012.cu(155): Warning: Cannot tell what pointer points to, assuming global memory space
./test_10_15_2012.cu(155): Warning: Cannot tell what pointer points to, assuming global memory space
./test_10_15_2012.cu(186): Warning: Cannot tell what pointer points to, assuming global memory space
./test_10_15_2012.cu(208): Warning: Cannot tell what pointer points to, assuming global memory space
./test_10_15_2012.cu(208): Warning: Cannot tell what pointer points to, assuming global memory space
./test_10_15_2012.cu(208): Warning: Cannot tell what pointer points to, assuming global memory space
./test_10_15_2012.cu(208): Warning: Cannot tell what pointer points to, assuming global memory space

多谢多谢。