python - 为什么我不能用 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 。您还可以阅读该算法的更现代版本，该版本使用 CUDA 的更新功能(warp shuffle 指令和原子事务)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 时会看到的那样。