基于GPU的密码S盒代数性质评估方法

2022-09-25 08:42蔡婧雯韦永壮刘争红
计算机应用 2022年9期
关键词:均匀度差分线程

蔡婧雯,韦永壮*,刘争红

(1.广西密码学与信息安全重点实验室(桂林电子科技大学),广西桂林 541004;2.广西无线宽带通信与信号处理重点实验室(桂林电子科技大学),广西桂林 541004)

0 引言

密码S 盒作为对称密码算法的核心部件,主要提供了必要的非线性变换,其代数性质往往决定着密码算法的安全强度。伴随着超级计算机其计算能力的迅速提升,特别是抵抗未来的量子计算攻击[1],高强度密码算法设计中对S 盒的输入及输出规模提出了新的要求,比如基于非线性反馈移位寄存 器(Nonlinear Feedback Shift Register,NFSR)或ARX(Addition-Rotation-XOR)操作部件等方法构造出16 比特或32 比特的大状态密码S 盒。美国国家标准与技术研究院(National Institute of Standards and Technology,NIST)发起轻量级密码算法公开征集[2],最终入围算法中SPARKLE[3]及GIFT-COFB(COmbined FeedBack)[4]等算法均采用了32 比特或者64 比特大状态密码S 盒。同年,在中国密码学会举办的全国密码算法设计竞赛[5]中,徐洪等[6]基于16 级NFSR 迭代构造了16 比特S 盒;田甜等[7]基于NFSR 设计了32 比特S 盒;2020 年美洲密码会上,Beierle 等[8]基于ARX 结构构造了64比特S 盒。注意到,密码S 盒的安全性与其安全性指标息息相关,传统安全性指标包括差分均匀度[9]、非线性度[10-11]、透明阶(Revised Transparency Order,RTO)[12]、飞来器连接表[13](Boomerang Connectivity Table,BCT)等。这些指标与相应的密码攻击密切相关,如差分均匀度、非线性度和透明阶分别刻画了S 盒抵抗差分密码分析[14]、线性密码分析[15]及差分功耗攻击(Differential Power Attack,DPA)[16]的能力。另一方面,对于n比特输入及n比特输出的密码S 盒,当n比较大时(如n>15 时)评估S 盒的各个安全性指标则较为困难,比如传统求解密码S 盒的差分均匀度、非线性度及透明阶时间复杂度分别约为O(23n)、O(23n)及O(23n·n2)。这些求解因搜索空间大,从而导致花销时间太长等问题。如何快速评估密码S 盒的代数性质是目前研究的热点之一。

为了解决计算资源瓶颈,图形处理器(Graphics Processing Unit,GPU)应运而生。GPU 主要应用于图像处理、视频音频处理、计算生物学等领域上。而利用GPU 解决密码学问题最早工作是Kedem 等[17]使用PixelFlow 图像引擎快速破解了UNIX 系统密钥;Manavski[18]利用计算统一设备架构(Compute Unified Device Architecture,CUDA)进行高级加密标准(Advanced Encryption Standard,AES)加速;Cheong等[19]在具有Kepler 架构上提出了加速分组密码算法国际数据加密算法(International Data Encryption Algorithm,IDEA),进一步提高加密吞吐量;Yeoh 等[20]提出了一种基于GPU 的分支定界算法。如何基于GPU 快速评估密码S 盒的安全强度仍有待进一步研究。

本文基于CPU-GPU 异构结构,对密码S 盒的差分均匀度、非线性度及透明阶提出求解优化方法,实现多线程并行计算,提出一种快速求解差分均匀度、非线性度及透明阶的方法。测试结果表明,与基于中央处理器(Central Processing Unit,CPU)实现相比,基于CPU-GPU 异构结构实现效率得到大幅度提升。本文方法利用单块GPU 分别计算差分均匀度、非线性度及透明阶所花销的时间与传统方法相比节省了90.28%、78.57%、60%。

1 预备知识

