CN103973431B

CN103973431B - 一种基于OpenCL的AES并行化实现方法

Info

Publication number: CN103973431B
Application number: CN201410153285.2A
Authority: CN
Inventors: 龚征; 袁宇恒; 何振忠; 温雅敏
Original assignee: South China Normal University
Current assignee: South China Normal University
Priority date: 2014-04-16
Filing date: 2014-04-16
Publication date: 2017-04-05
Anticipated expiration: 2034-04-16
Also published as: CN103973431A

Abstract

本发明公开了一种基于OpenCL的AES并行化实现方法，根据这个方案，能使AES在基于OpenCL的AMD GPU上并行运行时获得最佳性能。该方法包括如下步骤：S1、确定明/密文及轮密钥分组的数量，准备好明/密文及轮密钥的数据；S2、确定AES的执行模式，为编写内核函数做准备；S3、编写内核函数；S4、编写OpenCL程序主文件，为OpenCL程序的执行做准备；S5、设置程序运行参数，分配内存空间，运行程序；S6、获取加解密结果，释放资源。本发明主要通过对AES并行运行时，数据在内存中的合理分配以及并行粒度的合理选择来提升运行性能，可用于快速加解密或口令破解机。

Description

一种基于OpenCL的AES并行化实现方法

技术领域

本发明涉及密码算法的技术领域，特别涉及基于OpenCL的AMD GPU上的AES的快速实现方法。

背景技术

GPU原本是专门为处理图形图像数据而设计的，因而它具有高度并行的结构。现今，经过十多年的发展，GPU的并行运算能力已经远远超越了CPU，利用GPU对各种计算进行并行加速也成为了当今的研究热点。而在信息安全方面，GPU的一项重要应用，就是通过并行化来实现快速加解密。

事实上，现在实现加解密并行化有CPU、GPU、FPGA和OPENCL四种并行实现机制，它们也各有优劣。CPU单个核心的性能虽然比GPU的要高，但是不具备高度的并行化结构，并行能力不如GPU，而且受到结构限制，并行规模扩充比FPGA难；FPGA的并行规模扩充虽然比CPU和GPU都容易，然而硬件一旦升级换代，就需要改动大量的代码，不利于开发与维护；GPU则同样受到结构限制，并行规模扩充比FPGA难，单个核心的运算能力也不及CPU，显存与主机内存之间的数据交换会造成大量的延时，这种I/O的耗费成了GPU并行性能主要的瓶颈；OpenCL则是专为异构平台设计的，能够综合利用CPU、GPU及其他类型的处理器来提供并行计算，因而也日渐受到青睐。

OpenCL全称Open Computing Language（开放计算语言），是由Khronos Group维护的为异构平台提供编写程序（尤其是并行程序）的开放的框架标准。OpenCL由编写内核程序的语言和定义并控制平台的API两部分组成，可以在多核CPU或者GPU上编译运行。通过使用OpenCL，软件开发人员便能够高效利用各种异构处理平台，从高性能计算服务器，到家用计算机再到手持设备，都被OpenCL所支持，并且在OpenCL帮助下，能够组合工作。

发明内容

本发明的目的在于克服现有技术的缺点与不足，提供一种基于OpenCL的AES并行化实现方法。

本发明的目的通过下述技术方案实现：

一种基于OpenCL的AES并行化实现方法，包括下述步骤：

S1、确定明/密文及轮密钥分组的数量，准备好明/密文及轮密钥数据的数据；

S2、确定AES的执行模式，为编写内核函数做准备；

S3、编写内核函数；

S4、编写OpenCL程序主文件，为OpenCL程序的执行做准备；

S5、设置程序运行参数，分配内存空间，运行程序；

S6、获取加解密结果，释放资源。

优选的，步骤S1中，轮密钥的数据在主机端预先生成好。

优选的，步骤S2中，根据AES的设计，将每一轮的不同操作转化为仅用T表和异或来实现；其具体实现方式是：

