CUDA 中的矩阵乘法运行内存不足

Question

我尝试使用脚本计算矩阵乘法：

import numpy as np
import math
from timeit import default_timer as timer
from numba import cuda
from numba import *


def mult(a,b):
    return a*b
mult_gpu=cuda.jit(restype=float32,argtypes=[float32,float32],device=True)(mult)

@cuda.jit(argtypes=[float32[:,:],float32[:,:],float32[:,:,:]])
def mult_kernel(a,b,c):
    Ni=c.shape[0]
    Nj=c.shape[1]
    Nk=c.shape[2]

    startX,startY,startZ=cuda.grid(3)
    gridX=cuda.gridDim.x*cuda.blockDim.x
    gridY=cuda.gridDim.y*cuda.blockDim.y
    gridZ=cuda.gridDim.z*cuda.blockDim.z

    for i in range(startX,Ni,gridX):
        for j in range(startY,Nj,gridY):
            for k in range(startZ,Nk,gridZ):
                c[i,j,k]=mult_gpu(a[i,k],b[j,k])

def main():
    A=np.ones((20,50000),dtype=np.float32)
    B=np.ones((3072,50000),dtype=np.float32)
    C=np.ones((20,3072,50000),dtype=np.float32)
    (Ni,Nj,Nk)=C.shape

    my_gpu=cuda.get_current_device()
    thread_ct=my_gpu.WARP_SIZE
    block_ct_x=int(math.ceil(float(Ni)/thread_ct))
    block_ct_y=int(math.ceil(float(Nj)/thread_ct))
    block_ct_z=int(math.ceil(float(Nk)/thread_ct))

    blockdim=thread_ct,thread_ct,thread_ct
    griddim=block_ct_x,block_ct_y,block_ct_z
    print "Threads per block:",blockdim
    print "Blocks per grid:",griddim

    start=timer()
    Cg=cuda.to_device(C)
    mult_kernel[griddim,blockdim](A,B,Cg)
    Cg.to_host()
    dt=timer()-start
    print "Computation done in %f s"%(dt)

    print 'C[:3,1,1] = ',C[:3,1,1]
    print 'C[-3:,1,1] = ',C[-3:,1,1]


if __name__=='__main__':
    main()

执行此操作会产生错误：

numba.cuda.cudadrv.driver.CudaAPIError: [2] Call to cuMemAlloc results in CUDA_ERROR_OUT_OF_MEMORY

我该如何解决这个内存问题？

然而，使用更小的矩阵

    A=np.ones((20,500),dtype=np.float32)
    B=np.ones((372,500),dtype=np.float32)
    C=np.ones((20,372,500),dtype=np.float32)

还是有错误：

numba.cuda.cudadrv.driver.CudaAPIError: [1] Call to cuLaunchKernel results in CUDA_ERROR_INVALID_VALUE

我受到Mandelbrot Example的启发，实现了上面的计算。

EDIT1

为了解决大家的困惑，这实际上是一个3D矩阵乘法的3D矩阵：

    A=np.ones((20,1,50000),dtype=np.float32)
    B=np.ones((1,3072,50000),dtype=np.float32)
    C=np.ones((20,3072,50000),dtype=np.float32)

我在 A 和 B 中跳过了一个维度，因为计算不需要它。

EDIT2

我的 GPU 是：

    In [1]: from numba import cuda

    In [2]: gpu=cuda.get_current_device()

    In [3]: gpu.name
    Out[3]: 'GeForce GT 750M'

EDIT3

根据我GPU的显存（2GB），我将每个维度的大小减少了2：

    dimx=10
    dimy=1536
    dimz=25000
    A=np.ones((dimx,dimz),dtype=np.float32)
    B=np.ones((dimy,dimz),dtype=np.float32)
    C=np.ones((dimx,dimy,dimz),dtype=np.float32)

但我仍然收到 CUDA_ERROR_OUT_OF_MEMORY 错误。如何解释这一点？

计算得出 3 个矩阵的大小约为 1.7GB： (10*1536*25000*4.+10*25000*4+1536*25000*4.)/(10**9)=1.6906

Answer 1

关于第一个问题，您运行内存不足。一个主要原因是这不是人们通常进行矩阵乘法的方式。通常，当您将行和列元素相乘时，您会保留运行总和，然后将该总和存储在乘积（结果）矩阵中的适当位置。这将使 c 矩阵的尺寸小得多，即。它只需要 2 维，而不是 3 维。您可能希望只研究矩阵-矩阵乘法的线性代数定义。当您将 2D 矩阵乘以 2D 矩阵时，结果是 2D 矩阵，而不是 3D 矩阵。

简而言之，是这样的：

for i in range(startX,Ni,gridX):
    for j in range(startY,Nj,gridY):
        c[i,j] = 0
        for k in range(startZ,Nk,gridZ):
            c[i,j]= c[i,j] + mult_gpu(a[i,k],b[j,k])

并相应地调整你的 c 形状。

如果您确实需要 3D 形式的单个产品，就像您在此处所做的那样，那么除了您需要缩放问题以适应您使用的任何 GPU 的 GPU 内存大小之外，我无话可说使用。

关于第二个问题，你这里有问题：

thread_ct=my_gpu.WARP_SIZE
...

blockdim=thread_ct,thread_ct,thread_ct

WARP_SIZE 是 32（大概）所以你要求一个尺寸为 32*32*32 = 32K 线程的 3D 块。但是 CUDA 线程块限制为最多 1024 个线程，该限制是各个维度的乘积。

如果您将 thread_ct 更改为 8，例如：

thread_ct=8

您应该能够解决这个特定问题。

CUDA 中的矩阵乘法运行内存不足

Matrix multiplication in CUDA running out of memory

python

cuda

gpu

numba

CUDA 中的矩阵乘法 运行 内存不足

Matrix multiplication in CUDA running out of memory

python

cuda

gpu

numba

CUDA 中的矩阵乘法运行内存不足