GPU 并行编程- Inclusive Scan 全面扫描算法
定义
扫描操作采用二进制关联运算符 ⨁\bigoplus⨁(circle plus)和一个有n个元素的数组组成。
数组为
[x0,x1,......,xn−1][x_0,x_1,......,x_{n-1} ][x0,x1,......,xn−1]
返回值为
[x0,(x0⨁x1),......,(x0⨁x1⨁....⨁xn−1)][x_0,(x_0 \bigoplus x_1 ),......,(x_0\bigoplus x_1\bigoplus ....\bigoplus x_{n-1} ) ][x0,(x0⨁x1),......,(x0⨁x1⨁....⨁xn−1)]
例子:
如果 ⨁\bigoplus⨁代表加法,那么扫描操作
输入数组
A[]=[3 1 7 0 4 1 6 3],
返回值为
B[]=[3 4 11 11 15 16 22 25]
实例
问题
假设有一个 100 米长的三明治需要喂给10个人,
假设我们知道每个人需要的长度为:
A=[3 5 2 7 28 4 3 0 8 1]
问题1:我们如何快速分割三明治?
问题2:三明治会剩下多少?
解决方法
方法1
按顺序分割:第一次切3米,第二次切5米,第三次切2米…依此类推。
方法2
计算前缀和(prefix sum)
[3, 8, 10, 17, 45, 49,52,52,60,61] (剩下39)
在3米,8米,10米。。。。的位置分割就会得到每一段的长度
Scan的典型应用
Scan是一个简单并且有用的并行构筑模块
- Convert recurrences from sequential:
for(int i=1;i<n;i++)output[i]=output[i-1]+input[i];
- Into parallel
forall(i){temp[i]=f(i) }; scan(out, temp);
- Convert recurrences from sequential:
受益的并行算法:
- Radix sort
- Quick sort
- String comparison
- Lexical analysis
- Stream compaction
- Polynomial evaluation
- Solving recurrence
- Tree operations
- Histograms
其他应用
- Assigning camping spots
- Assigning Farmer‘s Market spaces
- Allocating memory to parallel threads
- Allocating memory bugger space for communication channels
- …
Inclusive Sequential Addition Scan
全面顺序扫描算法
Input: A sequence [x0,x1,x2....][x_0,x_1,x_2....][x0,x1,x2....]
output: [y0,y1,y2.....][y_0, y_1, y_2.....][y0,y1,y2.....]
y0=x0y_0=x_0y0=x0
y1=x0+x1y_1=x_0+x_1y1=x0+x1
y2=x0+x1+x2y_2=x_0+x_1+x_2y2=x0+x1+x2
递归定义:
yi=yi−1+xiy_i=y_{i-1}+x_iyi=yi−1+xi
“Work Efficient C Implementation”
y[0]=x[0]
for(int i=1; i < max_i; i++)y[i]=y[i-1]+x[i]
Computationally efficient:计算效率高
N个元素需要N次加法运算,O(N)
Naïve Inclusive Parallel Scan
- 给每个y的计算分配一个thread(线程)
- 让每个thread将计算这个y所需要的x相加
y0=x0y_0=x_0y0=x0
y1=x0+x1y_1=x_0+x_1y1=x0+x1
y2=x0+x1+x2y_2=x_0+x_1+x_2y2=x0+x1+x2
“并行计算在不考虑性能的时候十分简单”
一个更好的Parallel Scan 算法
- 从 device global memory 中将数据读取到 shared memory
- 迭代 log(n)log(n)log(n) 次,从1到n-1 跨步(stride),每一次迭代跨步翻倍
s | X | 3 | 1 | 7 | 0 | 4 | 1 | 6 | 3 |
---|---|---|---|---|---|---|---|---|---|
1 | T1T_{1}T1 | 3 | 1+3 | 7+1 | 0+7 | 4+0 | 1+4 | 6+1 | 6+3 |
1 | Y1Y_{1}Y1 | 3 | 4 | 8 | 7 | 4 | 5 | 7 | 9 |
2 | T2T_{2}T2 | 3 | 4 | 8+3 | 7+4 | 4+8 | 5+7 | 7+4 | 9+5 |
2 | Y2Y_{2}Y2 | 3 | 4 | 11 | 11 | 12 | 12 | 11 | 14 |
4 | T3T_{3}T3 | 3 | 4 | 11 | 11 | 12+3 | 12+4 | 11+11 | 14+11 |
4 | Y3Y_{3}Y3 | 3 | 4 | 11 | 11 | 15 | 16 | 22 | 25 |
处理依赖项
- 在每次迭代时,每个 thread 都会重写另一个 thread 的输入
- barrier synchronization 屏障同步:确保每一个输入都正确的发生
- 每个线程都必须保证能被重写输入项
- 需要屏障同步确保所有线程确保其输入的安全
- 所有线程执行加法和写输出操作
Work-Inefficient Scan Kernel
__global__ void work_inefficient_scan_kernel(float *X, float *Y, int InputSize)
{__shared__ float XY[SECTION_SIZE];int i = blockIdx.x * blockDim.x + threadIdx.x;if (i < InputSize) {XY[threadIdx.x] = X[i];}// the code below performs iterative scan on XYfor (unsigned int stride = 1; stride <= threadIdx.x; stride *= 2) {__syncthreads();float in1 = XY[threadIdx.x + stride];__syncthreads();XY[threadIdx.x] += in1;}__ syncthreads();If (i < InputSize) {Y[i] = XY[threadIdx.x];}
}
Work Efficiency Consideration
- Scan 操作log(n)log(n)log(n)次并行迭代
- 每一次迭代进行(n-1),(n-2),…,(n-n/2)次加法
- 总加法数量为:n∗log(n)−(n−1)n*log(n)-(n-1)n∗log(n)−(n−1) → O(n∗log(n))O(n*log(n))O(n∗log(n))
- 这种 Scan 算法不是 work efficient的
- 顺序 scan 算法进行 n 次加法,
- 乘上一个 log(n)log(n)log(n) 后运算次数将翻十倍,如果有1024 个元素。
- 当执行资源因工作效率低而饱和时,一个并行的算法会比一个顺序的算法更慢
改进 Work Efficient
- Balanced Trees
- 在输入数据上形成平衡的二叉树,并将其扫到根与根之间
- 树不是实际的数据结构,而是确定每个线程在每个步骤中做什么的概念
- For Scan
- 从叶子向下遍历到树中内部节点的部分根总和
- 根拥有所有叶子的总和
- 遍历备份树以建立部分和的输出
- 从叶子向下遍历到树中内部节点的部分根总和
GPU 并行编程- Inclusive Scan 全面扫描算法相关推荐
- openacc的Linux安装教程,科学网—opensuse 13.1 系统 openACC编译器使用及GPU并行编程环境配置 - 马小军的博文...
本文讲述opensuse13.1系统openACC编译器使用及GPU并行编程环境配置. 这里以笔记本显卡驱动为NVIDIA为例 在安装前,请确保系统已经安装kernel-devel ,kernel-s ...
- matlab gpuarray是什么,MATLAB GPU并行编程
gpuArray MATLAB中的gpuArray表示存储在GPU上的数据.使用gpuArray函数可以将数据从MATLAB工作空间传送到GPU.例如: A = data(10); G = gpuAr ...
- 使用 OpenCL.Net 进行 C# GPU 并行编程
在 初探 C# GPU 通用计算技术 中,我使用 Accelerator 编写了一个简单的 GPU 计算程序.也简单看了一些 Brahma 的代码,从它的 SVN 最新代码看,Brahma 要转移到使 ...
- 1.2.7存储结构-磁盘管理:磁盘移臂调度算法、先来先服务(FCFS)、最短寻道时间优先(SSTF)、扫描算法(SCAN)、循环扫描(CSCAN)
1.2.7存储结构-磁盘管理:磁盘移臂调度算法.先来先服务(FCFS).最短寻道时间优先(SSTF).扫描算法(SCAN).循环扫描(CSCAN) 先来先服务(FCFS) 最短寻道时间优先(SSTF) ...
- 求解VRP问题的节约里程法、sweep扫描算法和λ互换法
第05章 求解容量约束车辆路径问题的启发式算法 Edited by Jiannywang@163.com 目 录 5.1 节约里程法. 1 5.1.1 C-W节约算法简介. 1 5.1.2 C-W ...
- 基于正向扫描的并行区间连接平面扫描算法(IEEE论文)
作者: Panagiotis Bouros ∗ Department of Computer Science Aarhus University, Denmark pbour@cs.au.dk Nik ...
- 操作系统 先来先服务算法(FCFS)、最短寻到时间优先算法(SSTF)、扫描算法(电梯算法,SCAN)、循环扫描算法(CSCAN)
操作系统 先来先服务算法(FCFS).最短寻到时间优先算法(SSTF).扫描算法(电梯算法,SCAN).循环扫描算法(CSCAN)和N步扫描算法(NStepScan)的程序实现 复制到本地即可运行 # ...
- 操作系统之文件管理:9、磁盘的结构与磁盘调度算法(先来先服务FCFS、最短寻找时间优先SSTF、扫描算法SCAN、循环扫描算法C-SCAN、LOOK调度算法、C-LOOK调度算法)
9.磁盘的结构 磁盘结构 思维导图 磁盘.磁道.扇区.盘面.柱面 如何在磁盘中读/写数据? 磁盘的物理地址 一次磁盘读/写操作需要的时间 磁盘调度算法 1.先来先服务FCFS 2.最短寻找时间优先SS ...
- 最短寻道时间优先算法c语言程序,操作系统先来先服务、最短寻道时间优先(SSTf)、扫描算法(SCAN)、循环扫描算法(CSCAN)的c++实现.doc...
先来先服务(FCFS).最短寻道时间优先(SSTf).扫描算法(SCAN).循环扫描算法(CSCAN) 最后有运行截图. #include #include #include #include int ...
最新文章
- 华为年薪200万招募的“天才少年”,一句话让我陷入了深思
- c winform 上传文件到mysql_Winform下如何上传图片并显示出来。同时保存到数据库...
- hive 集群初探,查询比较
- wget镜像网站并且下载到指定目录 2012-06-20 19:40:56
- 博弈-sg函数的原理和优化(hdu-1536)
- pythonsuper继承规则,深入理解Python中的super()方法
- java 二进制文件修改_Java读写二进制文件操作
- hdu 5064 Find Sequence
- 计算机老提示安全证书到期,安全证书过期,教您怎么解决网站安全证书过期
- 计算机网络(考研)第二章 物理层
- spring boot项目使用ojdbc8连接oracle 12c(12.2.0.1.0),解决启动极慢问题!
- 黑科技解密!实现socket进程间迁移!
- Python使用Pillow库进行图像
- C语言的结构体前置声明,?C语言的不完整类型和前置声明
- wisp5学习日记1
- 16ava基础---Lambda,Stream流编程题
- 隐藏受保护的系统文件的注册表键值
- 更新并关机怎么关闭计算机,win7如何去掉“安装更新并关机”选项
- 7-55 查询水果价格 (15 分) 给定四种水果,分别是苹果(apple)、梨(pear)、桔子(orange)、葡萄…PTA:中M2021春C、Java入门练习第I段——变量、表达式、分支、循环
- 小区选择 html5,5G(NR)网络中终端的小区选择