为什么我不能用 numba (cuda python) 得到一维数组的正确总和？

Question

我尝试将 cuda python 与 numba 一起使用。代码是如下计算一维数组的总和，但我不知道如何得到一个值结果而不是三个值。

python3.5 与 numba + CUDA8.0

import os,sys,time
import pandas as pd
import numpy as np
from numba import cuda, float32

os.environ['NUMBAPRO_NVVM']=r'D:\NVIDIA GPU Computing Toolkit\CUDA\v8.0\nvvm\bin\nvvm64_31_0.dll'
os.environ['NUMBAPRO_LIBDEVICE']=r'D:\NVIDIA GPU Computing Toolkit\CUDA\v8.0\nvvm\libdevice'

bpg = (1,1) 
tpb = (1,3) 

@cuda.jit
def calcu_sum(D,T):
    ty = cuda.threadIdx.y
    bh = cuda.blockDim.y
    index_i = ty
    L = len(D)
    su = 0
    while index_i<L:
        su +=D[index_i]
        index_i +=bh
    print('su:',su)
    T[0,0]=su
    print('T:',T[0,0])


D = np.array([ 0.42487645,0.41607881,0.42027071,0.43751907,0.43512794,0.43656972,
               0.43940639,0.43864551,0.43447691,0.43120232], dtype=np.float32)
T = np.empty([1,1])
print('D: ',D)

stream = cuda.stream()
with stream.auto_synchronize():
    dD = cuda.to_device(D, stream)
    dT= cuda.to_device(TE, stream)
    calcu_sum[bpg, tpb, stream](dD,dT)

输出为：

D:  [ 0.42487645  0.41607881  0.42027071  0.43751907  0.43512794  0.43656972
  0.43940639  0.43864551  0.43447691  0.43120232]
su:  1.733004
su:  1.289852
su:  1.291317
T: 1.733004
T: 1.289852
T: 1.291317

为什么我不能得到输出“4.31417383”而不是“1.733004 1.289852 1.291317”？ 1.733004+1.289852+1.291317=4.314173.

我是 numba 新手，阅读了 numba 文档，但不知道如何操作。有人可以给点建议吗？

Answer 1

您没有得到预期总和的原因是您没有编写代码来生成该总和。

基本的 CUDA 编程模型（无论您使用 CUDA C、Fortran 还是 Python 作为您的语言）是您编写由每个线程执行的内核代码。您已经为每个线程编写了代码来读取和求和部分输入数组。您还没有为这些线程编写任何代码来共享它们各自的部分总和并将其汇总为最终总和。

有一个非常描述得很好的算法可以做到这一点——它被称为并行缩减。您可以在每个版本的 CUDA 工具包示例中附带的 PDF 中找到该算法的介绍，或者下载有关它的演示文稿 here. You can also read a more modern version of the algorithm which uses newer features of CUDA (warp shuffle instructions and atomic transactions) here。

学习完缩减算法后，您需要将标准 CUDA C 内核代码改编为 Numba Python 内核方言。至少，像这样：

tpb = (1,3) 

@cuda.jit
def calcu_sum(D,T):

    ty = cuda.threadIdx.y
    bh = cuda.blockDim.y
    index_i = ty
    sbuf = cuda.shared.array(tpb, float32)

    L = len(D)
    su = 0
    while index_i < L:
        su += D[index_i]
        index_i +=bh

    print('su:',su)

    sbuf[0,ty] = su
    cuda.syncthreads()

    if ty == 0:
        T[0,0] = 0
        for i in range(0, bh):
            T[0,0] += sbuf[0,i]
        print('T:',T[0,0])

可能会做你想做的事，尽管距离最佳的并行共享内存减少还有很长的路要走，正如你在阅读我提供的 material 链接时会看到的那样。

为什么我不能用 numba (cuda python) 得到一维数组的正确总和？

why can't I get the right sum of 1D array with numba (cuda python)?

python

cuda

numba

numba-pro