假设a为一轮的输入，长度为128比特，此处看成是4×4字节的矩阵，按照32比特划分为4份，即a0，a1，a2，a3，1份相当于原有矩阵中的1行，如a_0,j代表原有矩阵中的第一行第j列的元素，d为一轮的输出，那么，d可以表示为：

d_j=T₀[a_0,j]⊕T₁[a_1,j+1]⊕T₂[a_2,j+2]⊕T₃[a_3,j+3]⊕k_j；

其中T0,T1,T2,T3为4个T表，kj为第j列要与之异或的子轮密钥；经过上述处理后，原本需要经过字节替换、行移位、列混合以及轮密钥加四个操作才能获取的一轮的输出，现在就转化为仅依靠查表和异或就能获取到；在完成上述操作前需要预先准备好T表并将其写在内核文件中；

T表的类型可设为4字节的静态无符号整型私有值、4字节的无符号整型常量或静态无符号整型常量类型，需要注意，T表应写在内核函数之外，因为内核函数之内不允许使用静态类型，而且将T表写在内核函数中会导致程序运行时性能的下降。

优选的，步骤S2中，AES的执行模式选用电子密码本ECB模式、计时器CTR模式以及输出反馈OFB模式三种执行模式中的一种。

优选的，步骤S3中，编写内核参数的具体方法为：

S31、该函数有四个参数，分别是输入数据的地址、输出数据的地址、轮密钥的地址、轮数，其中，输入和输出数据的地址类型均为无符号字符向量全局指针globaluchar4*，轮密钥的地址类型为无符号字符向量常量指针constant uchar4*，轮数的类型为整型；该内核函数的参数可以根据实际的需求增加或减少；

S32、存放加/解密结果的中间变量设为私有无符号字符向量privateuchar4类型，其中uchar4为OpenCL中特有的类型，用来表示向量，类型后的数字表示该向量包含多少个值，uchar4就是包含4个无符号字符类型的向量；

S33、由于程序实际运行时，一个工作节点单独负责一个内核程序的执行，然而各个工作节点的输入输出数据也是独立的，所以编写内核函数的时候，应确保当前执行的工作节点访问到正确的数据；可以通过使用get_global_id()函数来标识当前工作节点在工作空间中的位置，利用该位置来限定各个工作节点所访问数据的位置；

S34、根据上述选择的AES的执行模式，结合上述的对AES进行优化的方法，将AES的加/解密算法写进内核函数，待加/解密的数据应利用输入数据地址从输入数据的空间之中读取，加/解密完成后，利用输出数据的地址将加/解密的结果写入到输出数据的空间之中。

优选的，步骤S4的具体步骤为：

S41、根据输入数据分组和轮密钥分组的数量设置参数global_work_size以及local_work_size的值，global_work_size和local_work_size分别指定执行内核时工作空间中每个维度工作节点的数量和工作组中每个维度的工作节点的数量，工作组的维度必须和整个工作空间的维度相同，并且工作组中的工作节点数量不应超过输入数据分组的数量，所有工作节点都将执行相同的内核程序；global_work_size的值根据输入数据以及轮密钥分组的数量来确定，local_work_size的值根据显卡可支持的最大工作节点数量来设置；也可只指定global_work_size的值，程序执行的时候，会自动划分一个工作组包含的工作节点的个数；

S42、配置OpenCL的运行环境，包括以下步骤：获取可用的计算平台，选择使用AMD的平台，获取AMD平台下的设备列表，选择GPU作为计算设备并创建上下文环境，创建内核程序对象并进行编译、创建命令队列。

优选的，步骤S5具体为：

S51、为输入、输出、轮密钥等数据创建内存空间，并将这些内存空间的地址告知内核程序，使内核函数能知道从哪里能取出其运行时所需要的参数；

S52、将待加密的明文或待解密的密文以及相应使用到的轮密钥复制到显存中；

S53、最后将内核程序放入命令队列中执行，由于之前已经指定了执行该内核程序的工作节点的数量，OpenCL会将内核程序交由各个工作节点独立执行，所有工作节点均是并行地运作的，所以工作节点的规模可以看成是加/解密执行时的并行规模。

