CUDA編譯器如何知道經線的分歧行爲？

Question

的CUDA編程指南（V4.1）描述了這個約斷言指令在5.4.2節：

編譯器取代了分支指令只如果由控制指令的數量預測 說明所述 分支條件是小於或等於一定的閾值：如果 編譯器確定的條件有可能產生許多 發散經線，該閾值是7，否則爲4

1. 條件如何產生很多發散扭曲？一個給定的條件可以 只分裂成兩塊。 很多是什麼意思？
2. 即使上述說法有道理，編譯器如何知道經紗的運行時間散度行爲？

Answer 1

經紗永遠不會「分裂」。它們要麼需要「有條件的執行」（所以執行時要使用非參與線程屏蔽）來爲有條件的發散代碼路徑服務，要麼它們不需要。

對於一個情況可能產品的多個分支warp如何，請考慮以下人爲的例子：

if (threadIdx.x < 128) { 
    // Only first four warps process here 
    int modthirtytwo = threadIdx.x % 32; 

    if (modthirtytwo == 0) { 
     // Action A only first thread in the warp 
    } else { 
     // Action B for the other threads in the warp 
    } 
} 

這裏，代碼可以產生多個不同的經紗，而編譯器應該能夠行爲在編譯模型時間。如果啓動界限指定給內核的編譯器，則更好。將這種情況與僅使用一個warp的共享內存減少進行比較。

if (threadIdx.x < 32) { 
    if (threadIdx.x < 16) shm[threadIdx.x] += shm[threadIdx.x+16]; 
    if (threadIdx.x < 8) shm[threadIdx.x] += shm[threadIdx.x+8]; 
    if (threadIdx.x < 4) shm[threadIdx.x] += shm[threadIdx.x+4]; 
    if (threadIdx.x < 2) shm[threadIdx.x] += shm[threadIdx.x+2]; 
    if (threadIdx.x == 0) shm[0] += shm[1]; 
} 

這裏的分歧限於每塊單個翹曲。所有這段文字說的是，這兩種情況下的編譯器行爲可能不同。

似乎「新」編譯器（它已用於OpenCL幾年）對於在分支變得更經濟之前應該使用多少謂詞指令具有heurisitics。看起來，指令管道中的許多分支對性能不利，因此，當編譯器可以計算出代碼將產生更高的「分支密度」時，它將更喜歡更多的預測指令而不是分支。