pycuda编程之并行前缀算法

持续创作，加速成长！这是我参与「掘金日新计划 · 10 月更文挑战」的第27天，点击查看活动详情

前面已经介绍过了许多关于cuda编程方面的知识，对于并行前缀算法却没有提及，对于这个算法，刚刚看到确实会有一些蒙圈，不知道什么是并行前缀算法，其实并行前缀算法也就是扫描设计模式，不过呢，不了解很可能就是不了解，扫描设计模式是什么，也不清楚，不要紧，继续往后看。

在前面的博客中我们使用过InclusiveScanKernel 和 ReductionKernel，其实这两个模板就是并行前缀算法的最简单的例子。

并行前缀算法需要做的是给定一种二元的运算符，可以表示为$⊕$，这个符号的意思就是处理两个输入值并给定一个输出值，例如+、x，最大值和最小值等。下面以朴素并行前缀算法对其实际的使用环境进行举例。

之所以称之为朴素，是因为假设输入的元素为n，且假设n可以对2的k次方进行整除操作，然后在n个处理器，此处可以理解为n个线程上对算法进行并行的运行。只不过这算法对n的元素进行了限制，不过算法的运行效率非常的可观。

内核函数编写如下

kernal = SourceModule(
    """
    __global__ void nt_profix(double *vec, double *out)
    {
        __shared__ double sum_buf[1024];
        int tid = threadIdx.x;
        sub_buf[tid] = vec[tid];
        int iter = 1;
        for(int i=0;i<10;i++)
        {
            __syncthreads();
            if(tid>=iter)
            {
                sum_buf[tid] = sum_buf[tid] + sum_buf[tid-iter];
            }
            iter *= 2;
        }
        __syncthreads();
        out[tid]=sum_buf[tid];
        __syncthreads();
    }
    """
)

由于之前的博客中已经给出过许多调用示例，本博客不再分享调用示例，用之前的代码替代即可，说一下这其中的含义，由于我们不关心我们的输入是什么，只关心最后的输出，所以最终通过引用形式得到的out，我们其最后一个值即可，通过索引-1得到。

不过如果对编程的理解没有经过前面博客逐步深入的过程，这个核函数看起来仍然区别不大，如果对齐仍然不理解的话，可以通过对上述的核函数进行自己修改并对照之前的代码进行理解。