CUDA：示例代碼具有300％多GPU縮放/性能優化

我正在將一個NBody問題解決代碼移植到具有CUDA的GPU上。CUDA：示例代碼具有300％多GPU縮放/性能優化

Nvidia公司提供了來與CUDA SDK的NBody解決模擬樣品英寸它在樣本/ 5_Simulations/nbody中。

我不是C++或CUDA方面的專家，很難理解他們的代碼分佈在多個文件中，因此我最終決定編寫我自己的實現，他們的算法在http://docs.nvidia.com/cuda/samples/5_Simulations/nbody/doc/nbody_gems3_ch31.pdf中描述。

我的執行很快成功的，我能計算在GTX泰坦每秒25間十億雙精度的相互作用。當我用-benchmark -fp64運行他們的實現時，我獲得了相同的性能。這讓我很驚訝，因爲在上面鏈接的文章中，他們每秒在200Gflop（單精度）卡上達到100億次單精度交互。 GTX Titan大約是1.3 Tflops（雙精度），因此我預計雙精度每秒可以處理650億次交互。

爲了增加神祕感，當我運行他們與-benchmark -fp64 -numdevices =執行6運算速度快18倍。這相當於300％的縮放？通過設置-numbodies來爲GPU增加更多工作量，比默認值大10倍450％縮放？

我要補充，我正在從CUDA SDK 5.5版本nbody實施，該系統是基於6個相同的GTX巨頭。

我已經嘗試了很多事情來獲得我的單GPU代碼上的250億次交互，但實際上它似乎處於峯值，因爲根據nvvp的內存訪問效率接近100％，佔用率爲50％（所有嘗試增加它對性能沒有影響）。

是否有任何可以讓GPU具有300％縮放比例的機制？

是我的代碼真正優化的，當我期待它運行快2-3倍觸發器判斷，可以嗎？

是否提供了與cuda竊聽的示例？

萬一這裏是內核的簡單版本：

__global__ 
void Force_GPU(double4 *d_r,double3 *d_F){ 
    unsigned int idx=threadIdx.x; 
    unsigned int index=blockIdx.x*blockDim.x+idx; 
    __shared__ double4 sharedpos[tilesize]; 
    double4 position=d_r[index]; 
    double3 temp_r,force={0.0,0.0,0.0}; 
    double temp_d; 
    #pragma unroll 
    for(int tile=0;tile<numtiles;tile++){ 
     sharedpos[idx]=d_r[tile*tilesize+idx]; 
     __syncthreads(); 
     #pragma unroll 
     for(int j=0;j<tilesize;j++){ 
      temp_r.x = position.x -sharedpos[j].x; 
      temp_r.y = position.y -sharedpos[j].y; 
      temp_r.z = position.z -sharedpos[j].z; 
      temp_d = temp_r.x * temp_r.x + temp_r.y * temp_r.y + temp_r.z * temp_r.z+1e-23; 
      temp_d = rsqrt(temp_d); 
      temp_d *= temp_d*temp_d; 
      temp_d *=sharedpos[j].w; 
      force.x += temp_r.x * temp_d; 
      force.y += temp_r.y * temp_d; 
      force.z += temp_r.z * temp_d; 
     } 
     __syncthreads(); 
    } 
    d_F[index]=force; 
}

我有更多的迷惑那些對內存流，擺脫軟化參數的1E-23，但性能的影響是針對無existant通過擺脫軟化參數（這需要更復雜的控制流程以避免計算粒子本身的力）來實現內存優化（不是我認爲的大量數據）和小（但清晰）。正如我所說，我也試圖增加佔用率，但它受限於寄存器，每個SM在2048個線程中有1024個線程。強制使用較低的寄存器使用率會產生可憐的性能，並且還需要調整瓦片以減少共享內存使用量。

任何幫助和意見將非常感激。

來源

2014-04-01 Agade

是具備CUDA樣本竊聽？

我相信通過CUDA 5.5的nbody示例代碼中存在一個錯誤，該錯誤在多GPU設置（甚至可能是2個GPU或更多）中開始與4個GPU或更多GPU一起生效。你可以通過運行cuda-memcheck的nbody示例代碼來得到一些指示，我想。

它可能是固定在當前可用的CUDA 6 RC drop中，我沒有檢查過它。

來源

2014-04-01 18:23:13

我試着用2,4和6顆GPU運行cuda-memcheck，它沒有報告任何錯誤。我也試着用選項--flush-to-disk yes和--leak-check full來運行它。 – Agade

嘗試使用CUDA 6 RC提供的nbody，它是[可供公開下載]（https://developer.nvidia.com/cuda-pre-production）。特別是，'bodysystemcuda.cu'中的重要源代碼更改與多設備使用相關，這與我所指的錯誤有關。 –

CUDA：示例代碼具有300％多GPU縮放/性能優化

回答

相關問題