结构材料辐照损伤的多GPU分子动力学模拟方法

摘要

公开了一种结构材料辐照损伤的多GPU分子动力学模拟方法，包括：进行初始化；动态地为每个节点划分网格；进行节点间通信；在GPU上建立排序的元胞列表；更新时间步长；根据粒子的坐标，找到粒子与网格编号的对应关系；预测粒子的位移、速度和加速度；计算每个粒子的受力；利用受力对粒子的位移、速度和加速度进行校正；根据系综修正速度，保证系统恒温；利用周期性边界条件，修正粒子位置；存储当前计算结果；迭代执行上述步骤直到设定步数。利用本方法，能高效快捷的模拟更大时空尺度上的材料辐照损伤过程，从微观尺度解释辐照损伤的长时间演化规律。

说明书

技术领域

本公开涉及分子动力学(MD)模拟仿真技术领域，具体涉及一种 使用多个图形处理单元(GPU)加速卡来模拟结构材料的辐照损伤过程 的分子动力学方法。

背景技术

结构材料的稳定性是保证反应堆安全运行的基础，反应堆内的极端 环境(高能粒子辐照)会导致其核心部件微结构严重损伤，进而引发其性 能降级。结构材料的辐照损伤是一个复杂的多参数过程，涉及多种缺陷 的形成、扩散、复合和聚集，以及嬗变氢、氦杂质与离位损伤尤其碰撞 级联损伤的关联，也是一个从微观、介观到宏观的多尺度材料问题。而 初级损伤事件发生在皮秒(ps)时间尺度内，开展实验研究费用高昂和 无法对缺陷进行在线观察，因此，分子动力学模拟方法成为唯一可行的 方案。为了更好的与实验结果对比，分子动力学方法需要模拟千万乃至 数亿量级的原子体系，这将对传统串行分子动力学模拟方法带来很大挑 战。因而，发展高性能并行计算方法势在必行。

在高性能计算并行初期，均是基于MPI和OPENMP的并行技术， 属于粗粒度的并行模式。近年来随着GPU技术的日益成熟，人们尝试将 GPU并行技术引入到分子动力学模拟方法中，其效果显著，这种属于细 粒度的并行模式。相对于MPI多节点并行，GPU具有较高的每瓦特性能、 每平方英尺性能和性能/价格比特性。然而，当计算模拟体系达到近上千 万时其计算效率会因受到GPU上流处理器数目的限制而下降。部分应用 目前采用单个节点上多个GPU并行的方案来扩大计算规模，但是这种方 法计算量还是相对有限，其计算模拟结果仍无法与实验结果比较。

发明内容

本公开实施例将MPI和GPU结合起来，公开了一种使用多个GPU 加速的结构材料辐照损伤的MD模拟方法，可有效地模拟更大时空尺度 上的材料辐照损伤过程，从微观尺度解释辐照损伤的演化规律，预测材 料辐照性能。在一个实施例中，提出了一种结构材料辐照损伤的多GPU 分子动力学模拟方法，其中多个GPU位于多个节点上并且能够并行运 行，所述方法包括：

a.在所述多个节点上进行初始化，包括空间分配、参数读入、材 料结构的生成、以及材料粒子初始位置和速度的设定；

b.动态地为每个节点划分网格，每个网络容纳一个或多个粒子；

c.进行节点间通信，将超出每个节点的网格区域的粒子及重叠区 粒子发送到对应的其他节点；

d.在所述多个节点的GPU上建立排序的元胞列表；

e.在所述多个节点的GPU上更新时间步长；

f.在所述多个节点的GPU上根据粒子的坐标，找到粒子与网格编 号的对应关系；

g.在所述多个节点的GPU上预测粒子的位移、速度和加速度；

h.在所述多个节点的GPU上根据步骤d得到的排序的元胞列表， 计算每个粒子的受力；

i.在所述多个节点的GPU上利用步骤h得到的受力对粒子的位移、 速度和加速度进行校正；

j.在所述多个节点的GPU上根据系综修正速度，保证系统恒温；

k.在所述多个节点的GPU上利用周期性边界条件，修正粒子位 置；

l.在所述多个节点上存储当前计算结果；

m.迭代执行步骤b到l，直到设定步数。

