如何使用 numba 在 GPU 上泛化快速矩阵乘法

Posted 2023-03-23

【中文标题】如何使用 numba 在 GPU 上泛化快速矩阵乘法

【英文标题】：How to generalize fast matrix multiplication on GPU using numba

【发布时间】：2021-01-19 16:18:06

【问题描述】：

最近我一直在尝试使用 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 * 的 if 语句TPB，但这些都不起作用。

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

    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 = 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)

我想编写一些不依赖于矩阵 A、B 和 C 的代码，这些矩阵的维度可以被 TPB 完全整除，因为这些有时不受我的控制。我知道 GPU 只有在处理非常大的矩阵时才能更快地进行矩阵乘法，但我想使用小示例来检查答案是否正确，然后再将其应用于实际数据。

【问题讨论】：

欢迎来到 ***！你能告诉我们你已经尝试过什么吗？

【参考方案1】：

that posted code 可以说至少有两个错误：

这不可能是一个正确的范围检查：

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

为了让我们确定网格中的特定线程不执行任何加载活动，我们要求要么x 超出范围或y 超出范围。 and 应该是 or。

如果块中的所有线程都不能参与该语句，则在条件代码中使用cuda.syncthreads() 是illegal。上面第 1 项中的先前 return 语句（即使从 and 更正为 or）几乎可以保证问题大小不能被线程块大小整除的这种非法行为。

因此，要解决这些问题，我们不能只对越界线程使用简单的return 语句。相反，在加载点，我们必须只允许线程从全局加载到共享内存，如果计算的全局加载索引（对于A 或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 和y 的含义（以及tx 和ty 的用法）以修复上述代码中的性能问题。原始发布的文档代码中也存在相同的性能问题。

同样，没有声称没有缺陷。此外，我确信可以得出“更优化”的代码。然而，优化矩阵乘法是一个应该很快导致使用库实现的练习。在 GPU 方法中使用cupy 应该是一种在 GPU 上利用高质量矩阵乘法例程的相当简单的方法。

编辑：正如所讨论的here OP 的代码（而且似乎是doc example）在设置tmp 变量时也存在性能问题。将其更改为适当的 32 位浮点变量会产生重要的性能差异。

【讨论】：

哇，非常感谢您的回答，这很有帮助。我以前使用 CuPy 进行矩阵乘法，但我想更好地了解它是如何工作的，所以我决定寻找一些教程。

看起来and 在文档中更新为or 介于0.52 和0.53 之间，尽管Google 仍然在列表（开发分支）中排名靠前保留and的例子。我注意到注意 URL 很重要，例如而不是numba.pydata.org/numba-doc/dev/cuda/examples.html，我们想要numba.pydata.org/numba-doc/latest/cuda/examples.html，正如本问答中给出的那样。它似乎仍然存在第二个问题，即退货声明的非法性。也许我会就此提交 PR。很好的答案@罗伯特