定义1设n比特输入、m比特输出的S 盒记为S(x)=(f1(x),f2(x),…,fm(x)):→,其中fi(x) 为→F2的布尔函数,记为密码S 盒的第i个分量函数(i=1,2,…,m)。本 文考虑n=m=16 的16 比特S 盒。

定义2差分均匀度[9]。设n比特输入、n比特输出的S盒记为S,对任意的输入差分α∈和输出差分β∈,其中差分对解的个数为:

则差分均匀度定义为:

当差分均匀度越小,S 盒的差分分布更均匀,安全性更好。

定义3非线性度[10-11]。一个n×n的S 盒的非线性度为S 盒的所有分量函数的非零线性组合中最小的非线性度,即:

定义4透明阶(RTO)[12]。对于任意n比特输入、n比特输出的S 盒,该S 盒的透明阶记为:

当透明阶越小时,抵抗差分功耗攻击的能力越强,安全性越好。

定义5S 盒的差分概率[21]。对于一个n比特输入、n比特输出的密码S 盒,对于输入差分α∈和输出差分β∈,存在差分概率

则称输入差分α经过S 盒后将以概率PS(α→β)得到输出差分β。

定义6S 盒的线性概率[21]。定义一个S 盒S:→,定义NS(θ,λ)=#{x∈:θ·x=λ·S(x)}构造密码S 盒的线性逼近表,其中,固定输入掩码θ∈,得到输出掩码λ的概率为:

即固定输入掩码θ,输出掩码λ,随机给定输入x,则θ·x=λ·S(x)以概率PLS(θ→λ)成立得到。

2 CUDA的并行计算

2.1 GPU与CUDA

随着人工智能、大数据等计算领域的不断发展,计算复杂度越来越大,图像处理器(GPU)从帮助CPU 做图像和图形运算转至海量数据处理上,涉及云计算、生物计算、天文学等多个领域,逐步成为了计算领域的研究热点。

由于CPU 与GPU 所应用的场景不同,两者的架构大不相同。从图1 中CPU 与GPU 的架构相比可知,CPU 与GPU 有如下区别:

图1 CPU与GPU的架构区别Fig.1 Difference between CPU and GPU architectures

1)GPU 采用了若干个算术逻辑单元(Arithmetic and Logic Unit,ALU)和超长的流水线,可以同时处理多个线程;然而CPU 拥有少量但很强大的算术逻辑单元,可以在很少的时间周期内完成运算。

2)CPU 具有强大的控制逻辑单元,在运算过程中提供逻辑预测能力降低延时;而GPU 控制能力稍逊色于CPU,运算过程中可以将多个访问合并成较少的访问。

3)CPU 有大量的缓存空间降低计算延时,而GPU 只有少量的缓存空间,与CPU 的缓存空间功能不同,GPU 的缓存空间将线程所需要访问的相同数据合并访问动态随机存取存储器。

从以上而可知,CPU 适合于需要较大的缓存空间且复杂控制逻辑的通信密集型运算;GPU 则适合于逻辑分支简单、计算量大的计算密集型运算。另外,从图2 可以得知,在计算密集型任务时,GPU 所需耗时占CPU 所需耗时的22%;而在计算通信密集的任务时,CPU 计算所需耗时较少,约为GPU 计算所需耗时的30%。在一些具有计算密集要求且逻辑分支较为简单的计算任务中,GPU 的处理能力比CPU 具有更大的优势。

图2 CPU与GPU任务耗时比较Fig.2 Task time consumption comparison of CPU and GPU

GPU 的造价和功耗与相同计算能力的CPU 相比,GPU 的造价和功耗相对较低。在计算领域中构建CPU 集群的超级计算机,造价昂贵。根据摩尔定律可以得知当CPU 计算速度达到一定程度时提升空间受限,GPU 的出现满足了需要计算大量数据而无法使用巨型计算机的用户需求。

