1 回答

TA貢獻1804條經驗 獲得超2個贊
可以說該發布的代碼中至少有兩個錯誤:
這不可能是正確的范圍檢查:
if?x?>=?C.shape[0]?and?y?>=?C.shape[1]:
為了讓我們決定網格中的特定線程不執行任何加載活動,我們要求要么超出
x
范圍,要么y
超出范圍。應該and
是一個or
.如果塊中的所有線程都無法參與該語句,則在條件代碼中使用是非法的。
cuda.syncthreads()
上面第 1 項中的前一個return
語句(即使從and
to更正or
)幾乎保證了對于問題大小不能被線程塊大小整除的這種非法行為。
因此,要解決這些問題,我們不能僅使用簡單的return
語句來處理越界線程。相反,在加載點,如果計算出的全局加載索引(forA
或B
)在邊界內(根據定義,共享索引在邊界內),我們必須只允許線程從全局加載到共享內存。此外,在寫入結果時,我們必須只寫入在 的范圍內的計算結果C
。
以下代碼修復了這些項目。對于您給定的測試用例,它似乎可以正常工作:
$ cat t49.py
from numba import cuda, float32
import numpy as np
import math
@cuda.jit
def fast_matmul(A, B, C):
? ? # Define an array in the shared memory
? ? # The size and type of the arrays must be known at compile time
? ? sA = cuda.shared.array(shape=(TPB, TPB), dtype=float32)
? ? sB = cuda.shared.array(shape=(TPB, TPB), dtype=float32)
? ? x, y = cuda.grid(2)
? ? tx = cuda.threadIdx.x
? ? ty = cuda.threadIdx.y
? ? bpg = cuda.gridDim.x? ? # blocks per grid
? ? # Each thread computes one element in the result matrix.
? ? # The dot product is chunked into dot products of TPB-long vectors.
? ? tmp = float32(0.)
? ? for i in range(bpg):
? ? ? ? # Preload data into shared memory
? ? ? ? sA[tx, ty] = 0
? ? ? ? sB[tx, ty] = 0
? ? ? ? if x < A.shape[0] and (ty+i*TPB) < A.shape[1]:
? ? ? ? ? sA[tx, ty] = A[x, ty + i * TPB]
? ? ? ? if y < B.shape[1] and (tx+i*TPB) < B.shape[0]:
? ? ? ? ? sB[tx, ty] = B[tx + i * TPB, y]
? ? ? ? # Wait until all threads finish preloading
? ? ? ? cuda.syncthreads()
? ? ? ? # Computes partial product on the shared memory
? ? ? ? for j in range(TPB):
? ? ? ? ? ? tmp += sA[tx, j] * sB[j, ty]
? ? ? ? # Wait until all threads finish computing
? ? ? ? cuda.syncthreads()
? ? if x < C.shape[0] and y < C.shape[1]:
? ? ? ? C[x, y] = tmp
#%%
x_h = np.arange(16).reshape([4,4])
y_h = np.ones([4,4])
z_h = np.zeros([4,4])
x_d = cuda.to_device(x_h)
y_d = cuda.to_device(y_h)
z_d = cuda.to_device(z_h)
TPB = 3
threadsperblock = (TPB, TPB)
blockspergrid_x = math.ceil(z_h.shape[0] / threadsperblock[0])
blockspergrid_y = math.ceil(z_h.shape[1] / threadsperblock[1])
blockspergrid = (blockspergrid_x, blockspergrid_y)
fast_matmul[blockspergrid, threadsperblock](x_d, y_d, z_d)
z_h = z_d.copy_to_host()
print(z_h)
print(x_h@y_h)
$ cuda-memcheck python t49.py
========= CUDA-MEMCHECK
[[ 6.? 6.? 6.? 6.]
?[22. 22. 22. 22.]
?[38. 38. 38. 38.]
?[54. 54. 54. 54.]]
[[ 6.? 6.? 6.? 6.]
?[22. 22. 22. 22.]
?[38. 38. 38. 38.]
?[54. 54. 54. 54.]]
========= ERROR SUMMARY: 0 errors
$
(請注意,在邊界測試中使用and此處是正確的。與測試一組索引是否越界相比,測試一組索引是否為入界在布爾意義上是不同的。在入界測試中,我們要求兩者都在界內。在出界測試中,任何一個索引越界都是不合格的)。
我并不是說上面的代碼沒有缺陷或適合任何特定目的。它旨在演示我發現的問題的可能修復方法。正如您所發現的,讓共享內存平鋪矩陣乘法在每個可以想象的配置中工作并非易事,并且除了此處顯示的內容之外,我還沒有對它進行測試。(例如,如果您決定使 TPB 大于 32,您會遇到其他問題。此外,原始發布的代碼僅用于方陣乘法,這在一般非方陣情況下不起作用。)
如上所述,發布的代碼和上面帶有“修復”的代碼將無法正確處理一般的非方形情況。我相信一些簡單的修改將使我們能夠處理非方形的情況。簡而言之,我們必須將網格設置得足夠大,以處理兩個輸入矩陣的維度,同時仍然只寫入輸出矩陣的界內值的結果。這是一個經過簡單測試的示例:
$ cat t49.py
from numba import cuda, float32
import numpy as np
import math
@cuda.jit
def fast_matmul(A, B, C):
? ? # Define an array in the shared memory
? ? # The size and type of the arrays must be known at compile time
? ? sA = cuda.shared.array(shape=(TPB, TPB), dtype=float32)
? ? sB = cuda.shared.array(shape=(TPB, TPB), dtype=float32)
? ? x, y = cuda.grid(2)
? ? tx = cuda.threadIdx.x
? ? ty = cuda.threadIdx.y
? ? bpg = cuda.gridDim.x? ? # blocks per grid
? ? # Each thread computes one element in the result matrix.
? ? # The dot product is chunked into dot products of TPB-long vectors.
? ? tmp = float32(0.)
? ? for i in range(bpg):
? ? ? ? # Preload data into shared memory
? ? ? ? sA[ty, tx] = 0
? ? ? ? sB[ty, tx] = 0
? ? ? ? if y < A.shape[0] and (tx+i*TPB) < A.shape[1]:
? ? ? ? ? sA[ty, tx] = A[y, tx + i * TPB]
? ? ? ? if x < B.shape[1] and (ty+i*TPB) < B.shape[0]:
? ? ? ? ? sB[ty, tx] = B[ty + i * TPB, x]
? ? ? ? # Wait until all threads finish preloading
? ? ? ? cuda.syncthreads()
? ? ? ? # Computes partial product on the shared memory
? ? ? ? for j in range(TPB):
? ? ? ? ? ? tmp += sA[ty, j] * sB[j, tx]
? ? ? ? # Wait until all threads finish computing
? ? ? ? cuda.syncthreads()
? ? if y < C.shape[0] and x < C.shape[1]:
? ? ? ? C[y, x] = tmp
#%%
x_h = np.arange(115).reshape([5,23])
y_h = np.ones([23,7])
z_h = np.zeros([5,7])
x_d = cuda.to_device(x_h)
y_d = cuda.to_device(y_h)
z_d = cuda.to_device(z_h)
#TPB must be an integer between 1 and 32
TPB = 32
threadsperblock = (TPB, TPB)
grid_y_max = max(x_h.shape[0],y_h.shape[0])
grid_x_max = max(x_h.shape[1],y_h.shape[1])
blockspergrid_x = math.ceil(grid_x_max / threadsperblock[0])
blockspergrid_y = math.ceil(grid_y_max / threadsperblock[1])
blockspergrid = (blockspergrid_x, blockspergrid_y)
fast_matmul[blockspergrid, threadsperblock](x_d, y_d, z_d)
z_h = z_d.copy_to_host()
print(z_h)
print(x_h@y_h)
$ cuda-memcheck python t49.py
========= CUDA-MEMCHECK
[[ 253.? 253.? 253.? 253.? 253.? 253.? 253.]
?[ 782.? 782.? 782.? 782.? 782.? 782.? 782.]
?[1311. 1311. 1311. 1311. 1311. 1311. 1311.]
?[1840. 1840. 1840. 1840. 1840. 1840. 1840.]
?[2369. 2369. 2369. 2369. 2369. 2369. 2369.]]
[[ 253.? 253.? 253.? 253.? 253.? 253.? 253.]
?[ 782.? 782.? 782.? 782.? 782.? 782.? 782.]
?[1311. 1311. 1311. 1311. 1311. 1311. 1311.]
?[1840. 1840. 1840. 1840. 1840. 1840. 1840.]
?[2369. 2369. 2369. 2369. 2369. 2369. 2369.]]
========= ERROR SUMMARY: 0 errors
$
x我還重新排序了and的含義(以及andy的用法)以修復上述代碼中的性能問題。原始發布的文檔代碼中也存在相同的性能問題。txty
再次強調,不聲稱無缺陷。此外,我確信可以得出“更優化”的代碼。然而,優化矩陣乘法是一項應該很快就會導致使用庫實現的練習。此處使用cupy
GPU 方法應該是一種相當簡單的方法,可以在 GPU 上利用高質量的矩陣乘法例程。
添加回答
舉報