根据本公开实施例的方法，利用GPU超强的浮点计算能力、高带 宽及轻量的计算核心的特点，采用例如NIVIDIA公司的CUDA并行架 构及MPI通信机制来实现该算法，此外在算法结构设计中，结合GPU 的硬件结构，设计出符合GPU并行模型的算法结构，采用了连续访问内 存、共享内存、寄存器、指令优化等一系列的优化策略，提高运行效率。 实施例-金属铁辐照损伤实验仿真结果表明该方法在保证计算精度的前 提下模拟更大时空尺度上的材料辐照损伤过程，进而从微观尺度解释辐 照损伤的长时间演化规律，同时降低了计算设备能耗和维修成本。

附图说明

现在将参考附图仅通过示例的方式描述本公开的实施例，附图中：

图1是根据本公开实施例的实现结构材料辐照损伤的多GPU分子 动力学模拟方法的系统结构示意图；以及

图2是根据本公开实施例的结构材料辐照损伤的多GPU分子动力 学模拟方法的示意流程图。

应注意，在整个附图中，相似的附图标记用于描述相同或相似的要 素、特征和结构。

具体实施方式

本公开可以具有多个实施例，并且可以在其中进行多种修改和变 型。因此，将参考附图中所示的特定实施例来详细描述本公开。然而， 应当理解的是，本公开并不限于具体实施例，而是在本公开的精神和范 围内包括所有修改、等同物和/或替换物。在对附图的描述中，相似的附 图标记用于表示相似的元件。

在本公开各实施例中使用的词语“包括”、“可以包括”和其他词 形表示存在相应公开的功能、操作和组成元件，而不限制一个或更多个 附加的功能、操作或组成元件。此外，如本公开的各种实施例中使用的 术语“包括”、“具有”和其词形变化旨在仅表示特定特性、数字、步 骤、操作、元件、组件或者其组合，而不应解释为首先排除一个或更多 个其他特性、数字、步骤、操作、元件、组件或者其组合的存在或可能 的添加。

在本公开的各个实施例中，表述“或”或“A或/和B中的至少一 个”包括所列词语的任意组合或全部组合。例如，表述“A或B”或“至 少A或/和B”可以包括A，可以包括B，或可以包括A和B二者。

在本公开中，包括诸如“第一”或“第二”等序数的表述可修饰各 个元件。然而，这种元件不被上述表述所限制。例如，上述表述并不限 制元件的顺序和/或重要性。上述表述仅用于将一个元件与其他元件进行 区分的目的。例如，在不脱离本公开的各种实施例的范围的情况下，第 一元件可以被称作第二元件，并且类似地，第二元件也可以被称作第一 元件。

在本公开的各个实施例中使用的术语仅用于描述一些实施例，而不 意图限制本公开。除非上下文另行明确指示，否则单数形式也意在包括 复数形式。此外，本文所用的所有术语(包括技术术语和科学术语)与 本公开所属技术领域的普通技术人员通常理解的含义相同。通用字典中 定义的此类术语应当被解释为具有与相关技术领域中的上下文含义相同 的含义，并且除非本公开的各个实施例中明确定义，否则这些术语不应 当被解释为具有理想或过度形式的意义。

以下结合附图对本公开的优选实例进行说明，应当理解，此处所描 述的优选实例仅用于说明和解释本公开，并不用于限定本公开。为了清 楚和简明起见，本公开省略了公知功能和结构的描述。

图1是根据本公开实施例的使用多个GPU并行实现的结构材料辐 照损伤的MD模拟方法的系统结构示意图。如图1所示，该系统包括客 户端1、包含GPU加速卡的高性能计算集群2以及用来传输信息的网络 3。虽然图中未示出，但是可以理解，在图1中，客户端1和计算集群2 可以采用任何适当的连接方式。例如，客户端1可以经由局域网或者广 域网等与计算集群20通信，远程进行仿真实验。

客户端1可以包含用于输入的硬件设备，例如键盘、触摸盘等，以 及可以远程管理计算集群的软件设备，例如PuTTY、XManager等。

