并行计算考前复习整理
(lwg老师会在最后一节课跟大家讲考点,考试考的东西不会在考点之外,这里面我整理的内容已经将考点全部囊括,最终100分)
一、向量求和函数
C语言的串行化实现
CUDA的并行化实现
1、问题一:i的计算
(考察点:给定代码,里面关于i的计算公式错了,让改正)
公式(牢记):int i = blockDim.x * blockIdx.x + threadIdx.x
解释:
· blockIdx.x:当前线程块block在网格grid中的索引。
· blockDim.x:每个线程块中的线程数量。
· threadIdx.x:当前现成在其所在线程块的索引。
具体来说:
· blockIdx.x * blockDim.x计算当前线程块的起始位置。
· threadIdx.x计算当前线程在其所在线程块中的偏移量。
· 将这两部分相加,就得到了当前线程在整个网格中的全局索引i。
2. 函数调用时传参的方式
(考察点:给出错误代码,如N/256,让改正,或许还让解释为什么这么传参)
传参方式(牢记):vecAdd <<<(N+255)/256,256>>>(a,b,c,N);
解释:
· 这两个参数分别是线程块的数量和每个线程块中线程的数量。
· 假设向量长度是N,如果直接计算N/256,可能会有一些剩余元素无法被处理到,而公式(N+255)/256确保有足够的线程块来覆盖所有的N个元素。通过加上255,再除以256,相当于向上取整,确保最后一个线程块也能处理剩余的元素。
二、线程同步指令__synthreads()
注意点:
·__synthreads()用来同步线程块内的所有线程,不能用于同步线程块之间的线程。
·一般在kernal中完成对共享内存(__shared__)和全局内存的写操作以后,调用__synthreads()函数来保证写操作正确反映在内存中。
例如下面的代码中,T被定义为一个共享内存数组(__shared__),那么在完成了对T的赋值之后,需要调用一次__synthreads()函数。
·一个块内的所有线程必须执行同一个__synthreads(),否则会导致GPU被挂起。
三、CUDA Reduce
考察方式:给出下面的CUDA Reduce的原始实现,让指出该实现目前在性能上存在什么问题?只需要指出第一层可优化点,并且要能够对其进行优化
答案:
1、存在2个性能问题:
(1)%运算导致高度分支的warp,效率非常低。
原因/解释:%会计算当前线程索引tid与步长2*s的取模结果,这会导致不同的线程可能得到不同的结果,使得同一个warp中的线程有不同的执行路径,也就是warp分支,GPU需要按顺序执行每个不同的路径,导致warp无法高效并行执行。
改进:将if判断语句改为index < blockDim.x,其中index = 2 * s * tid,这样就会使得同一个warp中的线程统一执行或者不执行该语句。
(2)%操作运算效率非常慢
原因/解释:%操作符在CUDA中被编译成20多条指令,运算速度非常慢。
改进:将%操作符改为乘法运算。
2、for循环改进后的代码
for(unsigned int s = 1;s < blockDim.x;s *= 2){int index = 2 * s * tid;if(index < blockDim.x){sdata[index] += sdata[index+s];}__synthreads();
}
四、矩阵向量乘
1、给出了下面的串行算法以及并行的实现,在这个并行实现中,通常假设m和n,即矩阵的行数和列数,都能够被t整除,t是线程的个数。但是,如果m和n不满足能被t整除的条件,那么用什么公式来分配数据?
串行算法(题目给)
并行实现(题目给)
答案:
quotient = n/p;
remainder = n % p;
if(my_rank < remainder){my_n_count = quotient + 1;my_first_i = my_rank * my_n_count;
}else{my_n_count = quotient;my_first_i = my_rank * my_n_count + remainder;
}
my_last_i = my_first_i + my_n_count – 1;
五、MPI_Reduce
考察点:MPI_Reduce中的实现只和目的进程(dest_process)上的调用顺序有关。
题目会给出MPI_Reduce的函数头,并给出几个进程,指明了操作方式为MPI_SUM,指明了dest_process是哪个进程,其中有一个非目的进程的调用顺序是错的,问最后的计算结果是什么。比如下面,Processs0是dest_process:
对于这3个进程,明显Process1的调用顺序是错的,看起来最终的计算结果是b=3,d=6,但是
MPI_Reduce的计算只和dest_process,也就是Process0的调用顺序有关,即使Process1调用顺序错了,也会把c当作a,把d当作b,因此最终的计算结果应该是b = a + c + a = 4,d = c + a + c = 5。
(如果dest_process变了,计算过程相应改变即可)
六、在冯·诺伊曼系统中加入缓存和虚拟内存改变了它作为SISD系统的类型吗?如果加入流水线呢?多发射或硬件多线程呢?
(1)加人缓存和虚内存不会改变系统的SISD类型。因为缓存和虚内存的加入没有改变一次可以执行的指令数或者一次可以操作的数据量,因此系统仍为SISD类型。
(2)加入流水线会使得系统变为SIMD类型。因为加入流水线相当于将一个指令应用于多个数据项,也就是SIMD。
(3)多发射和硬件多线程会使得系统变为MIMD类型。因为多发射相当于一个指令流在多个ALU上同时启动一部分,分裂为多个指令流,多个指令流用于不同的数据项,就是MIMD。硬件多线程是指当前任务被阻塞时,系统尝试切换到其他线程继续工作,这允许了多线程的存在,多个指令流用于不同的数据项,也是MIMD。
七、Cache命中问题
题目中给出算矩阵向量乘时的两种取数方式,分别是按行取数和按列取数,同时给出了Cache的情况,让你指出那种方式更加高效并且给出解释。
答案:
按行访问方式更快
原因:在按行访问时,由于内层循环是按行序问的,这与内存的存储方式一致(行优先存储),能够充分利用缓存行,提高缓存的命中率,速度更快。而按列访问时,每次访问都是跨行的,不是相邻的内存位置,不能充分利用缓存行,缓存命中率低,速度慢。
八、流水线计算
题目:当讨论流水线加法时,我们简单地假设每个功能单元都花费相同的时间。如果每个取命令和存命令都耗费2纳秒,其余每个操作均耗费1纳秒,如上图(考试也许不给图?)。
(1)在上述假设下,每个浮点数加法要耗费多少时间?
(2)非流水线1000对浮点数的加法要耗费多少时间?
(3)流水线1000对浮点数加法要耗费多少时间?
答案:
(1)9纳秒
(2)9 * 1000 = 9000纳秒
(3)9 + 999 * 2 = 2007纳秒
解释:流水线建立时间为9纳秒,剩下的999条指令,每条指令都只需2纳秒。
九、划分方法
题目:给定一定数量的元素数目(比如14),在给定进程数目(比如4),问使用块划分、循环划分、块循环划分的结果分别是什么?
(1)块划分
块划分首先计算分配给每个进程的元素数目,如果除不开,把剩余的元素数目依次分配给前几个进程。
(2)循环划分
循环划分时将元素依次分配给每个进程,直到分配完。
(3)块循环划分(BlockSize=2)
块循环划分首先将所有元素按照BlockSize划分成不同的块,再将这些块依次分配给每个进程。
最终将三个合并起来即为下面的表格形式
十、集合通信和点对点通信的异同?
(1)通信子中的所有进程必须调用相同的集合通信函数。
(2)每个进程传递给MPI通信函数的参数必须是相容的。
(3)参数output_data_p只用在dest_process上,但是所有的进程要传递一个与之相对应的实际参数,即使参数的值是NULL也可以。
(4)点对点通信是通过标签和通信子来匹配的。集合通信不使用标签,只通过通信子和调用的顺序来匹配。
十一、Cache一致性
题目1:Cache一致性的本质是什么?
Cache一致性的本质是在多处理器系统中保证所有处理器核心的Cache中存储的相同数据项的一致性,避免数据不一致问题。
题目2:解决Cache一致性问题有哪两种策略?各自的原理是什么?有什么异同?
解决Cache一致性问题的策略有监听Cache一致性和基于目录的Cache一致性。
(1)监听Cache一致性的原理:
· 所有的处理器核心共享一条总线。
· 在总线上传输的任何信号都会被连接到这条总线上的所有处理器核心看到。
· 当核心0更新其Cache中存储的x的副本时,会将此信息广播到总线上。
· 如果核心1监听总线,它将看到x被更新,并且将其缓存中x的副本标记为无效。
(2)基于目录的Cahce一致性的原理:
· 使用一个名为目录的数据结构来记录每个Cache缓存行的状态。、
· 当一个变量被更新时,会查询当前目录,目录会通知那些缓存了这个变量的缓存行的核心的缓存控制器,使其缓存行标记为无效。
(3)二者的异同
相同点:
· 都是为了在多处理器系统中保持缓存的一致性。
· 处理器核心之间都需要进行通信来确保缓存数据的同步。
不同点:
· 监听缓存一致性依赖于广播机制,而基于目录的缓存一致性依赖点对点通信。
· 监听缓存一致性依赖总线,扩展性差,而基于目录的缓存一致性扩展性好。
· 基于目录的缓存一致性需要维护一个目录,实现较为复杂,而监听缓存一致性实现较为简单。
十二、Cache伪共享问题
1、什么是Cache伪共享?
线程之间没有共享任何变量,但是共享了同一个缓存行,它们访问主存的行为看起来好像它们共享了一个变量,这种情况称为伪共享。
2、为什么会出现伪共享问题?
不同线程可能会操作各自独立的数据,但如果这些数据在内存位置上相邻,导致位于同一个缓存行中,当一个线程修改自己的数据时,会导致该缓存行在其他线程中的缓存失效。尽管这些线程没有共享变量,但它们的缓存行为表现得像是在共享数据,导致性能上的下降。
3、怎么解决Cache伪共享问题?
字节填充法。通过在数据结构中插入额外的字节,使得不同线程访问的数据不会位于同一个缓存行中。
十三、矩阵向量乘中的Cache问题
题目:给出了下面的原始的矩阵向量乘的代码,存在Cache伪共享问题,让写出优化后的代码(包括从OMP的初始化开始写)。
答案:
# pragma omp parallel num_threads(thread_count) \default(none) private(i,j,my_rank,sum) shared(A,x,y,m,n){my_rank = omp_get_thread_num();
# pragma omp forfor(i = 0; i < m; i++){sum = 0.0;for(j = 0; j < n; j++)sum += A[i*n+j]*x[j];y[i] = sum;
# ifdef DEBUGprintf(“调试信息”);
# endif}}
十四、循环依赖问题
题目:给出了下面的循环
在这个循环中显然有循环依赖,因为在计算a[i]前必须先算a[i-1]的值。请你找到一种方法消除循环依赖,并且并行化这个依赖。
答案:
(1)首先是消除依赖后的代码
for(i = 0; i < n; i++)
a[i] = i * (i+1)/2;
(2)并行化后的代码
# pragma omp parallel for num_threads(thread_count) \ default(none) private(i) shared(a,n)for(i = 0; i < n; i++)a[i] = i * (i+1)/2;
十五、打印宏_OPENMP
题目:如果已经定义了宏_OPENMP,它是一个int类型的十进制数。编写一个程序打印它的值。这个值的意义是什么?
答案:
#ifdef _OPENMP
#include<omp.h>
#endifvoid Usage(char progname[]);int main(int argc, char* argv[]){
# ifdef _OPENMPprintf(“_OPENMP = %d\n”,_OPENMP);
# elseprintf(“_OPENMP is not defined\n”);
# endifreturn 0;
}
_OPENMP的值是一个具有yyyymm形式的日期,它是已实现的OpenMP标准版本的年份和月。
十六、奇偶排序问题
题目:给定下面并行化的奇偶排序代码,让指出代码中存在的性能问题,并给出改进后的代码。
存在的性能问题:
(1)线程创建和销毁开销大:循环的每个阶段都会创建和销毁OpenMP线程,这增加了额外的开销。
(2)负载不均衡:循环变量i每次增量为2,可能导致线程之间的负载不均衡,特别是在元素数量不是2的倍数时。
改进后的代码:
# pragma omp parallel num_threads(thread_count) \default(none) shared(a,n) private(i,tmp,phase)for(phase = 0; phase < n; phase++){if(phase % 2 == 0)
# pragma omp forfor(i = 1; i < n;i += 2){if(a[i-1] > a[i]){tmp = a[i-1];a[i-1] = a[i];a[i] = tmp;}}else
# pragma omp forfor(i = 1; i < n-1; i += 2){if(a[i] > a[i+1]){tmp = a[i+1];a[i+1] = a[i];a[i] = tmp;}}
十七、Cache映射
Cache的三种映射方式:
(1)全相联映射:一个内存行能够映射到Cache的任何一个位置。
(2)直接映射:一个内存行只能映射到Cache中的一个唯一的位置。
(3)n路组相联映射:把Cache中的n个位置打包成一个组,一个内存行只能映射到一个组中的任何一个位置。
题目:给定一组Memory line,同时给定Cache的大小情况,让写出三种不同映射方式下的映射关系。如下表所示:
给定了15个Memory line,Cache的大小是4。对于全相联映射,很简单;直接映射则是用Memory Index % 4;最后是2路组相联映射,也就是Cache line 0和1作为一组,Cache line 2和3作为一组,用Memory index % 2映射到组后再在组内全相联即可。
十八、加速比和效率
1、加速比公式
2、效率公式
3、题目:一个串行程序只能够并行化90%的部分,假设我们使用的核心数是p,并且并行化是“完美的”,串行时间 = 20s,求它的加速比。
解答:
首先计算出可并行化部分的并行化运行时间:
再求出不可并行化部分的运行时间:
因此并行化后的总的运行时间为:
故加速比为
十九、超线性(Superlinear)
1、什么是超线性?
超线性加速是指并行程序的加速比超过了线性加速,即使用p个处理器时,加速比S > p。
2、给出几个超线性的例子。
(1)单核程序缓存有限,多核程序可以将所有数据放入缓存。
解释:在单核程序中,数据量可能超过单个核心的缓存容量,导致频繁的缓存未命中,从而增加内存访问时间,而在多核程序中,多个核心共享总缓存容量增加,可以将更多数据放入缓存,减少缓存未命中次数,提高访存效率。由于缓存效应,多核程序的实际执行时间可能比预期更短,导致超线性加速。
(2)并行DFS可以快速搜索目标结果。
在DFS中,并行化可以让多个线程同时搜索空间的不同部分,如果其中一个线程提前找到目标结果,那么其他线程可以立即停止搜索,从而节省了大量时间,这种情况下,实际加速效果可能超过线性加速。