2012-02-28 37 views
1

我有如下操作的工作流:CUDA 2D陣列 - 在分配之前考慮音高?

  1. 負荷初始值
  2. 過程值的中間結果甲
  3. 方法A中間體結果乙
  4. 方法B中間體結果Ç
  5. 過程C和B至中間結果D和E
  6. 總和部分D至最終結果F

我所有中間結果的自然結構都是2D數組,我用cudaMallocPitch()分配它。

不幸的是,我的算法要求我一次將D,E,C和B保存在內存中,並且單獨地,內存比B大4倍。由於我的處理中有另一個限制(迭代在存儲器中的圖形結構上),A或B的尺寸受D和E的最大尺寸限制,而D和E的最大尺寸又由初始值+ B的內存消耗+ C的內存消耗的內存使用量決定。這種依賴性是因爲我正在從主機中「分頁」到/從設備內存的中間結果部分(以容納非常大的問題集),並且我無法開始第4步,直到完成了整個步驟1-3問題集。

一旦我有B中所有的問題集,我可以刪除A.

我目前確定d + E的最大尺寸用下面的函數:

int gpuCalculateSimulPatterns(int lines, int patterns) { 
    // get free memory 
    size_t free_mem, total_mem; 
    int allowed_patterns; 
    cudaMemGetInfo(&free_mem, &total_mem); 
    allowed_patterns = (free_mem - (lines*sizeof(int)))/(lines*(sizeof(int)*2.5) + lines*sizeof(char)*1.5); 
    return min(patterns, allowed_patterns -(allowed_patterns % 32)); 
} 

它「作品」 ,但僅僅是因爲我高估了D或E(它們的尺寸和內存使用率相同)的大小25%,並且使B的預期大小加倍。即使如此,我仍遇到邊緣情況,因爲它的內存分配失敗內存不足。我想更有效地利用卡上的內存保持對齊,因爲我的內核對全局內存進行多次讀寫。

不,使用共享內存不是一個選項,因爲我在多個塊之間使用多個內核,並且塊內的線程根本不相互作用。

我發現cudaMallocPitch()只返回已成功分配的內存使用的音調。有沒有辦法給司機一個二維內存分配請求,只是要求它分配的音調?

我想搭建一個試驗/錯誤優化程序,但A,B,D和E之間維度的鏈接依賴性(CI計算先驗,因爲它沒有分配點距線性)使得這個解決方案很糟糕,它需要重新計算每個問題集。

有沒有人有更好的方法,可以讓我確定適合任意數量的設備內存的適當大小的中間數據集?

編輯:

爲中間體A的存儲器被重複使用,我的邊界的計算使得假設C + d + E + B >>初始+ A + B(它憑藉這樣的事實爲真甲& B是相同尺寸的1字節的字符,而C,d,E是整數)和這樣我只需要保證有一個爲B + C + d + E.

我只使用足夠的空間計算能力2.x卡可以用(Quadro 2000,Tesla C2075,GTX460)進行測試。

+0

是間距分配浪費了多少內存?你有多少例子?是否足夠擔心?如果是這樣,你可能會考慮做線性內存分配,並將你的索引線性化。 (另外,從你的依賴鏈看起來中間A的內存可以被重用)。 – harrism 2012-02-29 01:20:11

+0

我使用音調線性內存的主要原因是利用了合併,並且因爲一個進程中位置(X,Y)的結果在另一個進程中再次用於(X,Y)。 音高分配浪費了可變數量的內存。我認爲它保持到最接近的512字節,至少對於我正在測試的其中一張卡片。 – 2012-02-29 15:06:43

+0

你不需要間距線性來合併。您只需要(在Fermi上)查找經線中每個負載的地址,並將其放入單個對齊的128字節高速緩存行中。你永遠不會跨進程獲得緩存重用,所以我不確定爲什麼這是相關的。也許我誤解了。如果經線內的局部性模式是2D(並且具有良好的空間局部性),那麼使用紋理訪問會更好,因爲硬件紋理緩存和紋理內存佈局針對2D局部性進行了優化,而間距線性仍然是,線性的。 – harrism 2012-02-29 23:16:19

回答

1

節距以字節爲單位的計算是這樣的:

pitch_size = DIV_UP(width_in_bytes, prop.textureAlignment); 

凡DIV_UP舍入的第一個參數到第二個參數的下一個較高的倍數。

您可以通過致電cudaGetDeviceProperties()查詢prop.textureAlignment

+0

雖然我在cuda文檔中找不到DIV_UP,但它與宏的效果相同: '#define DIV_UP(x,y)(((x%y)> 0)?((x/y)+((x%y)> 0))* y:x)' – 2012-03-07 18:28:25

+0

嗯,我不會使用模數,這是一個相當昂貴的方式來做到這一點。我將使用'#define DIV_UP(x,y)((y)*(((x)+(y)-1)/(y)))' – harrism 2012-03-08 00:23:43

+0

我沒有在GPU代碼中使用宏,代碼謹慎使用。謝謝你的提醒。 – 2012-03-09 01:21:39

0

不應該以字節爲單位的間距爲:

pitch_size = DIV_UP(width_in_bytes, prop.textureAlignment)*prop.textureAlignment; 

代替:

pitch_size = DIV_UP(width_in_bytes, prop.textureAlignment);