c - 避免 CUDA 字符串搜索中的分支发散

Question

我想知道如何避免在使用 CUDA 进行字符串搜索时出现分支分歧，以及是否有好的方法可以做到这一点。

目前我试图让 Knuth Morris Pratt 适应 GPU，但我相信存在很多分歧，因为每个线程都在寻找 N 个字母并每次比较这些字母是否对应于单词 I' 的第一个字母我正在寻找。

int tid = blockDim.x * blockIdx.x + threadIdx.x;
int startId = tid * 64;
int x = 0;
for(int i = 0; i < 64; i++){
    if(array[startId + i] == 'C'){
        x++;
    }
}

如果我使用这个虚拟代码来查找字母“C”，但我也可以再看一遍以搜索更多字母。

Answer 1

您可以尝试将比较结果直接添加到值中，如下所示:

x+= (数组[startId + i] == 'C');

但我相信这可能仍然存在分支。我的解决方案是将 block 中的数组值存储到共享内存，然后为 block 中的每个线程分配一个所需的字符，并将结果放入它们自己的共享内存空间，然后减少。

__shared__ char l_array[BLOCK_SIZE];
__shared__ char l_results[BLOCK_SIZE];

int bid = blockDim.x * blockIdx.x;
int lid = threadIdx.x;
int tid = bid + lid;
int x=0;

char desired_char = get_character(lid);


l_array[lid] = -1;


//Store global values in shared memory
if(tid < array_size){
    l_array[lid] = array[tid];       
}

__syncthreads();

//Check local memory for desired character
for(int i = 0; i < BLOCK_SIZE; i++)
   x+=(l_array[i] == desired_char);

//Store results into shared memory
l_results[lid] = x;

__syncthreads();
//Then reduce (poorly)
if(lid==0){
    for(int i = 0; i < BLOCK_SIZE; i++)
        x+= l_results[i];
}

虽然我不知道算法本身，但我只是在猜测，但这里的一些内容可能会帮助您解决这个问题。