CUDA性能：分支和共享內存

我想請問性能兩個問題。我一直無法創建簡單的代碼來說明。CUDA性能：分支和共享內存

問題1：如何昂貴的非發散分支？在我的代碼中，似乎它甚至超過了4個非fma FLOPS的等價值。請注意，我正在談論BRA PTX代碼，其謂詞已經被計算出來

問題2：我一直在閱讀大量關於共享內存的性能和一些文章，如a Dr Dobbs article甚至說它可以像寄存器一樣快只要訪問得好）。在我的代碼中，塊內warps內的所有線程訪問相同的共享變量。我相信在這種情況下共享內存是以廣播方式訪問的，不是嗎？它是否應以這種方式達到寄存器的性能？是否有任何特殊的事情需要考慮使其發揮作用？

編輯：我已經能夠建造一些簡單的代碼，讓更多的有識之士爲我的查詢

這

#include <math.h> 
#include <stdlib.h> 
#include <stdio.h> 
#include <string.h> 
#include <float.h> 
#include "cuComplex.h" 
#include "time.h" 
#include "cuda_runtime.h" 
#include <iostream> 
using namespace std; 

__global__ void test() 
{ 
__shared__ int t[1024]; 
    int v=t[0]; 
    bool b=(v==-1); 
    bool c=(v==-2); 
    int myValue=0; 
    for (int i=0;i<800;i++) 
    { 
#if 1 
      v=i; 
#else 
      v=t[i]; 
#endif 

#if 0 
      if (b) { 
        printf("abs"); 
      } 
#endif 
      if (c) 
      { 
        printf ("IT HAPPENED"); 
        v=8; 
      } 
      myValue+=v; 

    } 
    if (myValue==1000) 
      printf ("IT HAPPENED"); 



} 
int main(int argc, char *argv[]) 
{ 
    cudaEvent_t event_start,event_stop; 
    float timestamp; 
float4 *data; 
    // Initialise 
    cudaDeviceReset(); 
    cudaSetDevice(0); 
dim3 threadsPerBlock; 
dim3 blocks; 
threadsPerBlock.x=32; 
threadsPerBlock.y=32; 
threadsPerBlock.z=1; 
blocks.x=1; 
blocks.y=1000; 
blocks.z=1; 
cudaEventCreate(&event_start); 
cudaEventCreate(&event_stop); 
cudaEventRecord(event_start, 0); 
test<<<blocks,threadsPerBlock,0>>>(); 
    cudaEventRecord(event_stop, 0); 
    cudaEventSynchronize(event_stop); 
    cudaEventElapsedTime(&timestamp, event_start, event_stop); 
    printf("Calculated in %f", timestamp); 
}

我運行在GTX680這個代碼。

現在結果如下..

如果運行，因爲它是它需要5.44毫秒

如果更改所述第一條件的#if 0（這將使從共享存儲器讀出），它會採取6.02ms ..沒有更多的，但對我來說還是

如果我能第二＃如果條件（插入一個分支，永遠不會評估爲真）它9.647040ms運行不夠。性能下降非常大。原因是什麼？可以做些什麼？

我也改變了略微的代碼，以使進一步檢查具有共享存儲器

代替

__shared__ int t[1024]

我確實

__shared__ int2 t[1024]

何予訪問T []我剛訪問噸[]。X。在性能進一步下降到10毫秒..（另一個400微秒）爲什麼會發生這種情況？

問候丹尼爾

來源

2012-11-05 Daniel

是否確定了，如果你的內核計算綁定或內存約束？如果你的內核是計算綁定的，你的第一個問題是最相關的，而如果你的內核是內存綁定的，第二個問題是最相關的。如果你假設一個，你可能會得到令人困惑或難以重現的結果，而另一個則是。

（1）我不認爲一個分支的成本已經出版。您可能只能通過實驗確定您的架構。 CUDA編程指南確實說沒有「分支預測並且沒有推測性執行」。

（2）你是對的，當你從所有線程沿經線訪問共享內存中一個32位值，該值是廣播。但是我的猜測是，只要不產生銀行衝突，從所有線程訪問單個值的成本與訪問任意值的組合的成本相同。所以你最終會從共享內存中獲取單個提取的延遲。我不認爲延遲週期的數量已經發布。它很短，通常很容易隱藏。

來源

2012-11-05 23:24:33

絕對很混亂我會說。我試圖製作一些很好的簡單代碼來說明問題..會回來一次希望成功 – Daniel

剛剛發佈了一些代碼 – Daniel

您需要記住編譯器是高度優化的。所以如果你註釋掉分支，你也會消除條件的評估，不管你是否把它留在源代碼中。因此，四個指令的區別似乎對你的例子非常合理：
1. 負載-1，
2. 比較v它（結果存儲到b）
3. 測試b，
4. 分支，
雖然我還沒有編譯你的例子，並看着代碼（這是你應該做的 - 在012上運行cuobjdump -sass你的二進制文件並查看機器碼的實際差異。
使用只有.x構成元素的int2，讓你從銀行衝突免費獲得去2路庫衝突，導致在你的榜樣輕微的進一步放緩改變了佈局共享內存。 IIRC共享內存訪問的延遲約爲30個週期，通常很容易被其他線程隱藏（正如Roger已經提到的那樣）。

來源

2012-11-08 22:00:47 tera

感謝您的幫助..我仍然消化所有這些..但你爲什麼說我創建了一個使用int2的.x組件時銀行衝突...？所有線程仍然訪問相同的元素 – Daniel

是否在運行它之前，我讀過您的示例和對其所做更改的描述，線程不訪問相同的元素。發佈您計時的確切代碼，我們可以看看發生了什麼。 – tera

CUDA性能：分支和共享內存

回答

相關問題