目前应用较为广泛的GPU 并行编程平台有CUDA、OpenCL 等。2006 年NVIDIA 公司推出并行开发平台CUDA,支持C、C++、Java、Python 等多种主流编程语言,便于各领域进行并行计算操作。CUDA 使用了具有很强的并行计算特点的单指令多线程(Single Instruction Multiple Thread,SIMT)的执行模型,模型在执行过程中,构建CPU 与GPU 异构架构,其中:CPU 主要负责串行计算工作,完成较为复杂的逻辑控制及通信密集的运算;GPU 主要负责并行计算工作,完成运算量大且计算任务较为简单的计算密集型工作。经过测试及多方考量,本文选用CUDA 作为本文的并行开发平台,构造CPU-GPU 异构架构进行测试。

2.2 线程块的分配

为了发挥GPU 的最大并行计算效率,在执行内核函数过程中,需要合理配置线程块数量及每块线程块中的线程数量。线程块的数目由配置参数所划分的网格所决定,通常为32 的倍数最佳。通过测试可知,将线程块的线程数量为512时具有最大的计算能力。

2.3 并行计算的影响因素

总结前文所述的GPU 与CUDA 的计算特点,影响并行计算效率的因素可以总结为以下3 点:

1)减少CPU 与GPU 之间的数据传输。由于CPU 与GPU使用不同的内存空间,CPU 与GPU 在数据交换过程中需要通过计算机总线,造成了额外的时间花销,因此在利用GPU 进行并行计算时,应尽量避免减少CPU 与GPU 之间的数据传输。

2)减少访问GPU 全局内存。为了减少内存访问产生的时延和消耗,在GPU 中应尽量减少过多的跳跃式访问,最大限度减少因对GPU 内存访问而造成的延迟。

3)合理的资源配置。为了提高并行计算的效率,合理设置线程块内的线程数量,最大限度地利用线程处理计算任务,减少资源的浪费。

3 NBC算法

NBC 算法为中国密码学会举办的全国密码算法设计竞赛分组算法[11]第二轮入选算法之一,其采用广义Feistel 结构[12],算法加密共有三种模式,具体如表1 所示。

表1 NBC算法的三种模式Tab.1 Three modes of NBC algorithm

本文使用的算法是数据分组长度为128 比特、密钥长度为128 比特的NBC 算法。设第i轮的输入为Xi=输出为NBC-128/128 算法结构如图3 所示。

图3 八分支的1轮NBC-128/128结构Fig.3 One-round NBC-128/128 structure with 8 branches

NBC-128/128 算法的S 盒采用16 级NFSR来构造16 比特S 盒,S 盒构造图如图4 所示。设S 盒的16 比特的输入为S0S1…S15,当全体内部状态经过迭代20 轮后形成S 盒输出。

图4 NBC-128算法的S盒构造Fig.4 S-box structure of NBC-128 algorithm

算法设计者称构造出来的S 盒的差分均匀度Diff(S)=22,非线性度NL(S)=31 982,透明阶RTO=15.982 6。

4 基于GPU的16比特密码S盒代数性质评估

由于在CPU 下求解差分均匀度、非线性度及透明阶的算法效率较低,在本章中,将传统求解密码S 盒代数性质评估方法进行优化,分别讨论基于单GPU 模式和多GPU 模式下将内核函数切片至多线程中,实现多线程并行化计算。

4.1 单GPU对16比特密码S盒性质评估

根据共享式内存的结构特点和对S 盒性质评估的求解流程,本文提出了单块GPU 环境下的CPU-GPU 异构模式,并行架构如图5 所示。

图5 CPU-GPU异构并行流程Fig.5 CPU-GPU heterogeneous parallel flowchart

程序在运行时控制一块GPU,创建多个线程共同完成计算任务。具体步骤如下所示:

1)检测显卡设备。函数cudaSetDevice()表示检测主机设备的显卡个数,当检测到主机存在可使用的显卡时,将对算法进行CUDA 并行计算做好准备;