计算集群2可以包括多个计算节点4-1，4-2，…，4-n(n是大于等于2 的整数)，每个计算节点可以包含CPU和GPU，以及可选地，还可以包 括存储设备等。作为示例，计算节点中的GPU可以是具有GF110核心 以上的NIVIDIA通用计算卡，包含NIVIDIA公司开发的CUDA架构的 GPU并行编程、以及用于节点间通信的接口协议，例如MPI接口协议。 利用MPI加CUDA的技术，可以模拟大规模的辐照损伤程序。多个计 算节点可以并行执行计算，每个节点上的每个GPU包括进行并行处理的 多个流处理器，每个流处理器对应处理一个粒子。MPI会逻辑上赋予每 个节点一个编号，通常从0开始，例如若有四个节点，那么节点编号是 0，1，2，3。在运行过程中默认节点0为主节点。

以上系统仅仅是一种示例的实现方式。本领域技术人员可以理解， 可以采用其他形式的系统架构，例如上述各个部件的功能可以进行再分 配和组合，以形成其他的系统架构。

下面结合图1的系统架构，参照图2来描述根据本公开实施例的结 构材料辐照损伤的多GPU分子动力学模拟方法。

如图2所示，在步骤202中在多个节点上进行初始化，包括例如空 间分配、参数读入、材料结构的生成、以及材料粒子初始位置和速度的 设定。例如，根据示例实施例，空间分配可以包括：

(1)读取节点个数nnodes和总粒子数nm，确定每个节点中的CPU 和GPU变量空间的大小，计算公式为空间长度nmPerNode＝2* ((nm+nnodes-1)/nnodes)；

(2)在每个节点上动态为CPU和GPU的变量数组开辟空间，包 括粒子位置x0、y0、z0，速度x1、y1、z1，加速度x2、y2、z2，加速 度的导数x3、y3、z3，粒子类型ispec，粒子的受力fx、fy、fz等，空间 长度为nmPerNode；

(3)对于每个节点，从存储于硬盘的文件中读入相关物理参数， 包括材料信息如粒子类型、粒子量、晶格类型、晶格常数、盒子大小、 系综、反应步数、构建晶格结构所需参数及开关量；从势参数文件中读 入对应势参数；从初始碰撞原子PKA文件中读入PKA个数、位置、方 向、能量，开关量；

(4)可选地，进行单位转换，使得整体单位一致；

(5)在主节点0的CPU上动态分配用于存储所有粒子x、y、z方 向的位置和速度的变量空间，分配长度为nm。

(6)计算结合能，用于测试；

(7)在主节点0上初始化晶体结构。若是弛豫，按照晶格结构生 成粒子初始位置，若是级联，直接将弛豫后粒子的坐标读入。晶体结构 是按照右手定则，x、y、z方向顺序设定，因为任务划分是按照线性网 格划分，这样可以确保尽可能多的粒子在对应的节点上而减少通信。在 程序运行时需要在弛豫与级联之间选择，是并行的。在物理上弛豫和级 联是两个不同的物理过程，先读入弛豫的输入文件让程序跑指定步数， 得到平衡系统下的粒子信息包括粒子位置、速度、加速度等，然后读入 级联的输入文件以及弛豫时得到的粒子信息进行辐照损伤模拟。

(8)若是弛豫，在主节点0上随机产生粒子(-0.5～0.5)的速度， 用方法修正保证这个系统动量为0，并且温度不变。若是级联，直接读 入弛豫得到的速度、加速度及粒子类型等。

在步骤204，分配任务，即动态地为每个节点划分网格，每个网络 容纳一个或多个粒子。根据示例实施例，分配任务可以包括：

(1)将粒子的位置和速度平均分配到每个节点的CPU上，并将计算 算得的张量广播到各节点；

(2)将每个节点上CPU端位移、速度、加速度拷贝到GPU，本公 开实施例的所有数组变量都采用一维向量形式；

(3)计算每个节点i对应子区域网格的起始位置cellPerNodeStart[i] 和结束位置cellPerNodeEnd[i]，初始值为平均为每个节点分配的网格 数。在一个实施例中，任务划分采用线性划分方法，初始时若总共的网 格数为Ncell，节点数为NG，那么每个节点是可计算的网格近似为 Ncell/NG，节点0则计算0～Ncell/NG号网格内的粒子，节点1则计算 Ncell/NG+1～2×Ncell/NG号网格内的粒子，以此类推；循环中考虑到级 联碰撞时粒子会在某个节点聚集，采用动态平衡的方法(将在下文中描 述)，在每个循环步内调节计算在每个节点上的网格数，以达到粒子数 近似平衡。在节点间(粗粒度的)任务划分时选用线性的任务划分方法， 相对于现有方法，例如Lammps和HOOMD等采用的“最小面积法”，根 据本公开实施例的方法，节点数可任意多、只需要和相邻节点通信、同 时便于为每个节点调节计算任务。具体地，Lammps和HOOMD都采用 “最小面积法”划分法(Lammps和HOOMD是国际通用分子动力学软件， 也采用MPI+cuda的并行模式，Lammps对部分计算量大的代码进行了并 行，HOOMD则是完全基于GPU的分子动力学通用软件)，“最小面 积法”划分法通信量最少，但是需要判断周围26个节点与本节点的重叠 区域的通信，并且对于节点数要求2ⁿ，对于负载平衡调节也比较麻烦。 而根据本公开实施例的方法相对于最小面积法可以只需要和上下两个相 邻的节点通信，动态负载平衡算法的实现也很简单并且对节点数要求也 比较低。

