3

我尝试将 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 文档，但不知道该怎么做。有人可以给建议吗？

4

1 回答 1

3

你没有得到你期望的总和的原因是因为你没有编写代码来产生这个总和。

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

有一个描述得非常好的算法可以做到这一点——它被称为并行归约。您可以在每个版本的 CUDA 工具包的示例中找到该算法的介绍，或在此处下载有关它的演示文稿。您还可以在此处阅读使用 CUDA（warp shuffle 指令和原子事务）的更新功能的算法的更现代版本。

在学习了归约算法之后，您需要将标准 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])

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

于 2017-03-29T08:54:52.863 回答