优选的，步骤S6中，各个工作节点执行内核程序后的结果会存放在先前分配的输出数据的内存空间中，但这时这些数据还在显存当中，因而，等待命令队列执行完毕后，需要将程序运行的结果从显存中读取出来，结果读取完毕后，需要释放程序中用到的各种资源。

本发明相对于现有技术具有如下的优点及效果：

1、本发明根据OpenCL与AES算法轮函数的特点，给出了基于OpenCL的AES高速并行化的编程实现；

2、本发明在实现AES并行化时采用的并行粒度（即1个工作节点负责1个AES分组的处理）以及数据分配方案（如T表、轮密钥、明密文及加/解密过程中需要用到的中间变量等）均是在理论分析和实际测试相结合后得出的最佳搭配，能确保实际实现时能发挥出AMDGPU的最佳性能。

3、本发明还具有易于实现、维护，性价比高的特点。

4、本发明提供的并行化快速加解密实现方案还有着极大的实用价值。它能广泛应用于云存储、服务器大型计算的数据安全性保护上。另外，由于目前大多数应用的口令或数据保护均是采用AES算法，本发明的实现能大大提高破解效率，为国家相关部门提供安全工具。

附图说明

图1本发明电子密码本加密模式；

图2本发明AES加密流程图；

图3本发明的流程图。

具体实施方式

下面结合实施例及附图对本发明作进一步详细的描述，但本发明的实施方式不限于此。

实施例

本实施例基于OpenCL给出了一种AES并行化加密的实现方法。根据该实施例的步骤，稍加修改也可用于并行化解密或并行化加解密的实现。

如图3所示，本发明基于OpenCL的AES并行化实现方法，包括下述步骤：

S1、确定明/密文及轮密钥分组的数量，准备好明/密文及轮密钥的数据；

S2、确定AES的执行模式，为编写内核函数做准备；

S3、编写内核函数；

S4、编写OpenCL程序主文件，为OpenCL程序的执行做准备；

S5、设置程序运行参数，分配内存空间，运行程序；

S6、获取加解密结果，释放资源。

下面结合具体实际的操作方式，对上述的步骤S1-S6做进一步的说明：

1、确定明（密）文及轮密钥分组的数量，准备好相应的数据。

本实施例中，待加密的明文分组个数为1048576（即1024*1024个分组，其中，一个分组为128比特），轮密钥分组个数为1。待加密的明文及加密需要使用的轮密钥数据均在主机（Host）端预先准备好。

2、确定AES的执行模式，为编写内核函数做准备。

本实施例中，AES的执行模式为电子密码本（ECB）模式，见图1。AES的算法将128比特的数据看成4×4字节的矩阵，由10轮操作组成（本实施例使用的密钥长度为128位，因而轮数为10轮），每一轮都包含四个操作：字节替换(SubBytes)、行移位(ShiftRows)、列混合(MixColumns)和轮密钥加(AddRoundKey)。当然，最后一轮的列混合操作是可选的，由用户自己决定是否执行，缺省是不执行的。加密流程如图2所示。为了提高并行实现的效率，根据它的设计，可以将每一轮的不同操作转化为仅用查表和异或来实现。所查的表称为T表（T-Box）。假设a为一轮的输入，长度为128比特，按照32比特划分为4份，即a0，a1，a2，a3，一份即为原有矩阵中一行的数据，如a_0,j表示原矩阵中第一行第j列的元素，d为一轮的输出，那么，d可以表示为：

d_j=T₀[a_0,j]⊕T₁[a_1,j+1]⊕T₂[a_2,j+2]⊕T₃[a_3,j+3]⊕k_j

其中T0,T1,T2,T3为4个T表，kj为第j列要与之异或的子轮密钥。

本实施例中，T表预先计算好并写在内核文件当中，用于存储T表的类型为4字节的静态无符号整型常量（static constant u32）。

3、编写内核函数

本实施例中，内核函数参数有四个，分别是存放明文数据的数组的地址（输入数据）、存放加密结果的数组的地址（输出数据）、存放轮密钥数据的数组的地址以及加密轮数。

