CUPY：矩阵向量乘法比向量向量乘法和小尺寸的 l2norm 更快
CUPY: matrix-vector multiplication is faster than vector-vector multiplication and l2norm for small sizes
我正在将我的 CPU 代码传输到 GPU。
在优化它的过程中，我发现了一个有争议的性能行为：
考虑计算向量的 L2 范数的简单任务。对于具有大量元素的向量，我的性能按预期缩放，但是对于少量元素 (256)，它不是：
import cupy as cp
a=cp.random.rand(256)
%timeit cp.linalg.norm(a)
32.3 µs ± 159 ns per loop (mean ± std. dev. of 7 runs, 10000 loops each)
现在，让我们将它与矩阵向量点积进行比较：
b=cp.random.rand(256,256)
%timeit cp.dot(a,b)
8.36 µs ± 80.1 ns per loop (mean ± std. dev. of 7 runs, 100000 loops each)
你可以看到矩阵-向量乘积出人意料地快了 4 倍。为什么会这样？
我开始研究这个玩具问题。首先，我创建了我的自定义缩减内核：
l2norm = cp.ReductionKernel('T x','T y', 'x * x','a + b', 'y = sqrt(a)', '0', 'l2norm')
使用这个内核，我的执行时间约为 17 微秒，比 linalg.norm 快两倍，但仍然比矩阵向量点积差两倍。我确信这个内核优化得很好，所以 C++ Thurst 实现不会做得更好。
我也试过用 cp.sqrt(a.dot(a)) 计算范数。我发现这是非常低效的，因为向量-向量点积 a.dot(a) 比矩阵-向量积 a.dot(b) 花费更长的时间！！！
我确实明白，对于这个小问题，性能是带宽有限的，所以大部分时间可以花在创建数组、copying/fetching 数据上，而不是算术上。但即使在这种情况下，我也希望 L2 范数比矩阵向量乘积快一点，因为它只需要 O(N) 次操作和提取，并且结果是一个数字。在矩阵向量乘积的情况下，我什至没有预先分配结果，我做了 N^2 次操作并从内存中获取 O(N^2) 个数。
对于大量元素（>1000 个元素），性能按预期扩展。
Ubuntu 18.05，anaconda 分布，python 3.8.3，cupy 8.2.，nvcc 11.0
首先，您只测量 CPU 时间，内核是异步执行的，您的测量只包括准备内核启动的部分时间，但您没有等待实际的内核执行。
如果我们更改代码以通过使用 cupyx.time.repeat 进行测量来考虑这一点，我们会得到
import cupy as cp
import cupyx
a = cp.random.rand(256)
cp.linalg.norm(a)
print(cupyx.time.repeat(cp.linalg.norm, (a,)))
b = cp.random.rand(256, 256)
print(cupyx.time.repeat(cp.dot, (a, b)))
c = cp.zeros(())
l2norm = cp.ReductionKernel(
"T x", "T y", "x * x", "a + b", "y = sqrt(a)", "0", "l2norm"
)
print(cupyx.time.repeat(l2norm, (a, c)))
结果是
norm : CPU: 32.077 us +/- 2.206 (min: 30.961 / max: 64.160) us GPU-0: 36.275 us +/- 2.223 (min: 34.880 / max: 68.512) us
dot : CPU: 9.572 us +/- 0.261 (min: 9.235 / max: 15.934) us GPU-0: 13.640 us +/- 0.347 (min: 12.896 / max: 21.440) us
l2norm : CPU: 10.216 us +/- 0.578 (min: 9.847 / max: 23.790) us GPU-0: 14.396 us +/- 0.591 (min: 13.504 / max: 27.936) us
cupy.linalg.norm 正在启动几个内核来计算范数，因此高 CPU 时间为 32 us，累积的 GPU 时间为 36 us。这里的数组大小非常小，这主要是添加几个内核的恒定开销。
dot 仅调用 cublas 函数，因此 cpu 时间大大减少，GPU 时间相当快，但随着大小的减少，这纯粹是开销。
由于生成实际内核所需的步骤，最终你的还原内核有更多的 cpu 时间，但 gpu 执行与点积大致相同。
如果我们将数组大小增加到 4096，结果如下：
norm : CPU: 31.637 us +/- 2.200 (min: 30.487 / max: 62.955) us GPU-0: 35.741 us +/- 2.215 (min: 34.336 / max: 67.008) us
dot : CPU: 9.547 us +/- 3.753 (min: 9.051 / max: 370.309) us GPU-0: 244.535 us +/- 3.791 (min: 241.952 / max: 598.624) us
l2norm : CPU: 10.170 us +/- 0.542 (min: 9.845 / max: 17.006) us GPU-0: 16.106 us +/- 0.725 (min: 15.168 / max: 29.600) us
请注意，GPU 执行时间仅针对点积发生变化，这与您的观察结果一致:)。对于其他内核，与初始开销相比，实际内核执行时间的大小仍然太小。
我正在将我的 CPU 代码传输到 GPU。 在优化它的过程中，我发现了一个有争议的性能行为：
考虑计算向量的 L2 范数的简单任务。对于具有大量元素的向量，我的性能按预期缩放，但是对于少量元素 (256)，它不是：
import cupy as cp
a=cp.random.rand(256)
%timeit cp.linalg.norm(a)
32.3 µs ± 159 ns per loop (mean ± std. dev. of 7 runs, 10000 loops each)
现在，让我们将它与矩阵向量点积进行比较：
b=cp.random.rand(256,256)
%timeit cp.dot(a,b)
8.36 µs ± 80.1 ns per loop (mean ± std. dev. of 7 runs, 100000 loops each)
你可以看到矩阵-向量乘积出人意料地快了 4 倍。为什么会这样？
我开始研究这个玩具问题。首先，我创建了我的自定义缩减内核：
l2norm = cp.ReductionKernel('T x','T y', 'x * x','a + b', 'y = sqrt(a)', '0', 'l2norm')
使用这个内核，我的执行时间约为 17 微秒，比 linalg.norm 快两倍，但仍然比矩阵向量点积差两倍。我确信这个内核优化得很好，所以 C++ Thurst 实现不会做得更好。
我也试过用 cp.sqrt(a.dot(a)) 计算范数。我发现这是非常低效的，因为向量-向量点积 a.dot(a) 比矩阵-向量积 a.dot(b) 花费更长的时间！！！
我确实明白，对于这个小问题，性能是带宽有限的，所以大部分时间可以花在创建数组、copying/fetching 数据上，而不是算术上。但即使在这种情况下，我也希望 L2 范数比矩阵向量乘积快一点，因为它只需要 O(N) 次操作和提取，并且结果是一个数字。在矩阵向量乘积的情况下，我什至没有预先分配结果，我做了 N^2 次操作并从内存中获取 O(N^2) 个数。
对于大量元素（>1000 个元素），性能按预期扩展。
Ubuntu 18.05，anaconda 分布，python 3.8.3，cupy 8.2.，nvcc 11.0
首先，您只测量 CPU 时间，内核是异步执行的，您的测量只包括准备内核启动的部分时间，但您没有等待实际的内核执行。
如果我们更改代码以通过使用 cupyx.time.repeat 进行测量来考虑这一点，我们会得到
import cupy as cp
import cupyx
a = cp.random.rand(256)
cp.linalg.norm(a)
print(cupyx.time.repeat(cp.linalg.norm, (a,)))
b = cp.random.rand(256, 256)
print(cupyx.time.repeat(cp.dot, (a, b)))
c = cp.zeros(())
l2norm = cp.ReductionKernel(
"T x", "T y", "x * x", "a + b", "y = sqrt(a)", "0", "l2norm"
)
print(cupyx.time.repeat(l2norm, (a, c)))
结果是
norm : CPU: 32.077 us +/- 2.206 (min: 30.961 / max: 64.160) us GPU-0: 36.275 us +/- 2.223 (min: 34.880 / max: 68.512) us
dot : CPU: 9.572 us +/- 0.261 (min: 9.235 / max: 15.934) us GPU-0: 13.640 us +/- 0.347 (min: 12.896 / max: 21.440) us
l2norm : CPU: 10.216 us +/- 0.578 (min: 9.847 / max: 23.790) us GPU-0: 14.396 us +/- 0.591 (min: 13.504 / max: 27.936) us
cupy.linalg.norm 正在启动几个内核来计算范数，因此高 CPU 时间为 32 us，累积的 GPU 时间为 36 us。这里的数组大小非常小，这主要是添加几个内核的恒定开销。
dot 仅调用 cublas 函数，因此 cpu 时间大大减少，GPU 时间相当快，但随着大小的减少，这纯粹是开销。
由于生成实际内核所需的步骤，最终你的还原内核有更多的 cpu 时间，但 gpu 执行与点积大致相同。
如果我们将数组大小增加到 4096，结果如下：
norm : CPU: 31.637 us +/- 2.200 (min: 30.487 / max: 62.955) us GPU-0: 35.741 us +/- 2.215 (min: 34.336 / max: 67.008) us
dot : CPU: 9.547 us +/- 3.753 (min: 9.051 / max: 370.309) us GPU-0: 244.535 us +/- 3.791 (min: 241.952 / max: 598.624) us
l2norm : CPU: 10.170 us +/- 0.542 (min: 9.845 / max: 17.006) us GPU-0: 16.106 us +/- 0.725 (min: 15.168 / max: 29.600) us
请注意，GPU 执行时间仅针对点积发生变化，这与您的观察结果一致:)。对于其他内核，与初始开销相比，实际内核执行时间的大小仍然太小。