时间戳

首页 > 解决方案 > 即使对于巨型矩阵，NUMBA CUDA 也比并行 CPU 慢

python - 即使对于巨型矩阵，NUMBA CUDA 也比并行 CPU 慢

问题描述

网上只有几个关于将 cuda 用于 numba 的示例，我发现它们都比并行 CPU 方法慢。带有 CUDA 目标和模板的矢量化甚至更糟，所以我尝试创建一个自定义内核。您到处都能找到的一篇博文是https://gist.github.com/mrocklin/9272bf84a8faffdbbe2cd44b4bc4ce3c。这个例子是一个简单的模糊过滤器：

import numpy as np
import time
from numba import njit, prange,cuda
import timeit
import numba.cuda


@numba.cuda.jit
def smooth_gpu(x, out):
    i, j = cuda.grid(2)
    n, m = x.shape

    if 1 <= i < n - 1 and 1 <= j < m - 1:
        out[i, j] = (x[i - 1, j - 1] + x[i - 1, j] + x[i - 1, j + 1] +
                    x[i    , j - 1] + x[i    , j] + x[i    , j + 1] +
                    x[i + 1, j - 1] + x[i + 1, j] + x[i + 1, j + 1]) / 9

x_gpu = np.ones((10000, 10000), dtype='float32')
out_gpu = np.zeros((10000, 10000), dtype='float32')

threadsperblock = (16, 16)
blockspergrid_x = math.ceil(x_gpu.shape[0] / threadsperblock[0])
blockspergrid_y = math.ceil(x_gpu.shape[1] / threadsperblock[1])
blockspergrid = (blockspergrid_x, blockspergrid_y)

# run on gpu
smooth_gpu[blockspergrid, threadsperblock](x_gpu, out_gpu) # compile before measuring time
start_time = time.time()
smooth_gpu[blockspergrid, threadsperblock](x_gpu, out_gpu)
print("GPU Time: {0:1.6f}s ".format(time.time() - start_time))

和 CPU 版本：

x_cpu = np.ones((10000, 10000), dtype='float32')
out_cpu = np.zeros((10000, 10000), dtype='float32')


@njit(nopython=True,parallel=True)
def smooth_cpu(x, out_cpu):

    for i in prange(1,np.shape(x)[0]-1):
        for j in range(1,np.shape(x)[1]-1):
            out_cpu[i, j] =  (x[i - 1, j - 1] + x[i - 1, j] + x[i - 1, j + 1] + x[i    , j - 1] + x[i    , j] + x[i    , j + 1] +x[i + 1, j - 1] + x[i + 1, j] + x[i + 1, j + 1]) / 9

# run on cpu
smooth_cpu(x_cpu, out_cpu) # compile before measuring time
start_time = time.time()
smooth_cpu(x_cpu, out_cpu)
print("CPU Time: {0:1.6f}s ".format(time.time() - start_time))

GPU 版本约 500 毫秒，CPU 版本约 50 毫秒。到底是怎么回事？

标签： pythonperformanceparallel-processingcudanumba

解决方案