为了确保工作节点能访问到正确的数据，这里使用两个整型变量：idx和idy，来记录工作节点在工作空间中的位置，idx标识所在行，idy标识所在列。idx和idy的范围由主程序文件中的global_work_size参数来控制。存放加密流程产生的中间结果的变量均为私有无符号字符向量（private uchar4）。get_global_id(x)函数可以获取当前执行的工作节点在工作空间中指定维度x上的索引位置。那么，

int idx=get_global_id(0);

int idy=get_global_id(1);

根据idx和idy的值就可以确定每个工作节点加密的数据以及用于加密的轮密钥数据的位置。需要注意的是，在本实施例中，一个工作节点负责一个明文分组的加密，但一个明文分组128比特，而用于存放明文数据的数组一个元素只有32比特（存放明文数据的数组是一个uchar4向量数组，uchar4向量包含4个无符号字符（uchar），其中，一个无符号字符占8比特，4个就是32比特），所以需要4个元素才能表示一个明文分组。我们用两个整型变量index_M和index_K来记录明文及轮密钥数据的索引，整型变量nrounds记录当前加密所需要的轮数，那么

int index_M=idx*4;

int index_K=idy*(nrounds+1)*4;

这样，index_M至index_M+3就为当前工作节点要处理的明文分组在存储明文数据的数组里的索引范围，而index_K至index_K+43就为当前工作节点要使用到的扩展密钥在存储轮密钥数据的数组中索引的范围。（在本实施例中，密钥长度为128比特，因而轮数为10轮，由图2可知，对于128比特明文分组，与轮密钥异或的次数为11次，而具体实现时一个明文分组又拆为4个32比特的分组，轮密钥亦同样要拆成4个32比特的分组才能对应与之进行异或，因而用于进行异或操作的轮密钥数据的索引为4*11=44个。）

4、编写OpenCL程序主文件，为OpenCL程序的执行做准备。

首先要设置global_work_size以及local_work_size两个参数的值：

本实施例中，明文分组大小为1048576，轮密钥分组大小为1，GPU支持的一个维度上工作节点的数量上限为256，所以：

size_t global_work_size[2]={1048576,1}

size_t local_work_size[2]={256,1}

也就是，执行该内核程序的工作空间为1维，1维中有1048576工作节点，其中每256个工作节点组成一个工作组。

然后要配置OpenCL的运行环境：

先使用clGetPlatformIDs和clGetPlatformInfo两个命令来获取可用的计算平台，并选择AMD环境作为计算平台；然后使用clCreateContextFromType命令生成上下文环境，命令的第二个参数设为CL_DEVICE_TYPE_GPU表示使用GPU作为计算设备；再使用clGetContextInfo命令获取计算设备的信息，检查所选设备是否正确；接着使用clCreateProgramWithSource命令创建程序对象，使用clBuildProgram编译程序对象，确保程序对象能被正常执行；然后使用clCreateKernel命令创建内核对象；最后使用clCreateCommandQueue命令创建命令队列。

5、设置程序运行参数，分配内存空间，运行程序

首先使用clCreateBuffer命令为输入数据、输出数据及轮密钥数据创建内存空间；然后使用clSetKernelArg命令为内核程序设置相应的参数，需要设置的参数个数与内核函数的参数个数一致；再使用clEnqueueWriteBuffer命令将待加密的明文以及要使用到的轮密钥复制到显存中；最后使用clEnqueueNDRangeKernel命令将内核程序放入命令队列中执行。

6、获取加密结果，释放资源

首先使用clFinish命令等待命令队列执行完毕；然后使用clEnqueueReadBuffer命令从显存中读取加密结果；最后使用clRelease的一系列命令释放资源。

实验结果

本实例运行环境为：CPU型号为Core i33120，内存6G，操作系统为Win7（64位），GPU型号为ATI Mobility Radeon HD7670m，显存1G，所使用的SDK版本为AMD APP SDKv2.9。

本实例在不同输入数据大小下所获得的最高吞吐率（不包含I/O耗费）：

