CN117473212A - Ntt算法的gpu加速方法、装置、设备及存储介质 - Google Patents
Ntt算法的gpu加速方法、装置、设备及存储介质 Download PDFInfo
- Publication number
- CN117473212A CN117473212A CN202311814455.2A CN202311814455A CN117473212A CN 117473212 A CN117473212 A CN 117473212A CN 202311814455 A CN202311814455 A CN 202311814455A CN 117473212 A CN117473212 A CN 117473212A
- Authority
- CN
- China
- Prior art keywords
- matrix
- gpu
- polynomial
- value
- algorithm
- 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.)
- Granted
Links
- 238000000034 method Methods 0.000 title claims abstract description 70
- 230000001133 acceleration Effects 0.000 title claims abstract description 61
- 239000011159 matrix material Substances 0.000 claims abstract description 406
- 238000004364 calculation method Methods 0.000 claims abstract description 67
- 230000009466 transformation Effects 0.000 claims abstract description 41
- 238000000354 decomposition reaction Methods 0.000 claims description 18
- 230000008569 process Effects 0.000 description 9
- 230000000694 effects Effects 0.000 description 7
- 238000004891 communication Methods 0.000 description 6
- HPTJABJPZMULFH-UHFFFAOYSA-N 12-[(Cyclohexylcarbamoyl)amino]dodecanoic acid Chemical compound OC(=O)CCCCCCCCCCCNC(=O)NC1CCCCC1 HPTJABJPZMULFH-UHFFFAOYSA-N 0.000 description 2
- 238000004458 analytical method Methods 0.000 description 2
- 238000006243 chemical reaction Methods 0.000 description 2
- 238000010586 diagram Methods 0.000 description 2
- 230000006870 function Effects 0.000 description 2
- 239000011800 void material Substances 0.000 description 2
- 238000009825 accumulation Methods 0.000 description 1
- 238000007792 addition Methods 0.000 description 1
- 230000008859 change Effects 0.000 description 1
- 125000004122 cyclic group Chemical group 0.000 description 1
- 238000005516 engineering process Methods 0.000 description 1
- 230000003287 optical effect Effects 0.000 description 1
- 238000005457 optimization Methods 0.000 description 1
- 238000000844 transformation Methods 0.000 description 1
- 238000011144 upstream manufacturing Methods 0.000 description 1
Classifications
-
- G—PHYSICS
- G06—COMPUTING; CALCULATING OR COUNTING
- G06F—ELECTRIC DIGITAL DATA PROCESSING
- G06F17/00—Digital computing or data processing equipment or methods, specially adapted for specific functions
- G06F17/10—Complex mathematical operations
- G06F17/14—Fourier, Walsh or analogous domain transformations, e.g. Laplace, Hilbert, Karhunen-Loeve, transforms
- G06F17/141—Discrete Fourier transforms
- G06F17/142—Fast Fourier transforms, e.g. using a Cooley-Tukey type algorithm
-
- G—PHYSICS
- G06—COMPUTING; CALCULATING OR COUNTING
- G06F—ELECTRIC DIGITAL DATA PROCESSING
- G06F17/00—Digital computing or data processing equipment or methods, specially adapted for specific functions
- G06F17/10—Complex mathematical operations
- G06F17/16—Matrix or vector computation, e.g. matrix-matrix or matrix-vector multiplication, matrix factorization
-
- 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
- Physics & Mathematics (AREA)
- Engineering & Computer Science (AREA)
- Mathematical Physics (AREA)
- General Physics & Mathematics (AREA)
- Computational Mathematics (AREA)
- Mathematical Analysis (AREA)
- Mathematical Optimization (AREA)
- Pure & Applied Mathematics (AREA)
- Data Mining & Analysis (AREA)
- Theoretical Computer Science (AREA)
- Algebra (AREA)
- Databases & Information Systems (AREA)
- Software Systems (AREA)
- General Engineering & Computer Science (AREA)
- Discrete Mathematics (AREA)
- Computing Systems (AREA)
- Complex Calculations (AREA)
Abstract
本发明公开了一种NTT算法的GPU加速方法、装置、设备及计算机可读存储介质,涉及信息安全技术领域,方法包括:获取多项式的系数序列和模数,获取第一数值和第二数值;根据第一数值和第二数值计算目标容量;在多项式的模数小于阈值且显存容量大于目标容量时,根据第一数值和第二数值将旋转因子矩阵拆分成多个子矩阵,对系数序列和子矩阵采用GPU矩阵乘法操作计算得到多项式的点值序列;在多项式的模数大于或等于预设阈值或显存容量小于或等于目标容量时,对系数序列采用GPU多线程并行执行蝴蝶变换,以计算得到多项式的点值序列。本发明提出一种NTT算法的GPU加速方案,提高全同态加密算法中NTT算法的计算效率。
Description
技术领域
本发明涉及信息安全技术领域,尤其涉及一种NTT算法的GPU加速方法、装置、设备及计算机可读存储介质。
背景技术
快速数论变换NTT(Number Theoretic Transform)算法是一种基于数论的离散傅里叶变换(DFT)的变种,常用于数字信号处理、多项式乘法以及在同态加密等领域中的应用。NTT算法的本质思想跟FFT是一致的,但它以数论为基础的具有循环卷积性质的快速数论变换(NTT),用有限域上的单位根来取代复平面上的单位根。NTT算法特别在同态加密中具有重要作用,可以将多项式从系数序列转换为向量点值序列形式的方法。这样做可以有效加速同态加密算法中多项式乘法的计算过程。常见CKKS算法,BFV/BGV算法都涉及到大量的NTT和INTT操作。但是,目前,基于CPU来实现BFV、BGV、CKKS等全同态加密算法中NTT算法,一方面由于CPU硬件的特性,另一方面由于全同态加密算法中NTT算法的使用较频繁,所以,导致基于CPU的实现方案计算效率差,耗时长。
发明内容
本发明的主要目的在于提供一种NTT算法的GPU加速方法、装置、设备及计算机可读存储介质,旨在提出一种NTT算法的GPU加速方案,提高BFV、BGV、CKKS等全同态加密算法中NTT算法的计算效率。
为实现上述目的,本发明提供一种NTT算法的GPU加速方法,包括以下步骤:
获取NTT算法中多项式的系数序列和所述多项式的模数,以及获取基于所述多项式的模数分解得到的第一数值和第二数值;
根据所述第一数值和所述第二数值计算GPU显存的目标容量;
在所述多项式的模数小于预设阈值,且GPU的显存容量大于所述目标容量的情况下,根据所述第一数值和所述第二数值将NTT算法中的旋转因子矩阵拆分成多个子矩阵,对所述系数序列和所述子矩阵采用GPU矩阵乘法操作计算得到所述多项式的点值序列;
在所述多项式的模数大于或等于所述预设阈值,或所述显存容量小于或等于所述目标容量的情况下,对所述系数序列采用GPU多线程并行执行NTT算法中的蝴蝶变换,以计算得到所述多项式的点值序列。
可选地,所述根据所述第一数值和所述第二数值计算GPU显存的目标容量的步骤包括:
计算所述第一数值的平方、所述第二数值的平方以及所述第一数值和所述第二数值的乘积,将三个计算结果相加后与预设位长和预设模数相乘,得到GPU显存的目标容量。
可选地,所述根据所述第一数值和所述第二数值将NTT算法中的旋转因子矩阵拆分成多个子矩阵的步骤包括:
将所述旋转因子矩阵拆分成第一子矩阵、第二子矩阵和第三子矩阵,其中,所述第一子矩阵的维度根据所述第一数值确定,所述第二子矩阵的维度根据所述第一数值和所述第二数值确定,所述第三子矩阵的维度根据所述第二数值确定。
可选地,所述对所述系数序列和所述子矩阵采用GPU矩阵乘法操作计算得到所述多项式的点值序列的步骤包括:
将所述系数序列转换为系数矩阵,所述系数矩阵的维度根据所述第一数值和所述第二数值确定;
通过GPU第一矩阵乘法操作对所述第一子矩阵与所述系数矩阵进行计算,得到第一中间矩阵,通过GPU第二矩阵乘法操作对所述第一中间矩阵与所述第二子矩阵进行计算,得到第二中间矩阵,通过GPU第一矩阵乘法操作对所述第二中间矩阵与所述第三子矩阵进行计算,得到第三中间矩阵;
根据所述多项式的模数对所述第三中间矩阵进行模运算,得到点值矩阵,将所述点值矩阵转换得到所述多项式的点值序列。
可选地,所述通过GPU第一矩阵乘法操作对所述第一子矩阵与所述系数矩阵进行计算,得到第一中间矩阵,通过GPU第二矩阵乘法操作对所述第一中间矩阵与所述第二子矩阵进行计算,得到第二中间矩阵,通过GPU第一矩阵乘法操作对所述第二中间矩阵与所述第三子矩阵进行计算,得到第三中间矩阵的步骤包括:
将所述第一子矩阵与所述系数矩阵进行GPU矩阵叉乘操作得到第一中间矩阵,将所述第一中间矩阵与所述第二子矩阵进行GPU矩阵点乘操作得到第二中间矩阵,将所述第二中间矩阵与所述第三子矩阵进行GPU矩阵叉乘操作得到第三中间矩阵。
可选地,所述将所述第一子矩阵与所述系数矩阵进行GPU矩阵叉乘操作得到第一中间矩阵,将所述第一中间矩阵与所述第二子矩阵进行GPU矩阵点乘操作得到第二中间矩阵的步骤包括:
根据各个GPU线程所在线程块的块索引和在各自线程块中的线程索引,为各所述GPU线程分配叉乘运算任务的第一任务数据和点乘运算任务的第二任务数据;
通过各所述GPU线程根据所分配的第一任务数据执行叉乘运算获得第一计算结果,其中,所述第一任务数据来自于所述第一子矩阵和所述系数矩阵;
通过各所述GPU线程根据所分配的第二任务数据执行点乘运算获得第二计算结果,其中,所述第二任务数据来自于所述第一中间矩阵和所述第二子矩阵;
基于各所述GPU线程获得的所述第一计算结果获得所述第一中间矩阵,基于各所述GPU线程获得的所述第二计算结果获得所述第二中间矩阵。
可选地,所述对所述系数序列采用GPU多线程并行执行NTT算法中的蝴蝶变换,以计算得到所述多项式的点值序列的步骤包括:
针对各个GPU线程,分别执行与所述多项式的模数相对应的次数的循环;
其中,在每一次循环中,根据所述GPU线程的线程编号和当前循环的轮次,为所述GPU线程分配蝴蝶变换操作任务的任务数据,通过所述GPU线程根据所分配的任务数据执行NTT算法中的蝴蝶变换获得计算结果,其中,所述任务数据来自于所述系数序列;
基于各所述GPU线程在各次循环中执行蝴蝶变换得到的计算结果,获得所述多项式的点值序列。
为实现上述目的,本发明还提供一种NTT算法的GPU加速装置,所述NTT算法的GPU加速装置包括:
获取模块,用于获取NTT算法中多项式的系数序列和所述多项式的模数,以及获取基于所述多项式的模数分解得到的第一数值和第二数值;
第一计算模块,用于根据所述第一数值和所述第二数值计算GPU显存的目标容量;
第二计算模块,用于在所述多项式的模数小于预设阈值,且GPU的显存容量大于所述目标容量的情况下,根据所述第一数值和所述第二数值将NTT算法中的旋转因子矩阵拆分成多个子矩阵,对所述系数序列和所述子矩阵采用GPU矩阵乘法操作计算得到所述多项式的点值序列;
第三计算模块,用于在所述多项式的模数大于或等于所述预设阈值,或所述显存容量小于或等于所述目标容量的情况下,对所述系数序列采用GPU多线程并行执行NTT算法中的蝴蝶变换,以计算得到所述多项式的点值序列。
为实现上述目的,本发明还提供一种NTT算法的GPU加速设备,所述NTT算法的GPU加速设备包括:存储器、处理器及存储在所述存储器上并可在所述处理器上运行的NTT算法的GPU加速程序,所述NTT算法的GPU加速程序被所述处理器执行时实现如上所述的NTT算法的GPU加速方法的步骤。
此外,为实现上述目的,本发明还提出一种计算机可读存储介质,所述计算机可读存储介质上存储有NTT算法的GPU加速程序,所述NTT算法的GPU加速程序被处理器执行时实现如上所述的NTT算法的GPU加速方法的步骤。
在本发明实施例中,提供了两种GPU加速NTT算法的方法:算法1是针对NTT算法蝴蝶变换的实现场景的GPU加速方法,考虑了分配多线程并行计算的方式实现高效利用GPU,从而实现通过GPU加速NTT算法的效果;算法2是针对NTT算法矩阵向量乘法的实现场景的GPU加速方法,在旋转因子矩阵较大的情况下对其进行矩阵分解,以解决GPU内存限制的问题,并利用GPU矩阵乘法提高GPU利用率,从而实现通过GPU加速NTT算法的效果;并且,针对实现GPU加速NTT算法的方法的特点,提出通过计算多项式的模数和GPU硬件显存大小的关系,动态决策使用算法1还是算法2以实现GPU加速NTT算法,通过对多项式的模数和GPU显存容量大小的分析,决策使用算法1或算法2以使得GPU加速NTT算法的方法更优,从而提高了BFV、BGV、CKKS等全同态加密算法中NTT算法的计算效率。
附图说明
图1为本发明实施例方案涉及的硬件运行环境的结构示意图;
图2为本发明NTT算法的GPU加速方法第一实施例的流程示意图;
图3为本发明一可行实施方式中涉及的实现GPU加速NTT算法的整体流程图;
图4为本发明NTT算法的GPU加速装置较佳实施例的功能模块示意图。
本发明目的的实现、功能特点及优点将结合实施例,参照附图做进一步说明。
具体实施方式
应当理解,此处所描述的具体实施例仅仅用以解释本发明,并不用于限定本发明。
如图1所示,图1是本发明实施例方案涉及的硬件运行环境的设备结构示意图。
需要说明的是,本发明实施例NTT算法的GPU加速设备,所述NTT算法的GPU加速设备可以是个人计算机、服务器、智能手机等设备,在此不做具体限制。
如图1所示,该NTT算法的GPU加速设备可以包括:处理器1001,例如CPU,网络接口1004,用户接口1003,存储器1005,通信总线1002。其中,通信总线1002用于实现这些组件之间的连接通信。用户接口1003可以包括显示屏(Display)、输入单元比如键盘(Keyboard),可选用户接口1003还可以包括标准的有线接口、无线接口。网络接口1004可选的可以包括标准的有线接口、无线接口(如WI-FI接口)。存储器1005可以是高速RAM存储器,也可以是稳定的存储器(non-volatile memory),例如磁盘存储器。存储器1005可选的还可以是独立于前述处理器1001的存储装置。
本领域技术人员可以理解,图1中示出的设备结构并不构成对NTT算法的GPU加速设备的限定,可以包括比图示更多或更少的部件,或者组合某些部件,或者不同的部件布置。
如图1所示,作为一种计算机存储介质的存储器1005中可以包括操作系统、网络通信模块、用户接口模块以及NTT算法的GPU加速程序。操作系统是管理和控制设备硬件和软件资源的程序,支持NTT算法的GPU加速程序以及其它软件或程序的运行。在图1所示的设备中,用户接口1003主要用于与客户端进行数据通信;网络接口1004主要用于与服务器建立通信连接;而处理器1001可以用于调用存储器1005中存储的NTT算法的GPU加速程序,并执行以下操作:
获取NTT算法中多项式的系数序列和所述多项式的模数,以及获取基于所述多项式的模数分解得到的第一数值和第二数值;
根据所述第一数值和所述第二数值计算GPU显存的目标容量;
在所述多项式的模数小于预设阈值,且GPU的显存容量大于所述目标容量的情况下,根据所述第一数值和所述第二数值将NTT算法中的旋转因子矩阵拆分成多个子矩阵,对所述系数序列和所述子矩阵采用GPU矩阵乘法操作计算得到所述多项式的点值序列;
在所述多项式的模数大于或等于所述预设阈值,或所述显存容量小于或等于所述目标容量的情况下,对所述系数序列采用GPU多线程并行执行NTT算法中的蝴蝶变换,以计算得到所述多项式的点值序列。
在一可行实施方式中,所述根据所述第一数值和所述第二数值计算GPU显存的目标容量的操作包括:
计算所述第一数值的平方、所述第二数值的平方以及所述第一数值和所述第二数值的乘积,将三个计算结果相加后与预设位长和预设模数相乘,得到GPU显存的目标容量。
在一可行实施方式中,所述根据所述第一数值和所述第二数值将NTT算法中的旋转因子矩阵拆分成多个子矩阵的操作包括:
将所述旋转因子矩阵拆分成第一子矩阵、第二子矩阵和第三子矩阵,其中,所述第一子矩阵的维度根据所述第一数值确定,所述第二子矩阵的维度根据所述第一数值和所述第二数值确定,所述第三子矩阵的维度根据所述第二数值确定。
在一可行实施方式中,所述对所述系数序列和所述子矩阵采用GPU矩阵乘法操作计算得到所述多项式的点值序列的操作包括:
将所述系数序列转换为系数矩阵,所述系数矩阵的维度根据所述第一数值和所述第二数值确定;
通过GPU第一矩阵乘法操作对所述第一子矩阵与所述系数矩阵进行计算,得到第一中间矩阵,通过GPU第二矩阵乘法操作对所述第一中间矩阵与所述第二子矩阵进行计算,得到第二中间矩阵,通过GPU第一矩阵乘法操作对所述第二中间矩阵与所述第三子矩阵进行计算,得到第三中间矩阵;
根据所述多项式的模数对所述第三中间矩阵进行模运算,得到点值矩阵,将所述点值矩阵转换得到所述多项式的点值序列。
在一可行实施方式中,所述通过GPU第一矩阵乘法操作对所述第一子矩阵与所述系数矩阵进行计算,得到第一中间矩阵,通过GPU第二矩阵乘法操作对所述第一中间矩阵与所述第二子矩阵进行计算,得到第二中间矩阵,通过GPU第一矩阵乘法操作对所述第二中间矩阵与所述第三子矩阵进行计算,得到第三中间矩阵的操作包括:
将所述第一子矩阵与所述系数矩阵进行GPU矩阵叉乘操作得到第一中间矩阵,将所述第一中间矩阵与所述第二子矩阵进行GPU矩阵点乘操作得到第二中间矩阵,将所述第二中间矩阵与所述第三子矩阵进行GPU矩阵叉乘操作得到第三中间矩阵。
在一可行实施方式中,所述将所述第一子矩阵与所述系数矩阵进行GPU矩阵叉乘操作得到第一中间矩阵,将所述第一中间矩阵与所述第二子矩阵进行GPU矩阵点乘操作得到第二中间矩阵的操作包括:
根据各个GPU线程所在线程块的块索引和在各自线程块中的线程索引,为各所述GPU线程分配叉乘运算任务的第一任务数据和点乘运算任务的第二任务数据;
通过各所述GPU线程根据所分配的第一任务数据执行叉乘运算获得第一计算结果,其中,所述第一任务数据来自于所述第一子矩阵和所述系数矩阵;
通过各所述GPU线程根据所分配的第二任务数据执行点乘运算获得第二计算结果,其中,所述第二任务数据来自于所述第一中间矩阵和所述第二子矩阵;
基于各所述GPU线程获得的所述第一计算结果获得所述第一中间矩阵,基于各所述GPU线程获得的所述第二计算结果获得所述第二中间矩阵。
在一可行实施方式中,所述对所述系数序列采用GPU多线程并行执行NTT算法中的蝴蝶变换,以计算得到所述多项式的点值序列的操作包括:
针对各个GPU线程,分别执行与所述多项式的模数相对应的次数的循环;
其中,在每一次循环中,根据所述GPU线程的线程编号和当前循环的轮次,为所述GPU线程分配蝴蝶变换操作任务的任务数据,通过所述GPU线程根据所分配的任务数据执行NTT算法中的蝴蝶变换获得计算结果,其中,所述任务数据来自于所述系数序列;
基于各所述GPU线程在各次循环中执行蝴蝶变换得到的计算结果,获得所述多项式的点值序列。
基于上述的结构,提出NTT算法的GPU加速方法的各个实施例。
参照图2,图2为本发明NTT算法的GPU加速方法第一实施例的流程示意图。
本发明实施例提供了NTT算法的GPU加速方法的实施例,需要说明的是,虽然在流程图中示出了逻辑顺序,但是在某些情况下,可以以不同于此处的顺序执行所示出或描述的步骤。在本实施例中,NTT算法的GPU加速方法的执行主体可以是部署有GPU硬件环境的设备,在本实施例中并不做限制,以下为便于描述,省略执行主体进行各实施例的阐述。在本实施例中,所述NTT算法的GPU加速方法包括步骤S10~S40:
步骤S10,获取NTT算法中多项式的系数序列和所述多项式的模数,以及获取基于所述多项式的模数分解得到的第一数值和第二数值。
NTT算法在同态加密中具有重要作用,可以将多项式从系数序列转换为点值序列形式,这样做可以有效加速同态加密算法中多项式乘法的计算过程。NTT算法有蝴蝶变换和矩阵向量乘法两种实现方式,均可以实现将多项式从系数序列转换为点值序列形式。若要加速实现同态加密算法场景,加速优化NTT算法是最关键的一个环节。
本实施例中,为解决通过CPU实现NTT算法的方案计算速度慢的问题,提出基于GPU加速NTT算法的方案。具体地,提供了两种GPU加速NTT算法的方法:一种是针对NTT算法蝴蝶变换的实现场景的GPU加速方法,也即,对系数序列采用GPU多线程并行执行NTT算法中的蝴蝶变换,以计算得到多项式的点值序列(以下称为算法1),考虑了分配多线程并行计算的方式实现高效利用GPU;另一种是针对NTT算法矩阵向量乘法的实现场景的GPU加速方法,也即,根据第一数值和第二数值将NTT算法中的旋转因子矩阵拆分成多个子矩阵,对系数序列和子矩阵采用GPU矩阵乘法操作计算得到多项式的点值序列(以下称为算法2),在旋转因子矩阵(也可称为NTT转换矩阵)较大的情况下进行矩阵分解,以解决GPU内存限制的问题;并且,通过预估模数和GPU显存容量大小可以决策使用哪种GPU加速NTT算法的方法更优。
多项式(polynomial)是指由变量、系数以及它们之间的加、减、乘、幂运算(非负整数次方)得到的表达式。多项式可以通过系数序列和点值序列表示。
多项式的系数序列和模数是已知量,可以从上游计算模块中获取,本实施例对其获取方式并不做限制。
可以将多项式的模数分解为第一数值和第二数值。在一可行实施方式中,分解得到的第一数值和第二数值相乘可以是等于该多项式的模数的。例如,采用X表示多项式的模数,X1表示第一数值,X2表示第二数值,那么X=X1*X2。在将多项式的模数分解为两个数值相乘的方式有很多种的情况下,可以将各种分解方式以其中较大的数值进行大小排序,将最小数值所对应的分解方式作为最终的分解方式,获得第一数值和第二数值。例如,X=16,那么有2*8和4*4两种分解方式,将两种分解方式以其中较大的数值进行大小排序,也即将8和4进行排序,最小的数值为4,那么将4*4这种分解方式作为最终的分解方式,获得第一数值和第二数值均为4,由于方阵(行列维度相等的矩阵)的计算速度比非方阵(行列维度不等的矩阵)的计算速度更快,因此一般选择方阵作为最终分解方式。
步骤S20,根据所述第一数值和所述第二数值计算GPU显存的目标容量。
根据第一数值和第二数值计算目标容量,目标容量即采用算法2实现加速NTT算法所需的GPU的显存容量。根据第一数值和第二数值计算目标容量的方式在本实施例中并不做限制。例如,在一可行实施方式中,所述步骤S20包括步骤S201:
步骤S201,计算所述第一数值的平方、所述第二数值的平方以及所述第一数值和所述第二数值的乘积,将三个计算结果相加后与预设位长和预设模数相乘,得到GPU显存的目标容量。
采用X1表示第一数值、X2表示第二数值、假设预设位长设置为64bit,那么,目标容量=k*(X1*X1+X1*X2+X2*X2)*64bit。k表示预设模数根据中国剩余定理进行分解得到的模数个数。
步骤S30,在所述多项式的模数小于预设阈值,且GPU的显存容量大于所述目标容量的情况下,根据所述第一数值和所述第二数值将NTT算法中的旋转因子矩阵拆分成多个子矩阵,对所述系数序列和所述子矩阵采用GPU矩阵乘法操作计算得到所述多项式的点值序列。
旋转因子矩阵是NTT算法矩阵向量乘法的实现场景中所使用的用于将系数序列转换为点值序列的转换矩阵。
算法2更适合多项式的模数较小的场景。预设阈值可以预先根据需要设置,例如设置为1024,当多项式的模数小于预设阈值时,表示采用算法2实现加速NTT算法更优。当多项式的模数小于预设阈值,且GPU的显存容量大于目标容量的情况下,采用算法2实现加速NTT算法,也即:根据第一数值和第二数值将维度较大的旋转因子矩阵拆分成多个小矩阵(以下称为子矩阵),从而避免采用较大的旋转因子矩阵导致GPU显存空间会超过限制;对系数序列和子矩阵采用GPU矩阵乘法操作计算得到多项式的点值序列。GPU矩阵乘法操作的具体实施方式有很多种,在本实施例中并不做限制。GPU矩阵乘法操作通过多线程并行处理乘法操作任务,所以能够提高GPU利用率,从而实现通过GPU加速NTT算法的效果。
步骤S40,在所述多项式的模数大于或等于所述预设阈值,或所述显存容量小于或等于所述目标容量的情况下,对所述系数序列采用GPU多线程并行执行NTT算法中的蝴蝶变换,以计算得到所述多项式的点值序列。
在多项式的模数大于或等于预设阈值的情况下,或在GPU的显存容量小于或等于目标容量的情况下,说明采用算法1实现加速NTT算法更优。采用算法1实现加速NTT算法,考虑到了GPU多线程的优势,通过利用GPU多线程来并行执行NTT算法中的蝴蝶变换,从而实现通过GPU加速NTT算法的效果。
在本实施例中,提供了两种GPU加速NTT算法的方法:算法1是针对NTT算法蝴蝶变换的实现场景的GPU加速方法,考虑了分配多线程并行计算的方式实现高效利用GPU,从而实现通过GPU加速NTT算法的效果;算法2是针对NTT算法矩阵乘法的实现场景的GPU加速方法,在旋转因子矩阵较大的情况下对其进行矩阵分解,以解决GPU内存限制的问题,并利用GPU矩阵乘法提高GPU利用率,从而实现通过GPU加速NTT算法的效果;并且,针对实现GPU加速NTT算法的方法的特点,提出通过计算多项式的模数和GPU硬件显存大小的关系,动态决策使用算法1还是算法2以实现GPU加速NTT算法,通过对多项式的模数和GPU显存容量大小的分析,决策使用算法1或算法2以使得GPU加速NTT算法的方法更优,从而提高了BFV、BGV、CKKS等全同态加密算法中NTT算法的计算效率。上述两种GPU加速NTT算法的方法对比CPU硬件下实现NTT算法的方法,在计算效率提升上有显著效果,算法1在GPU上实现NTT算法比在CPU硬件上提速20倍+,算法1在多项式模数较大的情况下表现更优且未超过GPU的显存限制;算法2在多项式模数较小的情况下表现更优,算法2在原始旋转因子矩阵为32768*32768维度时,可将其分解成三个小矩阵进行存储,未超过GPU内存使用限制。
在一可行实施方式中,可以按照如图3所示的流程实现GPU加速NTT算法。
基于上述第一实施例,提出本发明NTT算法的GPU加速方法第二实施例。在本实施例中,提出一种拆分旋转因子矩阵的可行实施方式,所述步骤S30中根据所述第一数值和所述第二数值将NTT算法中的旋转因子矩阵拆分成多个子矩阵的步骤包括S301。
步骤S301,将所述旋转因子矩阵拆分成第一子矩阵、第二子矩阵和第三子矩阵,其中,所述第一子矩阵的维度根据所述第一数值确定,所述第二子矩阵的维度根据所述第一数值和所述第二数值确定,所述第三子矩阵的维度根据所述第二数值确定。
在CPU实现NTT算法的场景下,NTT算法使用矩阵乘法公式如下:
其中,旋转因子矩阵表示为N,大小为X*X,其中矩阵每个位置的值nij如上公式所示,等于单位根的幂次。e表示多项式的系数序列,eT表示系数序列的转置。M表示最终NTT变换后结果,即为多项式的点值矩阵。
在本实施例中,可以将旋转因子矩阵拆分成三个子矩阵,分别称为第一子矩阵、第二子矩阵和第三子矩阵。其中,第一子矩阵的维度可以是根据第一数值确定的,第二子矩阵的维度可以根据第一数值和第二数值确定的,第三子矩阵的维度可以是根据第二数值确定的。
例如,在一可行实施方式中,第一子矩阵的行数和列数均为第一数值,第二子矩阵的行数和列数分别为第一数值和第二数值,第三子矩阵的行数和列数均为第二数值。第一子矩阵表示为N1,第二子矩阵表示为N2,第三子矩阵表示为N3,第一数值和第二数值分别表示为X1和X2,那么,N1、N2、N3三个子矩阵的维度大小分别表示为X1*X1、X1*X2、X2*X2。
在本实施例中,还提出一种基于拆分的子矩阵和系数序列采用GPU矩阵乘法操作计算多项式点值序列的可行实施方式,所述步骤S30中对所述系数序列和所述子矩阵采用GPU矩阵乘法操作计算得到所述多项式的点值序列的步骤包括S302~S304。
步骤S302,将所述系数序列转换为系数矩阵,所述系数矩阵的维度根据所述第一数值和所述第二数值确定。
多项式的系数序列e的长度为X,将系数序列e转换为系数矩阵,系数矩阵表示为m,系数矩阵的行数和列数可以是分别根据第一数值和第二数值确定的。例如,在一可行实施方式中,系数矩阵的行数和列数分别为第一数值和第二数值,那么m的大小为X1*X2,其中序列对应位置的值不变化,按照顺序排布为矩阵形式,X1、X2分别为系数矩阵的行数和列数。
步骤S303,通过GPU第一矩阵乘法操作对所述第一子矩阵与所述系数矩阵进行计算,得到第一中间矩阵,通过GPU第二矩阵乘法操作对所述第一中间矩阵与所述第二子矩阵进行计算,得到第二中间矩阵,通过GPU第一矩阵乘法操作对所述第二中间矩阵与所述第三子矩阵进行计算,得到第三中间矩阵。
在一可行实施方式中,步骤S303包括:
步骤S3031,将所述第一子矩阵与所述系数矩阵进行GPU矩阵叉乘操作得到第一中间矩阵,将所述第一中间矩阵与所述第二子矩阵进行GPU矩阵点乘操作得到第二中间矩阵,将所述第二中间矩阵与所述第三子矩阵进行GPU矩阵叉乘操作得到第三中间矩阵。
GPU第一矩阵乘法操作可以是GPU矩阵叉乘操作,GPU第二矩阵乘法操作可以是GPU矩阵点乘操作,也即,第一子矩阵与系数矩阵的矩阵叉乘可以通过GPU矩阵叉乘操作实现,第一中间矩阵与第二子矩阵的矩阵点乘可以通过GPU矩阵点乘操作实现,第二中间矩阵与第三子矩阵的矩阵叉乘可以通过GPU矩阵叉乘操作实现。GPU矩阵点乘操作和GPU矩阵叉乘操作的原理是利用GPU多线程并行计算,其具体实现方式有很多种,在本实施方式中并不做限制。
步骤S304,根据所述多项式的模数对所述第三中间矩阵进行模运算,得到点值矩阵,将所述点值矩阵转换得到所述多项式的点值序列。
在将系数序列转换为系数矩阵,将旋转因子矩阵转换为三个小矩阵后,即可将NTT算法变换为:
其中,N1、N2、N3分别为第一子矩阵、第二子矩阵和第三子矩阵,m为系数矩阵,M为点值矩阵,M的大小为X1*X2。
将点值矩阵转换为点值序列即将点值矩阵的第一行与第二行拼接,再与第三行拼接,以此类推,得到一个序列,该过程不改变点值矩阵中的数值。
在本实施例中,还提出一种通过矩阵乘法操作计算矩阵叉乘的可行实施方式,步骤S3031中将所述第一子矩阵与所述系数矩阵进行GPU矩阵叉乘操作得到第一中间矩阵的步骤包括S30311~S30312。
步骤S30311,根据各个GPU线程所在线程块的块索引和在各自线程块中的线程索引,为各所述GPU线程分配叉乘运算任务的第一任务数据,通过各所述GPU线程根据所分配的第一任务数据执行叉乘运算获得第一计算结果,其中,所述第一任务数据来自于所述第一子矩阵和所述系数矩阵。
根据第一子矩阵和系数矩阵的大小,可以确定GPU线程的数量,并将各个GPU线程按照块划分,为各个GPU线程块给定块索引,以及为各个GPU线程给定其在线程块中的线程索引。根据各个GPU线程所在的线程块的块索引和在各自线程块中的线程索引,可以从第一子矩阵和系数矩阵中为各个GPU线程分配叉乘运算任务的第一任务数据,然后即可利用各个GPU线程根据各自被分配的任务数据并行地执行叉乘运算,以获得各自的第一计算结果。
为各个GPU线程分配叉乘运算任务的任务数据的实现方式有很多种,在本实施方式中并不做限制。
步骤S3032,基于各所述GPU线程获得的第一计算结果获得所述第一中间矩阵。
将各个GPU线程获得的第一计算结果进行拼接即可得到最终的叉乘结果,也即,得到将第一子矩阵和系数矩阵进行矩阵叉乘得到的第一中间矩阵。
在一可行实施方式中,可以采用上述步骤S3031~S3032中类似的步骤,实现通过GPU矩阵乘法操作,将所述第二中间矩阵与所述第三子矩阵进行矩阵叉乘得到第三中间矩阵,区别是为各个GPU线程分配的任务数据来自于第二中间矩阵和第三子矩阵。
在一可行实施方式中,可以采用如下代码所展示的原理来实现通过GPU矩阵叉乘操作计算第一子矩阵与系数矩阵的矩阵叉乘,或第二中间矩阵和第三子矩阵的矩阵叉乘。
__global__ void matrixMultiply(int* a,int* b,int* c,int rows,intcols,int shared_dim)
{
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row<rows&&col<cols)
{
int sum = 0;
for (int i = 0; i<shared_dim; ++i)
{
sum += a[row * shared_dim + i] * b[i * cols + col];
}
c[row * cols + col] = sum;
}
}
其中,__global__是CUDA核函数的标记,表示这个函数在GPU上执行。
a和b是输入矩阵,也即需要进行矩阵叉乘的两个矩阵,可以是第一子矩阵与系数矩阵,或第二中间矩阵和第三子矩阵。c是输出矩阵,也即矩阵叉乘得到的结果。a、b、c在设备内存上。
rows、cols和shared_dim是输入矩阵和输出矩阵的维度信息,其中,rows是a的行数和c的行数,shared_dim是a的列数和b的行数,cols是b的列数和c的列数。
int row = blockIdx.y * blockDim.y + threadIdx.y用于计算当前线程的行索引。int col= blockIdx.x * blockDim.x + threadIdx.x用于计算当前线程的列索引。blockIdx.y 和 blockIdx.x是当前线程块的索引,blockDim.y和blockDim.x 是线程块的维度,threadIdx.y和threadIdx.x是当前线程在线程块中的索引。这些索引被用于计算线程在整个矩阵中的位置。
if (row<rows&&col<cols){...}用于确保当前线程在矩阵的有效范围内。
int sum = 0用于初始化一个累加变量,保存叉乘的结果。
for (int i = 0; i<shared_dim; ++i){...}用于循环遍历矩阵的共享维度,执行叉乘运算。
sum += a[row * shared_dim + i] * b[i * cols + col]用于执行叉乘运算,并将结果累加到sum中。
c[row * cols + col] = sum用于将叉乘的结果存储到输出矩阵c中。
这个核函数通过并行计算输出矩阵的每个元素的方式,利用GPU的并行性来加速矩阵叉乘运算。
在本实施例中,还提出一种通过矩阵乘法操作计算矩阵点乘的可行实施方式,步骤S3031中通过GPU矩阵乘法操作,将所述第一中间矩阵与所述第二子矩阵进行矩阵点乘得到第二中间矩阵的步骤包括S30313~S30314。
步骤S30313,根据各个GPU线程所在线程块的块索引和在各自线程块中的线程索引,为各所述GPU线程分配点乘运算任务的第二任务数据,通过各所述GPU线程根据所分配的第二任务数据执行点乘运算获得第二计算结果,其中,所述第二任务数据来自于所述第一中间矩阵和所述第二子矩阵。
根据第一中间矩阵和第二子矩阵的大小,可以确定GPU线程的数量,并将各个GPU线程按照块划分,为各个GPU线程块给定块索引,以及为各个GPU线程给定其在线程块中的线程索引。根据各个GPU线程所在的线程块的块索引和在各自线程块中的线程索引,可以从第一子矩阵和系数矩阵中为各个GPU线程分配点乘运算任务的第二任务数据,然后即可利用各个GPU线程根据各自被分配的任务数据并行地执行点乘运算,以获得各自的第二计算结果。
为各个GPU线程分配点乘运算任务的任务数据的实现方式有很多种,在本实施方式中并不做限制。
步骤S30314,基于各所述GPU线程获得的第二计算结果获得所述第二中间矩阵。
将各个GPU线程获得的第二计算结果进行拼接即可得到最终的点乘结果,也即,得到将第一中间矩阵和第二子矩阵进行矩阵点乘得到的第二中间矩阵。
在一可行实施方式中,可以采用如下代码所展示的原理来实现通过GPU矩阵点乘操作计算第一子矩阵与系数矩阵的矩阵点乘,或第二中间矩阵和第三子矩阵的矩阵点乘。
__global__ void matrixElementWiseMultiply(const int* a, const int* b,int* c, int rows, int cols)
{
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row<rows&&col<cols)
{
c[row * cols + col] = a[row * cols + col]* b[row * cols + col];
}
}
其中,__global__是CUDA核函数的标记,表示这个函数在GPU上执行。
a和b是输入矩阵,也即需要进行矩阵叉乘的两个矩阵,可以是第一中间矩阵与第二子矩阵。c是输出矩阵,也即矩阵叉乘得到的结果。a、b、c在设备内存上。
rows是a、b、c的行数,cols是a、b、c的列数。
int row = blockIdx.y * blockDim.y + threadIdx.y用于计算当前线程的行索引。int col= blockIdx.x * blockDim.x + threadIdx.x用于计算当前线程的列索引。blockIdx.y和blockIdx.x是当前线程块的索引,blockDim.y和blockDim.x是线程块的维度,threadIdx.y和threadIdx.x是当前线程在线程块中的索引。这些索引被用于计算线程在整个矩阵中的位置。
if (row<rows&&col<cols){...}用于确保当前线程在矩阵的有效范围内。
c[row * cols + col] = a[row * cols + col]* b[row * cols + col]用于执行矩阵的点乘运算,并将结果存储到输出矩阵c中。这里的乘法是对应位置上的元素相乘。
这个核函数通过并行计算输出矩阵的每个元素的方式,利用GPU的并行性来加速矩阵点乘运算。
基于上述第一和/或第二实施例,提出本发明NTT算法的GPU加速方法第三实施例。在本实施例中,提出一种采用GPU多线程并行执行蝴蝶变换的可行实施方式,所述步骤S40中对所述系数序列采用GPU多线程并行执行NTT算法中的蝴蝶变换,以计算得到所述多项式的点值序列的步骤包括S401~S402。
步骤S401,针对各个GPU线程,分别执行与所述多项式的模数相对应的次数的循环;其中,在每一次循环中,根据所述GPU线程的线程编号和当前循环的轮次,为所述GPU线程分配蝴蝶变换操作任务的任务数据,通过所述GPU线程根据所分配的任务数据执行NTT算法中的蝴蝶变换获得计算结果,其中,所述任务数据来自于所述系数序列。
要想在GPU中实现NTT算法,需要考虑分配GPU多线程以实现高的利用率,为了获得最佳性能,应尽可能让每个线程负载相等。
在本实施方式中,根据多项式的模数大小,可以确定GPU线程的数量,并为各个GPU线程给定其对应的线程编号。为实现采用多线程并行执行NTT算法中三层循环的蝴蝶变换操作,可以针对各个GPU线程,分别执行与多项式的模数相对应的次数的循环,并在每一次循环过程中,根据GPU线程的线程编号和当前循环的轮次,从系数序列中分配蝴蝶变换操作任务的任务数据,然后通过该GPU线程根据所分配的任务数据执行蝴蝶变换,从而获得计算结果。
为各个GPU线程分配蝴蝶变换操作任务的任务数据的实现方式有很多种,在本实施例中并不做限制。
步骤S402,基于各所述GPU线程在各次循环中执行蝴蝶变换得到的计算结果,获得所述多项式的点值序列。
在各个GPU线程分别在各次循环中执行蝴蝶变换得到计算结果后,将各个计算结果进行拼接,即可获得多项式的点值序列。
在一可行实施方式中,可以采用如下流程来实现通过GPU多线程并行执行蝴蝶变换,以将多项式的系数序列转换为点值序列。
用n表示多项式的模数,对于length从1到n的循环,执行以下步骤:
在每次迭代中,length的值翻倍,表示多项式长度的增加。
计算tid,它表示GPU线程的全局索引,也即GPU线程的线程编号。
计算step=(n/length)/2,step用于控制多项式分组的步长。
计算psi_step=tid / step,psi_step是一个辅助变量,用于确定GPU线程在多项式中的位置。
计算target_idx=(psi_step * step* 2)+(tid mod step),target_idx是线程需要处理的多项式元素的索引,用于确定GPU线程要处理的多项式元素的位置。
计算step-group=length + psi_step,step-group用于确定GPU线程所在的多项式分组。
获取psi值,psi是一个与step-group相关的常量,也即单位根的幂,用于NTT算法中的数论变换步骤。
获取U值,U是多项式中target_idx处的元素。
获取V值,V是多项式中target_idx+step处的元素。
在此时,线程的任务分配已经完成。每个线程都有了它的任务数据,包括U、V和psi值。
等待每个线程完成其任务,确保每个线程都有了相应的值,结束循环。
此外,本发明实施例还提出一种NTT算法的GPU加速装置,参照图4,所述NTT算法的GPU加速装置包括:
获取模块10,用于获取NTT算法中多项式的系数序列和所述多项式的模数,以及获取基于所述多项式的模数分解得到的第一数值和第二数值;
第一计算模块20,用于用于根据所述第一数值和所述第二数值计算GPU显存的目标容量;
第二计算模块30,用于在所述多项式的模数小于预设阈值,且GPU的显存容量大于所述目标容量的情况下,根据所述第一数值和所述第二数值将NTT算法中的旋转因子矩阵拆分成多个子矩阵,对所述系数序列和所述子矩阵采用GPU矩阵乘法操作计算得到所述多项式的点值序列;
第三计算模块40,用于在所述多项式的模数大于或等于所述预设阈值,或所述显存容量小于或等于所述目标容量的情况下,对所述系数序列采用GPU多线程并行执行NTT算法中的蝴蝶变换,以计算得到所述多项式的点值序列。
在一可行实施方式中,所述第一计算模块20还用于:
计算所述第一数值的平方、所述第二数值的平方以及所述第一数值和所述第二数值的乘积,将三个计算结果相加后与预设位长和预设模数相乘,得到GPU显存的目标容量。
在一可行实施方式中,所述第二计算模块20还用于:
将所述旋转因子矩阵拆分成第一子矩阵、第二子矩阵和第三子矩阵,其中,所述第一子矩阵的维度根据所述第一数值确定,所述第二子矩阵的维度根据所述第一数值和所述第二数值确定,所述第三子矩阵的维度根据所述第二数值确定。
在一可行实施方式中,所述第二计算模块20还用于:
将所述系数序列转换为系数矩阵,所述系数矩阵的维度根据所述第一数值和所述第二数值确定;
通过GPU第一矩阵乘法操作对所述第一子矩阵与所述系数矩阵进行计算,得到第一中间矩阵,通过GPU第二矩阵乘法操作对所述第一中间矩阵与所述第二子矩阵进行计算,得到第二中间矩阵,通过GPU第一矩阵乘法操作对所述第二中间矩阵与所述第三子矩阵进行计算,得到第三中间矩阵;
根据所述多项式的模数对所述第三中间矩阵进行模运算,得到点值矩阵,将所述点值矩阵转换得到所述多项式的点值序列。
在一可行实施方式中,所述第二计算模块20还用于:
将所述第一子矩阵与所述系数矩阵进行GPU矩阵叉乘操作得到第一中间矩阵,将所述第一中间矩阵与所述第二子矩阵进行GPU矩阵点乘操作得到第二中间矩阵,将所述第二中间矩阵与所述第三子矩阵进行GPU矩阵叉乘操作得到第三中间矩阵。
在一可行实施方式中,所述第二计算模块20还用于:
根据各个GPU线程所在线程块的块索引和在各自线程块中的线程索引,为各所述GPU线程分配叉乘运算任务的第一任务数据和点乘运算任务的第二任务数据;
通过各所述GPU线程根据所分配的第一任务数据执行叉乘运算获得第一计算结果,其中,所述第一任务数据来自于所述第一子矩阵和所述系数矩阵;
通过各所述GPU线程根据所分配的第二任务数据执行点乘运算获得第二计算结果,其中,所述第二任务数据来自于所述第一中间矩阵和所述第二子矩阵;
基于各所述GPU线程获得的所述第一计算结果获得所述第一中间矩阵,基于各所述GPU线程获得的所述第二计算结果获得所述第二中间矩阵。
在一可行实施方式中,所述第三计算模块20还用于:
针对各个GPU线程,分别执行与所述多项式的模数相对应的次数的循环;
其中,在每一次循环中,根据所述GPU线程的线程编号和当前循环的轮次,为所述GPU线程分配蝴蝶变换操作任务的任务数据,通过所述GPU线程根据所分配的任务数据执行NTT算法中的蝴蝶变换获得计算结果,其中,所述任务数据来自于所述系数序列;
基于各所述GPU线程在各次循环中执行蝴蝶变换得到的计算结果,获得所述多项式的点值序列。
本发明NTT算法的GPU加速装置的具体实施方式的拓展内容与上述NTT算法的GPU加速方法各实施例基本相同,在此不做赘述。
此外,本发明实施例还提出一种计算机可读存储介质,所述存储介质上存储有NTT算法的GPU加速程序,所述NTT算法的GPU加速程序被处理器执行时实现如下所述的NTT算法的GPU加速方法的步骤。
本发明NTT算法的GPU加速设备和计算机可读存储介质各实施例,均可参照本发明NTT算法的GPU加速方法各个实施例,此处不再赘述。
需要说明的是,在本文中,术语“包括”、“包含”或者其任何其他变体意在涵盖非排他性的包含,从而使得包括一系列要素的过程、方法、物品或者装置不仅包括那些要素,而且还包括没有明确列出的其他要素,或者是还包括为这种过程、方法、物品或者装置所固有的要素。在没有更多限制的情况下,由语句“包括一个……”限定的要素,并不排除在包括该要素的过程、方法、物品或者装置中还存在另外的相同要素。
上述本发明实施例序号仅仅为了描述,不代表实施例的优劣。
通过以上的实施方式的描述,本领域的技术人员可以清楚地了解到上述实施例方法可借助软件加必需的通用硬件平台的方式来实现,当然也可以通过硬件,但很多情况下前者是更佳的实施方式。基于这样的理解,本发明的技术方案本质上或者说对现有技术做出贡献的部分可以以软件产品的形式体现出来,该计算机软件产品存储在一个存储介质(如ROM/RAM、磁碟、光盘)中,包括若干指令用以使得一台终端设备(可以是手机,计算机,服务器,空调器,或者网络设备等)执行本发明各个实施例所述的方法。
以上仅为本发明的优选实施例,并非因此限制本发明的专利范围,凡是利用本发明说明书及附图内容所作的等效结构或等效流程变换,或直接或间接运用在其他相关的技术领域,均同理包括在本发明的专利保护范围内。
Claims (10)
1.一种NTT算法的GPU加速方法,其特征在于,所述NTT算法的GPU加速方法包括以下步骤:
获取NTT算法中多项式的系数序列和所述多项式的模数,以及获取基于所述多项式的模数分解得到的第一数值和第二数值;
根据所述第一数值和所述第二数值计算GPU显存的目标容量;
在所述多项式的模数小于预设阈值,且GPU的显存容量大于所述目标容量的情况下,根据所述第一数值和所述第二数值将NTT算法中的旋转因子矩阵拆分成多个子矩阵,对所述系数序列和所述子矩阵采用GPU矩阵乘法操作计算得到所述多项式的点值序列;
在所述多项式的模数大于或等于所述预设阈值,或所述显存容量小于或等于所述目标容量的情况下,对所述系数序列采用GPU多线程并行执行NTT算法中的蝴蝶变换,以计算得到所述多项式的点值序列。
2.如权利要求1所述的NTT算法的GPU加速方法,其特征在于,所述根据所述第一数值和所述第二数值计算GPU显存的目标容量的步骤包括:
计算所述第一数值的平方、所述第二数值的平方以及所述第一数值和所述第二数值的乘积,将三个计算结果相加后与预设位长和预设模数相乘,得到GPU显存的目标容量。
3.如权利要求1所述的NTT算法的GPU加速方法,其特征在于,所述根据所述第一数值和所述第二数值将NTT算法中的旋转因子矩阵拆分成多个子矩阵的步骤包括:
将所述旋转因子矩阵拆分成第一子矩阵、第二子矩阵和第三子矩阵,其中,所述第一子矩阵的维度根据所述第一数值确定,所述第二子矩阵的维度根据所述第一数值和所述第二数值确定,所述第三子矩阵的维度根据所述第二数值确定。
4.如权利要求3所述的NTT算法的GPU加速方法,其特征在于,所述对所述系数序列和所述子矩阵采用GPU矩阵乘法操作计算得到所述多项式的点值序列的步骤包括:
将所述系数序列转换为系数矩阵,所述系数矩阵的维度根据所述第一数值和所述第二数值确定;
通过GPU第一矩阵乘法操作对所述第一子矩阵与所述系数矩阵进行计算,得到第一中间矩阵,通过GPU第二矩阵乘法操作对所述第一中间矩阵与所述第二子矩阵进行计算,得到第二中间矩阵,通过GPU第一矩阵乘法操作对所述第二中间矩阵与所述第三子矩阵进行计算,得到第三中间矩阵;
根据所述多项式的模数对所述第三中间矩阵进行模运算,得到点值矩阵,将所述点值矩阵转换得到所述多项式的点值序列。
5.如权利要求4所述的NTT算法的GPU加速方法,其特征在于,通过GPU第一矩阵乘法操作对所述第一子矩阵与所述系数矩阵进行计算,得到第一中间矩阵,通过GPU第二矩阵乘法操作对所述第一中间矩阵与所述第二子矩阵进行计算,得到第二中间矩阵,通过GPU第一矩阵乘法操作对所述第二中间矩阵与所述第三子矩阵进行计算,得到第三中间矩阵的步骤包括:
将所述第一子矩阵与所述系数矩阵进行GPU矩阵叉乘操作得到第一中间矩阵,将所述第一中间矩阵与所述第二子矩阵进行GPU矩阵点乘操作得到第二中间矩阵,将所述第二中间矩阵与所述第三子矩阵进行GPU矩阵叉乘操作得到第三中间矩阵。
6.如权利要求5所述的NTT算法的GPU加速方法,其特征在于,所述将所述第一子矩阵与所述系数矩阵进行GPU矩阵叉乘操作得到第一中间矩阵,将所述第一中间矩阵与所述第二子矩阵进行GPU矩阵点乘操作得到第二中间矩阵的步骤包括:
根据各个GPU线程所在线程块的块索引和在各自线程块中的线程索引,为各所述GPU线程分配叉乘运算任务的第一任务数据和点乘运算任务的第二任务数据;
通过各所述GPU线程根据所分配的第一任务数据执行叉乘运算获得第一计算结果,其中,所述第一任务数据来自于所述第一子矩阵和所述系数矩阵;
通过各所述GPU线程根据所分配的第二任务数据执行点乘运算获得第二计算结果,其中,所述第二任务数据来自于所述第一中间矩阵和所述第二子矩阵;
基于各所述GPU线程获得的所述第一计算结果获得所述第一中间矩阵,基于各所述GPU线程获得的所述第二计算结果获得所述第二中间矩阵。
7.如权利要求1至6中任一项所述的NTT算法的GPU加速方法,其特征在于,所述对所述系数序列采用GPU多线程并行执行NTT算法中的蝴蝶变换,以计算得到所述多项式的点值序列的步骤包括:
针对各个GPU线程,分别执行与所述多项式的模数相对应的次数的循环;
其中,在每一次循环中,根据所述GPU线程的线程编号和当前循环的轮次,为所述GPU线程分配蝴蝶变换操作任务的任务数据,通过所述GPU线程根据所分配的任务数据执行NTT算法中的蝴蝶变换获得计算结果,其中,所述任务数据来自于所述系数序列;
基于各所述GPU线程在各次循环中执行蝴蝶变换得到的计算结果,获得所述多项式的点值序列。
8.一种NTT算法的GPU加速装置,其特征在于,所述NTT算法的GPU加速装置包括:
获取模块,用于获取NTT算法中多项式的系数序列和所述多项式的模数,以及获取基于所述多项式的模数分解得到的第一数值和第二数值;
第一计算模块,用于根据所述第一数值和所述第二数值计算GPU显存的目标容量;
第二计算模块,用于在所述多项式的模数小于预设阈值,且GPU的显存容量大于所述目标容量的情况下,根据所述第一数值和所述第二数值将NTT算法中的旋转因子矩阵拆分成多个子矩阵,对所述系数序列和所述子矩阵采用GPU矩阵乘法操作计算得到所述多项式的点值序列;
第三计算模块,用于在所述多项式的模数大于或等于所述预设阈值,或所述显存容量小于或等于所述目标容量的情况下,对所述系数序列采用GPU多线程并行执行NTT算法中的蝴蝶变换,以计算得到所述多项式的点值序列。
9.一种NTT算法的GPU加速设备,其特征在于,所述NTT算法的GPU加速设备包括:存储器、处理器及存储在所述存储器上并可在所述处理器上运行的NTT算法的GPU加速程序,所述NTT算法的GPU加速程序被所述处理器执行时实现如权利要求1至7中任一项所述的NTT算法的GPU加速方法的步骤。
10.一种计算机可读存储介质,其特征在于,所述计算机可读存储介质上存储有NTT算法的GPU加速程序,所述NTT算法的GPU加速程序被处理器执行时实现如权利要求1至7中任一项所述的NTT算法的GPU加速方法的步骤。
Priority Applications (1)
Application Number | Priority Date | Filing Date | Title |
---|---|---|---|
CN202311814455.2A CN117473212B (zh) | 2023-12-27 | 2023-12-27 | Ntt算法的gpu加速方法、装置、设备及存储介质 |
Applications Claiming Priority (1)
Application Number | Priority Date | Filing Date | Title |
---|---|---|---|
CN202311814455.2A CN117473212B (zh) | 2023-12-27 | 2023-12-27 | Ntt算法的gpu加速方法、装置、设备及存储介质 |
Publications (2)
Publication Number | Publication Date |
---|---|
CN117473212A true CN117473212A (zh) | 2024-01-30 |
CN117473212B CN117473212B (zh) | 2024-04-16 |
Family
ID=89633324
Family Applications (1)
Application Number | Title | Priority Date | Filing Date |
---|---|---|---|
CN202311814455.2A Active CN117473212B (zh) | 2023-12-27 | 2023-12-27 | Ntt算法的gpu加速方法、装置、设备及存储介质 |
Country Status (1)
Country | Link |
---|---|
CN (1) | CN117473212B (zh) |
Citations (10)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
CN102609393A (zh) * | 2012-02-08 | 2012-07-25 | 浪潮(北京)电子信息产业有限公司 | 一种线性方程组的数据处理方法及装置 |
US20140002617A1 (en) * | 2012-06-27 | 2014-01-02 | The Board Of Trustees Of The University Of Illinois | Particle tracking system and method |
CN109753682A (zh) * | 2018-11-29 | 2019-05-14 | 浙江大学 | 一种基于gpu端的有限元刚度矩阵模拟方法 |
CN111078398A (zh) * | 2019-11-29 | 2020-04-28 | 苏州浪潮智能科技有限公司 | 一种gpu的分配方法、设备以及存储介质 |
CN113032007A (zh) * | 2019-12-24 | 2021-06-25 | 阿里巴巴集团控股有限公司 | 一种数据处理方法及装置 |
CN113541921A (zh) * | 2021-06-24 | 2021-10-22 | 电子科技大学 | 一种全同态加密gpu高性能实现方法 |
CN115393172A (zh) * | 2022-08-26 | 2022-11-25 | 无锡砺成智能装备有限公司 | 基于gpu实时提取光条纹中心的方法及设备 |
WO2022266920A1 (en) * | 2021-06-24 | 2022-12-29 | Intel Corporation | METHODS AND APPARATUS TO PERFORM MIXED RADIX FAST FOURIER TRANSFORM (FFT) CALCULATIONS ON GRAPHICS PROCESSING UNITS (GPUs) |
US20230163945A1 (en) * | 2021-11-24 | 2023-05-25 | Electronics And Telecommunications Research Institute | Method and apparatus for hardware-based accelerated arithmetic operation on homomorphically encrypted message |
CN116308989A (zh) * | 2022-12-09 | 2023-06-23 | 杭州后量子密码科技有限公司 | 一种全同态快速数论变换的gpu加速方法 |
-
2023
- 2023-12-27 CN CN202311814455.2A patent/CN117473212B/zh active Active
Patent Citations (10)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
CN102609393A (zh) * | 2012-02-08 | 2012-07-25 | 浪潮(北京)电子信息产业有限公司 | 一种线性方程组的数据处理方法及装置 |
US20140002617A1 (en) * | 2012-06-27 | 2014-01-02 | The Board Of Trustees Of The University Of Illinois | Particle tracking system and method |
CN109753682A (zh) * | 2018-11-29 | 2019-05-14 | 浙江大学 | 一种基于gpu端的有限元刚度矩阵模拟方法 |
CN111078398A (zh) * | 2019-11-29 | 2020-04-28 | 苏州浪潮智能科技有限公司 | 一种gpu的分配方法、设备以及存储介质 |
CN113032007A (zh) * | 2019-12-24 | 2021-06-25 | 阿里巴巴集团控股有限公司 | 一种数据处理方法及装置 |
CN113541921A (zh) * | 2021-06-24 | 2021-10-22 | 电子科技大学 | 一种全同态加密gpu高性能实现方法 |
WO2022266920A1 (en) * | 2021-06-24 | 2022-12-29 | Intel Corporation | METHODS AND APPARATUS TO PERFORM MIXED RADIX FAST FOURIER TRANSFORM (FFT) CALCULATIONS ON GRAPHICS PROCESSING UNITS (GPUs) |
US20230163945A1 (en) * | 2021-11-24 | 2023-05-25 | Electronics And Telecommunications Research Institute | Method and apparatus for hardware-based accelerated arithmetic operation on homomorphically encrypted message |
CN115393172A (zh) * | 2022-08-26 | 2022-11-25 | 无锡砺成智能装备有限公司 | 基于gpu实时提取光条纹中心的方法及设备 |
CN116308989A (zh) * | 2022-12-09 | 2023-06-23 | 杭州后量子密码科技有限公司 | 一种全同态快速数论变换的gpu加速方法 |
Non-Patent Citations (1)
Title |
---|
张岩;: "CPU-OpenMP和GPU-CUDA并行计算技术对矩阵乘法运算的加速效果分析", 《科技视界》, no. 26, 30 September 2017 (2017-09-30), pages 50 - 52 * |
Also Published As
Publication number | Publication date |
---|---|
CN117473212B (zh) | 2024-04-16 |
Similar Documents
Publication | Publication Date | Title |
---|---|---|
CN111062472B (zh) | 一种基于结构化剪枝的稀疏神经网络加速器及其加速方法 | |
US10691996B2 (en) | Hardware accelerator for compressed LSTM | |
CN108170639B (zh) | 基于分布式环境的张量cp分解实现方法 | |
CN108205702B (zh) | 一种多输入多输出矩阵卷积的并行处理方法 | |
CN106846235B (zh) | 一种利用NVIDIA Kepler GPU汇编指令加速的卷积优化方法及系统 | |
CN112200300B (zh) | 卷积神经网络运算方法及装置 | |
Peng et al. | GLU3. 0: Fast GPU-based parallel sparse LU factorization for circuit simulation | |
CN110390075B (zh) | 矩阵预处理方法、装置、终端及可读存储介质 | |
CN109885407B (zh) | 数据处理方法和装置、电子设备、存储介质 | |
CN112668708B (zh) | 一种提高数据利用率的卷积运算装置 | |
CN110377875B (zh) | 矩阵求逆方法、装置、设备及计算机可读存储介质 | |
CN112199636B (zh) | 适用于微处理器的快速卷积方法及装置 | |
CN111639701B (zh) | 一种图像特征提取的方法、系统、设备及可读存储介质 | |
CN110782009B (zh) | 基于ARMv8体系的计算内核优化方法 | |
CN114138231B (zh) | 执行矩阵乘法运算的方法、电路及soc | |
CN106682258B (zh) | 一种高层次综合工具中的多操作数加法优化方法及系统 | |
JP6357345B2 (ja) | ビデオデータ処理時に空間領域と周波数領域との間の変換を実行するためのデータ処理装置および方法 | |
CN118193914A (zh) | 面向分布式平台的lu分解方法、装置、设备及存储介质 | |
CN117473212B (zh) | Ntt算法的gpu加速方法、装置、设备及存储介质 | |
EP2538345A1 (en) | Fast fourier transform circuit | |
CN115130672B (zh) | 一种软硬件协同优化卷积神经网络计算的方法及装置 | |
CN113836481B (zh) | 矩阵计算电路、方法、电子设备及计算机可读存储介质 | |
CN111178505B (zh) | 卷积神经网络的加速方法和计算机可读存储介质 | |
CN116821576B (zh) | 用于基于risc-v加速n:m稀疏网络的方法和装置 | |
CN111915014B (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 |