將「for」循環轉換爲CUDA並行代碼

Question

我正在轉換CUDA並行化代碼中最初編寫的C代碼的過程。 還是一個新手，我把大部分代碼轉換成了CUDA，但是我的一些內核沒有正確地完成這項工作。

這裏是我的內核：

__global__ void kernel(long int *neighbour, double *f, double *r, double *b, double *fn, double *rn, double *bn, int nfluidsite){ 

int ns = blockDim.x * blockIdx.x + threadIdx.x; 

    if(ns<nfluidsite) 
{ 
    double tempr = r[ns]; 
    double tempb = b[ns]; 
    rn[ns]=tempr; 
    bn[ns]=tempb; 
    for(int q=1;q<Q;++q) 
    { 
    double confr=r[q*NSITE+ns]; 
    double confb=b[q*NSITE+ns]; 
    __syncthreads(); 
    int ns1=neighbour[q*NTOTAL+ns]; 
    __syncthreads(); 
    rn[q*NSITE+ns1]=confr; 
    bn[q*NSITE+ns1]=confb; 
    } 
} 

if(ns<NSITE) 
{ 
    for(int q=0;q<Q;++q) 
    { 
     double rqns = rn[q*NSITE+ns]; 
     double bqns = bn[q*NSITE+ns]; 
    __syncthreads(); 
    r[q*NSITE+ns]=rqns; 
    b[q*NSITE+ns]=bqns; 
    f[q*NSITE+ns]=rqns+bqns; 
    } 
} 

} 

所以，這個代碼工作正常（雖然它不是在所有優化），但我也希望進行並行內的循環上q。所以，我是這樣的：

int ns = blockIdx.x; 
    int q = threadIdx.x; 

，我開始了我的內核如下：

blocksPerGrid = NSITE; 
threadsPerBlock = Q; 
kernel<<<blocksPerGrid,threadsPerBlock>>>(neighbourCu, fCu, rCu, bCu, fnCu, rnCu, bnCu, nfluidsite); 

而且它不會在所有的工作，CUDA不會對數組返回任何錯誤，但操作是隨機的...我在完全並行版本中添加了__syncthreads()命令，但它並沒有解決這些差異。

而且，我不爲什麼，但如果我使用超過1024個線程，在我的內核中的指令也運行隨機...

嗯，我一直不解兩週內，如果有人看到我需要做的，請給我一個提示！

Answer 1

我認爲你不需要__syncthreads();你已經放置它們，你正在分配和讀取線程本地（寄存器或本地內存）變量。

在另一邊肯定需要調用__threadfence();的代碼塊前開始：

if(ns<NSITE) { 
    ... 
} 

的違規聲明：

rn[q*NSITE+ns1]=confr; 
    bn[q*NSITE+ns1]=confb; 

前：

double rqns = rn[q*NSITE+ns]; 
    double bqns = bn[q*NSITE+ns]; 

之後。

由於一個線程寫入一些全局內存和一個不同讀取它，你需要至少兩個放在中間：

__threadfence(); 
__syncthreads(); 

（更多信息here）。只有rn和bn被同一個塊中的線程修改，這才能正確工作。如果碰巧塊外部的線程可以修改它們，這還不夠：在繼續之前，需要保證所有塊都已經到達該點（注意：__syncthreads()只保證同一塊內的線程已經達到了這一點 - 它是塊局部障礙）。你有三種選擇：

1. 在這一點上將內核拆分到兩個不同的內核中 - 單獨的內核調用被隱式地同步。我會從這開始，看看它是否有效。

2. 如果NSITE只是一個線程塊，就像看起來一樣，您可以再次遵循全局總和here的示例：只有最後一個塊會執行 - 其餘所有其他塊都會跳過。

3. 在這一點上實現了一個全球性障礙（這不是一件容易的事情 - 但如果你在谷歌周圍，你會發現很好的參考）。