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

system · 2012 年10 月 29 日 18:20

我想用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之间的同步会很复杂。十分感谢。

system · 2012 年10 月 30 日 01:35

调用顺序
kernel1
kernel2
kernel3
kernel4

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

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

system · 2012 年10 月 30 日 02:48

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自行根据算法思考决定。

system · 2012 年10 月 30 日 02:55

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

system · 2012 年10 月 30 日 03:07

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

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

system · 2012 年10 月 30 日 04:38

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

system · 2012 年11 月 1 日 19:25

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

system · 2012 年11 月 6 日 22:27

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得到的。我写的代码不对，用问号标注出了不解的地方。求指导。十分感谢。

system · 2012 年11 月 7 日 03:51

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

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

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

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

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

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

system · 2012 年11 月 7 日 08:20

嗯…死循环了…

system · 2012 年11 月 7 日 18:32

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

system · 2012 年11 月 8 日 01:17

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才能实现。所以我建议还是用一维数组吧，然后进行一维、二维互相映射的方式解决二维的访问形式。

system · 2012 年11 月 8 日 03:14

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

system · 2012 年11 月 8 日 03:32

此外，我补充一下ICE版主的帖子：

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

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

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

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

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

system · 2012 年11 月 8 日 23:11

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

system · 2012 年11 月 9 日 04:47

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

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

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

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

祝您编程愉快~

system · 2012 年11 月 19 日 23:48

这是我的两个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

多谢多谢。