2)读取数据并复制入GPU。采用cudaMalloc()函数在设备Device 中开辟计算中所需要参数的空间。由于GPU 在计算过程中,无法直接读取CPU 内存中的数据,故在计算前需要在设备Decive 开辟相应的空间。

3)当Device 中开辟了相应的空间大小后,使用cudaMemcpy()函数将所需要的参数S 盒复制进入GPU 内。

4)计算内核函数。伪代码中存在3 个不同的内核函数,分别为differenceUniformity()、degreeOfNolinearity()及calculateRTO(),其中:differenceUniformity()为计算差分均匀度的内核函数;degreeOfNolinearity()为计算非线性度的内核函数;calculateRTO()为计算透明阶的内核函数。

5)在内核函数中,<<<Block,Thread>>>表示在启动内核函数时,分配Block个线程组,每个线程组中分配Thread个线程,故共分配Block*Thread线程总数。通过合理设置线程组和线程数量,才能更好地发挥GPU 的计算能力。本文使用的是一个线程处理一个分组,例如当处理100 组数据时,需要在GPU 内分配100 个线程,故本文计算16 比特S 盒的密码性质中,共需要处理65 536 个分组数据,使用了128 个线程块,其中每个线程块512 个线程。

6)检查并返回结果。当每一个线程完成了内核函数中的计算任务时,使用cudaGetLastError()函数检查内核函数在计算过程中是否存在错误:若存在错误,将错误返回至CPU中;若不存在,利用函数cudaMemcpy()将计算结果返回至CPU 中,计算结束。

求解复杂度分析如下:

1)由差分均匀度的定义可以得知:针对n比特输入、n比特输出的密码S 盒,传统求解差分均匀度需要遍历输入差分α∈、输出差分β∈及x∈三个变量,时间复杂度约为O(23n)。根据GPU 并行计算的特性,使用切片技术对求解差分均匀度的最外层循环分解到各个线程中并行,即除最外层循环外部分设为内核函数,此时求解差分均匀度的时间复杂度降低至O(22n)。为了进一步提高效率,减少计算逻辑分支数,再将遍历的输出差分β循环放置内核函数外,此在GPU 内计算的内核函数的时间复杂度将降低至O(2n)。

2)对求解非线性度及透明阶进行求解分析。传统求解非线性度及透明阶的时间复杂度为O(23n)、O(23n·n2)。利用相同的切片技术,将求解最外层循环分解到各个线程中,求解过程中利用线程索引对应最外层循环所遍历的值,此时求解非线性度及透明阶的时间复杂度降低至O(22n)、O(22n·n2)。另外再将一层循环放在内核函数外,最终GPU 内计算非线性度及透明阶的内核函数的时间复杂度将降低至O(2n)、O(2n·n2),与传统求解方法相比,该方法的时间复杂度降低了两个指数级,节省了求解时间花销。

算法1 测试主程序。

输入 S 盒;

输出 差分均匀度,非线性度,透明阶。

4.2 多GPU对大状态S盒性质评估

4.1 节分析了在CPU-GPU 异构计算结构下,对16 比特S盒安全性指标测评比在传统CPU 计算下所具有的时间优势,在相同的实验条件下,使用单块GPU 构建的CPU-GPU 异构计算比传统CPU 计算时间节省90.28%。但对于n比特输入、n比特输出的密码S 盒,当n比较大时(如n>15 时),由于计算搜索空间大,运算量大,单GPU 计算时间仍然较长,故提出在多GPU 环境下,对评估NBC 算法的16 比特S 盒的差分均匀度、非线性度等安全性指标方案并行化研究,对计算过程中涉及的数据传输过程进行研究与优化。分析并行化计算所遇到的瓶颈主要在数据传输过程,在结果保证正确性的基础上,调整程序的传输方式,由同步传输调整至异步传输,且利用多流技术与异步传输相结合逐步提高加速比。最后通过实现分析说明基于多GPU 架构下对大状态S 盒的安全性指标计算性能。

