我正在分析一些代码,但无法找出性能差异 . 我试图在两个数组之间进行简单的元素添加(就地) . 这是使用numba的CUDA内核:
from numba import cuda
@cuda.jit('void(float32[:], float32[:])')
def cuda_add(x, y):
ix = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
stepSize = cuda.gridDim.x * cuda.blockDim.x
while ix < v0.shape[0]:
y[ix] += x[ix]
ix += stepSize
我认为性能很好,但后来我把它与cuBLAS方法进行了比较:
from accelerate.cuda.blas import Blas
blas = Blas()
blas.axpy(1.0, X, Y)
对于大型阵列(20M元件),BLAS方法的性能大约快25% . 这是在"warming up"内核之前通过调用它来编译的PTX代码已经被缓存(不确定这是否重要,但只是为了确保这不是问题) .
我可以理解3级矩阵矩阵运算的性能差异,但这是一个简单的补充 . 我能做些什么来从cuda.jit代码中挤出更多性能?我问,因为我想要优化的真实代码是一个二维数组,不能传递给blas.axpy .
EDIT 执行代码和其他所需的包:
import numpy as np
def main():
n = 20 * 128 * 128 * 64
x = np.random.rand(n).astype(np.float32)
y = np.random.rand(n).astype(np.float32)
## Create necessary GPU arrays
d_x = cuda.to_device(x)
d_y = cuda.to_device(y)
## My function
cuda_add[1024, 64](d_x , d_y)
## cuBLAS function
blas = Blas()
blas.axpy(1.0, d_x , d_y)
1 回答
很简短的答案是否定的 . CUBLAS利用许多东西(纹理,矢量类型)来改善numba CUDA方言目前不支持的内存绑定代码的性能 .
我在CUDA中删掉了这个:
我的基准测试显示它在CUBLAS的5%左右,但我不相信你现在可以在numba中做到这一点 .
顺便说一句,我不明白你关于无法在2D数组上运行
saxpy
的说法 . 如果数组在内存中是连续的(因为我怀疑它们必须是)并且具有相同的布局(即不尝试添加转置),那么您可以在2D数组上使用saxpy
.