CN102708009B - 一种基于cuda实现多任务共享gpu的方法 - Google Patents
一种基于cuda实现多任务共享gpu的方法 Download PDFInfo
- Publication number
- CN102708009B CN102708009B CN201210115719.0A CN201210115719A CN102708009B CN 102708009 B CN102708009 B CN 102708009B CN 201210115719 A CN201210115719 A CN 201210115719A CN 102708009 B CN102708009 B CN 102708009B
- Authority
- CN
- China
- Prior art keywords
- task
- block
- dimension
- piecemeal
- gpu
- 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.)
- Expired - Fee Related
Links
Images
Landscapes
- Multi Processors (AREA)
Abstract
本发明公开了一种基于CUDA实现多任务共享GPU的方法。包括:在Global Memory中建立映射表,确定合并后的Kernel中,每个Block执行的任务编号和任务分块编号;一次用一个Kernel启动N个Block,N等于所有任务的任务分块数目之和;用标记和阻塞等待的方法,满足原有任务之间的约束关系;对于Shared Memory,采用预申请和静态分配的方式进行多任务共享。通过本发明,可以简便地实现在现有GPU硬件架构上实现多任务共享,可以简化实际应用中的编程工作,并在一定情况下取得良好的性能。
Description
技术领域
本发明涉及一种多任务共享GPU的实现方法,具体涉及在NVIDA的CUDA架构中合并多个任务,实现任务并行的方法,属于GPGPU计算领域。
背景技术
GPGPU(General-purpose computing on graphics processing units),是利用GPU来进行大规模计算的技术。CUDA是NVIDA公司提供的GPGPU架构。CUDA自从推出开始,就成为广泛应用的众核并行计算形式。
GPU具有远远高于CPU的浮点运算能力和内存带宽(附图1),同时由于其高度的并行性,非常适合于大规模数据处理。
然而,由于GPU的硬件设计,GPU上的编程和CPU上的并行编程有所不同。一个显著的区别就是,GPU不支持多任务共享:每个任务在GPU上的运行都是对GPU的硬件资源独占的,不允许有其他的Kernel也在执行。例如,当多个任务都要使用GPU的时候,只能够一个一个的顺序地执行,而不能一起同时在GPU上运行。这一点和CPU上允许进程之间进行切换是有很大的不同。
目前,尚未发现有专利或者文献针对GPU上的多任务共享进行讨论。
发明内容
本发明所使用的一些术语定义如下:
Kernel:CUDA架构中,GPU一次运行所执行的代码。
Thread,Block,Grid:CUDA架构中,一个大的Kernel被划分为了很多小的基本单位,称为线程(Thread)。
所有Thread组织成为了两级结构。(附图2)
首先,这些Thread划分为了若干个线程块(Block),每个Block包含相同数目的Thread。Thread是以Block为单位分发到硬件资源上来执行的。在Block内部,线程的ID编号可以采用一维、二维或者三维,这称为Block维数。在某一个确定的Kernel里面,每个Block都含有相同的维数,以及每一维上相同的大小(这称为维度)。维数和维度都是在GPU启动Kernel之前由程序员指定好的,执行过程中是不能改变的。
然后,所有的Block组成的整体(也就是所有的Thread)称为线程网格(Grid)。和Block的维数定义类似,Block的ID编号,可以组织成一维、二维或者三维,称为Grid的维数,每一维上的大小称为维度。
SM(Streaming Multiprocessor):GPU上的硬件单元,包括运算单元、寄存器、片上存储器。每个Block都会被分配到SM上进行执行。一个Block只能在一个SM上执行,一个SM上可以同时执行一个或者多个Block(取决于单个Block消耗的寄存器、片上存储器资源的数目)。
Global Memory:显卡上的存储器,属于片外存储器,GPU可以从中读取,但是速度相对比较慢。
Shared Memory:GPU上的一种片上存储器,每个SM中拥有一定数量的Shared Memory,通常为KB级大小,可以视作一种可供程序员维护的Cache,在CUDA架构中有非常重要的意义。
任务分块:单个任务划分为多个子任务,每个子任务称为任务分块。一个任务分块将会对应到一个Block中完成计算。
本发明旨在提供一种方法使得在CUDA架构上实现多任务共享GPU,解决现有GT200架构不支持多任务共享GPU的问题。
本发明的原理包括三点:
1.对每个Block执行的任务分块,在GPU运行之初由程序员通过对一个映射表的赋值来确定。通常,每个Block执行的任务分块,在算法设计的时候就已经固定。本发明通过增加一个映射表,使得能够在GPU启动之前,程序员可以对每个Block执行的任务分块进行重排。
2.利用GPU高并行度的特点,把多个任务的Block合并到的一个Kernel中去。
CUDA架构中,可以同时启动成百上千的线程(Thread),进行并行度非常高的计算任务。所有的Thread被按照一定数目组织成为若干的Block。本发明让一个GPU在一个Kernel中同时启动大量的Block,执行所有任务的任务分块。
3.考虑到任务之间可能具有约束性,利用GPU动态调度Block到SM上的特性,处理约束关系。
CPU上一种通常的做法是,如果有约束关系,那么在需要同步的地方,直接使用原子操作即可。虽然GPU也支持原子操作,但是原子操作的代价是非常大的,会带来性能上的严重损失,并且极易造成死锁。
本发明充分考虑到了GPU调度Block的规律性。
GPU通常拥有数十到上百个SM。GPU上的Block在启动的时候并不会拥有硬件资源,而是在执行的过程中,动态地把Block分发给SM。但SM的数量一般相对于Block少很多,所以每个时刻,在SM上执行的Block只是一部分,其他的Block只能等待。一旦有SM上Block的任务计算结束,GPU回收得到空闲资源的时候,就会从尚未执行的Block中选取一定的Block分发到有空闲资源的SM上去。
为Block分配资源的时候,是有一定的顺序性的,这种顺序性表现为:
1.总是优先分发ID编号较小的Block到SM上去。例如ID编号为0的Block分发到SM上去的时间总是不晚于ID编号为1的Block被分发的时间。
2.相对的,回收Block资源的时候,也是优先回收ID编号较小的Block的空闲资源。例如ID编号为0的Block尚未结束,ID编号为10的Block执行结束,但此时GPU不会回收编号为10的Block的资源,因为ID编号更小,ID编号为0的Block的资源并未被回收。
根据这种顺序性,本发明先把需要优先执行的任务,在映射表中适当排序:
1.被其他任务依赖的任务,使ID编号较小的Block来执行它的任务分块,这样它会优先获得资源,被调度到SM上去进行先执行;
2.依赖于其他任务的任务,使ID编号较大的Block来执行它的任务分块,同时,辅以适当的阻塞等待操作,保证其依赖的任务已经完全执行结束。
完整的技术方案如下(流程参见图4):
一种基于CUDA实现多任务共享GPU的方法,包括如下步骤:
1)在Global Memory中建立映射表,确定合并后的Kernel中,每个Block执行的任务编号和任务分块编号;
2)一次用一个Kernel启动N个Block,N等于所有任务的任务分块数目之和;
3)用标记和阻塞等待的方法,满足原有任务之间的约束关系;
4)对于Shared Memory,采用预申请和静态分配的方式进行多任务共享。
其中,步骤1)的优选实现方法如下:
1.1)映射表要给出Block到任务和任务分块的映射关系,即确定某个Block执行哪个任务的哪个分块;
1.2)任务在映射表中的排布要求满足约束条件的拓扑顺序:如果任务A依赖于任务B,那么,执行任务A的所有Block的ID编号应该大于所有执行任务B的Block的ID编号;
1.3)在满足步骤1.2)所述的约束条件的情况下,其他无约束关系任务在映射表中以任意的方式进行排布。
步骤2)的优选实现方法如下:
2.1)将原有的任务的grid维数转换为一维:原有的任务可能是有不同的grid维数和维度,这里可以统一选取一维;如果原来的任务是多维的,只需要进行一维到多维的换算即可;
2.2)若Block的维数不一致,则将之统一转换为一维;若转成一维之后Block的维度不一致,则统一选取一个最大的维度,让其他较小的Block添加空线程补足,这样所有的任务都采用了相同的Block维数和维度;
2.3)所有的Block在开始执行计算之前,首先从映射表中获取该Block需要执行的任务编号和任务分块编号;
2.4)根据步骤2.3)中读取的任务编号,选择执行不同的任务代码;将步骤2.3)中读取的任务分块编号,使用到具体的任务计算中去。
步骤3)的优选实现方法如下:
3.1)给每个任务设置标志位mark,每个任务的每个任务分块设置标记数组tag[],用以标记任务和任务分块的执行情况;
3.2)对于任务的每一个任务分块,在该任务分块完成之后,将对应的标记位tag置位,表明该任务分块已经执行结束;
3.3)用每个任务的最后一个Block,在计算返回前阻塞,循环检查同一任务的其他任务分块的对应tag标记位是否已经被置位,一旦全部被置位,则对标志位mark置位,表明该任务已经结束;
3.4)如果任务A需要依赖于任务B,那么在任务A计算开始之前阻塞,循环检测B的任务标志位mark,直到其置位。
步骤4)的优选实现方法如下:
4.1)预先在Kernel开始的时候,申请一个足够大的Shared Memory数组,其大小至少应该等于每个任务所需要的Shared Memory用量的最大值;
4.2)每个任务单独写成一个函数,将Shared Memory数组地址传给这个函数,函数中需要使用Shared Memory的时候,直接在这个数组中静态分配使用。
通过本发明,可以简便地实现在现有GPU硬件架构上实现多任务共享,可以简化实际应用中的编程工作,并在一定情况下取得良好的性能。
附图说明
图1GPU的浮点运算能力存储器带宽同CPU的比较(图片来源:NVIDIA CUDAProgramming Guide Version 2.3)
图2GPU中的Thread,Block,Grid结构(图片来源:NVIDIA CUDA Programming GuideVersion 2.3)
图3两种映射表的排布。图3(a)实施方案中的3个任务使用图形表示(约束关系用箭头表示,任务1依赖于任务0);图3(b)一种合法的映射表排布;图3(c)一种不合法的映射表排布(任务1依赖于任务0,却有Block排布到了任务0前面)。
图4本发明所述方法的流程图。
具体实施方式
以下以一个具体的例子,对本发明做进一步的说明。但是需要注意的是,公布实施例的目的在于帮助进一步理解本发明,但是本领域的技术人员可以理解:在不脱离本发明及所附的权利要求的精神和范围内,各种替换和修改都是可能的。因此,本发明不应局限于实施例所公开的内容,本发明要求保护的范围以权利要求书界定的范围为准。
具体的例子是:3个计算任务(具体任务内容此处并无影响)。
任务存在以下的约束关系:任务1必须在任务0完成后才能进行,因为任务1需要使用任务0的结果,而任务2同任务0和任务1没有任何约束关系。(附图3(a),圆圈代表任务,箭头代表依赖关系)
使用技术方案中所述的方法实现任务并行。
为了方便叙述,定义以下device函数分别完成3个计算任务,分别称为任务0,1和2。
任务0 | _device_void compute0(...); |
任务1 | _device_void compute1(...); |
任务2 | _device_void compute2(...); |
同时假设每个任务都进行了相同大小的任务分块,各自的计算任务划分为了N个任务分块。
实施过程分为以下步骤:
A.建立映射表
开辟两个一维的数组,长度等于所有任务的任务分块数之和,此例中为3*N。两个数组的具体含义如下
1.task_id[],取值为0、1或者2。这个数组给出了Kernel中的Block需要执行的任务。例如,图3(b),task_id[0]到task_id[N-1]的值都为0,代表了Block ID从0到N-1的Block都要执行任务0;其他值为1和2的元素,意义类似。
2.block_id[],取值为0到N-1。这个数组给出了Kernel中Block需要执行的分块编号。例如,图3(b),block_id[N]等于0,代表了编号为N的Block需要执行编号为0的任务分块(任务编号由1中的task_id指定)。
编号的顺序需要满足具体问题所要求的任务之间的约束关系。任务1依赖于任务0,所以执行任务1的分块的那些Block,ID编号应该大于执行任务0的Block ID编号。
例如简单地做法,把任务0和任务2的block排布到前端,任务1的block排布到后端,如图3(b)所示。
而图3(c)的排布是不满足约束关系的拓扑序的,任务1的分块被排布到了任务0前面,在GPU的调度过程中会首先执行,造成错误,所以是是不合法的。
确定好顺序之后,可以在Kernel启动之前直接对task_id[]和block_id[]进行赋值即可。
B.启动Kernel
定义一个合并后的Kernel,Block的数目等于原有所有任务的任务分块数之和,此例中为3*N。这个Kernel的参数列表应该传入三个任务需要的所有参数。
合并的Kernel应该先获取自身的Block ID号bidx
int bidx=blockIdx.x+gridDim.x*blockIdx.y; |
利用bidx从task_id和block_id数组中第bidx个位置读取任务和分块编号信息taskid和bid。
int bid=block_id[bidx],taskid=task_id[bidx]; |
然后根据taskid相应地选择执行不同的device函数,device函数需要做简单修改,把bid作为参数传给device函数使用。
如下面代码所示。
C.处理约束
Kernel启动前为每个任务在Global Memory中各自开辟一个标记数组tag0[]、tag1[]和tag2[],每个元素取值0或者1,用以标记每个任务的Block是否执行结束。例如tag0[]的第10个元素为0,表明任务0的第10个Block尚未执行结束。
此外,为每个任务在Global Memory中各自设置一个标志变量mark0、mark1、mark2,取值0或者1,以标记任务的执行状态。例如,mark2=0代表任务2已经执行结束。
tag0[]、tag1[]和tag2[]的所有元素,以及mark0、mark1和mark2,一开始都被设置为0。
某个任务的编号为bid任务分块执行结束,置对应的标记数组第bid个位置1。例如,任务0的第10个任务分块执行结束,就把tag0[]的第10个元素设置为1。
每个任务的最后一个Block(执行此任务的ID编号最大的Block),在退出前,检查当前任务的其他任务分块的标记,查看是否都已经变成1,如果是,则对应任务的标志位为1,否则循环等待所有标记都变为1。
例如,对于任务0,任务0的最后一个Block,也就是bid=N-1的Block在调用了device函数返回之后,检查任务0其他任务分块,循环阻塞,等待tag0[]全部变为1。
在需要满足约束的地方,对对应的mark进行阻塞检查。例如,任务1依赖于任务0,所以在任务1的device函数之前,需要满足约束。于是增加一个循环等待语句,阻塞等待任务0的标志位mark0变为1。
D.共享Shared Memory
预先在Kernel开始的地方,申请出一块足够大小的Shared Memory数组shared,大小应该是三个device函数使用的Shared Memory的最大值。
_shared_float shared[MAX]; |
修改device函数,把shared传给device函数以作为Shared Memory使用。Device函数中所有用到Shared Memory的地方,都从传入的shared数组上,进行静态分配。
Claims (5)
1.一种基于CUDA实现多任务共享GPU的方法,包括如下步骤:
1)在Global Memory中建立映射表,确定合并后的Kernel中,每个Block执行的任务编号和任务分块编号,任务在映射表中的排布要求满足约束条件的拓扑顺序;
2)一次用一个Kernel启动N个Block,N等于所有任务的任务分块数目之和;
3)用标记和阻塞等待的方法,满足原有任务之间的约束关系;
4)对于Shared Memory,采用预申请和静态分配的方式进行多任务共享。
2.如权利要求1所述的方法,其特征是,步骤1)的实现方法如下:
1.1)映射表要给出Block到任务和任务分块的映射关系,即确定某个Block执行哪个任务的哪个分块;
1.2)任务在映射表中的排布要求满足约束条件的拓扑顺序:如果任务A依赖于任务B,那么,执行任务A的所有Block的ID编号应该大于所有执行任务B的Block的ID编号;
1.3)在满足步骤1.2)所述的约束条件的情况下,其他无约束关系任务在映射表中以任意的方式进行排布。
3.如权利要求1所述的方法,其特征是,步骤2)的实现方法如下:
2.1)将原有的任务的grid维数转换为一维;
2.2)若Block的维数不一致,则将之统一转换为一维;若转成一维之后Block的维度不一致,则统一选取一个最大的维度,让其他较小的Block添加空线程补足,这样所有的任务都采用了相同的Block维数和维度;
2.3)所有的Block在开始执行计算之前,首先从映射表中获取该Block需要执行的任务编号和任务分块编号;
2.4)根据步骤2.3)中读取的任务编号,选择执行不同的任务代码;将步骤2.3)中读取的任务分块编号,使用到具体的任务计算中去。
4.如权利要求1所述的方法,其特征是,步骤3)的实现方法如下:
3.1)给每个任务设置标志位mark,每个任务的每个任务分块设置标记数组tag[],用以标记任务和任务分块的执行情况;
3.2)对于任务的每一个任务分块,在该任务分块完成之后,将对应的标记位tag置位,表明该任务分块已经执行结束;
3.3)用每个任务的最后一个Block,在计算返回前阻塞,循环检查同一任务的其他任务分块的对应tag标记位是否已经被置位,一旦全部被置位,则对标志位mark置位,表明该任务已经结束;
3.4)如果任务A需要依赖于任务B,那么在任务A计算开始之前阻塞,循环检测B的任务标志位mark,直到其置位。
5.如权利要求1所述的方法,其特征是,步骤4)的实现方法如下:
4.1)预先在Kernel开始的时候,申请一个足够大的Shared Memory数组,其大小至少应该等于每个任务所需要的Shared Memory用量的最大值;
4.2)每个任务单独写成一个函数,将Shared Memory数组地址传给这个函数,函数中需要使用Shared Memory的时候,直接在这个数组中静态分配使用。
Priority Applications (1)
Application Number | Priority Date | Filing Date | Title |
---|---|---|---|
CN201210115719.0A CN102708009B (zh) | 2012-04-19 | 2012-04-19 | 一种基于cuda实现多任务共享gpu的方法 |
Applications Claiming Priority (1)
Application Number | Priority Date | Filing Date | Title |
---|---|---|---|
CN201210115719.0A CN102708009B (zh) | 2012-04-19 | 2012-04-19 | 一种基于cuda实现多任务共享gpu的方法 |
Publications (2)
Publication Number | Publication Date |
---|---|
CN102708009A CN102708009A (zh) | 2012-10-03 |
CN102708009B true CN102708009B (zh) | 2014-04-02 |
Family
ID=46900823
Family Applications (1)
Application Number | Title | Priority Date | Filing Date |
---|---|---|---|
CN201210115719.0A Expired - Fee Related CN102708009B (zh) | 2012-04-19 | 2012-04-19 | 一种基于cuda实现多任务共享gpu的方法 |
Country Status (1)
Country | Link |
---|---|
CN (1) | CN102708009B (zh) |
Families Citing this family (10)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
CN103440660B (zh) * | 2013-09-03 | 2016-03-30 | 四川大学 | 一种基于gpu的集成成像微图像阵列快速生成方法 |
CN104866297B (zh) * | 2014-02-26 | 2018-05-29 | 华为技术有限公司 | 一种优化核函数的方法和装置 |
CN104102513B (zh) * | 2014-07-18 | 2017-06-16 | 西北工业大学 | 一种基于Kepler架构的CUDA运行时参数透明优化方法 |
WO2016041126A1 (zh) * | 2014-09-15 | 2016-03-24 | 华为技术有限公司 | 基于gpu的数据流处理方法和装置 |
WO2019000435A1 (zh) * | 2017-06-30 | 2019-01-03 | 华为技术有限公司 | 任务处理方法、装置、介质及其设备 |
CN107329818A (zh) * | 2017-07-03 | 2017-11-07 | 郑州云海信息技术有限公司 | 一种任务调度处理方法及装置 |
CN110825514B (zh) * | 2018-08-10 | 2023-05-23 | 昆仑芯(北京)科技有限公司 | 人工智能芯片以及用于人工智能芯片的指令执行方法 |
CN110825530B (zh) * | 2018-08-10 | 2022-12-23 | 昆仑芯(北京)科技有限公司 | 用于人工智能芯片的指令执行方法和装置 |
CN110413408A (zh) * | 2019-06-29 | 2019-11-05 | 苏州浪潮智能科技有限公司 | 一种深度学习框架的显存控制方法、设备以及存储介质 |
CN113407333B (zh) * | 2020-12-18 | 2023-05-26 | 上海交通大学 | Warp级别调度的任务调度方法、系统、GPU及设备 |
Citations (2)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
CN101937425A (zh) * | 2009-07-02 | 2011-01-05 | 北京理工大学 | 基于gpu众核平台的矩阵并行转置方法 |
CN102298522A (zh) * | 2011-09-13 | 2011-12-28 | 四川卫士通信息安全平台技术有限公司 | 一种使用gpu实现sha-1算法的方法 |
-
2012
- 2012-04-19 CN CN201210115719.0A patent/CN102708009B/zh not_active Expired - Fee Related
Patent Citations (2)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
CN101937425A (zh) * | 2009-07-02 | 2011-01-05 | 北京理工大学 | 基于gpu众核平台的矩阵并行转置方法 |
CN102298522A (zh) * | 2011-09-13 | 2011-12-28 | 四川卫士通信息安全平台技术有限公司 | 一种使用gpu实现sha-1算法的方法 |
Non-Patent Citations (4)
Title |
---|
CUDA并行程序的内存访问优化技术研究;邹岩 等;《计算机测量与控制》;20091231;第17卷(第12期);全文 * |
付娟.信息安全算法的GPU高速实现.《南昌大学硕士学位论文》.2010, |
信息安全算法的GPU高速实现;付娟;《南昌大学硕士学位论文》;20100106;全文 * |
邹岩 等.CUDA并行程序的内存访问优化技术研究.《计算机测量与控制》.2009,第17卷(第12期), |
Also Published As
Publication number | Publication date |
---|---|
CN102708009A (zh) | 2012-10-03 |
Similar Documents
Publication | Publication Date | Title |
---|---|---|
CN102708009B (zh) | 一种基于cuda实现多任务共享gpu的方法 | |
CN102902512B (zh) | 一种基于多线程编程及消息队列的多线程并行处理方法 | |
Venkataraman et al. | Presto: distributed machine learning and graph processing with sparse matrices | |
Song et al. | Dynamic task scheduling for linear algebra algorithms on distributed-memory multicore systems | |
TWI525540B (zh) | 具有橫跨多個處理器之平行資料執行緒的映射處理邏輯 | |
US8990827B2 (en) | Optimizing data warehousing applications for GPUs using dynamic stream scheduling and dispatch of fused and split kernels | |
KR101559090B1 (ko) | 이종 코어를 위한 자동 커널 마이그레이션 | |
Cho et al. | PARADIS: An efficient parallel algorithm for in-place radix sort | |
CN104050032A (zh) | 用于有条件的屏障和急迫的屏障的硬件调度的系统和方法 | |
CN103870213B (zh) | 一种生成性能数据的图形处理管线及计算系统 | |
US20180329753A1 (en) | Scheduling heterogenous computation on multithreaded processors | |
CN103996216A (zh) | 用于曲面细分和几何着色器的电力高效属性处置 | |
CN103870309A (zh) | 用于集群多级寄存器堆的寄存器分配 | |
Sbîrlea et al. | Bounded memory scheduling of dynamic task graphs | |
CN115934102A (zh) | 通用寄存器动态分配方法、装置、计算机设备和存储介质 | |
Fang et al. | Aristotle: A performance impact indicator for the OpenCL kernels using local memory | |
Zhang et al. | Optimization of N-queens solvers on graphics processors | |
CN103699363A (zh) | 一种用于在多核平台下优化关键临界区的方法 | |
Valero et al. | Towards a more efficient use of gpus | |
Huynh et al. | TP-PARSEC: A task parallel PARSEC benchmark suite | |
CN116775265A (zh) | 协作组阵列 | |
CN116774914A (zh) | 分布式共享存储器 | |
JP4950325B2 (ja) | モンテカルロ法の効率的な並列処理手法 | |
Siddiqui et al. | Design space exploration of embedded applications on heterogeneous cpu-gpu platforms | |
Jeannot | Performance analysis and optimization of the tiled cholesky factorization on numa machines |
Legal Events
Date | Code | Title | Description |
---|---|---|---|
C06 | Publication | ||
PB01 | Publication | ||
C10 | Entry into substantive examination | ||
SE01 | Entry into force of request for substantive examination | ||
C14 | Grant of patent or utility model | ||
GR01 | Patent grant | ||
CF01 | Termination of patent right due to non-payment of annual fee |
Granted publication date: 20140402 Termination date: 20150419 |
|
EXPY | Termination of patent right or utility model |