输入数据大小（MB）	明文分组个数	吞吐率（Gbps）
			1	64*1024	4.51
4	256*1024	4.89
			8	512*1024	4.98
16	1024*1024	5.03
			128	810241024	5.04

上述实施例为本发明较佳的实施方式，但本发明的实施方式并不受上述实施例的限制，其他的任何未背离本发明的精神实质与原理下所作的改变、修饰、替代、组合、简化，均应为等效的置换方式，都包含在本发明的保护范围之内。

Claims

1.一种基于OpenCL的AES并行化实现方法，其特征在于，包括下述步骤：

S2、确定AES的执行模式，为编写内核函数做准备，根据AES的设计，将每一轮的不同操作转化为仅用T表和异或来实现，其具体实现方式是：

其中T0,T1,T2,T3为4个T表，k_j为第j列要与之异或的子轮密钥；经过上述T表异或处理后，原本需要经过字节替换、行移位、列混合以及轮密钥加四个操作才能获取的一轮的输出，现在就转化为仅依靠查表和异或就能获取到；在完成上述T表异或操作前需要预先准备好T表并将其写在内核文件中；

T表的类型可设为4字节的静态无符号整型私有值、4字节的无符号整型常量或静态无符号整型常量类型，需要注意，T表应写在内核函数之外，因为内核函数之内不允许使用静态类型，而且将T表写在内核函数中会导致程序运行时性能的下降；

S3、编写内核函数；

S4、编写OpenCL程序主文件，为OpenCL程序的执行做准备；

S5、设置程序运行参数，分配内存空间，运行程序；

S6、获取加解密结果，释放资源。

2.根据权利要求1所述的基于OpenCL的AES并行化实现方法，其特征在于，步骤S1中，轮密钥的数据在主机端预先生成好。

3.根据权利要求1所述的基于OpenCL的AES并行化实现方法，其特征在于，步骤S2中，AES的执行模式选用电子密码本ECB模式、计时器CTR模式以及输出反馈OFB模式三种模式中的一种。

4.根据权利要求1所述的基于OpenCL的AES并行化实现方法，其特征在于，步骤S3中，编写内核函数的具体方法为：

S31、该函数有四个参数，分别是输入数据的地址、输出数据的地址、轮密钥的地址、轮数，其中，输入和输出数据的地址类型均为无符号字符向量全局指针global uchar4*，轮密钥的地址类型为无符号字符向量常量指针constant uchar4*，轮数的类型为整型；该内核函数的参数可根据实际的需求增加或减少；

S32、存放加/解密结果的中间变量设为私有无符号字符向量private uchar4类型，其中uchar4为OpenCL中特有的类型，用来表示向量，类型后的数字表示该向量包含多少个值，uchar4就是包含4个无符号字符类型的向量；

S33、由于程序实际运行时，一个工作节点单独负责一个内核程序的执行，然而各个工作节点的输入输出数据也是独立的，所以编写内核函数的时候，应确保当前执行的工作节点访问到正确的数据；可通过使用get_global_id()函数来标识当前工作节点在工作空间中的位置，利用该位置来限定各个工作节点所访问数据的位置；

S34、根据上述步骤S2确定的AES的执行模式，结合上述步骤S31-S33中对AES进行优化的方法，将AES的加/解密算法写进内核函数，待加/解密的数据应利用输入数据地址从输入数据的空间之中读取，加/解密完成后，利用输出数据的地址将加/解密的结果写入到输出数据的空间之中。

5.根据权利要求1所述的基于OpenCL的AES并行化实现方法，其特征在于，步骤S4的具体步骤为：

6.根据权利要求1所述的基于OpenCL的AES并行化实现方法，其特征在于，步骤S5具体为：

7.根据权利要求1所述的基于OpenCL的AES并行化实现方法，其特征在于，步骤S6中，各个工作节点执行内核程序后的结果会存放在先前分配的输出数据的内存空间中，但这时这些数据还在显存当中，因而，等待命令队列执行完毕后，需要将程序运行的结果从显存中读取出来，结果读取完毕后，需要释放程序中用到的各种资源。