根据示例实施例，多个节点间的通信采用MPI通信协议，多个节点 分为主节点和从节点，主节点负责数据的分发和收集，按照负载平衡原 则动态地为每个节点划分网格。

在步骤206，通过MPI通信协议将移出本节点对应区域的粒子及重 叠区粒子发送到对应的其他节点，在步骤208，每个节点在GPU上生成 各自排序的元胞列表，步骤206和208是交互的。在一个示例实施例 中，这两个步骤可以包括：

(1)在每个GPU上利用流处理器对每个粒子进行并行处理，根据 粒子坐标，计算每个粒子对应的网格编号；

(2)在每个GPU上利用流处理器对每个粒子进行并行处理，找出 不在本节点的粒子并存储到缓冲区，同时记录相应的粒子数，将不在本 节点且网格编号小于本节点网格的起始位置的粒子存放到缓冲区，将网 格编号大于本节点网格的结束位置的粒子存放到另一缓冲区，并将该粒 子对应的粒子编号和网格编号赋最大值；

每个流处理器对应一个线程，为了确保各个线程都执行完成，可 以用同步机制对各线程进行同步；

(3)计算出与本节点相邻的两个节点的编号，将不在本节点的粒 子发送到对应节点的缓冲区，从缓冲区读取发送过来的信息，在传输之 前，需要预通信，让接收的节点知道需要接收的粒子数；

当节点数大于2时，为了避免缓冲区读写死锁，让主节点0先接收 缓冲区信息，而其他节点发送信息，当节点数等于2时，两个节点可以 直接发送信息；

(4)将接收到的缓冲区信息拷贝到GPU，在GPU上对缓冲区中的 粒子并行存储到对应变量的最后面；

(5)在每个节点上对粒子按照其对应的网格编号排序，例如可以 利用CUDA的thrust库函数sort_by_key；

(6)为了保证GPU线程访问地址的连续性，提高执行效率，在每 个节点上的GPU上按照排序的网格编号对粒子对应的位置、速度、加 速度排序，这时因之前不在本节点的网格编号都被赋了最大值，对应的 粒子信息也排在了最后面；

(7)由于在后续计算中需要遍历网格周围的26个网格，则相邻节 点间网格需要通信形成重叠(overlap)区，由于是线性划分任务以及周 期边界方法，为了减少通信量，一次性将包裹本节点的网格的全部信息 传输到本节点，向前传的网格数为cellPerNodeStart[i]+startoffset， startoffset＝nlcx*nlcy+(nlcx*nlcy-cellPerNodeStart[i]％(nlcx*nlcy))％(nlcx* nlcy)，nlcx、nlcy分别是x、y方向上的网格数，向后传的网格数为 cellPerNodeEnd[i]-endoffset，endoffset＝nlcx*nlcy+cellPerNodeEnd[i]％(nlcx *nlcy)，信息传发送和接收及存储与(3)和(4)相同，存储时将信息 拷贝到有效粒子信息后面，覆盖以前非本网格的信息；

(8)在每个节点的GPU上利用共享内存，按照排序的网格编号记录 每个网格中粒子的起始和结束位置分别存储在数组cellStart和cellEnd 中。

在步骤210，动态地更新时间步长，其为最大时间步长、最大能量 截断、最远距离截断中所需时间最小的那个。通过先在每个节点上求得 极值，然后求得节点间的极值，来获得最大时间步长、最大能量截断以 及最远距离截断中所需时间最小的那个。在一个示例实施例中，时间步 长更新如下：

