关于核函数线程分配问题

遇到问题如下:
需求:
一共有1024个数据包,每个数据包都包含一个长度为64的字符串str[i],我用包中的字符串去匹配一个已知的长度为64的字符串string,将匹配上的包记录在一个数组中
方案:
按包分配线程,一共开启1024个线程,每个线程负责一个包与已知字符串的对比,在核函数中:
#define N 1024
int tid = threadIdx.x + blockIdx.x * blockDim.x;
while (tid < N)
{实现字符串匹配}

问题:
现在我想开启64*1024个线程,每64个线程负责一个包,每个线程负责匹配字符串中的一位,请问核函数内如何实现以64个线程为一组的1024次循环??不能用while ( tid < 64 * N)了呀,请关注。。。。。。。。

LZ您好:

您的1024个数据包是否是连续保存在显存中的,如果是的话,那这只是一个访存问题,并且很好解决。

您可以安排block的时候,安排为一行64个线程,这样一行线程处理一个字符串,这一行的线程threadIdx.y是相同的,然后threadIdx.y不同的线程处理不同行,用threadIdx.y辅助寻址找到每个string的开始位置即可。

您的grid在安排的时候,也可以安排为x和z方向为1,y方向排开的形式,这样用blockIdx.y和threadIdx.y就能完成全局的寻址,具体为glb_y_index=blockIdx.y*blockDim.y+threadIdx.y,由此可以找到每个字符串的开始位置(当然您的字符串需要是连续存储的)。

大致如此,供您参考。

祝您编码顺利~

以及稍微补充一下,CUDA中有提供Warp Vote Function,如果您的字符串比较仅仅是比较相同还是不同,那么这些函数可能会有帮助。

祝您好运~

版主您好!

感谢您的回复,对我很有用,O(∩_∩)O谢谢!

版主,字符串就是单纯的比较相同或者不同,您说的这个函数是类似于CPU中的strcmp(),直接调用就可以了么?

ICE不是这个意思。

warp vote是__all/any/ballot这样的函数。

他可以用来在您的warp内快速交换比较结果,这样可以辅助您决策是否下一步warp需要继续下一次比较操作。

例如您在比较字符串的过程中,可以warp快速比较是否所有的32个4B都比较成立(或者有任何一个不成立),这样您可以知道需要继续比较(或者不需要继续比较)。

这些函数还有一个大版本,block的,分别是:
__syncthreads_and/or/count()

他们和上面的warp vote函数类似,但可以为您提供block级别的决策支持。

感谢您的来访。