在使用多GPU 构架中,选择单个节点连接到高速串行计算机扩展总线标准(Peripheral Component Interconnect express,PCIe)总线上,具体架构如图6 所示。程序在运行时使用函数cudaSetDecice()对GPU 设备组上的各设备进行绑定,使得每个线程管理一个GPU,实现多个GPU 并行工作。与单GPU 结构相比,多GPU 结构可以开辟更多的线程,运算速度得到进一步提升。

图6 多GPU节点架构Fig.6 Multi-GPU node architecture

由于同步传输的并行化计算中,传输数据占用了大量的时间。本节利用多流技术与异步技术相结合,在计算过程中使计算过程与数据传输两个步骤进行重叠,从而减少一部分时间的开销。有无重叠优化的时间开销对比如图7 所示。在无重叠优化时,由于默认只有一个流队列,此时所有的计算过程皆为串行执行。先对数据传输至GPU 的全局内存内,传输完毕后再进行数据计算,等待GPU 内的所有线程计算完毕后再将结果复制回CPU 内。作为对比,在使用重叠技术进行时间优化后,当一个流队列在计算部分数据的同时,另一个流队列可以对剩下数据进行传输至GPU 内等待计算。当一个流上的数据计算完毕后,利用另一个流队列传输回CPU 内,下一个流队列等待数据传输。计算与传输时间重叠技术的优化既能保持计算任务仍按照串行执行,又能掩盖GPU 与CPU 数据传输之间所带来的大量时间开销,从而进一步减少程序所需要的执行时间,提高并行效率。

图7 有无重叠优化的时间开销对比Fig.7 Comparison of time cost with and without overlapping optimization

利用多GPU 对大状态S 盒进行评估过程具体如下:

1)在CPU 端获取已有的GPU 设备数量和每个GPU 设备信息。利用CUDA 中自带的函数cudaGetDeviceCount(&ngpus),读取已有的GPU 设备数量,并将GPU 数量信息存储在变量ngpus中,可通过设备号dev进行选择使用GPU设备。

2)在同一节点上的GPU 设备构成GPU 设备组,GPU 设备组内的GPU 设备直接进行通信和数据传输。

3)在CPU 端进一步准备计算所需要的数据集,根据GPU设备组的数量,将数据平分至各GPU 设备上,另外在CPU 端设置S 盒,在CPU 端将S 盒以结构体的形式传输至GPU 设备组中对应的常量存储区中,S 盒参数都在对应的GPU 设备运行过程中将会被核函数多次调用。

4)在CPU 端设置循环遍历所有的GPU 设备,将GPU 设备组中的GPU 分别置于对应并行流上,通过对工作流在不同时间下的操作和阻塞实现GPU 设备的异步,设置CUDA 工作流Steam 的异步操作隐藏了部分访问延迟和实现了任务的并发执行,减少了数据处理时间。

5)明确每个核函数分配的变量和变量空间,利用函数cudaMemory()将数据以异步方式传输至GPU 设备组中对应的GPU 上。

6)为了确保核函数运行时有较好的性能,延用上一节对GPU 的线程数分配,使用了128 个线程组,其中每个线程组512 个线程,共计65 536 个线程数。

7)核函数完成线程配置后,数据根据“分而治之”思想,将输入数据划分成多个子集分别复制。由于每个问题都是独立的,所以分别安排在不同的并行流中进行计算,不同的流之间输出传输于另一个流的核计算进行重叠。

8)当线程中循环遍历完所有的块,完成内核函数Kernel的计算后,利用重叠流的思想保证每个线程计算后优先传输至CPU 内。

9)当所有线程都计算完毕后,CPU 端对GPU 设备组的各GPU 设备返回的结果进行统一归总,并按照规定的格式进行输出。

算法2 多GPU 测试主程序。

输入 S 盒;

输出 差分均匀度,非线性度,透明阶。

5 测试与结果分析

5.1 测试环境

