没有合适的资源?快使用搜索试试~ 我知道了~
10通过组合同步提高并发GPU B+树的性能和QoS0张伟华 �†‡¶ , 赵传雷 � , 彭璐 § , 林宇哲 � , 张丰哲 � , 陆云平 †¶0� 计算机科学学院,复旦大学 † 大数据研究所,复旦大学 ‡ 数学工程与先进计算国家重点实验室 ¶ 复旦大学并行处理研究所 § 图拉尼大学计算机科学系{zhangweihua,clzhao20,yzlin14,fzzhang}@fudan.edu.cn,lpeng3@tulane.edu,luyping@sina.com0摘要0并发B+树已经广泛应用于许多系统中。随着数据请求规模呈指数增长,系统面临着巨大的性能压力。GPU已经显示出加速并发B+树性能的潜力。在处理许多并发请求时,应该检测和解决冲突。以前的方法通过基于锁或软件事务内存(STM)的方法来保证并发GPUB+树的正确性。然而,这些方法使请求处理逻辑复杂化,增加了内存访问次数并引入了执行路径分歧。它们导致性能下降和响应时间的方差增加。此外,以前的方法不能保证并发请求之间的线性化。在本文中,我们设计了一种基于组合的并发控制框架,称为Eirene,用于GPUB+树,以减少冲突检测和解决的开销。首先,设计了一种基于组合的同步方法来组合和发出请求。它将具有相同键的请求组合在一起,构建它们的依赖关系,决定发出的请求,并确定它们的返回值。由于每个键只发出一个请求,键冲突被消除。然后,使用一种乐观的STM方法来减少结构冲突。查询和更新请求被分成不同的内核。对于更新内核,只有在重试次数达到阈值时才涉及STM。最后,提出了一种局部感知的warp重新组织优化来改善内存行为和0未经费用许可,允许个人或课堂使用本作品的全部或部分内容进行数字或硬拷贝,前提是不得为盈利或商业优势而制作或分发副本,并且副本上必须带有本通知和第一页的完整引用。必须尊重本作品中其他作者拥有的组成部分的版权。允许带有信用的摘要。要复制其他内容,或重新发布,或发布到服务器上,或分发到列表上,需要事先获得特定的许可和/或费用。请向permissions@acm.org申请权限。PPoPP '23,2023年2月25日至3月1日,加拿大蒙特利尔©2023版权归所有者/作者所有。出版权由ACM许可。ACM ISBN979-8-4007-0015-6/23/02...$15.00https://doi.org/10.1145/3572848.35774740通过利用请求之间的局部性来减少冲突。在NVIDIA A100GPU上的评估表明,Eirene是高效的(每秒24亿个吞吐量),并且可以保证线性化。与最先进的GPUB+树相比,它可以实现7.43倍的加速,并将响应时间方差从36%降低到5%。0CCS概念:•计算理论→数据结构设计与分析;•软件及其工程→并发控制。0关键词:GPU B+树,并发性01 引言作为最受欢迎的索引结构之一,并发B+树已经被广泛应用于不同领域,如文件系统[36],关系数据库[20,27,33]和键值存储系统[22,24]。如今,在诸如大数据之类的领域中,并发请求的数量呈指数增长。例如,亚马逊数据库在2020年PrimeDay上每秒处理超过8000万个请求[6]。底层的B+树面临着为如此大量的并发请求提供服务的巨大压力。由于其高计算能力和大内存带宽,GPU已经成为最受欢迎的加速器之一[18,25,42]。它们已经在数据中心环境中被采用,以提高许多应用程序的核心算法的性能[3,15]。GPU还显示出提高并发B+树性能的潜力。为了充分利用GPU上的并行计算资源,首先将并发请求缓冲在主机内存中,然后传输到GPU进行处理[43,44]。为了支持在GPU上进行并发请求处理,一些并发控制方法,如锁或软件事务内存(STM),被用于检测和解决请求之间的冲突1。然而,通过分析现有的并发GPUB+树设计[4,19],我们意识到并发冲突的检测和解决使01冲突可以分为键冲突(由相同键的请求引起)和结构冲突(由更新触发的树平衡操作引起)。20PPoPP '23,2023年2月25日至3月1日,加拿大蒙特利尔,魁北克省,张等人。0处理逻辑,从而增加内存访问和分支控制指令。这严重影响了GPUB+树的性能和服务质量(QoS),并导致响应时间的显著变化。而这些并发GPU B+树中没有一个保证线性化[13, 16,17],这对于并发请求的公平性至关重要。在上述分析的启发下,我们设计和实现了一种基于合并的并发控制框架,称为Eirene,希腊和平女神,用于GPUB+树,以最小化冲突检测和解决的开销。首先,设计了一种基于合并的同步方法来消除请求之间的键冲突。它将具有相同键的请求合并并通过它们的逻辑时间戳构建它们之间的依赖关系。时间戳指的是请求的到达顺序,在线性化语义中确定请求的结果。对于具有相同键的请求,只发出一个请求,并且其他请求的返回值由它们的依赖关系确定。为了减少结构冲突,使用了一种乐观的STM方法。根据请求的类型,将请求分为不同的内核(查询内核和更新内核)。对于查询内核,请求在没有STM保护的情况下进行处理。对于更新内核,只有在重试达到阈值时才涉及STM。最后,提出了一种基于局部性感知的warp重组优化,通过利用合并请求之间的局部性来提高内存性能并减少冲突。Eirene的这些设计可以实现更好的性能和QoS,并通过消除键冲突、减少结构冲突的可能性和利用局部性来保持线性化。我们使用YCSB基准测试[11]评估键值存储性能。评估结果显示,相对于现有的并发GPU B+树(STMGB-树[19]和LockGB-树[4]),Eirene在性能和响应时间方面都有显著改进。对于95%查询/5%更新的并发输入,它的吞吐量达到每秒24亿次。与最先进的LockGB-树[4]相比,它的性能提升了约7.43倍,响应时间的方差从36%降低到5%。总之,这项工作的贡献包括以下内容:0• 我们对并发GPUB+树的先前工作进行了全面分析,并发现了几个性能瓶颈。0•我们提出了一种基于合并的并发控制框架Eirene,这是第一个用于并发GPUB+树的基于合并的同步方法。Eirene可以减少冲突和冲突检测与解决的开销,以实现更好的性能和QoS,并保持0•结果显示,Eirene每秒可以达到24亿次的吞吐量,仅有5%的响应时间方差。0本文的剩余部分组织如下。第2节介绍了背景并讨论了我们的动机。第3节调查了相关工作。第4节介绍了基于合并的并发控制方法。第5节提出了基于局部性感知的warp重组优化。第7节介绍了Eirene的实现。第8节展示了实验结果,第9节总结了这项工作。02 背景和动机0本节首先介绍了并发GPUB+树。然后,我们讨论了现有并发GPU B+树的问题。02.1 并发GPU B+树0B+树是一种用于大规模数据存储的流行索引结构,在不同领域广泛应用,例如数据库[20, 27,33]和文件系统[36]。它是B树的一种变体,其中内部节点仅存储键和子节点引用。叶节点包含键和值。在实际应用中,B+树的并发请求(包括查询请求和更新请求,其中包括更新、插入和删除)同时进行处理。随着请求数量呈指数增长,GPU显示出提高并发B+树性能的潜力[4,41]。一个GPU包含多个流多处理器(SMs)。多个连续的线程(例如32个线程)组织成一个warp,并以锁步方式执行相同的指令。由多个warp组成的线程块构成GPU内核,即在所有SM上执行的并行函数。为了充分利用GPU上的并行计算资源,首先将并发的键值请求缓冲在主机内存中,然后传输到GPU进行处理[43,44]。当并发请求并行处理时,必须检测和解决它们之间的冲突以保证正确性。冲突可以使用软件事务内存(STM)[8, 19,40]或基于锁的方法[4, 43, 44]来检测和解决。02.2 动机0对于并发的GPUB+树,无论是基于锁还是基于STM的方法,冲突检测和解决都会降低请求处理的效率。冲突检测和解决是在树遍历和节点操作期间进行的。为了检测和解决冲突,引入了大量的条件语句和分支,这对GPUSIMT架构不友好,并且显著增加了内存访问次数。此外,当发生冲突时,请求必须重新执行或停顿,这增加了处理延迟。因此,冲突检测和解决对请求处理的吞吐量和响应时间产生严重影响[23, 37,45]。为了说明这种影响,我们收集了内存访问(������_����)和控制流指令(�������_����)的数据。 10 100 1000 10000 0 0.2 0.4 0.6 0.8 1 1.2 1.430PPoPP '23,2023年2月25日至3月1日,加拿大蒙特利尔0使用Nsight Compute [31]对基于STM的GPUB+tree[19](简称STM GB-tree)和基于锁的GPUB-tree[4](简称LockGB-tree)进行性能分析。性能分析环境与第8节中描述的相同。0memory_inst control_inst0每个请求的指令数(对数)0没有并发控制的GB-tree0STM GB-tree0Lock GB-tree0图1. STM GB-tree和LockGB-tree的性能分析。注意y轴是对数刻度。0如图1所示的数据,STMGB-tree每个请求有209次内存访问,LockGB-tree每个请求有79次内存访问。STMGB-tree每个请求有8562条控制流指令,LockGB-tree每个请求有5445条控制流指令。与没有并发控制的GPUB+树(图1中每组的第一根柱子,可以视为理想条件)相比,STM GB-tree和LockGB-tree的内存访问分别增加了2.98倍和1.12倍,控制流指令分别增加了4.49倍和2.85倍。此外,STM GB-tree和LockGB-tree的内存访问分别占总指令数的18%和10%。STMGB-tree和LockGB-tree的控制流指令分别占总指令数的30%和25%。为了说明响应时间的情况,我们还收集了STM GB-tree和LockGB-tree的最大响应时间、最小响应时间和平均响应时间。为了比较性能,所有平均结果都归一化为STMGB-tree的平均响应时间。不同方法的最大时间和最小时间分别归一化为它们的平均响应时间,以计算方差。如图2中数据的左两根柱子所示,上述方法的响应时间方差较大(STMGB-tree为40%,LockGB-tree为36%)。线性化[17]对于并发处理的正确性和公平性至关重要。如果并发处理的结果与按照它们的先前顺序(通常是实时顺序)从顺序执行中获得的结果相同,则并发处理是线性化的。线性化可以为并发请求提供重要的正确性条件和公平性,这在大多数并发系统中是必需的[13]。在STMGB-tree或LockGB-tree的设计中,主要关注的是利用GPU中的并行性。它们都不能保证线性化。0STM GB树 锁 GB树 Eirene0标准化每个请求的时间0图2. 标准化每个请求的时间。03 相关工作0为了更高效地处理并发请求,许多工作致力于通过软件或硬件加速来提高索引结构的性能。本节讨论了与索引结构加速、并发GPUB+树和合并同步技术相关的工作。索引结构加速。随着并发请求的数量呈指数增长,人们努力通过并行硬件(如FPGA用于哈希表[7, 26, 39],GPU用于哈希表[43, 44]和跳表[29,30],硬件事务内存(HTM)用于B+树[38]以及GPUHTM[9,10])来加速并发索引结构,并且还提出了一些改进在GPU上提高索引结构的查询性能[35,41]。与先前的工作相比,Eirene专注于使用GPU加速B+树,并且不仅关注查询性能,还关注并发(混合查询和更新)性能。并发GPUB+树。Holey等人[19]提出了一种轻量级STM设计,具有简单的STM接口,可以很好地适应许多GPU线程。Awad等人[4]提出了一种并发GPUB树设计。它使用细粒度锁进行并发控制,并针对GPUSIMT进行了优化,以实现合并内存访问并避免分支分歧。然而,这两个工作无法实现令人满意的性能和QoS,并且不支持线性化语义。Yan等人[41]通过几种优化提高了GPU上B+树的查询吞吐量。然而,它不支持并发处理。在我们的工作中,我们使用[19]提供的STM接口来保证更新请求的树遍历的正确性,该接口提供了轻量级的STM原语,开销很小。基于此,我们提出了一种乐观策略,用于B+树上的STM,以最小化由于虚假事务冲突引起的可能回退开销。合并同步技术。合并同步是一种用于减少多核环境中同步冲突的技术。一个或多个线程成为合并器,将其他线程的具有相同键的请求合并为大请求,并由合并器处理这些合并的请求。实验证明,基于合并的实现方式R1R5R5R4W4,aW4,bR4U(5,e)U(5,f)Q1Q4U(4,a)Q4U(4,b)T7Q1T8Q1U(5,f)U(4,b)Q1Q4U(4,a)Q4U(5,e)T7Q1U(5,f)U(4,b)Q1Q4U(4,a)Q4U(5,e)T7a40PPoPP '23,2023年2月25日至3月1日,加拿大蒙特利尔,魏翔等人0优于细粒度锁实现。一些工作[5, 14, 28,34]使用合并同步技术来提高各种数据结构(如堆、队列和红黑树)的并发性能。Ak-senov等人[2]提出了一种将扁平合并应用于批处理的并行合并方法,并展示了显式同步在并发场景中的性能优势。然而,这些工作是为CPU上的并发处理而设计的,而不是GPU。与这些工作相比,Eirene是第一个将合并同步技术应用于并发GPU B+树的系统。04 基于合并的并发控制在本节中,我们介绍基于合并的并发控制方法的设计,它减少了冲突的影响并实现了线性化。04.1 基于合并的同步合并同步是一种低开销的并发同步技术,其中合并线程将具有相同键的并发请求合并并逐个处理,以减少冲突开销[2, 14,16]。然而,先前的合并设计主要集中在多核CPU上,不能应用于GPU。04.1.1基于合并的同步。当应用合并同步以减少冲突开销时,合并器需要能够检测具有相同键的请求(找到键冲突)。朴素地,可以通过比较任意两个请求之间的键来实现。由于不同请求的输入键是随机分布的,这种检测是耗时的。然而,如果按照键对请求进行排序,检测开销将大大减小,因为每个请求只需要与其相邻的请求进行比较。在Eirene中,首先按照目标键对请求进行排序,然后按照它们的逻辑时间戳对具有相同键的请求进行排序。注意,由于这些请求在到达GPU时按时间顺序缓冲,因此我们将缓冲区中的顺序视为它们的逻辑时间戳。为了合并具有相同键的请求,首先扫描排序后的请求。如果存在具有相同键的多个请求,并且它们都是查询请求或更新请求,那么这些请求将被合并。只有具有最大时间戳的请求将被发出以遍历B+树。如果特定键的所有请求都是查询请求,则未使用的查询请求可以与发出的查询请求共享结果。如果特定键的所有请求都是更新请求,则忽略未发出的更新请求,因为它们的值将被最后一个更新请求覆盖。如果对于同一键存在混合类型的请求,即查询和更新请求,则发出具有最大时间戳的更新请求,并通过发出的请求检索叶节点中的旧(原始)值。对于一个查询请求0在混合请求中,如果在其之前有一个删除请求而没有其他更新请求在它们之间,其结果应该为null。如果在其之前有一个更新请求而没有删除请求在它们之间,它使用此更新请求的值作为其结果。否则,它使用由发出的请求检索到的旧(原始)值作为其结果。因此,所有未发出的请求可以在不遍历B+树的情况下计算出它们的结果。0T2 T3 T5 T6 时间戳0(a) 原始请求0T8 T2 T3 T5 时间戳 T6 T4 T10(b) 排序后的请求01_值0T8 T2 T3 T5 时间戳 T4 T6 T104_旧值0(c) 请求合并0图3. 合并方法的工作示例。0这里给出一个例子来说明基于合并的同步方法的工作原理。如图3所示,��表示一个目标键为�的查询请求。�(�,�)表示一个更新请求,用于将键�更新为值�。��表示一个时间戳,其中较小的�对应较小的时间戳,即发生在较早的时间。图3(a)中的请求是带有时间戳的原始请求。如图3(b)所示,具有相同键的排序请求在时间戳顺序中是相邻的。在这个例子中,目标键1的请求都是查询,因此发出了�1和�8的�1,并且�1和�4共享结果(1_���)。目标键5的请求都是更新。因此,只发出了最后一个请求(�(5,�))。具有目标键4的请求是混合类型的。合并后,发出了最后一个更新操作(�(4,�)),并返回旧(原始)值(4_���_���)而不是值a或b。没有时间戳小于�4和�2的更新,因此其结果是4_���_���。�4和�5获取更新�(4,�)的值(�),该值是它之前最近的更新操作的值。排序和合并后,消除了请求之间的键冲突,因为每个键只发出一个请求。同时,所有其他未发出的具有相同键的请求可以根据它们的时间戳获取它们的结果,从而实现了线性化[17]。04.1.2 范围查询。R(3,5)U(4,e)......T2T10U(4,b)T1T104a3 4 5ceR(3,5)U(4,e)T2443 4 5 63650PPoPP '23,2023年2月25日至3月1日,加拿大蒙特利尔0范围查询是并发B+树的一种基本操作,用于查询范围内的数据。由于范围查询获取的是一系列键而不是一个键,范围查询不能直接通过合并同步方法进行处理。此外,如果范围查询仍然按照其原始方式进行处理,它们可能会得到错误的结果,因为其范围内的一些更新请求可能会与其他请求合并。0(a) 原始请求0结果0旧值0应该是‘b’0发出的请求0(b) 范围错误0图4. 原始范围查询在合并方法下会出错的示例。0图4显示了一个例子。�(�,�)表示一个具有下界�和上界�的范围查。图4(a)显示了原始请求序列。图中显示了所有针对键4的请求。如图4(b)所示,具有时间戳�2的�(3,5)应该获取键4的值�。然而,在合并键4的请求之后,只发出了�(4,�)。因此,�(3,5)只能获取键4的旧值或由�(4,�)写入的值�,这是错误的值。我们引入了一种机制,以确保范围查询可以与基于合并的同步方法正确工作。当处理缓冲请求时,所有范围查询也根据它们的下界值与其他请求一起排序。然后,扫描排序后的请求。如果一个键在范围查询的范围内,并且存在一些更新该键的请求,则为该键生成一个带有范围查询的时间戳的人工查询,并根据其时间戳将其插入到该键的请求序列中。范围查询首先以原始方式获取其结果。对于每个具有人工查询的键,其值将使用相应人工查询的结果进行更新。图5给出了范围查询的示例。图5(a)0显示原始请求。�(3,6)是一个范围查询,它查询从键3到键6的键。图5(b)显示了�(3,6)的处理过程。首先根据键和时间戳对请求进行排序。键4和键6在�(3,6)的范围内。此外,还有它们的更新请求。因此,为�(3,6)生成了两个人工查询(��4和��6),时间戳为�2,并将其插入到相应的位置。这请求被0R(3,6)0T20Q40T40U(4,e)0T50U(4,b)0T10Q30T30U(6,a)0T60(a)排序和合并之前的请求0Q4 R(3,6) U)0T2 T4 T50U(4,b)0T10Q30T30U(6,a)0T60Q r 40T20Q r 60T20结果03_v5_val a03_val 6_val0b 6_val0update0al0e0a0b 6_val0(b)范围查询处理0图5.范围查询处理示例0处理,��4和��6获得它们的结果(�和6_���)。之后,键4的范围查询结果被键��4的结果(�)更新,键6的范围查询结果被键��6的结果(6_���)更新。04.2基于合并的并发控制0在并发请求中,存在键冲突和结构冲突。在我们的并发控制设计中,我们首先使用基于合并的同步方法消除键冲突。然后,我们使用基于STM的方法处理结构冲突以保证正确性。基于合并的并发控制方法的整个工作流程如算法1所示。当并发请求被传输到GPU时,它们首先通过基于合并的同步方法(第2行)进行处理,以消除键冲突。在将具有相同键的请求合并之后,已发出的请求之间只存在结构冲突。然后,我们根据请求的类型(第3行)将要发出的请求分区到不同的GPU内核中:查询内核(用于查询请求和���������请求)和更新内核(用于更新请求,包括更新、插入和删除)(第3-5行)[4,43]。我们根据以下观察进行这样的内核分区:对于查询请求,它们只需要获取值而不进行任何树结构修改。即使它们没有受到保护,也可以正确执行。不对查询请求应用同步保护可以减少结构冲突的可能性。首先启动查询内核,然后启动更新内核。这两个内核是分开执行的。由于已发出的请求之间没有键冲突,发出请求的执行顺序不会影响返回的结果。查询内核上没有同步保护,并且可以应用先前工作中的一些优化来进一步提高树遍历的性能(第19行)。对于更新内核(第23-46行),我们使用急切的STM[19]来保护更新请求,并使用多树区域分割方法[12,38]来3:4:QUERY_KERNEL(QueryReqs)5:UPDATE_KERNEL(UpdateReqs)6:RESULT_CAL(Reqs,QueryReqs,UpdateReqs))17:18:val = &QueryReqs[id].val19:leaf = findLeaf (root, key)20:val = findValInLeaf(leaf, key)25:29:31:X_BEGIN()32:leaf = findLeaf(root,key)33:TX_END()35:s = leaf.version39:41:etry_count++45:D()Warp 2Warp 3Warp 1Q8Q7Q6Q5Q4Q2Q1Q3Q9…123415116219322425……1558……2341560PPoPP '23,2023年2月25日至3月1日,加拿大蒙特利尔,QC,张等人0算法1 基于合并的并发控制0声明:0Reqs , QueryReqs , UpdateReqs :指向传入请求数组、查询请求和更新请求的指针。CombReqs :指向合并后的请求数组的指针。01: 过程 基于合并的并发控制 2: CombReqs =COMBINING(Reqs)07: 结束过程 8: 9: 过程COMBINING(Reqs) 10: Sort(Reqs)012: 返回 CombReqs013: 结束过程 14: 15: 过程QUERY_KERNEL(QueryReqs) 16: id = blockIdx.x* blockDim.x + threadIdx.x021: 结束过程 22: 23: 过程UPDATE_KERNEL(UpdateReqs) 24: id =blockIdx.x * blockDim.x + threadIdx.x027: RETRY: � 内部节点区域028: 如果 retry_count < THRESHOLD 则030: 否则034: 结束如果036: 37: TX_BEGIN() � 叶节点区域038: 如果 leafvers == leaf.version & key 在范围内(leaf) 则040: 否则043: 转到 RETRY044: 结束如果046: 结束过程0减少结构冲突。树遍历分为两部分:内部节点遍历(第 29行)和叶节点操作(第 39行)。只有叶节点操作始终由STMs保护(第 37 - 45行),而内部节点遍历则0首先是不受保护的(第 29行)。当请求处理过程中发生冲突时,会进行重试。如果重试次数达到阈值,对应请求的内部节点遍历将由STM保护(第 30 - 34行)。在内部节点遍历之后,请求获取到目标叶节点的引用。在请求获取到引用之后,可能会出现一致性问题,即在请求获取到引用之后,目标叶节点被分裂。为了解决这个问题,我们在内部节点遍历和叶节点操作之间使用了验证机制[ 38]。每个叶节点都有一个版本号,该版本号也由其父节点存储。一旦叶节点被分裂,其版本号会原子性地增加一,并更新存储在其父节点中的版本号。在获取到叶节点的引用时,也会获取到相应的版本号(第 35行)。在对叶节点进行操作之前,会将记录的版本号与叶节点的版本号进行比较。如果两个版本号不相等,则发生了冲突,请求处理将中止并重试(第 38 - 44行)。发出的请求完成后,它们将检索节点的旧值。如第 4.1小节所讨论的,未使用的请求根据其与发出的请求之间的依赖关系通过调用 ������ _ ��� 方法计算其结果(第 6行)。此校正过程在GPU上并行执行。05 本地感知的Warp重新组织0相同目标节点 相邻目标节点0RG1 RG2 RG30Q45 Q46 Q470远程目标节点0Warp 40g0叶节点0树高:50节点id0RG40(a) 原始warp0节点id RF值 时间0迭代warp0缓冲状态0null0RG 1 Q3 Q2 Q10Q6 Q5 Q4 RG 20Q9 Q8 Q7 RG 30Q47 Q46 Q45 RG 40(b) 迭代warp0图6. 本地感知warp重新组织方法的示例。70PPoPP ’23,2023年2月25日至3月1日,加拿大蒙特利尔0在请求排序和合并后,发出的相邻请求很可能访问相同或相邻的叶节点。当这些请求被分成请求组(RGs)并映射到不同的warp时,这些warp也可能访问相同或相邻的叶节点。图6(a)展示了这样一个例子,其中��表示具有键�的请求。相同颜色的多个请求意味着它们的目标键位于同一个叶节点中。例如,RG1中的�3和RG2中的�4具有相同的颜色(粉色),表示它们的目标键都在叶节点2中。在这个例子中,树的高度为5。在这些请求被分成不同的RGs之后,相邻的RGs将访问相同或相邻的叶节点。例如,RG2和RG1将访问相同的叶节点2,RG3将访问与RG2访问的叶节点3相邻的叶节点4。因此,如果一个RG能够从其相邻RG的目标叶节点以及与目标叶节点相连的叶节点中获取其目标叶节点,而不是从根节点遍历,与从根节点遍历相比,遍历步骤(需要遍历的节点数)可以减少。然而,在实际中,这些RG被映射到不同的warp之后,很难重用这样的局部性,因为warp在同一个SM中动态调度,调度顺序无法预先定义。为了利用相邻RG之间的局部性,提出了一种称为本地感知warp重新组织的方法。在这种方法中,几个相邻的RG被组织成一个warp,并以循环方式执行(称为迭代warp)。迭代warp中的RG被映射到一个warp,并逐个执行。每个迭代warp都有一个共享缓冲区来存储目标叶节点。当执行一个迭代warp时,它的第一个RG按照原始方式遍历树(从根到叶,也称为垂直遍历)。在一个RG完成执行后,最后访问的叶节点(可能是最接近下一组的叶节点)被存储在缓冲区中。接下来的RG将通过从缓冲节点开始遍历链接的叶节点来获取其目标叶节点(称为水平遍历)。对于更新请求,水平遍历使用锁或STM进行保护。如果发生冲突,它们将重试并通过垂直遍历遍历树。因为原始相邻RG之间可能存在结构冲突,这样的设计可以避免它们被合并为迭代warp后的一些结构冲突。如果缓冲左节点和目标叶节点之间的步骤或遍历的节点数小于树的高度,这样的设计可以表现更好。然而,如果步骤大于树的高度,性能将下降。为了避免这种情况,设计了一个机制来帮助Eirene决定是垂直遍历还是水平遍历。每个叶节点中都扩展了一个范围字段(RF),用于记录从该节点开始的遍历步骤为树高加一的节点的最小键。例如,假设树的高度为5,叶节点6的最小键为16;叶节点1的RF值将为16。这些RF值在构建树时进行初始化。0只有当叶节点是水平遍历的起始点且遍历步长大于树高时,叶节点的RF值才会更新。当叶节点信息存储在迭代warp缓冲区中时,其RF值也存储在缓冲区中。对于下一次迭代,它将首先检查其键是否大于缓冲的RF值。如果是,则下一次迭代将垂直遍历树。否则,树将水平遍历。由于一次迭代中的请求以SIMT方式执行,不同线程的执行步骤遵循最大步骤。因此,使用迭代中的最大目标键而不是最小目标键与存储的RF值进行比较。另一个设计考虑因素是warp中的迭代次数。较大的迭代次数会增加迭代之间的局部性。然而,它会牺牲warp之间可用的并行性。为了充分利用计算资源,RG均匀分布在GPU上的不同SM中。然后,它们被组织成在每个SM上执行的迭代warp。图6(b)展示了这种方法的工作原理。假设树高为5,不同叶节点的RF值如图6(b)中的表所示。四个RG被组织成一个迭代warp,而不是映射到不同的warp。在执行第一次迭代之前,缓冲区是空的。第一次迭代(RG1)中的请求从根节点遍历到叶节点。在RG1的请求完成后,缓冲区被其最大访问的叶节点(叶节点2)填充。当执行RG2时,它从缓冲区获取叶节点2,并检查其最大键(6)是否小于缓冲的RF值(19)。由于最大键小于RF值,它水平遍历树。在RG3完成后,叶节点4存储在缓冲区中。在执行RG4之前,它首先检查其最大键(47)是否大于缓冲的RF值(25)。由于最大键大于RF值,它垂直遍历树。06个证明概述本节概述了基于组合的并发控制的正确性和线性化的论证。主要的证明义务是展示设计可以正确工作,并且由它处理的并发请求的结果与按照它们的逻辑时间戳顺序处理这些请求的结果相同。证明概述。线性化[17]是并发的正确性条件。它可以定义为并发请求的执行结果与实时或时间戳顺序的顺序执行结果一致。因此,设计是否满足线性化取决于如何处理具有相同键的请求。如果具有相同键的请求的执行结果等同于按照时间戳顺序的顺序执行结果,则设计是线性化的。在我们的设计中,每个请求都有一个逻辑时间戳,根据其在缓冲区中的到达顺序确定。逻辑时间戳在请求到达GPU时已经确定。如果具有相同键的请求都是查询或范围查询请求,无论以什么顺序执行这些请求,它们总是会得到相同的结果。因此,在基于组合的并发控制中,请求可以获得与顺序执行相同的结果。对于一些具有混合请求类型的键,它们首先按照它们的时间戳排序,并且在我们的基于组合的同步方法中构建它们的依赖关系。然后只发出一个请求,其他未发出的请求将根据依赖关系获得它们的结果。因此,这样的设计可以实现线性化,因为具有相同键的请求按照它们的时间戳顺序顺序执行时得到一致的结果。在基于组合的并发控制之后,只存在结构冲突。因此,将这些请求分成不同的内核是安全的。对于一个更新请求,树遍历被分为内部节点遍历和叶节点操作。首先在没有任何STM保护的情况下遍历内部节点。当失败时,它会重试直到达到阈值。然后打开STM保护。此外,叶节点操作始终受到STM保护。这种方法的正确性已经在[38]中得到证明。因此,基于组合的并发控制设计可以正确工作。在局部感知的warp重新组织方法中,通过利用局部性选择不同的遍历方式来搜索目标叶节点。局部感知的warp重新组织不会影响并发请求的结果,因此对于线性化是无害的。由于水平遍历在叶节点上执行,并且始终对更新请求进行STM保护,所以局部感知的warp重新组织的正确性可以得到保证。• Can Eirene achieve better performance than state-of-80PPoPP ’23,2023年2月25日至3月1日,加拿大蒙特利尔,张等人。0每个请求都有一个逻辑时间戳,根据其在缓冲区中的到达顺序确定。逻辑时间戳在请求到达GPU时已经确定。如果具有相同键的请求都是查询或范围查询请求,无论以什么顺序执行这些请求,它们总是会得到相同的结果。因此,在基于组合的并发控制中,请求可以获得与顺序执行相同的结果。对于一些具有混合请求类型的键,它们首先按照它们的时间戳排序,并且在我们的基于组合的同步方法中构建它们的依赖关系。然后只发出一个请求,其他未发出的请求将根据依赖关系获得它们的结果。因此,这样的设计可以实现线性化,因为具有相同键的请求按照它们的时间戳顺序顺序执行时得到一致的结果。在基于组合的并发控制之后,只存在结构冲突。因此,将这些请求分成不同的内核是安全的。对于一个更新请求,树遍历被分为内部节点遍历和叶节点操作。首先在没有任何STM保护的情况下遍历内部节点。当失败时,它会重试直到达到阈值。然后打开STM保护。此外,叶节点操作始终受到STM保护。这种方法的正确性已经在[38]中得到证明。因此,基于组合的并发控制设计可以正确工作。在局部感知的warp重新组织方法中,通过利用局部性选择不同的遍历方式来搜索目标叶节点。局部感知的warp重新组织不会影响并发请求的结果,因此对于线性化是无害的。由于水平遍历在叶节点上执行,并且始终对更新请求进行STM保护,所以局部感知的warp重新组织的正确性可以得到保证。07 Eirene的实现 我们使用CUDAC++实现了Eirene的GPU并发B+树,代码量为2200行。在这里,我们描述了Eirene实现的细节。树结构是一个常规的B+树,其中内部节点包含键和子指针。子指针指向子节点的位置。整个树结构存储在GPU的全局内存中。我们在CPU端缓冲并发请求,并在数量达到可配置的阈值(当前实现中为100万)后将它们传输到GPU。在基于合并的同步方法中,使用CUB库[32]中的基数排序算法根据键和逻辑时间戳对请求进行排序。排序后,我们通过并行地扫描每个请求与其相邻的请求来合并这些请求,其中每个GPU线程扫描一个请求。然后,我们实现了两个CUDA核函数(query和update)来发出不同类型的请求,其中query核函数在update核函数之前启动。由于query核函数只处理不修改B+树的查询或范围查询请求,我们在处理这些请求时不使用任何STM保护,并使用[41]中的缩小线程组优化来加速查询性能。update核函数执行树遍历和更新操作,采用乐观的STM方法实现。我们使用渴望冲突检测[19]实现了STM。理论上,实现中可以使用除STM之外的其他同步方案,如细粒度锁。为了支持局部感知的warp重新组织方法,我们将多个连续的warp重新组织在一起形成一个迭代warp。在原始的warp中,每个线程只处理一个请求。在迭代warp中,每个线程将逐个处理来自多个原始warp的相同位置的请求。0• Eirene能否比LockGB-tree等最先进的系统实现更好的性能?08 评估在本节中,我们评估了Eirene的性能,并尝试回答以下问题:0• Eirene能否实现更好的QoS? •Eirene是否解决了第2节中的挑战? •每个设计选择对性能有何影响? •Eirene的范围查询性能如何?08.1 实验方法0我们的实验在一台64核服务器上进行(2个32核的AMDEPYC 7532s),配备了一块Ampere GPU(NVIDIAA100)。我们使用GCC 7.5.0和CUDA 11在Ubuntu18.04(内核5.4.0)上编译所有代码,并使用O3优化选项。我们按照第7节中描述的方式实现了Eirene的原型。性能基准使用了STM GB-tree [19]和Lock GB-tree [4]。LockGB-tree是开源的,我们根据论文[19]实现了STMGB-tree。为了展示设计的性能有效性,我们使用了广泛使用的键值存储基准测试Yahoo! Cloud ServingBenchmark(YCSB)[11]对Eirene进行评估。每个记录有一个32位键和一个32位值。我们使用具有不同查询/更新比例的请求数据集。默认的查询/更新比例是95%/5%。默认的分布是均匀分布。这些请求在不同的树大小上进行处理,包括2^23、2^24、2^25和2^26,其中默认的树大小是2^23。所有结果都是通过5次执行的平均值。Eirene的所有评估都包括基于合并的并发控制方法中的排序时间。在本文中,我们专注于GPU上B+树在并发请求下的性能。因此,结果不包括CPU和GPU之间的传输时间。虽然CPU到GPU之间的传输时间是不可忽略的,但我们在当前的设计和评估中没有考虑它,也没有与现有的工作进行比较。 0 300 600 900 1200 1500 1800 2100 240023242526STM GB-treeLock GB-treeEirene 0 1 2
下载后可阅读完整内容,剩余1页未读,立即下载
cpongm
- 粉丝: 4
- 资源: 2万+
上传资源 快速赚钱
- 我的内容管理 收起
- 我的资源 快来上传第一个资源
- 我的收益 登录查看自己的收益
- 我的积分 登录查看自己的积分
- 我的C币 登录后查看C币余额
- 我的收藏
- 我的下载
- 下载帮助
会员权益专享
最新资源
- zigbee-cluster-library-specification
- JSBSim Reference Manual
- c++校园超市商品信息管理系统课程设计说明书(含源代码) (2).pdf
- 建筑供配电系统相关课件.pptx
- 企业管理规章制度及管理模式.doc
- vb打开摄像头.doc
- 云计算-可信计算中认证协议改进方案.pdf
- [详细完整版]单片机编程4.ppt
- c语言常用算法.pdf
- c++经典程序代码大全.pdf
- 单片机数字时钟资料.doc
- 11项目管理前沿1.0.pptx
- 基于ssm的“魅力”繁峙宣传网站的设计与实现论文.doc
- 智慧交通综合解决方案.pptx
- 建筑防潮设计-PowerPointPresentati.pptx
- SPC统计过程控制程序.pptx
资源上传下载、课程学习等过程中有任何疑问或建议,欢迎提出宝贵意见哦~我们会及时处理!
点击此处反馈
安全验证
文档复制为VIP权益,开通VIP直接复制
信息提交成功