cuda麻木三维循环性能差

Question

我正在尝试建立一个3D循环来完成任务。

 C(i,j,k) = A(i,j,k) + B(i,j,k)

在我的GPU上使用Python。这是我的GPU：

http://www.geforce.com/hardware/desktop-gpus/geforce-gt-520/specifications

我所看到的/与之比较的来源是：

http://nbviewer.ipython.org/gist/harrism/f5707335f40af9463c43

python.ipynb

有可能我已经导入了比必要的更多的模块。这是我的密码：

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

@autojit
def myAdd(a, b):
  return a+b

myAdd_gpu = cuda.jit(restype=f8, argtypes=[f8, f8], device=True)(myAdd)

@cuda.jit(argtypes=[float32[:,:,:], float32[:,:,:], float32[:,:,:]])
def myAdd_kernel(a, b, c):
    tx = cuda.threadIdx.x
    ty = cuda.threadIdx.y
    tz = cuda.threadIdx.z
    bx = cuda.blockIdx.x
    by = cuda.blockIdx.y
    bz = cuda.blockIdx.z
    bw = cuda.blockDim.x
    bh = cuda.blockDim.y
    bd = cuda.blockDim.z
    i = tx + bx * bw
    j = ty + by * bh
    k = tz + bz * bd
    if i >= c.shape[0]:
      return
    if j >= c.shape[1]:
      return
    if k >= c.shape[2]:
      return
    for i in xrange(0,c.shape[0]):
      for j in xrange(0,c.shape[1]):
        for k in xrange(0,c.shape[2]):
          # c[i,j,k] = a[i,j,k] + b[i,j,k]
          c[i,j,k] = myAdd_gpu(a[i,j,k],b[i,j,k])

def main():
    my_gpu = numba.cuda.get_current_device()
    print "Running on GPU:", my_gpu.name
    cores_per_capability = {1: 8,2: 32,3: 192,}
    cc = my_gpu.compute_capability
    print "Compute capability: ", "%d.%d" % cc, "(Numba requires >= 2.0)"
    majorcc = cc[0]
    print "Number of streaming multiprocessor:", my_gpu.MULTIPROCESSOR_COUNT
    cores_per_multiprocessor = cores_per_capability[majorcc]
    print "Number of cores per mutliprocessor:", cores_per_multiprocessor
    total_cores = cores_per_multiprocessor * my_gpu.MULTIPROCESSOR_COUNT
    print "Number of cores on GPU:", total_cores

    N = 100
    thread_ct = my_gpu.WARP_SIZE
    block_ct = int(math.ceil(float(N) / thread_ct))

    print "Threads per block:", thread_ct
    print "Block per grid:", block_ct

    a = np.ones((N,N,N), dtype = np.float32)
    b = np.ones((N,N,N), dtype = np.float32)
    c = np.zeros((N,N,N), dtype = np.float32)

    start = timer()
    cg = cuda.to_device(c)
    myAdd_kernel[block_ct, thread_ct](a,b,cg)
    cg.to_host()
    dt = timer() - start
    print "Wall clock time with GPU in %f s" % dt
    print 'c[:3,:,:] = ' + str(c[:3,1,1])
    print 'c[-3:,:,:] = ' + str(c[-3:,1,1])


if __name__ == '__main__':
    main()

我运行这个程序的结果如下：

Running on GPU: GeForce GT 520
Compute capability:  2.1 (Numba requires >= 2.0)
Number of streaming multiprocessor: 1
Number of cores per mutliprocessor: 32
Number of cores on GPU: 32
Threads per block: 32
Block per grid: 4
Wall clock time with GPU in 1.104860 s
c[:3,:,:] = [ 2.  2.  2.]
c[-3:,:,:] = [ 2.  2.  2.]

当我在源代码中运行这些例子时，我看到了显著的加速。我不认为我的例子运行正常，因为挂钟的时间比我预期的要长得多。我主要是从第一个示例链接中的“更大的加速带cuda python”部分建模的。

我相信我的索引是正确和安全的。也许问题就在我的块上？还是鬼鬼祟祟？或者我用错类型的GPU了。我想我读到它们一定是某种类型的。我对此非常陌生，所以这个问题很可能是微不足道的！

所有的帮助都是非常感谢的！

Answer 1

您正在正确地创建您的索引，但是您忽略了它们。运行嵌套循环

for i in xrange(0,c.shape[0]):
    for j in xrange(0,c.shape[1]):
        for k in xrange(0,c.shape[2]):

强制所有线程循环遍历所有维度中的所有值，这不是您想要的结果。您希望每个线程计算一个块中的一个值，然后继续。

我觉得这样的事情应该更有效..。

i = tx + bx * bw
while i < c.shape[0]:
    j = ty+by*bh
    while j < c.shape[1]:
        k = tz + bz * bd
        while k < c.shape[2]:
            c[i,j,k] = myAdd_gpu(a[i,j,k],b[i,j,k])
            k+=cuda.blockDim.z*cuda.gridDim.z
        j+=cuda.blockDim.y*cuda.gridDim.y
    i+=cuda.blockDim.x*cuda.gridDim.x

试着编译并运行它。也要确保验证它，因为我没有。

Answer 2

我不认为您使用imshow，或show，所以没有必要导入这些。

您似乎没有使用数学的导入(我没有看到任何math.some_function调用)。

你从伦巴和麻雀那里进口的东西似乎是重复的。您的“从numba进口库达”超过您的“从麻木的进口库达”，因为它是后续的。你打给库达的电话用的是南巴的库达，而不是麻木瓜。当您调用“从numba import *”时，您从numba导入所有东西，而不仅仅是cuda，这似乎是您唯一使用的东西。而且，(我相信)导入numba.cuda等同于从numba.cuda。为什么不用一个“从numba进口的库达”来消除您从numba进口的所有进口产品。