(1)在每个节点的GPU上借助CUDA的库函数thrust：：min_element/ max_element完成对速度、加速度及其导数的极值；

(2)通过MPI通信协议中的MPI_Send()和MPI_Recv()函数将各节点的 极值收集到主节点0，然后求最大值和最小值；

(3)在主节点的在CPU上求出最小时间步长dtmax、最大能量截断 demax和最远距离截断dxmax中所需时间最小的那个；

(4)在GPU上根据更新(或修正)前后时间步长的比值修正粒子的 速度、加速度及其导数，若是NPT系综，同时修正张量。

在步骤212，在多个节点的GPU上根据粒子的坐标，找到粒子与网 格编号的对应关系。这实际上是动态地调节负载平衡，首先在主节点0 统计出每个网格的粒子数，以每个节点的平均粒子数作为标准，调节在 每个节点上网格的起始和结束位置，使得在每个节点上计算的粒子数近 似相等。

在步骤214，在每个GPU上利用流处理器对每个粒子进行并行处 理，例如用verlet方法预测粒子的位移、速度、加速度；

若是NPT系综，还需要在GPU上预测张量。

在步骤216，在所述多个节点的GPU上根据步骤208得到的排序的 元胞列表，计算每个粒子的受力。具体地，计算每个粒子的多体势，得 到每个粒子的受力。在一个示例实施例中，粒子的受力包括x、y、z方 向的受力，即粒子与周围所有粒子间对势、粒子与电子间嵌入势、修正 项相加得到粒子的总受力在x、y、z方向的投影，通过以下步骤计算粒 子的总受力：

(1)基于排序的数组，让GPU上的流处理器连续遍历本粒子所在的 网格及其相邻的26个网格中的所有粒子，分别计算出它们的距离进而算 出该粒子对应的电子密度及该粒子对应的电子密度非球对称部分的贡 献；

(2)通过计算得到的电子密度计算粒子的嵌入势及其导数，近而计算 粒子的修正项及其导数；

(3)基于排序的数组，让GPU上的流处理器连续遍历本粒子所在网 格及其相邻的26个中的所有粒子，分别计算出它们的距离，根据对势函 数计算两粒子间的对势及对势的导数；以及

(4)将粒子与周围所有粒子的对势、嵌入势和修正项相加得到粒子的 总受力，进而投影出x、y、z方向的受力。

在步骤218，在每个GPU上利用流处理器对每个粒子进行并行处 理，根据步骤216得到的受力，对粒子的位移、速度、加速度进行校 正。

若是NPT系综，对张量进行校正。

在步骤220，在多个节点的GPU上根据系综修正速度，保证系统恒 温。例如，在非NVE系综下，在GPU上修正速度，保证系统恒温。

在步骤222，在多个节点的GPU上利用周期性边界条件，修正粒子 位置。

在步骤224，在多个节点上存储当前计算结果。例如，根据存储需 求，将GPU端信息拷贝到CPU。

在步骤226，判断循环是否到达设定步数，若没有，则返回步骤 204，进行循环。

若判断循环到达设定步数，结束循环，释放开辟的存储空间以及 系统内存空间。

根据示例实施例，辐照损伤过程包括不同晶格结构和不同系综下材 料的辐照损伤过程，上述方法还包括在NPT系综下对于张量的调节处理。

可选地，为了进一步提高计算准确性，可以在步骤208与步骤210之 间加入以下一个或多个步骤：

初始加速度计算步骤：在每个节点的GPU上利用粒子位置计算每 个粒子的多体势，得到每个粒子的总受力，进而得到粒子的初始加速 度，包括：

(1)基于排序的数组，让每个节点GPU上的流处理器连续地遍历 本粒子所在的网格及其相邻的26个网格中的所有粒子(若是非本节点的 网格则到数组后面对应网格中读取)，分别计算出它们的距离进而算出 该粒子对应的电子密度及该粒子对应的电子密度非球对称部分的贡献；

(2)在每个节点GPU上对每个粒子进行并行处理，通过计算得到 的电子密度计算粒子的嵌入势及其导数，近而计算粒子的修正项及其导 数；

(3)基于排序的数组，让每个节点GPU上的流处理器连续地遍历 本粒子所在的网格及其相邻的26个网格中的所有粒子，分别计算出它 们的距离，根据对势函数计算两粒子间的对势，及对势的导数；将粒子 与周围所有粒子的对势、嵌入势和修正项相加得到粒子的总受力，进而 投影出x、y、z方向的受力；

