2016-01-10 36 views
2

我需要一種算法來計算數組的並行前綴和而不使用共享內存。 如果沒有別的方法可以使用共享內存,那麼解決衝突問題的最佳方法是什麼?並行前綴和CUDAfy

+0

這包含在[GPU Gems 3]中(http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html) – harold

回答

3

此鏈接包含並行前綴和的順序和並行算法的詳細分析。

http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html

它還包含用於並行前綴算法的執行,並用於避免共享存儲器衝突的詳細說明的C碼的一個片段。

您既可以將代碼移植到CUDAfy,也可以簡單地定義C的區域,並將它們用作應用程序中的非託管代碼。 但是CUDA C代碼有幾處錯誤。我正在Cudafy.NET中編寫代碼的修正版本

[Cudafy] 
public static void prescan(GThread thread, int[] g_odata, int[] g_idata, int[] n) 
{ 
    int[] temp = thread.AllocateShared<int>("temp", threadsPerBlock);//threadsPerBlock is user defined 
    int thid = thread.threadIdx.x; 
    int offset = 1; 
    if (thid < n[0]/2) 
    { 
     temp[2 * thid] = g_idata[2 * thid]; // load input into shared memory 
     temp[2 * thid + 1] = g_idata[2 * thid + 1]; 

     for (int d = n[0] >> 1; d > 0; d >>= 1)     // build sum in place up the tree 
     { 
      thread.SyncThreads(); 
      if (thid < d) 
      { 
       int ai = offset * (2 * thid + 1) - 1; 
       int bi = offset * (2 * thid + 2) - 1; 
       temp[bi] += temp[ai]; 
      } 
      offset *= 2; 
     } 
     if (thid == 0) 
     { 
      temp[n[0] - 1] = 0; 
     } // clear the last element 


     for (int d = 1; d < n[0]; d *= 2) // traverse down tree & build scan 
     { 
      offset >>= 1; 
      thread.SyncThreads(); 
      if (thid < d) 
      { 
       int ai = offset * (2 * thid + 1) - 1; 
       int bi = offset * (2 * thid + 2) - 1; 
       int t = temp[ai]; 
       temp[ai] = temp[bi]; 
       temp[bi] += t; 
      } 
     } 
     thread.SyncThreads(); 
     g_odata[2 * thid] = temp[2 * thid]; // write results to device memory 
     g_odata[2 * thid + 1] = temp[2 * thid + 1]; 
    } 
} 

您可以使用上述修改的代碼而不是鏈接中的代碼。