📚“九层之台,起于垒土”。学习技术须脚踏实地。
☀️你要做冲出的黑马,而不是坠落的星星。
🔑前言
本文介绍了CUDA并行编程的经典案例——CUDA归约求和的第一个版本。这个案例对于理解掌握CUDA并行原理很有帮助。想了解更多CUDA相关内容,与大牛相互讨论交流,可以用这款刷题、模拟面试神器,可助你斩获心仪大厂offer:点我免费刷题、模拟面试
CUDA 归约求和算法:
并行归约(Reduction)是一种基础的并行算法,简单来说,我们有N个输入数据,使用一个符合结合律的二元操作符作用其上,最终生成1个结果。这个二元操作符可以是求和、取最大、取最小、平方、逻辑与或等等。
由于加法的交换律和结合律,数组可以以任意顺序求和。所以我们会自然而然产生这样的思路:首先把输入数组划分为更小的数据块,之后用一个线程计算一个数据块的部分和,最后把所有部分和再求和得出最终结果。
上面演示了单个线程块内的归约求和过程,结合下面的代码解释:
- 13行第一轮循环,对应上图的第二行,数组中索引为偶数的值被对应部分和替代。
- 第二轮循环,对应上图第三行,数组中索引为 4 的整数倍的值被对应部分和替代。
- 第三轮循环,对应上图第四行,数组中索引为 8 的整数倍的值被对应部分和替代。
- …
- 直至 stride 的值为 blockDim.x 的一半,这一步块内和就计算完了,并存储在threadIdx.x 为 0 的线程中。
但对于不同块内的和,就需要将每个块内 threadIdx.x 为 0 的线程的和相加才能求得总的向量和。
下面给出代码:
#include<stdio.h>
#include<stdlib.h>
#include<cuda.h>
#define THREAD_LENGTH 256
__global__ void reduceSum(double *d_A, int n){
unsigned int t = threadIdx.x;
__shared__ double partialSum[THREAD_LENGTH<