跳转至

笔记

Parallel BFS

常规广度优先搜索的局限性

简单bfs是每次从队列中取出当前的访问节点后,之后就将它的邻接节点加入到队列中,这样明显不利于并行化。

实现一: 两队列

使用了两个队列,第一个队列存当前同一层的节点,第二个队列用来存储当前层节点的所有邻接节点(下一层节点),等到当前层的节点全部访问完毕后,再将第一个队列与第二个队列进行交换,即可。

这样做的优势在于便于以后的并行化。同一层的节点可以一起运行,不会受到下一层节点的干扰。

但是写有问题,并行计算要储存下一层节点时,push操作不是原子的,需要加锁。

q1.push(1);
   distance[1]=0;
   while(!q1.empty()){
      clear(q2);
      for(int i=0;i<q1.size();i++){
         int node=q1.front();
         q1.pop();
         cout<<node<<"->";
         int start_edge=g.outgoing_starts[node];
         int end_edge=(node==g.num_nodes-1)?g.num_edges:g.outgoing_starts[node+1];
         for(int j=start_edge;j<end_edge;j++){
               int outgoing=g.outgoing_edges[j];
               if(distance[outgoing]==-1){
                  distance[outgoing]=1;
                  q2.push(outgoing);
               }
         }
      }
      swap(q1, q2);
   }

实现二:PBFS

思想和实现一一样,但是对并行时竞争情况有处理。

思想:都是将并行处理的一层节点bag的相邻节点存到下一次处理的bag中。

竞争:

  1. 第 18 行的 v.dist 更新为 d+1。
  2. 幸运的是,由于都是设置为同一个值,该竞争不影响结果。
  3. 第19行,可能同时将同一个节点加入out-bag,导致额外的工作
  4. 但是这个bag数据结构是具有元素各异的性质,正好解决这个问题。
  5. 根本没有元素各异好吧。
  6. 第19行,可能同时将同一个节点加入out-bag,不会写同一个地方吗?insert

最后还是要加锁,屮。

辅助数据结构 - pennant

辅助数据结构,称为 "pennant", 是一个\(2^k\)个节点的特殊树。根只有个left子节点,该子节点有个完整是二叉树。

通过这个辅助数据结构,可以实现动态无序集的 bag 数据结构

注意,普通的k层完全二叉树,有\(\(1+2+4+2^{k-1}=2^k-1\)\)个节点。

PENNANT-UNION(x,y)
   1 y. right = x. left
   2 x. left = y
   3 return x
PENNANT-SPLIT(x)
   1 y = x. left
   2 x. left = y. right
   3 y. right = NULL
   4 return y

数据结构 - bags

bag是pennant的集合,其中没有相同大小的pennant。PBFS采用固定大小的数组S[0...r]称作 backbone 骨干。 数组中的第k元素S[k]为空或者保存着指向大小为\(2k\)的pennant。可知,整个数组S[0...r]最多保存\(2\)元素。 所以BAG-CREATE分配了保存空指针固定大小的数组backbone,会花费Θ(r)时间。

This bound can be improved to O(1) by keeping track of the largest nonempty index in the backbone.

BAG-Insert类似二进制计数器递增的过程

BAG-INSERT(S,x)
   1 k = 0
   2 while S[k] != NULL
      3 x = PENNANT-UNION(S[k],x)
      4 S[k++] = NULL
   5 S[k] = x                 #不执行循环 S[0] = x,不是在图7这种插入。
BAG-UNION(S1,S2)     
   1 y = NULL // The “carry” bit.
   2 for k = 0 to r
   3 (S1[k],y) = FA(S1[k],S2[k],y)
BAG-UNION(S1,S2)因为每一个PENNANT-UNION都需要固定的时间,计算FA(x,y,z)的值也需要恒定的时间。计算结果包主干中的所有条目需要Θ(r)时间。该算法可以改进为Θ(lgn),其中n是两个包中较小的包中的元素数,方法是保持每个包主干的最大非空索引,并将索引较小的包合并为索引较大的包

其中FA(x,y,z)相当于三个数的加法,c是进位,s是余位。

函数解释

函数时间开销

BAG-CREATE O(1) BAG-INSERT(n) O(1)-O(lgn) BAG-UNION(n) ~O(lgn)

函数时间开销

需要进一步的研究学习

暂无

遇到的问题

PBFS: 1. 11行为什么是lg呢,全部遍历不好吗? 1. 因为存的是\(2^r\)的元素

This bound can be improved to O(1) by keeping track of the largest nonempty index in the backbone.

开题缘由、总结、反思、吐槽~~

flood fill的需要

参考文献

https://www.cnblogs.com/JsonZhangAA/p/9016896.html

https://zhuanlan.zhihu.com/p/134637247

