cuda子矩陣

我有4個矩陣（64x64）的單精度數字。需要做的計算，如：

R = A * sin(B) + C * cos(D)

想法：

加快計算使用共享內存。由於每個線程塊都有（在我的GPU的情況下）16KB共享內存，浮點大小爲4，因此可以在共享內存中存儲4000個浮點數。所以每個矩陣使用1000個元素，每個維度31個元素。

所以每個矩陣在768,16 16子矩陣（16×16）被devided。

dim3 dimBlock(16, 16, 1) 
dim3 dimGrid(4, 4, 1)

內核：

int Tx = threadIdx.x; 
int Ty = threadIdx.y; 

int Bx = blockIdx.x; 
int By = blockIdx.y; 

int idx = Bx * blockDim.x + Tx; 
int idy = By * blockDim.y + Ty; 

__shared__ float s_A[16*16]; 
__shared__ float s_B[16*16]; 
__shared__ float s_C[16*16]; 
__shared__ float s_D[16*16]; 

// I am not sure how to write this part 

s_A[(Tx * blockDim.x + Ty + By) + Bx] = A[idx * 64 + idy]; 
s_B[(Tx * blockDim.x + Ty + By) + Bx] = B[idx * 64 + idy]; 
s_C[(Tx * blockDim.x + Ty + By) + Bx] = C[idx * 64 + idy]; 
s_D[(Tx * blockDim.x + Ty + By) + Bx] = D[idx * 64 + idy]; 

R[idx * 64 + idy] = s_A[(Tx * blockDim.x + Ty + By) + Bx] * sin(s_B[(Tx * blockDim.x + Ty + By) + Bx]) + s_C[(Tx * blockDim.x + Ty + By) + Bx] * cos(s_D[(Tx * blockDim.x + Ty + By) + Bx]);

如何devide原矩陣submatrixs所以每塊都有自己的小矩陣4，計算它們。

來源

2012-03-29 user1281071

當你寫'cos（D）'時，你是指D的矩陣餘弦還是D的所有元素的餘弦？這兩件事情非常不同。 – talonmies 2012-03-29 09:55:29

SRY，餘弦所有元素的d- – user1281071 2012-03-29 10:10:43

除非我誤解了你的問題，你不需要和不應該使用共享內存進行此項操作。共享內存對於在同一塊內的線程之間共享和重新發送數據以及促進內存訪問的共享非常有用。您的操作似乎不需要這些東西正常工作。在你提出的方式使用共享內存很可能是慢不是直接從全局內存只是閱讀。此外，因爲你只擔心元素的運算，內核的索引方案可以大大簡化 - 事實上，A，B，C和D是「矩陣」是風馬牛不相及的計算，因爲我明白你的問題。

其結果是，你的內核的接近最優版本，可以寫成簡單地

在這段代碼中，你將推出儘可能多的塊作爲將達到您的GPU的峯值吞吐量，每個線程如果數組的大小超過您啓動的最佳一維網格的大小，將會處理多個輸入/輸出值。當然，如果你總共只處理4096個元素，這是非常的理論 - 這可能大約有2個數量級，太小而不能從使用GPU獲得任何好處。

來源

2012-03-29 10:45:03 talonmies

我的原始內核是： 'INT IDX = blockIdx.x * blockDim.x + threadIdx.x;'' INT IDX = blockIdx.x * blockDim.x + threadIdx.x ;'012x'[idx * 64 + idy] = A [idx * 64 + idy] * sin（B [idx * 64 + idy]）;' 'if（idx> = 64 || idy> = 64）return;' ' idy * 64 + idy * sin（D [idx * 64 + idy]）;' 所以共享mem無法加速？ – user1281071 2012-03-29 10:56:24

請不要在評論中張貼代碼，請將其編輯進原來的問題。但答案是否定的。在這種情況下使用共享內存不會有任何性能優勢。 – talonmies 2012-03-29 14:09:28

在這裏你有一個問題，你的操作/傳輸比例爲1級。由於線程和全局內存之間的帶寬瓶頸，你可能很難從GPU獲得任何體面的速度，並且沒有辦法減少這一點。

共享存儲器解決方案通常是最好的，當有一些數據從全局存儲器重複調用。無需從低帶寬，高延遲的全局內存中重複加載此數據，只需從中加載一次，然後從更高帶寬，更低延遲的共享內存中進行後續加載。請注意，這是更高和更低，而不是高和低。使用共享內存仍然會帶來性能損失。

你的情況下，由於要素不叫從全局內存幾次，將它們存儲在共享內存，只會增加帶寬限制和延遲自帶的共享內存使用情況。所以，實際上，這個解決方案只會增加從共享內存訪問您的數據加載延遲。

現在，如果您有多個計算需要執行，並且其中也使用了其中的一些矩陣，那麼將它們組合到一個內核中可能會提高速度，因爲您可能可以爲整體加載一次而不是每次操作一次。如果情況並非如此，並且您無法提高您的操作/傳輸率，那麼您將難以獲得一些體面的速度，並且在CPU上執行這些計算可能會更好。

你甚至可以從CPU上的多線程獲得一些體面的結果。

來源

2012-04-05 01:17:11

回答

相關問題