CN102708009B - 一种基于cuda实现多任务共享gpu的方法 - Google Patents

一种基于cuda实现多任务共享gpu的方法 Download PDF

Info

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
Application number
CN201210115719.0A
Other languages
English (en)
Other versions
CN102708009A (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.)
Peking University
Huawei Technologies Co Ltd
Original Assignee
Peking University
Huawei Technologies Co Ltd
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 Peking University, Huawei Technologies Co Ltd filed Critical Peking University
Priority to CN201210115719.0A priority Critical patent/CN102708009B/zh
Publication of CN102708009A publication Critical patent/CN102708009A/zh
Application granted granted Critical
Publication of CN102708009B publication Critical patent/CN102708009B/zh
Expired - Fee Related legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

Landscapes

  • Multi Processors (AREA)

Abstract

本发明公开了一种基于CUDA实现多任务共享GPU的方法。包括:在Global Memory中建立映射表,确定合并后的Kernel中,每个Block执行的任务编号和任务分块编号;一次用一个Kernel启动N个Block,N等于所有任务的任务分块数目之和;用标记和阻塞等待的方法,满足原有任务之间的约束关系;对于Shared Memory,采用预申请和静态分配的方式进行多任务共享。通过本发明,可以简便地实现在现有GPU硬件架构上实现多任务共享,可以简化实际应用中的编程工作,并在一定情况下取得良好的性能。

Description

一种基于CUDA实现多任务共享GPU的方法
技术领域
本发明涉及一种多任务共享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函数使用。
如下面代码所示。
Figure BDA0000154914300000061
Figure BDA0000154914300000071
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。
Figure BDA0000154914300000072
在需要满足约束的地方,对对应的mark进行阻塞检查。例如,任务1依赖于任务0,所以在任务1的device函数之前,需要满足约束。于是增加一个循环等待语句,阻塞等待任务0的标志位mark0变为1。
Figure BDA0000154914300000073
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的时候,直接在这个数组中静态分配使用。
CN201210115719.0A 2012-04-19 2012-04-19 一种基于cuda实现多任务共享gpu的方法 Expired - Fee Related CN102708009B (zh)

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)

* Cited by examiner, † Cited by third party
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)

* Cited by examiner, † Cited by third party
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算法的方法

Patent Citations (2)

* Cited by examiner, † Cited by third party
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)

* Cited by examiner, † Cited by third party
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