性能推力::計數

Question

我有下面的代碼數據的在CUDA內核以後使用重組計劃的一部分：

thrust::device_ptr<int> dev_ptr = thrust::device_pointer_cast(dev_particle_cell_indices); 
int total = 0; 
for(int i = 0; i < num_cells; i++) { 
    particle_offsets[i] = total; 
    // int num = 0; 
    int num = thrust::count(dev_ptr, dev_ptr + num_particles, i); 
    particle_counts[i] = num; 
    total += num; 
} 

現在，如果我設置num爲0（取消對5日線，並註釋掉第6），應用程序運行速度超過30 fps，這是我的目標。但是，當我設置num等於thrust::count呼叫時，幀率下降到大約1-2幀/秒。爲什麼會發生？

我的理解是推力應該是一組利用GPU優勢的高度優化的算法，所以我很驚訝它會對我的程序性能產生這樣的影響。這是我第一次使用推力，所以我可能不知道一些重要的細節。

在循環中使用thrust::count會導致它運行得如此緩慢嗎？我如何優化我的使用情況？

舉一些數字，在我目前的測試情況下，num_particles是2000左右，而num_cells約爲1500

Answer 1

的thrust::count的表現就好了，這是你要使用它就是那樣性能問題。如果你有很多的粒子和只有幾個單元格，那麼你的實現使用thrust::count可能不是一個壞主意。你的問題是你有1500個單元。這意味着每次您想要執行計算時，將會有1500個調用count和1500個設備來主機傳輸內存。正如您所發現的，所有內核啓動和所有PCI-e總線副本的延遲都會導致性能下降。

了大量的細胞更好的方法可能是這樣的：

thrust::device_ptr<int> rawin = thrust::device_pointer_cast(dev_particle_cell_indices); 

// Sort a scratch copy of the cell indices by value 
thrust::device_vector<int> cidx(num_particles); 
thrust::copy(rawin, rawin+num_particles, cidx.begin()); 
thrust::sort(cidx.begin(), cidx.end()); 

// Use binary search to extract all the cell counts/offsets 
thrust::counting_iterator<int> cellnumber(0); 
thrust::device_vector<int> offsets(num_cells), counts(num_cells); 

// Offsets come from lower_bound of the ordered cell numbers 
thrust::lower_bound(cidx.begin(), cidx.end(), cellnumber, cellnumber+num_cells, offsets.begin()); 

// Counts come from the adjacent_difference of the upper_bound of the ordered cell numbers 
thrust::upper_bound(cidx.begin(), cidx.end(), cellnumber, cellnumber+num_cells, counts.begin()); 
thrust::adjacent_difference(counts.begin(), counts.end(), counts.begin()); 

// Copy back to the host pointer 
thrust::copy(counts.begin(), counts.end(), particle_counts); 
thrust::copy(offsets.begin(), offsets.end(), particle_offsets); 

在這裏，我們首先排序小區索引的本地副本，然後使用推力二進制搜索功能來執行相同的作爲代碼運行，但GPU內存中的數據通過的次數要少得多，並且只有兩臺設備才能託管內存副本，以便將所有結果返回給主機。

當我使用我在上面發佈的用於非平凡情況的代碼（在OS X上使用CUDA 4.1的GeForce 320M上的10000個隨機粒子和2000個單元格）的代碼進行基準測試時，我發現您的版本需要大約0.95秒運行，而排序/搜索版本運行約需0.003秒。所以如果你使用更有效的策略和更合適的算法，可能會有幾百倍的速度可以使用推力。

Answer 2

我必須坦率地說在這裏，這將是

推力

炫耀ArrayFire（其中我是一個核心開發者）的推力

批評

• 批評

  他們在優化向量輸入的並行算法方面做得很好。他們使用數據級並行性（除其他外）來標準化效果非常好的算法大,矢量輸入。但是他們沒有對它進行改進，並且一路去執行true數據級別的並行性。即一個大一些小問題。

  第二種情況在許多現實世界的應用中很有用，ArrayFire在這方面提供瞭解決方案（查看gfor，一個並行for循環）。

  爲​​

  插頭什麼應該是一個簡單的呼叫進行還原和掃描，而不是變爲4個算法（其中之一是一個昂貴排序）和在推力3個存儲器拷貝。

  下面是代碼是如何工作的ArrayFire：

  array cell_indices(num_particles, 1, dev_particle_cell_indices, afDevicePointer); 
  array particle_counts = zeros(num_cells); 
  
  gfor(array i, num_cells) // Parallel for loop 
       particle_counts(i) = sum(cell_indices == i); 
  
  array particle_offsets = accum(particle_counts); // Inclusive sum 
  

  -

  設置

  • 我使用talonmies代碼來對arrayfire基準。
  • 我在Linux 64（cuda 4.1/gcc 4.7）上使用類似的圖形卡（gts 360m）。
  • 您可以通過here找到完整的基準測試代碼。

  基準1

  隨着num_particles = 2000和num_cells = 1500（如原來的問題）

  $ ./a.out 
  Thrust time taken: 0.002384 
  ArrayFire time taken: 0.000131 
  

  ArrayFire更快倍

  基準2

  隨着num_particles = 10000和num_cells = 2000（像talonmies'測試用例）

  $ ./a.out 
  Thrust time taken: 0.002920 
  ArrayFire time taken: 0.000132 
  

  ArrayFire是更快倍

  基準3

  隨着num_particles = 50000和num_cells = 5000（只是一個較大的測試用例）

  $ ./a.out 
  Thrust time taken: 0.003596 
  ArrayFire time taken: 0.000157 
  

  ArrayFire是快倍

  注意

  • 推力要求你重寫代碼
  • 推力提供子320x〜一個加速過的原代碼
  • ArrayFire幾乎不需要重寫鱈魚E（變更用於向GFOR）
  • ArrayFire是18-23倍的速度（有效〜7300x超過原碼）
  • ArrayFire能更好地伸縮（運行時間增加了50％推力，用於ArrayFire 15％）

  結論

  推力確實提供了一個體面的速度，如果你可以重新寫你的問題。但是這並不總是可行的，對於更復雜的問題來說也不是微不足道的。這些數字表明，存在更高的性能範圍（因爲高度的數據並行性），而這種性能根本無法實現。

  ArrayFire以更有效的方式利用並行資源，時間表明gpu仍然不飽和。

  您可能想編寫自己的自定義cuda代碼或使用ArrayFire。我只想指出，有時使用推力不是一種選擇，因爲它在大量小問題中實際上毫無用處。

  編輯固定基準1次的結果（我用的是不正確的數字）