kernel之间如何进行运算

[thinkincuda123]

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

[linxxx3]

顺序调用四个kernel应该就行了吧，每个kernel把自己的全部工作都完成，再调下一个，用不着同步。
当然更精细的方法肯定是有的。

[angelababy88]

有具体点的做法吗？

[linxxx3]

从lz你的描述看，要完成功能，就按你写的，一步步调用kernel就可以了，用不着显式的同步。好好开发每个kernel内部的并行性，或许性能就不错了。

上面功能的描述写的已经算清楚了，lz编码上应该不会有问题了我猜，那还需要说明更具体的什么哈？或者lz执行之后，发觉性能有瓶颈了，咱们再讨论

[thinkincuda123]

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

[linxxx3]

你的代码跟你第一段的描述不一样啊，最后一个？行，S矩阵的值不仅依赖S0，还依赖S矩阵当前自身的值？你确定是这个逻辑么？这跟你第一段里说的“例如……”有本质的差别。 有依赖自身的值，意味着运算有内在的串行性。想要利用cuda并行，需要自己分析依赖关系，尽可能发掘可以并行的部分了，同步机制可能需要精细设计。抱歉没时间深入看了，lz要花些功夫了

[thinkincuda123]

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

[linxxx3]

可能是二级指针的原因，中间一级指针的值是不确定的，这就是warning提示的意思，编译器认为有风险。特别是如果主机是32位系统，这个可能要特别小心，因为CUDA的device指针是64位的，很可能在二级指针出问题。 保险的做法是，给kernel传参数，统一用一级指针，kernel内计算偏移地址。当然warning不一定意味着出错，能用就好了。 另，判断浮点数的相等最好不要用“==”，除非你十分确定里面存的是整型。

[thinkincuda123]

请问怎么用一级指针代替二级指针？怎么在kernel 内部计算偏移量？ 这两个kernel 函数的算法有问题么?还是只是二级指针的问题？ 多谢帮忙。！

[linxxx3]

先试试这个二维数组的方法结果是不是对的，对了就无所谓了。算法对不对是另外的问题了，当然也要花时间调。 如果确定是二维数组有问题，转一维的也容易，第一个元素的地址赋值给一级指针，后面二维数组的每一行的起始地址、每个元素的地址，都是可以计算出来的，lz自己推算下吧。我对这算法不熟，写出来可能就错了。