本文实验环境所使用的CPU 为Intel Xeon Silver 4210 2.20 GHz;GPU 为NVIDIA Quadro RTX 8000;在多GPU 环境下,共使用4 块相同型号的GPU,且显卡型号为NVIDIA Quadro RTX 8000;操作系统为Ubuntu 18.04.4 LTS,64 bits;编程环境为CUDA 7.0、GCC 7.5.0。本实验的CPU 代码用的C 语言进行编写,GPU 代码用CUDA C 进行编写。

5.2 测试结果

本次测试是针对NBC-128/128 算法的16 比特S 盒分别进行差分均匀度、非线性度和透明阶运算,其中测试可得NBC 算法的差分均匀度为Diff(S)=22,非线性度为NL=31 982,透明阶RTO=15.982 6,运行时间如图8 所示。

图8 对比CPU、单块GPU及多块GPU下的运行时间Fig.8 Comparison of running time under CPU,single GPU and multi-GPU

通过以上实验结果表明,在相同的实验条件下,使用GPU 测试16 比特S 盒差分均匀度所用时间比在CPU 测试16比特S 盒所用时间约减少90.28%;使用GPU 测试16 比特S盒的非线性度所用时间比在CPU 测试16 比特S 盒所用时间约减少78.57%;使用GPU 测试16 比特S 盒透明阶所用时间比在CPU 测试16 比特S 盒所用时间约减少60%。实验结果证明使用GPU 测试大比特S 盒性质所消耗时间明显少于使用CPU 测试大比特S 盒性质所用的时间。在使用多GPU 并行计算的架构下,在相同实验条件下,使用多GPU 测试差分均匀度所用时间比单GPU 测试所用时间约减少99.52%;使用多GPU 测试非线性度所用时间比单GPU 测试所用时间约减少91.67%;使用多GPU 测试透明阶所用时间比单GPU 测试所用时间约减少78.13%,使用多块GPU 并行计算的计算速率明显高于单块GPU 计算速率。

通过密码S 盒的差分概率定义可知,当输入尺寸n比较大时(如n>15 时),需要遍历输入差分α∈、输出差分β∈及x∈三个变量,所以求解差分概率所需的时间复杂度约为O(23n)。类似地,当求解线性概率时,同样需要遍历3 个变量,分别是输入掩码θ、输出掩码λ及输入x,即时间复杂度也约为O(23n)。注意到,利用切片技术对差分概率及线性概率的计算过程可以分解到各个线程中并行计算。因而,求解差分概率及线性概率与求解差分均匀度方法类似,预计所花销的时间大致相当,限于篇幅,本文不再赘述。

6 结语

本文基于CPU-GPU 结构,结合差分均匀度、非线性度等计算特征,将内核函数利用切片技术拆分至多线程上,实现多线程并行计算,并由此提出快速评估密码S 盒代数性质新方法。在单块GPU 及4 块GPU 环境下对NBC-128/128 密码算法的S 盒进行差分均匀度、非线性度及透明阶3 个性质计算,实验结果证实:与基于CPU 的实现环境相比,基于单块GPU 所构建的CPU-GPU 架构的实现效率得到了显著的提升,即计算差分均匀度、非线性度及透明阶分别节省了90.28%、78.57%、60%的时间。下一步的研究工作可以考虑针对32 比特、64 比特等大状态的密码S 盒,基于CPU-GPU 结构进行安全性评估。

猜你喜欢
均匀度差分线程
5G终端模拟系统随机接入过程的设计与实现
一类分数阶q-差分方程正解的存在性与不存在性(英文)
实时操作系统RT⁃Thread启动流程剖析
实时操作系统mbedOS 互斥量调度机制剖析
一个求非线性差分方程所有多项式解的算法(英)
一类caputo分数阶差分方程依赖于参数的正解存在和不存在性
蛋鸡育成期的饲养管理要点
喷头高度对防火林带喷淋效果的影响研究
基于差分隐私的数据匿名化隐私保护方法
PTT纤维纱线生产及在毛精纺面料中的应用