我要指出两点：

  1. 您在 GPU 版本的计时中包括将输入数组从主机传输到设备所需的时间，以及从设备到主机的结果。如果这是您比较的目的，那就这样吧；结论是 GPU 不适合这项任务（以一种有趣的方式）。

  2. GPU 代码在给出正确结果时并没有为获得良好的性能而组织。问题出在这里：

    i, j = cuda.grid(2)

    再加上这些索引用于访问数据的顺序：

    out[i, j] = (x[i - 1, j - 1] ...

    这会导致 GPU 访问效率低下。我们可以通过颠倒上面描述的两个顺序之一来解决这个问题。

考虑到上述两个问题，以下是您的代码稍作调整：

$ cat t29a.py
import numpy as np
import time
from numba import njit, prange,cuda
import timeit
import numba.cuda


x_cpu = np.ones((10000, 10000), dtype='float32')
out_cpu = np.zeros((10000, 10000), dtype='float32')


@njit(parallel=True)
def smooth_cpu(x, out_cpu):

    for i in prange(1,x.shape[0]-1):
        for j in range(1,x.shape[1]-1):
            out_cpu[i, j] =  (x[i - 1, j - 1] + x[i - 1, j] + x[i - 1, j + 1] + x[i    , j - 1] + x[i    , j] + x[i    , j + 1] +x[i + 1, j - 1] + x[i + 1, j] + x[i + 1, j + 1]) / 9

# run on cpu
smooth_cpu(x_cpu, out_cpu) # compile before measuring time
start_time = time.time()
smooth_cpu(x_cpu, out_cpu)
print("CPU Time: {0:1.6f}s ".format(time.time() - start_time))
$ python t29a.py
CPU Time: 0.161944s

$ cat t29.py
import numpy as np
import time
from numba import njit, prange,cuda
import timeit
import numba.cuda
import math

@numba.cuda.jit
def smooth_gpu(x, out):
    j, i = cuda.grid(2)
    m, n = x.shape

    if 1 <= i < n - 1 and 1 <= j < m - 1:
        out[i, j] = (x[i - 1, j - 1] + x[i - 1, j] + x[i - 1, j + 1] +
                    x[i    , j - 1] + x[i    , j] + x[i    , j + 1] +
                    x[i + 1, j - 1] + x[i + 1, j] + x[i + 1, j + 1]) / 9

x = np.ones((10000, 10000), dtype='float32')
out = np.zeros((10000, 10000), dtype='float32')
x_gpu = cuda.to_device(x)
out_gpu = cuda.device_array_like(out)
threadsperblock = (16, 16)
blockspergrid_x = math.ceil(x_gpu.shape[0] / threadsperblock[0])
blockspergrid_y = math.ceil(x_gpu.shape[1] / threadsperblock[1])
blockspergrid = (blockspergrid_x, blockspergrid_y)

# run on gpu
smooth_gpu[blockspergrid, threadsperblock](x_gpu, out_gpu) # compile before measuring time
cuda.synchronize()
start_time = time.time()
smooth_gpu[blockspergrid, threadsperblock](x_gpu, out_gpu)
cuda.synchronize()
print("GPU Time: {0:1.6f}s ".format(time.time() - start_time))
$ python t29.py
GPU Time: 0.021776s
$

所以我们看看如果我们针对所指出的两个问题进行调整，GPU（在我的例子中是 GTX 960）比 CPU 快大约 8 倍。这样的测量在某种程度上取决于用于比较的 CPU 和 GPU——你不应该假设我的测量与你的测量相当——你最好运行这些修改后的代码进行比较。但是，数据传输时间肯定比 GPU 计算时间要大很多，在我的情况下也超过 CPU 计算时间。这意味着（至少在我的情况下，在任何方面都不是特别快的系统）即使我们将 GPU 计算时间减少到零，传输数据的成本仍然会超过 CPU 计算时间成本。

因此，当你遇到这种情况时，是不可能获胜的。那时可以给出的唯一建议是“不要那样做”，即找到一个更有趣、更复杂的问题供 GPU 解决。如果我们在计算上让问题变得非常简单，比如这个，或者向量加法，这是你唯一想在 GPU 上做的事情，与在 CPU 上做的比较几乎从来都不是一个有趣的比较。希望您能看到让矩阵变大在这里并没有太大帮助，因为它也会影响数据传输时间/成本。

如果我们考虑数据传输成本（并且不要在我们的 GPU 代码中犯下严重的性能错误），根据我的测试，GPU 比 CPU 快。如果我们包括数据传输成本，对于这个非常简单的问题，GPU 很可能不可能比 CPU 更快（即使 GPU 计算时间减少到零）。

毫无疑问，可以做更多的事情来稍微改进 GPU 的情况（例如改变块形状、使用共享内存等），但我个人不希望把时间花在打磨无趣的东西上。

您可以在此处获得有关 Numba GPU 内存管理的更多描述。

与索引排序相关的内存效率问题的一般描述是here

推荐阅读