只使用Shuffle + ballot + popc從選定的車道寫入數據

我在嘗試保存一些共享內存以改善處理時遇到了一些問題。在這段代碼中，每個線程從主機中選擇一個模板索引（t_initdwl），並在下一個內核處理它之前將其擴展到上限。只使用Shuffle + ballot + popc從選定的車道寫入數據

這個想法是使用shuffle + ballot來「保留」適當的空間量，以便只寫入沒有超出上限的線程索引。由於它是一系列限制，因此首先測試每個列極限，然後寫入它。「擴展條目」的數量根據初始索引中的值而變化，但是，一旦線程超過限制，任何進一步的增量都無效，因此它設置play = false（僅用於安全）並返回。

問題是始終是trasnp_line + 1的bCol值，給出的想法是popc在此行中只能正常工作，因爲1只是lineID 0的正確值。我沒有收到任何錯誤，

預期的結果是寫入行的正確位置，考慮到仍在「播放」的行數（尚未返回）。

__global__ void dwLgen(const int maxthreads, short* __restrict__ kpL, int* nkpl, 
     const short* __restrict__ init_dwL, const short rloops){ 
    __shared__ short t_initdwl[1024][24]; 

    const int warpID(threadIdx.x/warpSize); 
    const int laneID(threadIdx.x % warpSize); 
    const int st(rloops + 2); 
    const int stb((p - kpMax + 1 + 1) * BUFFERSIZE_MAX); 
    const int idx = blockDim.x * blockIdx.x + threadIdx.x; 
    unsigned int cAlive, tAlive, bCol, bline, transp_line; 
    int i, j; 
    bool volatile play(true); 

    if (idx < maxthreads) { 
     for (j = 0; j < st; j++){ 
      t_initdwl[threadIdx.x][j] = init_dwL[idx + j * BUFFERSIZE_MAX]; 
     } 
     __syncthreads(); 
     for (i = 0; i < p; i++) { 
      for (j = 0; j < st; j++) 
       if ((t_initdwl[threadIdx.x][j] + i) > dwLt[j]) { 
        play = false; 
        return; 
       } 

      cAlive = __ballot((play == true)); 
      tAlive = __popc(cAlive); 
      bline = __ffs(cAlive) - 1; 
      if (laneID == bline) 
       transp_line = atomicAdd(nkpl, tAlive); 
      transp_line = __shfl(transp_line, bline); 

      tAlive = exp2f(laneID + 1); 
      bline = tAlive & cAlive; // the bline value is ok 
      bCol = transp_line + __popc(bline); // but __popc(bline) become always 1 


      for (j = 0; j < st; j++) 
       kpL[bCol + j * stb] = t_initdwl[threadIdx.x][j] + i; 
     } 
    } 
}

在此先感謝！

來源

2016-08-23 Thiago Conrado

我已經讀過兩次了，現在老老實實地不知道你在問什麼。看來你正在尋找調試幫助。如果是這樣，請將自包含的可構建和可運行代碼與樣本輸入以及預期輸出和實際輸出一起發佈：[MCVE]（http://stackoverflow.com/help/mcve） – njuffa

那麼，您是對的。當我一步一步檢查時，我看到'bline = tAlive＆cAlive'的期望值，但是'kpL [bCol + j * stb]'總是評估爲** transp_line + 1 ** ... I' m猜測我應該在'__popc（bline）;掩碼或重鑄或別的** bline **;' 可能會注意到數據中的錯誤結果（在下一個使用它的內核中）和當前內核在最後一行。 –

您似乎允許一些線程變爲非活動狀態（例如，通過你的'return'語句）。你有沒有考慮到不活動的線程時，各種instrinsics的行爲？例如，當目標經線不活動時，經線混洗的行爲不確定。沒有[mcve]我只能猜測。 –

感謝您的支持！

閱讀@njuffa並測試它的代碼變體的工作。

以前的方法來掩蔽laneID是不正確的。
使用exp2f只會產生第一條和最後一條車道的預期結果。

tAlive = exp2f(laneID + 1)

被取代：

tAlive = 0xffffffff >> (warpSize - laneID);

所以，現在它的工作原理，就好了。

編輯：今天我發現一個機會一個post in Parallel for ForALL，這很好地解釋瞭如何使用洗牌，巴洛特利POPC爲了過濾有助於主要結果線程。

我希望這個編輯可以幫助別人找到它，因爲我沒有過去。

Regards，

來源

2016-08-24 04:14:10

只使用Shuffle + ballot + popc從選定的車道寫入數據

回答

相關問題