因为在分子动力学中，粒子间作用势能函数的选取非常重要，针 对不同种类的材料，由于其结构不同，所选择的势能函数也不同。

当系综为NPT时，还需要对相应张量初始化。

(4)在GPU上对每个粒子进行并行处理，根据每个粒子的受力及 速度得到粒子的初始加速度。

PKA设置步骤：判断是否进行PKA设置，若需设置PKA，设置 PKA，PKA位置和速度的设置在主节点0处完成，可多次设置PKA， 并可以随机或者指定PKA的位置和方向，在多个GPU计算时，PKA应 该在哪个节点的哪个网格，首先根据PKA的位置，计算出对应的网 格，并找到对应的节点，将PKA信息发送到对应节点，遍历包括PKA 在内的27个网格，找到距离最小的那个，判断对应位置是不是在本节 点，如果不是，则发送到对应的其他节点，找到位置替换。PKA是在 循环开始时根据需求设置，人为设置的开关量，从输入文件读入的。

可选地，在步骤218与220之间，可以包括如下步骤：

判断是否需要执行淬火，如果需要，则执行。这可以根据输入文 件中的设置，用开关量iquen控制。

可选地，为了验证方法的正确性，在每一循环步骤中所有粒子的 电子能量和粒子间能量相加，并与结合能对比。例如，将在初始化阶段 利用公式算出结合能cohe与每次循环结束后算出的的能量总和pe比 较。当pe与cohe值相等时，说明正确。

可以用MPI中的函数MPI_Reduce()将各节点的pe值规约。

下表列出了实施例一金属铁在NVIDIA的多GPU上辐照损伤实验 仿真结果，该表是在不同尺寸下，运行步数为5000时CPU与GPU运行 时间对比，结果表明本公开在降低能耗的同时极大地提高了运算效率。

根据本公开的各种实施例的方法可以由存储在计算机可读记录介 质中的计算机可读代码来实现。计算机可读记录介质包括其中存储了可 由计算机系统读取的数据的所有类型的记录设备。对于这样的记录介质， 例如，可以使用ROM、RAM、光盘、磁带、软盘、硬盘或非易失性存 储器。另外，计算机可读记录介质可以存储有如下代码，所述代码分布 于通过网络连接的计算机系统中，从而可以通过分布式方法由计算机读 取和执行所述代码。

尽管已经参照本公开的各种实施例示出并描述了本公开，但是本领 域技术人员将理解，在不脱离由所附权利要求及其等同物限定的本公开 的范围的情况下，可以在其中进行形式和细节上的各种改变。

将理解，本公开的实施例可以被实现为硬件、软件或硬件和软件结 合的形式。任何这种软件可以被存储为易失性存储器或非易失性存储器 的形式(例如ROM之类的存储设备，不论是否是可擦除的或可重写的)， 或者被存储为存储器的形式(例如RAM、存储器芯片、装置或集成电路)， 或者被存储在光或磁可读介质上(例如，CD、DVD、磁盘或磁带等)。 将认识到，存储设备和存储介质是适于存储程序的机器可读存储装置的 实施例，所述程序包括在被执行时实施本公开实施例的指令。

贯穿本说明书的描述和权利要求，词语“包括”和“包含”及该词的变 型表示“包括但不限于”，不意在排除其他组件、整体或步骤。

贯穿本说明书的描述和权利要求，单数形式包含复数形式，除非上 下文另外要求。具体地，在使用不定冠词的情况下，本说明书将被理解 为考虑复数以及单数，除非上下文另外要求。

结合本公开的具体方面、实施例或示例所描述的特征、整体或特性 将被理解为适用于本文中描述的任何其他方面、实施例或示例，除非不 相容。

还将认识到，贯穿本说明书的描述和权利要求，“用于Y的X”的通 用形式的语言(其中，Y是某个动作、活动或步骤，X是某个用于执行 该动作、活动或步骤的装置)包含用于进行Y的专门(但不排他)布置 或改编的装置X。

此外，本说明书中公开的实施例对技术内容的描述和理解进行启 示，但并不限制本公开的范围。因此，本公开的范围应被解释为包括基 于本公开的技术思想而做出的所有修改或各种其他实施例。