避免CUDA字符串搜索中的分支分歧

我想知道如何避免使用CUDA進行字符串搜索時的分支分歧，以及是否有一種好方法。避免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」，但我還可以做第二次看，以尋找更多的字母爲好。

來源

2013-01-07 Anoracx

請顯示一些代碼。 – sgarizvi

@Anoracx：你的鏈接只包含一個簡單的串行實現。你的GPU代碼呢？ – talonmies

這不是我的代碼，我只是試圖在GPU上實現它。 – Anoracx

你可以嘗試添加comparisions的結果直接進入像這樣的值：

X + =（陣列[startId + 1] == 'C'）;

但我相信這可能仍然是分支。我的解決方案是將塊中的數組值存儲到共享內存中，然後將塊中的每個線程分配給所需的字符，並將結果放入它們自己的共享內存空間中，然後減少。

__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]; 
}

雖然我不知道算法本身，但我只是猜測，但沿着這裏的一些線可能會幫助你弄清楚這一點。

來源

2015-09-30 16:52:55 Primalfido

避免CUDA字符串搜索中的分支分歧

回答

相關問題