CUDA：發散式經線罰款細節

我想了解下一段代碼中經線發散的罰分是多少。我理解如何在原則上處理經線分歧（小分支的預測性指令，大分支的經線投票和分支 - 如果所有經線都一致，否則爲斷定的指令並且沒有分支，與小分支相同）。但是，我不明白具體細節 - 處理break/continue的while循環的方式。CUDA：發散式經線罰款細節

在下面的例子中，當scrapEverythingCondition（）評估爲真用於車道X，其下面的會發生：在內部循環

評價繼續進行，直到n ==可N_N，泳道X評估的NOP的整個時間，我增加了，現在所有的車道一起工作。
大家除了車道X確實someMoreWork（），而車道X評估空指令，車道X確實someCostlyInitialization（），加入N = 0，而其他人評估空指令，所有車道繼續共同評估內環（含明顯不同n值）。
其他我沒有想到的東西。

代碼：

__global__ void chainKernel() { 
    int i = threadIdx.x + blockIdx.x * blockDim.x; 
    while (i < N_I) { 
     someCostlyInitialization(); 
     for(int n = 0; n < N_N; ++n) { 
      someStatisticsComputations(n); 
      if (scrapEverythingCondition(n)) { 
       // Everything we did for current i is no good. Scrap and begin again 
       i -= BLOCKS*THREADS; 
       break; 
      } 
      someMoreWork(); 
     } 
     i += BLOCKS*THREADS; 
    } 
}

我試圖編譯成PTX並查看生成的代碼，但它太複雜，我:(

編輯：感謝馬庫的答案我也能夠使用老式的printf（）來查看答案，我可以看到哪些線程按照什麼順序排列，實際上選項1是正確的（X線路暫停，直到內部for循環已用盡）

來源

2013-07-02 user2412789

我在這個問題上發現了一個有趣的文檔：pdf

據我所知，控制流程語句（包括break）定義了線程的同步點。在你的情況在 i += BLOCKS*THREADS; 所以車道X葉for循環，並等待其他線程達到上述行的那樣。

來源

2013-07-02 09:02:26 Maku

的鏈接似乎被打破的「具有DOI號文件......」 – user2412789

shouold現在的工作 – Maku

至於我可以告訴當他們描述同步點時，他們在那裏引用他們的MCUDA框架。但是，他們確實試圖模仿CUDA架構的實際行爲，所以您的觀點似乎是有效的。謝謝！ – user2412789

在我的理解所有線程其中scrapEverythingCondition（n）爲真正在執行什麼的，如果塊內，並退出for循環。所有其他線程都被阻塞，直到該塊被執行。當這些線程退出for循環時，其他的trheads將執行someMoreWork（）;

試用NVidia Visual Profiler。這真的有助於分析這些問題。

這也是關於這個（第13頁 - 18）的一些信息：
http://mc.stanford.edu/cgi-bin/images/3/34/Darve_cme343_cuda_3.pdf

來源

2013-07-02 09:29:17 ipa

我從來沒有試過NVidia Visual Profiler，我現在試試，謝謝。我已經看到您鏈接到的文檔。不幸的是，就像我目前閱讀的所有內容一樣，它只討論if/else情況。 – user2412789

我試過探查器，但它沒有給我每個線程的信息，所以它沒有幫助解決這個問題。 – user2412789

CUDA：發散式經線罰款細節

回答

相關問題