亚洲在线久爱草,狠狠天天香蕉网,天天搞日日干久草,伊人亚洲日本欧美

為了賬號安全,請及時綁定郵箱和手機立即綁定
已解決430363個問題,去搜搜看,總會有你想問的

如何使用 numba 在 GPU 上推廣快速矩陣乘法

如何使用 numba 在 GPU 上推廣快速矩陣乘法

偶然的你 2023-10-06 11:05:53
最近,我一直在嘗試使用 Numba 庫在 Python 中進行 GPU 編程。我一直在使用那里的教程在他們的網站上閱讀它,目前我陷入了他們的示例,可以在這里找到:https: //numba.pydata.org/numba-doc/latest/cuda/examples。 html。我試圖稍微概括一下快速矩陣乘法的示例(其形式為 A*B=C)。在測試時,我注意到尺寸不能完全被每塊線程數 (TPB) 整除的矩陣不會產生正確的答案。我從https://numba.pydata.org/numba-doc/latest/cuda/examples.html的示例中復制了下面的代碼,并創建了一個帶有 4 x 4 矩陣的非常小的測試用例。如果我選擇 TPB=2 一切都很好,但是當我設置 TPB=3 時就會出錯。我知道代碼超出了矩陣的范圍,但我無法阻止這種情況發生(我在ty + i * TPB和tx + i * TPB上嘗試了一些 if 語句,但這些不起作用。from numba import cuda, float32import numpy as npimport [email protected] 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    if x >= C.shape[0] and y >= C.shape[1]:        # Quit if (x, y) is outside of valid C boundary        return    # Each thread computes one element in the result matrix.    # The dot product is chunked into dot products of TPB-long vectors.    tmp = 0.    for i in range(bpg):        # Preload data into shared memory        sA[tx, ty] = A[x, ty + i * TPB]        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()    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 整除的矩陣 A、B 和 C 的代碼,因為這些有時超出了我的控制范圍。據我所知,GPU 僅在對非常大的矩陣進行矩陣乘法時速度更快,但我想使用小示例來檢查答案是否正確,然后再將其應用于實際數據。
查看完整描述

1 回答

?
慕婉清6462132

TA貢獻1804條經驗 獲得超2個贊

可以說該發布的代碼中至少有兩個錯誤:

  1. 這不可能是正確的范圍檢查:

    if?x?>=?C.shape[0]?and?y?>=?C.shape[1]:

    為了讓我們決定網格中的特定線程不執行任何加載活動,我們要求要么超出x范圍,要么y超出范圍。應該and是一個or.

  2. 如果塊中的所有線程都無法參與該語句,則在條件代碼中使用是非法的。cuda.syncthreads()上面第 1 項中的前一個return語句(即使從andto更正or)幾乎保證了對于問題大小不能被線程塊大小整除的這種非法行為。

因此,要解決這些問題,我們不能僅使用簡單的return語句來處理越界線程。相反,在加載點,如果計算出的全局加載索引(forAB)在邊界內(根據定義,共享索引在邊界內),我們必須只允許線程從全局加載到共享內存。此外,在寫入結果時,我們必須只寫入在 的范圍內的計算結果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

再次強調,不聲稱無缺陷。此外,我確信可以得出“更優化”的代碼。然而,優化矩陣乘法是一項應該很快就會導致使用庫實現的練習。此處使用cupyGPU 方法應該是一種相當簡單的方法,可以在 GPU 上利用高質量的矩陣乘法例程。



查看完整回答
反對 回復 2023-10-06
  • 1 回答
  • 0 關注
  • 109 瀏覽
慕課專欄
更多

添加回答

舉報

0/150
提交
取消
微信客服

購課補貼
聯系客服咨詢優惠詳情

幫助反饋 APP下載

慕課網APP
您的移動學習伙伴

公眾號

掃描二維碼
關注慕課網微信公眾號