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

摘要

本发明公开了一种基于OpenCL的AES并行化实现方法，根据这个方案，能使AES在基于OpenCL的AMD GPU上并行运行时获得最佳性能。该方法包括如下步骤：S1、确定明/密文及轮密钥分组的数量，准备好明/密文及轮密钥的数据；S2、确定AES的执行模式，为编写内核函数做准备；S3、编写内核函数；S4、编写OpenCL程序主文件，为OpenCL程序的执行做准备；S5、设置程序运行参数，分配内存空间，运行程序；S6、获取加解密结果，释放资源。本发明主要通过对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、该函数有四个参数，分别是输入数据的地址、输出数据的地址、轮 密钥的地址、轮数，其中，输入和输出数据的地址类型均为无符号字符向量全 局指针global uchar4*，轮密钥的地址类型为无符号字符向量常量指针 constant uchar4*，轮数的类型为整型；该内核函数的参数可以根据实际的需 求增加或减少；

S32、存放加/解密结果的中间变量设为私有无符号字符向量private uchar4类型，其中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表、轮密钥、明密文及加/解密 过程中需要用到的中间变量等）均是在理论分析和实际测试相结合后得出的最 佳搭配，能确保实际实现时能发挥出AMD GPU的最佳性能。

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 8*1024*1024 5.04

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