CN110083488B - 一种面向gpgpu的细粒度低开销的容错系统 - Google Patents

一种面向gpgpu的细粒度低开销的容错系统 Download PDF

Info

Publication number
CN110083488B
CN110083488B CN201910320906.4A CN201910320906A CN110083488B CN 110083488 B CN110083488 B CN 110083488B CN 201910320906 A CN201910320906 A CN 201910320906A CN 110083488 B CN110083488 B CN 110083488B
Authority
CN
China
Prior art keywords
gpu
error
execution
kernel
fault
Prior art date
Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
Active
Application number
CN201910320906.4A
Other languages
English (en)
Other versions
CN110083488A (zh
Inventor
季振洲
郭明周
李金宇
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
Harbin Institute of Technology
Original Assignee
Harbin Institute of Technology
Priority date (The priority date is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the date listed.)
Filing date
Publication date
Application filed by Harbin Institute of Technology filed Critical Harbin Institute of Technology
Priority to CN201910320906.4A priority Critical patent/CN110083488B/zh
Publication of CN110083488A publication Critical patent/CN110083488A/zh
Application granted granted Critical
Publication of CN110083488B publication Critical patent/CN110083488B/zh
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

Classifications

    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F11/00—Error detection; Error correction; Monitoring
    • G06F11/07—Responding to the occurrence of a fault, e.g. fault tolerance
    • G06F11/14—Error detection or correction of the data by redundancy in operation
    • G06F11/1402—Saving, restoring, recovering or retrying
    • G06F11/1415—Saving, restoring, recovering or retrying at system level
    • G06F11/142—Reconfiguring to eliminate the error
    • G06F11/1428—Reconfiguring to eliminate the error with loss of hardware functionality
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F11/00—Error detection; Error correction; Monitoring
    • G06F11/07—Responding to the occurrence of a fault, e.g. fault tolerance
    • G06F11/14—Error detection or correction of the data by redundancy in operation
    • G06F11/1402—Saving, restoring, recovering or retrying
    • G06F11/1446—Point-in-time backing up or restoration of persistent data
    • G06F11/1458—Management of the backup or restore process
    • Y—GENERAL TAGGING OF NEW TECHNOLOGICAL DEVELOPMENTS; GENERAL TAGGING OF CROSS-SECTIONAL TECHNOLOGIES SPANNING OVER SEVERAL SECTIONS OF THE IPC; TECHNICAL SUBJECTS COVERED BY FORMER USPC CROSS-REFERENCE ART COLLECTIONS [XRACs] AND DIGESTS
    • Y02—TECHNOLOGIES OR APPLICATIONS FOR MITIGATION OR ADAPTATION AGAINST CLIMATE CHANGE
    • Y02D—CLIMATE CHANGE MITIGATION TECHNOLOGIES IN INFORMATION AND COMMUNICATION TECHNOLOGIES [ICT], I.E. INFORMATION AND COMMUNICATION TECHNOLOGIES AIMING AT THE REDUCTION OF THEIR OWN ENERGY USE
    • Y02D10/00—Energy efficient computing, e.g. low power processors, power management or thermal management

Landscapes

  • Engineering & Computer Science (AREA)
  • Theoretical Computer Science (AREA)
  • Quality & Reliability (AREA)
  • Physics & Mathematics (AREA)
  • General Engineering & Computer Science (AREA)
  • General Physics & Mathematics (AREA)
  • Hardware Redundancy (AREA)
  • Retry When Errors Occur (AREA)

Abstract

本发明提出了一种面向GPGPU的细粒度低开销的容错系统，其中包括任务划分模块，检查点备份模块，冗余执行与错误检测模块，错误修复模块。本发明可以实现对GPU计算部件瞬时故障的容错处理，并且可以解决GPU的传统软件容错方法中容错粒度大、错误修复代价高、容错系统性能差等问题。本发明的有益效果：能够把线程任务进行划分，减小kernel的计算规模，在检查点备份时只需对相对活跃变量进行备份，减少了存储带来的时空开销，在错误修复时只需把与错误相关的部分对象进行复算，减小了复算带来的容错代价，并且充分利用CPU‑GPU异构系统的异步机制，隐藏了因为数据传输带来的时间延迟，提高了系统的性能。

Description

一种面向GPGPU的细粒度低开销的容错系统
技术领域
本发明涉及计算机技术领域，具体而言，涉及一种面向GPGPU的细粒度低开销的容错系统。
背景技术
近年来，通用图形处理器单元(GPGPU)由于其卓越的计算能力、内存访问带宽和改进的可编程性而变得日益流行。利用GPU的计算能力进行高性能计算的异构并行计算机已经受到绝大多数科学领域的研究人员的青睐，包括金融分析，地震探测，高能量物理，量子化学，分子动力学，甚至药物设计。
由于GPU初始主要应用于图形图像处理领域，而该领域的应用本身就具有一定程度的容错特性，单个像素点计算结果的错误并不影响整个图像的显示效果，因此传统的图形处理器设计没有考虑计算部件的可靠性问题。即使NVIDIA公司在Fermi架构中为GPU的存储系统提供了ECC校验，解决了GPU存储部件的可靠性问题，但是计算逻辑单元中仍不提供支持软硬错误检测的硬件以及其他容错机制。由于高性能计算领域对计算结果的正确性要求极其严格，因此需要对GPU计算部件的瞬时故障进行容错处理。
当前针对解决GPU的容错问题所采用的技术可分为硬件冗余和软件冗余技术。增加冗余硬件的容错方法并不适用于目前被广泛使用的GPU，因此主要采用软件冗余技术解决GPU的容错问题。然而，GPU的软件容错方法仍处于起步阶段，存在容错粒度大，错误修复代价高，容错系统性能差，实现容错难度大等问题。
发明内容
基于背景技术存在的技术问题，本发明提出了一种面向GPGPU的细粒度低开销的容错系统。
本发明提出的一种面向GPGPU的细粒度低开销的容错系统，所述系统包括任务划分模块，检查点备份模块，冗余执行与错误检测模块，错误修复模块。
优选地，任务划分模块用于从总体上划分输入数据集合，然后将数据相关的计算kernel和数据传输归入单个流中，在系统级实现多个流的并行。
优选地，检查点备份模块用于在kernel执行前设置检查点标记，然后利用CPU-GPU异构系统的异步机制，分别对CPU、GPU上的数据进行备份保存并记录相应的内存地址，以便后续进行错误修复。
优选地，冗余执行与错误检测模块用于记录各条GPU语句的执行顺序，并且通过冗余副本执行来检测数据运算结果的正确性，当错误发生时，该模块需返回GPU计算结果的详细错误信息。
优选地，错误修复模块用于根据给定详细的错误信息定位到需要复算的对象，然后采用检查点备份模块所备份的数据恢复复算对象的初始值，最后复算这些对象，进而修复之前的错误数据。
本发明的技术特点及有益效果：
(1)本发明的任务划分模块利用GPGPU计算、GPU-CPU之间数据传输的异步性，可以实现计算与数据传输在时间轴上的重叠，从而隐藏因为数据传输带来的时间延迟，提高系统性能，并且由于每个kernel的计算规模减小，可以减少纠错时复算的计算量。
(2)本发明的检查点备份模块利用CPU-GPU异构系统的异步机制，分别对CPU、GPU进行检查点备份，可以有效地利用CPU的空闲时间，同时也避免同步模式下更多的存储时间开销。当进行检查点备份时，只需备份目标程序段中的相对活跃变量，以此减少备份所需的存储空间。
(3)本发明的错误修复模块不需要复算自上一检查点至错误发生位置的所有计算，只需复算部分与错误相关的对象，这样可以有效降低复算带来的系统开销。
附图说明
图1为本发明的任务划分阶段示意图。
图2为本发明异步执行检查点备份的流程示意图。
图3为本发明的冗余代理线程组织示意图。
图4为本发明的整体运行流程示意图。
具体实施方式
下面结合具体实施例对本发明作进一步解说。
实施例
参考图1，任务划分模块利用CUDA平台的流计算模式将输入数据集划分为N个数据子集，然后将数据相关的计算kernel和数据传输归入单个流中，在系统级实现N个流的并行执行。可以使用cudaStreamCreate()和cudaMemcpyAsync()这两个功能函数，实现流的创建和流中的数据异步传输的功能。这种方法利用GPGPU计算、GPU-CPU之间数据传输的异步性，可以实现计算与数据传输在时间轴上的重叠，从而隐藏因为数据传输带来的时间延迟，提高系统性能，并且由于每个kernel的计算规模减小，可以减少纠错时复算的计算量。
参考图2，由于kernel执行过程中无法中断，因此将检查点统一设置在kernel执行前的位置，并且利用CPU-GPU异构系统的异步机制，分别对CPU、GPU进行检查点备份，可以有效地利用CPU的空闲时间，同时也避免同步模式下更多的存储时间开销。当进行检查点备份时，在kernel中未被赋值的数据无需备份，只需备份目标程序段中的相对活跃变量，即那些在该目标程序段中被定过值，且首次操作为被引用的变量，其余的数据可以通过活跃变量恢复得到，以此减少备份所需的存储空间。所述方法包括如下步骤：
步骤1：在kernel执行前设置检查点标记；
步骤2：采用编译的定值-引用分析技术分析kernel中的变量，将其中的相对活跃变量备份保存到CPU的内存当中，并记录相应的内存地址。
步骤3：执行kernel调用；
步骤4：当GPU在进行kernel计算时，将CPU上的数据进行备份保存到内存当中，并记录相应的内存地址。
参考图3，本实施例利用CUDA计算平台上的线程执行层次，在线程block级上把原kernel程序和冗余代理融合在计算kernel的一次发射中并行执行，这样可以充分利用GPU中的计算资源，减少冗余执行带来的时间开销。冗余执行和错误检测包括如下步骤：
步骤1：CPU端调用kernel程序时，为每个GPU相关语句静态指定一个唯一的ID；
步骤2：将原kernel程序中每个Grid的线程块数目N修改为2*N,修改新增的N个线程块的线程号，使其与原来N个线程块的线程号相同；
步骤3：为原kernel程序中的每个输入变量、输出变量添加一个副本，将这些变量副本使用CudaMalloc分配在GPU端，然后把冗余代理的原输入和输出指针分别指向输入副本和输出副本；
步骤4：执行原kernel和冗余代理的计算任务，记录下GPU语句的执行顺序；
步骤5：比较原kernel程序和冗余代理的各个计算结果，如果两个计算结果不相同，则把详细的错误信息存放在错误位数组中；
步骤6：遍历步骤5中的错误位数组，如果错误位数组不空，则把该错误位数组和GPU语句的执行顺序发送到错误修复模块，否则，表明计算正确，返回计算结果，结束该阶段执行；
以上所述，仅为本发明较佳的具体实施方式，但本发明的保护范围并不局限于此，任何熟悉本技术领域的技术人员在本发明揭露的技术范围内，根据本发明的技术方案及其发明构思加以等同替换或改变，都应涵盖在本发明的保护范围之内。

Claims (4)

1.一种面向GPGPU的细粒度低开销的容错系统，其特征在于：能够把线程任务进行划分，在检查点备份时只需对具有修改标识的变量进行备份，在错误修复时只需把那些直接或间接影响了最终错误计算结果的对象进行复算，减小了复算带来的容错代价，实现了对GPGPU计算部件瞬时故障的容错处理；所述系统包括任务划分模块，检查点备份模块，冗余执行与错误检测模块，错误修复模块；
所述任务划分模块，利用CUDA平台的流计算模式将输入数据集划分为N个数据子集，然后将数据相关的计算kernel和数据传输归入单个流中，在系统级实现N个流的并行执行；
所述检查点备份模块，用于在kernel执行前设置检查点标记，然后利用CPU-GPU异构系统的异步机制，分别对CPU、GPU上的数据进行备份保存并记录相应的内存地址，以便后续进行错误修复；
所述冗余执行与错误检测模块，用于记录各条GPU语句的执行顺序，并且通过冗余副本执行来检测数据运算结果的正确性，当错误发生时，该模块需返回GPU计算结果的详细错误信息；
所述错误修复模块，用于根据给定详细的错误信息定位到需要复算的对象，然后采用检查点备份模块所备份的数据恢复复算对象的初始值，最后复算这些对象，进而修复之前的错误数据。
2.根据权利要求1所述的面向GPGPU的细粒度低开销的容错系统，其特征在于所述检查点备份模块；
由于kernel执行过程中无法中断，因此将检查点统一设置在kernel执行前的位置，并且利用CPU-GPU异构系统的异步机制，分别对CPU、GPU进行检查点备份；当进行检查点备份时，在kernel中未被赋值的数据无需备份，只需备份目标程序段中的相对活跃变量，即那些在该目标程序段中被定过值，且首次操作为被引用的变量，其余的数据可以通过活跃变量恢复得到，以此减少备份所需的存储空间；检查点备份包括如下步骤：
步骤1：在kernel执行前设置检查点标记；
步骤2：采用编译的定值-引用分析技术分析kernel中的变量，将其中的相对活跃变量备份保存到CPU的内存当中，并记录相应的内存地址；
步骤3：执行kernel调用；
步骤4：当GPU在进行kernel计算时，将CPU上的数据进行备份保存到内存当中，并记录相应的内存地址。
3.根据权利要求1所述的面向GPGPU的细粒度低开销的容错系统，其特征在于所述冗余执行与错误检测模块，为每个GPU相关语句静态指定一个唯一的ID，并且基于双模冗余检测技术，为GPU上kernel的运行添加一个完全相同的冗余代理，在执行的过程中，记录各条GPU语句的执行顺序，最后比较原kernel程序和冗余代理的计算结果；冗余执行与错误检测包括如下步骤：
步骤1：CPU端调用kernel程序时，为每个GPU相关语句静态指定一个唯一的ID；
步骤2：创建一个执行相同计算任务的冗余代理；
步骤3：为kernel程序中的每个输入变量、输出变量添加一个副本，将这些变量副本使用CudaMalloc分配在GPU端，然后把冗余代理的原输入和输出指针分别指向输入副本和输出副本；
步骤4：执行原kernel和冗余代理的计算任务，记录下GPU语句的执行顺序；
步骤5：比较原kernel程序和冗余代理的各个计算结果，如果两个计算结果不相同，则把详细的错误信息存放在错误位数组中；
步骤6：遍历步骤5中的错误位数组，如果错误位数组不空，则把该错误位数组和GPU语句的执行顺序发送到错误修复模块，否则，表明计算正确，返回计算结果，结束该阶段执行。
4.根据权利要求1所述的面向GPGPU的细粒度低开销的容错系统，其特征在于所述错误修复模块，通过逆向遍历GPU的执行轨迹和计算详细错误信息得到GPU执行轨迹中各个GPU相关语句所需要复算的对象，即那些直接或间接影响了最终错误计算结果的对象，然后采用检查点备份模块所备份的数据恢复复算对象的初始值，最后在原kernel程序中添加相应的掩码控制语句并调用该部分复算kernel复算这些对象即可完成错误修复。
CN201910320906.4A 2019-04-21 2019-04-21 一种面向gpgpu的细粒度低开销的容错系统 Active CN110083488B (zh)

Priority Applications (1)

Application Number Priority Date Filing Date Title
CN201910320906.4A CN110083488B (zh) 2019-04-21 2019-04-21 一种面向gpgpu的细粒度低开销的容错系统

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
CN201910320906.4A CN110083488B (zh) 2019-04-21 2019-04-21 一种面向gpgpu的细粒度低开销的容错系统

Publications (2)

Publication Number Publication Date
CN110083488A CN110083488A (zh) 2019-08-02
CN110083488B true CN110083488B (zh) 2023-03-17

Family

ID=67415785

Family Applications (1)

Application Number Title Priority Date Filing Date
CN201910320906.4A Active CN110083488B (zh) 2019-04-21 2019-04-21 一种面向gpgpu的细粒度低开销的容错系统

Country Status (1)

Country Link
CN (1) CN110083488B (zh)

Families Citing this family (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN111124691B (zh) * 2020-01-02 2022-11-25 上海交通大学 多进程共享的gpu调度方法、系统及电子设备
CN111309514B (zh) * 2020-02-21 2021-06-01 吉林大学 一种gpgpu寄存器的纠错码生成方法
CN112131034B (zh) * 2020-09-22 2023-07-25 东南大学 一种基于检测器位置的检查点软错误恢复方法

Citations (5)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US5948112A (en) * 1996-03-19 1999-09-07 Kabushiki Kaisha Toshiba Method and apparatus for recovering from software faults
CN102929738A (zh) * 2012-11-06 2013-02-13 无锡江南计算技术研究所 大规模异构并行计算的容错方法
CN105022673A (zh) * 2015-07-15 2015-11-04 南京师范大学 一种面向数据并行计算容错的快速并行复算方法
CN105677486A (zh) * 2016-01-08 2016-06-15 上海交通大学 数据并行处理方法及系统
CN107506261A (zh) * 2017-08-01 2017-12-22 北京丁牛科技有限公司 适应cpu、gpu异构集群的级联容错处理方法

Family Cites Families (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US9952932B2 (en) * 2015-11-02 2018-04-24 Chicago Mercantile Exchange Inc. Clustered fault tolerance systems and methods using load-based failover

Patent Citations (5)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US5948112A (en) * 1996-03-19 1999-09-07 Kabushiki Kaisha Toshiba Method and apparatus for recovering from software faults
CN102929738A (zh) * 2012-11-06 2013-02-13 无锡江南计算技术研究所 大规模异构并行计算的容错方法
CN105022673A (zh) * 2015-07-15 2015-11-04 南京师范大学 一种面向数据并行计算容错的快速并行复算方法
CN105677486A (zh) * 2016-01-08 2016-06-15 上海交通大学 数据并行处理方法及系统
CN107506261A (zh) * 2017-08-01 2017-12-22 北京丁牛科技有限公司 适应cpu、gpu异构集群的级联容错处理方法

Non-Patent Citations (3)

* Cited by examiner, † Cited by third party
Title
Sarah Azimi ; Boyang Du ; Luca Sterpone."Evaluation of Transient Errors in GPGPUs for Safety Critical Applications: An Effective Simulation-based Fault Injection Environment".《Journal of Systems Architecture》.2017, *
基于Charm++运行时环境的异构计算应用容错研究;孟晨等;《计算机工程与应用》;20160414(第13期);全文 *
异构计算环境下的三维Kirchhoff叠前深度偏移混合域并行算法;王一达等;《石油地球物理勘探》;20180601(第03期);全文 *

Also Published As

Publication number Publication date
CN110083488A (zh) 2019-08-02

Similar Documents

Publication Publication Date Title
US10884859B2 (en) Resiliency to memory failures in computer systems
US10838808B2 (en) Error-correcting code memory
CN110083488B (zh) 一种面向gpgpu的细粒度低开销的容错系统
US10776267B2 (en) Mirrored byte addressable storage
US9836354B1 (en) Automated error detection and recovery for GPU computations in a service environment
US7779294B2 (en) Power-safe disk storage apparatus, systems, and methods
KR101863406B1 (ko) 검증된 데이터 세트의 비휘발성 매체 저널링 기법
CN107278297B (zh) 用于软件测试的计算设备、方法以及介质
US9502139B1 (en) Fine grained online remapping to handle memory errors
US20140089760A1 (en) Storage of codeword portions
Pourghassemi et al. cudacr: An in-kernel application-level checkpoint/restart scheme for cuda-enabled gpus
US7302619B1 (en) Error correction in a cache memory
US20170123915A1 (en) Methods and systems for repurposing system-level over provisioned space into a temporary hot spare
US20210141697A1 (en) Mission-Critical AI Processor with Multi-Layer Fault Tolerance Support
EP4014119A1 (en) Data race analysis based on altering function internal loads during time-travel debugging
Fu et al. A stack-based single disk failure recovery scheme for erasure coded storage systems
US11409608B2 (en) Providing host-based error detection capabilities in a remote execution device
WO2021184901A1 (zh) 一种数据的写入方法、装置以及设备
CN110750385A (zh) 一种基于受限恢复的图迭代器及方法
Yang et al. The fault tolerant parallel algorithm: the parallel recomputing based failure recovery
Balaji et al. Fault tolerance techniques for scalable computing
US11907124B2 (en) Using a shadow copy of a cache in a cache hierarchy
US20240111623A1 (en) Extended protection storage system put operation
Jia et al. Hessenberg reduction with transient error resilience on gpu-based hybrid architectures
CN107193692B (zh) 基于检查点的计算机的容错方法

Legal Events

Date Code Title Description
PB01 Publication
PB01 Publication
SE01 Entry into force of request for substantive examination
SE01 Entry into force of request for substantive examination
GR01 Patent grant
GR01 Patent grant