A Work-Efficient Parallel Breadth-First Search Algorithm (or How to Cope with the Nondeterminism of Reducers) from ACM SPAA 2010

Flood fill Algorithm

算法介绍

泛洪算法——Flood Fill,(也称为种子填充——Seed Fill,其中基于栈/队列的显式实现(有时称为“Forest Fire算法”))是一种算法,用于确定连接到多维数组中给定节点的区域。用于填充具有不同颜色的连接的,颜色相似的区域。泛洪算法的基本原理就是从一个像素点出发,以此向周边的像素点扩充着色,直到图形的边界。

算法参数及特点

参数:

  1. 起始节点(start node)
  2. 目标颜色(target color)
  3. 替换颜色(replacement color)

数据结构:

明确地或隐式地使用队列或堆栈数据结构。

常见实现

四邻域泛洪

将像素点(x,y)周围的上下左右四个点分别进行着色。

八邻域泛洪

将一个像素点的上下左右,左上,左下,右上,右下都进行着色。

描绘线算法(Scanline Fill)

(我感觉就是满足访存局部性原理,然后就快一些

void floodFillScanline(int x, int y, int newColor, int oldColor){
    if(newColor==oldColor) return;
    if(screen[x][y]!=oldColor) return;
    emptyStack();
    int x1;
    bool spanAbove, spanBelow;
    if(!push(x,y)) return;
    while(pop(x,y)){
        x1=x;
        while(x1>0&&screen[x1][y]==oldColor) x1--;
        x1++;
        spanAbove = spanBelow = 0;
        while(x1<w&&screen[x1][y]==oldColor){
            screen[x1][y]=newColor;
            if(!spanAbove&&y>0&&screen[x1][y-1]==oldColor){ //这里判断出了上行的元素可以被染色,可能为了修改screen的访存连续性,所以这里没修改。而且你改了上行的值,却没考虑其四周,会有完备性的问题。
                if(!push(x1,y-1)) return;
                spanAbove=1;
            }
            else if(spanAbove&&y>0&&screen[x1][y-1]!=oldColor){
                spanAbove=0; //不压入重复过多的元素
            }
            if(!spanBelow&&y<h-1&&screen[x1][y+1]==oldColor){
                if(!push(x1,y+1)) return;
                spanBelow=1;
            }
            else if(spanBelow&&y<h-1&&screen[x1][y+1]!=oldColor){
                spanBelow=0;
            }
            x1++;
        }
    }
}
我感觉用队列更好,读上一行就是从左到右。

但是程序是上下行同时入栈/队列,写成两个,先pop一个,访存更好。

描绘线算法(Scanline Fill)的并行实现

直接将push变成用新线程重新调用程序task。

讨论数据冲突(IPCC)

每行各自修改自己行screen,不会冲突。 访问隔壁行screen,但是只在意oldColor,对是否染色不关心。?

记录修改过的数据,是每个线程每行中的连续一个范围,index1~index2。由于每个线程的任务变成每行,任务量多了,这里加个锁,应该不会影响性能,维护一个行数大小的数组,每个元素存每行被标记的index,要不要有序呢?这主要看cache页能不能存下一行。然后后面还原就可以对每行omp,保证访存连续性。但这不是最优的,详见ipcc10总结。

大规模行为(Large-scale behaviour)

以数据为中心(data-centric)

小并行,每个像素一次检查4个或者8个

以流程为中心(process-centric)

使用堆栈。使用邻接技术和堆栈作为其种子像素存储的4路泛洪填充算法产生具有“后续填充的间隙”行为的线性填充。 这种方法尤其适用于较旧的8位计算机游戏。

需要进一步的研究学习

最后的以流程为中心方法没懂,怎么实现

并行BFS(广度优先搜索)

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

IPCC 初赛 EnforceLabelConnectivity函数需要改成并行的图填充算法

参考文献

https://www.pianshen.com/article/172962944/

IPCC Preliminary SLIC Optimization 3

node6

因为例子太小,导致之前的分析时间波动太大。所以写了个了大一点的例子,而且给每个函数加上了时间的输出,好分析是否有加速。(Qrz,node5有人在用。

技术路线 描述 总时间 加速比 备注
Baseline 串行程序 207 s 1
simpleomp 两处omp 57s
more1omp maxlab 48s
more2omp sigma + delete maxxy 24.8s 8.35
more3omp DetectLabEdges + EnforceLabelConnectivity(该算法无法并行) 21.2s
icpc 13.4s
+ -O3 13.2s
+ -xHost 13.09s
+ -Ofast -xHost 基于icpc 12.97s
+ -ipo 12.73s 16.26
-no-prec-div -static -fp-model fast=2 14.2s 时间还多了,具体其他选项需要到AMD机器上试
### Baseline 207s
1. DoRGBtoLABConversion 10.4s
2. PerformSuperpixelSegmentation_VariableSandM 187.3s
1. core 15.3s
2. maxlab 1s
3. sigma 2.3s
### simpleomp 57s
1. DoRGBtoLABConversion 0.89s
2. PerformSuperpixelSegmentation_VariableSandM 46s
1. core 0.94-1.8s
2. maxlab 1s
3. sigma 2.3-2.6s
### more1omp 48s
1. DoRGBtoLABConversion 0.82s
2. PerformSuperpixelSegmentation_VariableSandM 37s
1. core 1-2.3s
2. maxlab 0.04-0.1s
3. sigma 2.3s
### more2omp 24.8s
1. DoRGBtoLABConversion 0.85s
2. PerformSuperpixelSegmentation_VariableSandM 13.5s
1. core 0.8-1.7s
2. maxlab 0.02-0.1s
3. sigma 0.1s
3. DetectLabEdges 3.7s
4. EnforceLabelConnectivity 5.2s

more2omp 21.2s

  1. DoRGBtoLABConversion 0.74s
  2. PerformSuperpixelSegmentation_VariableSandM 12.3s
  3. core 1.1s
  4. maxlab 0.02-0.1s
  5. sigma 0.1s
  6. DetectLabEdges 0.7s
  7. EnforceLabelConnectivity 5.8s (需要换算法
  8. PerformSuperpixelSegmentation_VariableSandM (vector声明的时间,可以考虑拿到外面去) 1.6s

icpc 13.4s

  1. DoRGBtoLABConversion 0.44s
  2. PerformSuperpixelSegmentation_VariableSandM 8.49s
  3. core 0.5-1.1s
  4. maxlab 0.04s
  5. sigma 0.05s
  6. DetectLabEdges 0.54s
  7. EnforceLabelConnectivity 2.79s (需要换算法
  8. PerformSuperpixelSegmentation_VariableSandM (vector声明的时间,可以考虑拿到外面去) 1.16s

12.7s

  1. DoRGBtoLABConversion 0.42s
  2. PerformSuperpixelSegmentation_VariableSandM 7.98s
  3. core 0.5-1.1s
  4. maxlab 0.04s
  5. sigma 0.05s
  6. DetectLabEdges 0.49s
  7. EnforceLabelConnectivity 2.69s (需要换算法
  8. PerformSuperpixelSegmentation_VariableSandM (vector声明的时间,可以考虑拿到外面去) 1.13s

IPCC AMD

技术路线 描述 总时间 加速比 备注
Baseline 串行程序 161.7s s 1
more3omp 前面都是可以证明的有效优化 omp_num=32 14.08s
more3omp 前面都是可以证明的有效优化 omp_num=64 11.4s
deletevector 把sz大小的3个vector,移到全局变量,但是需要提前知道sz大小/声明一个特别大的 10.64s 可以看出写成全局变量也不会影响访问时间
enforce_Lscan ipcc opt 4 8.49s
### Baseline 161.7s
1. DoRGBtoLABConversion 11.5s
2. PerformSuperpixelSegmentation_VariableSandM 143s
1. core 11.5s
2. maxlab 0.8s
3. sigma 1.7s
3. DetectLabEdges 2.74s
4. EnforceLabelConnectivity 3.34s
5. PerformSuperpixelSegmentation_VariableSandM 1.11s

more2omp 14.08s

  1. DoRGBtoLABConversion 0.69s
  2. PerformSuperpixelSegmentation_VariableSandM 8.08s
  3. core 0.73s
  4. maxlab 0.02s
  5. sigma 0.05s
  6. DetectLabEdges 0.37s
  7. EnforceLabelConnectivity 3.8s
  8. PerformSuperpixelSegmentation_VariableSandM 1.1s

more2omp 11.4s

  1. DoRGBtoLABConversion 0.61s
  2. PerformSuperpixelSegmentation_VariableSandM 5.86s
  3. core 0.53s
  4. maxlab 0.02s
  5. sigma 0.03s
  6. DetectLabEdges 0.33s
  7. EnforceLabelConnectivity 3.5s
  8. PerformSuperpixelSegmentation_VariableSandM 1.02s

deletevector 10.64s

  1. DoRGBtoLABConversion 0.59s
  2. PerformSuperpixelSegmentation_VariableSandM 5.75s
  3. core 0.53s
  4. maxlab 0.02s
  5. sigma 0.03s
  6. DetectLabEdges 0.41s
  7. EnforceLabelConnectivity 3.84s
  8. PerformSuperpixelSegmentation_VariableSandM 0s

enforce_Lscan 8.49s

  1. DoRGBtoLABConversion 0.56s
  2. PerformSuperpixelSegmentation_VariableSandM 5.52s
  3. core 0.53s
  4. maxlab 0.02s
  5. sigma 0.03s
  6. DetectLabEdges 0.31s
  7. EnforceLabelConnectivity 1.19s
  8. PerformSuperpixelSegmentation_VariableSandM 0.88s

需要进一步的研究学习

  1. 外面声明vector
  2. EnforceLabelConnectivity 换并行算法
  3. 数据结构要求:
    1. 保存已经染色区域的位置,之后可能要还原
      1. 可以无序,有序最好,会访存连续
      2. x,y或者index也行。还是xy好判断边界
    2. 是4分还是8分,既然有重复,记录来的方向/路径,只向某方向移动。4是符合理论的,8不和要求,2有情况不能全部遍历。
    3. 3分倒是可以,但是实现小麻烦
  4. flood fill 与 PBFS 特定结合
  5. openmp线程池+锁(sz 大小的两个数组存 x y,nlabels存新的分类结果)+计时声明与flood+把这些在sz声明放外面
  6. openmp线程池+队列(最后可以并行处理吧,要一个个pop?)+需要锁吗(这取决于队列的实现有没有靠计数器)
  7. openmpfor+双队列*4/2?+需要锁吗
  8. 扫描行实现 + 上下建线程,左右在线程里跑
    1. 多线程的访问存储连续性
  9. 队列/栈是怎么实现代码的,速度怎么样(写入读取push pop,还有size)
  10. 栈有size吗
  11. 在AMD机器加入MPI进行混合编程,运行2节点

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

Git Lfs

安装

mkdir git-lfs | cd git-lfs 
wget https://github.com/git-lfs/git-lfs/releases/tag/v2.13.3
tar -zxvf git
sudo ./install.sh

使用

git lfs install
git lfs track “*.rar” # 这个是要指定的大文件
git lfs track "*.txt" # 对一批,然后正常add commit
git add .gitattributes # 关联这个文件
git commit -m “aaa”

git 恢复

  • 工作区修改了文件(add之前),但是发现文件是你不想修改的,或者修改错误的,执行git checkout - 文件名,在工作区把文件恢复到修改之前的状态;

  • 工作区修改了文件,并且已经添加到缓存区(add之后,承之前),执行git reset HEAD文件名(HEAD表示最新的版本),此操作是把缓存区修改的内容返回到工作区,如果此时你还是不想修改此文件的话,就再次执行第一步操作,就可以恢复到文件修改前的状态;

  • 已经把文件提交给了分支(commit之后,推之前),执行git reset - hard HEAD ^(HEAD ^表示上一个版本),或者先用git log查看已经提交的版本号,执行git reset - -hard版本号的ID,就可以恢复到之前的版本,此时工作区和缓存区也是干净的;

  • 推的时候忽略文件的操作:(忽略大文件操作.gitignore不好使的时候),在commit提交之后push推之前,输入命令:

     git filter-branch --force --index-filter "git rm --cached --ignore-unmatch 有关文件"  --prune-empty --tag-name-filter cat -- --all # 如果git提示包含未提交的更改,需要再提交一下
    
     git commit --amend -CHEAD # 这个文件将会从你的提交记录里移除,并且以后commit都将不会再提交
    
     git push
    

需要进一步的研究学习

暂无

遇到的问题

很搞笑的是node5的IPCC/SLIC我就是弄不好,明明是按照步骤来的。

开题缘由、总结、反思、吐槽~~

大于100MB的文件上传不了github

参考文献

OpenMP Reductions

遇到竞争写情况怎么办

critical section

最简单的解决方案是通过声明一个critical部分来消除竞争。

double result = 0;
#pragma omp parallel num_threads(ndata)
{
  double local_result;
  int num = omp_get_thread_num();
  if (num==0)      local_result = f(x);
  else if (num==1) local_result = g(x);
  else if (num==2) local_result = h(x);
#pragma omp critical
  result += local_result;
}

double result = 0;
#pragma omp parallel
{
   double local_result;
#pragma omp for
   for (i=0; i<N; i++) {
    local_result = f(x,i);
#pragma omp critical
   result += local_result;
} // end of for loop
}

原子操作/加锁

性能是不好的,变串行了

#pragma omp atomic
 pi += sum;
static omp_lock_t lock;
void omp_init_lock(&lock):初始化互斥器
void omp_destroy_lock(omp_lock*):销毁互斥器
void omp_set_lock(omp_lock*):获得互斥器
void omp_unset_lock(omp_lock*):释放互斥器
void omp_test_lock(omp_lock*): 试图获得互斥器,如果获得成功则返回true,否则返回false

reduction clause 子句

将其添加到一个omp并行区域有如下效果。 * OpenMP将为每个线程制作一个reduction变量的副本,初始化为reduction操作的身份,例如\(1\)用于乘法。 * 然后,每个线程将其reduce到其本地变量中。 * 在并行区域结束时,本地结果被合并,再次使用reduction操作,合并到全局变量。

多个变量的情况

reduction(+:x,y,z)
reduction(+:array[:])

对于复杂结构体

如果代码过于复杂,还是建议复制全局变量来手工实现,最后再合并。

//错误示例
double result,local_results[3];
#pragma omp parallel
{
  int num = omp_get_thread_num();
  if (num==0)      local_results[num] = f(x)
  else if (num==1) local_results[num] = g(x)
  else if (num==2) local_results[num] = h(x)
}
result = local_results[0]+local_results[1]+local_results[2]
虽然上面这段代码是正确的,但它可能是低效的,因为有一个叫做虚假共享的现象。即使线程写到不同的变量,这些变量也可能在同一个缓存线上。这意味着核心将浪费大量的时间和带宽来更新对方的缓存线副本。

可以通过给每个线程提供自己的缓存线来防止错误的共享。

// 不是最好
double result,local_results[3][8];
#pragma omp parallel
{
  int num = omp_get_thread_num();
  if (num==0)      local_results[num][1] = f(x)
// et cetera
}
最好的方法给每个线程一个真正的局部变量,并在最后用一个critial部分对这些变量进行求和。
double result = 0;
#pragma omp parallel
{
  double local_result;
  local_result = .....
#pragam omp critical
  result += local_result;
}

默认的归约操作

Arithmetic reductions: \(+,*,-,\max,\min\)

Logical operator reductions in C: & && | || ^

归约变量的初始值

初始化值大多是不言而喻的,比如加法的0和乘法的1。对于min和max,它们分别是该类型的最大和最小可表示值。

用户自定义reduction的声明与使用

语法结构如下

#pragma omp declare reduction
    ( identifier : typelist : combiner )
    [initializer(initializer-expression)]
例子1: 取int最大
int mymax(int r,int n) {
// r is the already reduced value
// n is the new value
  int m;
  if (n>r) {
    m = n;
  } else {
    m = r;
  }
  return m;
}
#pragma omp declare reduction \
  (rwz:int:omp_out=mymax(omp_out,omp_in)) \
  initializer(omp_priv=INT_MIN)
  m = INT_MIN;
#pragma omp parallel for reduction(rwz:m)
  for (int idata=0; idata<ndata; idata++)
    m = mymax(m,data[idata]);

openmp减法归约浮点运算有精度损失

如何对vector归约

累加

#include <algorithm>
#include <vector>

#pragma omp declare reduction(vec_float_plus : std::vector<float> : \
                              std::transform(omp_out.begin(), omp_out.end(), omp_in.begin(), omp_out.begin(), std::plus<float>())) \
                    initializer(omp_priv = decltype(omp_orig)(omp_orig.size()))

std::vector<float> res(n,0);
#pragma omp parallel for reduction(vec_float_plus : res)
for(size_t i=0; i<m; i++){
    res[...] += ...;
}
编辑:原始initializer很简单:initializer(omp_priv = omp_orig)。但是,如果原始副本没有全零,结果将是错误的。因此,我建议使用更复杂的initializer,它总是创建零元素向量。

求最大值

#pragma omp declare reduction(vec_double_max : std::vector<double> : \
                          std::transform(omp_out.begin(), omp_out.end(), omp_in.begin(), omp_out.begin(), [](double a, double b) {return std::max(a,b);}))     \
                    initializer(omp_priv = decltype(omp_orig)(omp_orig.size()))

#pragma omp parallel for reduction(vec_double_max:maxlab)
for( int i = 0; i < sz; i++ )
{
   maxlab[klabels[i]] = max(maxlab[klabels[i]],distlab[i]);
}

std::transform

在指定的范围内应用于给定的操作,并将结果存储在指定的另一个范围内。

需要进一步的研究学习

  1. 对vector的归约

  2. 泥菩萨: 你这么改,开-g,在vtune里面看汇编

泥菩萨: 看有没有vmm指令

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

写IPCC发现:openmp没想象中简单,

参考文献

https://stackoverflow.com/questions/43168661/openmp-and-reduction-on-stdvector

https://pages.tacc.utexas.edu/~eijkhout/pcse/html/omp-reduction.html

http://www.cplusplus.com/forum/general/201500/

CPU vs GPU

GPU vs CPU

CPU: latency-oriented design

低延时的设计思路

  1. large L1 caches to reduce the average latency of data
  2. 时钟周期的频率是非常高的,达到3-4GHz
  3. Instruction-level parallelism to compute partial results ahead of time to further reduce latency
    1. 当程序含有多个分支的时候,它通过提供分支预测的能力来降低延时。
    2. 数据转发。 当一些指令依赖前面的指令结果时,数据转发的逻辑控制单元决定这些指令在pipeline中的位置并且尽可能快的转发一个指令的结果给后续的指令。

相比之下计算能力只是CPU很小的一部分。擅长逻辑控制,串行的运算。

GPU: throughput-oriented design

大吞吐量设计思路

  1. GPU采用了数量众多的计算单元和超长的流水线
  2. 但只有非常简单的控制逻辑
  3. 几乎省去了Cache。缓存的目的不是保存后面需要访问的数据的,减少cache miss。这点和CPU不同,而是为thread提高服务的。
  4. GPU “over-subscribed” threads: GPU运行任务会启动远超物理核数的thread,原因是借助极小的上下文切换开销,GPU能通过快速切换Threads/warps来隐藏访存延迟。
    1. GPU线程的创建与调度使用硬件而不是操作系统,速度很快(PowerPC创建线程需要37万个周期)1
    2. Cost to switch between warps allocated to a warp scheduler is 0 cycles and can happen every cycle.[^2]

对带宽大的密集计算并行性能出众,擅长的是大规模并发计算。

对比项 CPU GPU 说明
Cache, local memory 低延时
Threads(线程数)
Registers 多寄存器可以支持非常多的Thread,thread需要用到register,thread数目大,register也必须得跟着很大才行。
SIMD Unit 单指令多数据流,以同步方式,在同一时间内执行同一条指令

DRAM vs GDRAM

其实最早用在显卡上的DDR颗粒与用在内存上的DDR颗粒仍然是一样的。后来由于GPU特殊的需要,显存颗粒与内存颗粒开始分道扬镳,这其中包括了几方面的因素:

  1. GPU需要比CPU更高的带宽 GPU不像CPU那样有大容量二三级缓存,GPU与显存之间的数据交换远比CPU频繁,而且大多都是突发性的数据流,因此GPU比CPU更加渴望得到更高的显存带宽支持。位宽×频率=带宽,因此提高带宽的方法就是增加位宽和提高频率,但GPU对于位宽和频率的需求还有其它的因素。
  2. 显卡需要高位宽的显存显卡PCB空间是有限的,在有限的空间内如何合理的安排显存颗粒,无论高中低端显卡都面临这个问题。从布线、成本、性能等多种角度来看,显存都需要达到更高的位宽。 3090是384位。而内存则没有那么多要求,多年来内存条都是64bit,所以单颗内存颗粒没必要设计成高位宽,只要提高容量就行了,所以位宽一直维持在4/8bit。
  3. 显卡能让显存达到更高的频率显存颗粒与GPU配套使用时,一般都经过专门的设计和优化,而不像内存那样有太多顾忌。GPU的显存控制器比CPU或北桥内存控制器性能优异,而且显卡PCB可以随意的进行优化,因此显存一般都能达到更高的频率。而内存受到内存PCB、主板走线、北桥CPU得诸多因素的限制很难冲击高频率。由此算来,显存与内存“分家”既是意料之外,又是情理之中的事情了。为了更好地满足显卡GPU的特殊要求,一些厂商(如三星等)推出了专门为图形系统设计的高速DDR显存,称为“Graphics Double Data Rate DRAM”,也就是我们现在常见的GDDR。

内存频率

sudo dmidecode|grep -A16 "Memory Device"|grep "Speed"
   Speed: 2666 MT/s

显存等效频率

因为显存可以在一个时钟周期内的上升沿和下降沿同时传送数据,所以显存的实际频率应该是标称频率的一半。

从GDDR5开始用两路传输,GDDR6采用四路传输(达到类似效果)。

GDDR6X的频率估计应该至少从16Gbps(GDDR6目前的极限)起跳,20Gbps为主,这样在同样的位宽下,带宽比目前常见的14Gbps GDDR6大一半。比如在常见的中高端显卡256bit~384位宽下能提供512GB/s~768GB/s的带宽。

RTX 3090的GDDR6X显存位宽384bit,等效频率19Gbps到21Gbps,带宽可达912GB/s1006GB/s,达到T级。(384*19/8=912)

RTX 3090 加速频率 (GHz) 1.7, 基础频率 (GHz) 1.4

19/1.4 = 13.57
21/1.7 = 12.35

消费者设备 GDDR6x DDR4 的带宽对比

  • 上一小节 RTX 3090 带宽在912GB/s1006GB/s 附近
  • DRAM Types 一文里有分析,个人主机插满4条DDR4带宽" 3.2 Gbps * 64 bits * 2 / 8 = 51.2GB/s

可见两者差了20倍左右。

GPU / CPU workload preference

通过上面的例子,大致能知道: 需要高访存带宽和高并行度的SIMD的应用适合分配在GPU上。

最佳并行线程数

\[ 144 SM * 4 warpScheduler/SM * 32 Threads/warps = 18432 \]

参考文献

https://zhuanlan.zhihu.com/p/156171120?utm_source=wechat_session

https://www.cnblogs.com/biglucky/p/4223565.html

https://www.zhihu.com/question/36825227/answer/69351247

https://baijiahao.baidu.com/s?id=1675253413370892973&wfr=spider&for=pc

https://zhuanlan.zhihu.com/p/62234511

https://kknews.cc/digital/x6v69xq.html

AOCC

https://developer.amd.com/amd-aocc/

Install

cd <compdir>\
tar -xvf aocc-compiler-<ver>.tar
cd aocc-compiler-<ver>
bash install.sh
# It will install the compiler and displaythe AOCC setup instructions.

source <compdir>/setenv_AOCC.sh
# This will setup the shell environment for using AOCC C, C++, and Fortran compiler where the command is executed.

Using AOCC

Libraries

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

https://developer.amd.com/wp-content/resources/AOCC_57223_Install_Guide_Rev_3.1.pdf

AMD Epyc Compiler Options

AMD EPYC™ 7xx2-series Processors Compiler Options Quick Reference Guide

AOCC compiler (with Flang -Fortran Front-End)

Latest release: 2.1, Nov 2019

https://developer.amd.com/amd-aocc/Advanced

GNU compiler collection (gcc, g++, gfortran)

Intel compilers (icc, icpc, ifort)

amd prace guide

需要进一步的研究学习

  1. Amd uprof
  2. PGI compiler
  3. Numactl
  4. OMP_PROC_BIND=TRUE; OMP_PLACES=sockets

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

https://developer.amd.com/wordpress/media/2020/04/Compiler%20Options%20Quick%20Ref%20Guide%20for%20AMD%20EPYC%207xx2%20Series%20Processors.pdf

https://prace-ri.eu/wp-content/uploads/Best-Practice-Guide_AMD.pdf#page35

Prace guide

Intel Compile Options

Win与Linux的区别

选项区别

对于大部分选项,Intel编译器在Win上的格式为:/Qopt,那么对应于Lin上的选项是:-opt。禁用某一个选项的方式是/Qopt-和-opt-。

Intel的编译器、链接器等

在Win上,编译器为icl.exe,链接器为xilink.exe,VS的编译器为cl.exe,链接器为link.exe。

在Linux下,C编译器为icc,C++编译器为icpc(但是也可以使用icc编译C++文件),链接器为xild,打包为xiar,其余工具类似命名。

GNU的C编译器为gcc,C++编译器为g++,链接器为ld,打包为ar

并行化

-qopenmp

-qopenmp-simd

如果选项 O2 或更高版本有效,则启用 OpenMP* SIMD 编译。

-parallel

告诉自动并行程序为可以安全地并行执行的循环生成多线程代码。

要使用此选项,您还必须指定选项 O2 或 O3。 如果还指定了选项 O3,则此选项设置选项 [q 或 Q]opt-matmul。

-qopt-matmul

启用或禁用编译器生成的矩阵乘法(matmul)库调用。

向量化(SIMD指令集)

-xHost

必须至少与-O2一起使用,在Linux系统上,如果既不指定-x也不指定-m,则默认值为-msse2。

-fast

On macOS* systems: -ipo, -mdynamic-no-pic,-O3, -no-prec-div,-fp-model fast=2, and -xHost

On Windows* systems: /O3, /Qipo, /Qprec-div-, /fp:fast=2, and /QxHost

On Linux* systems: -ipo, -O3, -no-prec-div,-static, -fp-model fast=2, and -xHost

指定选项 fast 后,您可以通过在命令行上指定不同的特定于处理器的 [Q]x 选项来覆盖 [Q]xHost 选项设置。但是,命令行上指定的最后一个选项优先。

-march

必须至少与-O2一起使用,如果同时指定 -ax 和 -march 选项,编译器将不会生成特定于 Intel 的指令。

指定 -march=pentium4 设置 -mtune=pentium4。

-x

告诉编译器它可以针对哪些处理器功能,包括它可以生成哪些指令集和优化。

AMBERLAKE
BROADWELL
CANNONLAKE
CASCADELAKE
COFFEELAKE
GOLDMONT
GOLDMONT-PLUS
HASWELL
ICELAKE-CLIENT (or ICELAKE)
ICELAKE-SERVER
IVYBRIDGE
KABYLAKE
KNL
KNM
SANDYBRIDGE
SILVERMONT
SKYLAKE
SKYLAKE-AVX512
TREMONT
WHISKEYLAKE

-m

告诉编译器它可能针对哪些功能,包括它可能生成的指令集。

-ax

生成基于多个指令集的代码。

HLO

High-level Optimizations,高级(别)优化。O1不属于

-O2

更广泛的优化。英特尔推荐通用。

在O2和更高级别启用矢量化。

在使用IA-32体系结构的系统上:执行一些基本的循环优化,例如分发、谓词Opt、交换、多版本控制和标量替换。

此选项还支持:

内部函数的内联
文件内过程间优化,包括:
   内联
   恒定传播
   正向替代
   常规属性传播
   可变地址分析
   死静态函数消除
   删除未引用变量
以下性能增益功能:
   恒定传播
   复制传播
   死码消除
   全局寄存器分配
   全局指令调度与控制推测
   循环展开
   优化代码选择
   部分冗余消除
   强度折减/诱导变量简化
   变量重命名
   异常处理优化
   尾部递归
   窥视孔优化
   结构分配降低与优化
   死区消除

-O3

O3选项对循环转换(loop transformations)进行更好的处理来优化内存访问。

比-O2更激进,编译时间更长。建议用于涉及密集浮点计算的循环代码。

既执行O2优化,并支持更积极的循环转换,如Fusion、Block Unroll和Jam以及Collasing IF语句。

此选项可以设置其他选项。这由编译器决定,具体取决于您使用的操作系统和体系结构。设置的选项可能会因版本而异。

当O3与options-ax或-x(Linux)或options/Qax或/Qx(Windows)一起使用时,编译器执行的数据依赖性分析比O2更严格,这可能会导致更长的编译时间。

O3优化可能不会导致更高的性能,除非发生循环和内存访问转换。在某些情况下,与O2优化相比,优化可能会减慢代码的速度。

O3选项建议用于循环大量使用浮点计算和处理大型数据集的应用程序。

与非英特尔微处理器相比,共享库中的许多例程针对英特尔微处理器进行了高度优化。

-Ofast

-O3 plus some extras.

IPO

Interprocedural Optimizations,过程间优化。

典型优化措施包括:过程内嵌与重新排序、消除死(执行不到的)代码以及常数传播和内联等基本优化。

过程间优化,当程序链接时检查文件间函数调用的一个步骤。在编译和链接时必须使用此标志。使用这个标志的编译时间非常长,但是根据应用程序的不同,如果与-O*标志结合使用,可能会有明显的性能改进。

内联

内联或内联展开,简单理解,就是将函数调用用函数体代替,主要优点是省去了函数调用开销和返回指令的开销,主要缺点是可能增大代码大小。

PGO

PGO优化是分三步完成的,是一个动态的优化过程。

PGO,即Profile-Guided Optimizations,档案导引优化。

具体选项详解

-mtune=processor

此标志对特定的处理器类型进行额外的调整,但是它不会生成额外的SIMD指令,因此不存在体系结构兼容性问题。调优将涉及对处理器缓存大小、指令优先顺序等的优化。

为支持指定英特尔处理器或微体系结构代码名的处理器优化代码。

-no-prec-div

不启用 提高浮点除法的精度。

-static

不用动态库

-fp-model fast=2

自动向量化时按照固定精度,与OpenMP的选项好像有兼容性的问题

-funroll-all-loops

展开所有循环,即使进入循环时迭代次数不确定。此选项可能会影响性能。

-unroll-aggressive / -no-unroll-aggressive

此选项决定编译器是否对某些循环使用更激进的展开。期权的积极形式可以提高绩效。

此选项可对具有较小恒定递增计数的回路进行积极的完全展开。

falign-loops

将循环对齐到 2 的幂次字节边界。

-falign-loops[=n]是最小对齐边界的可选字节数。它必须是 1 到 4096 之间的 2 的幂,例如 1、2、4、8、16、32、64、128 等。如果为 n 指定 1,则不执行对齐;这与指定选项的否定形式相同。如果不指定 n,则默认对齐为 16 字节。

-O0 / -Od

关闭所有优化选项,-O等于-O2 (Linux and macOS)

-O1

在保证代码量不增加的情况下编译,

  1. 实现全局优化;这包括数据流分析、代码运动、强度降低和测试替换、分割生存期分析和指令调度。
  2. 禁用某些内部函数的内联。

遇到的问题

 icpc -dM -E -x c++ SLIC.cpp

https://stackoverflow.com/questions/34310546/how-can-i-see-which-compilation-options-are-enabled-on-intel-icc-compiler

parallel 与mpicc 或者mpiicc有什么区别呢

开题缘由、总结、反思、吐槽~~

讲实话,IPO PGO我已经晕了,我先列个list,之后再研究

参考文献

https://blog.csdn.net/gengshenghong/article/details/7034748

按字母顺序排列的intel c++编译器选项列表