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（大概），所以您要求一个 3D 块，尺寸为 32*32*32 = 32K 线程。但是 CUDA 线程块被限制为最多 1024 个线程，这个限制是各个维度的乘积。

如果你把你的thread_ct改为8，例如：

thread_ct=8

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