基于GPU并行枚举DES密钥Bitslice方法的实现
2019-12-11
(石化盈科信息技术有限责任公司,北京 100007)
1 研究背景
统一计算设备架构CUDA(Compute Unified Device Architecture),它是NVIDIA公司提出的一种通用并行计算平台和编程模型。通过使用GPU高性能并行多线程能力,解决大数据计算的模型已被广泛应用在密码破译[1-5]、模型分析[6-9]、图形图像[10-13]等领域。通过使用CUDA开发平台可以开发在CPU和GPU上运行的通用计算程序。并行编程的中心思想是分组数据:将大数据划分为一些小块数据组,再把这些小块数据组交给相应的处理单元并行处理。CUDA中的问题被粗粒度地划分为若干子问题(块Block),每个块都由一些线程组成,线程是CUDA中最小的处理单元。子问题进一步划分为若干更小的细粒度的线程问题以解决。CUDA线程数目通常能达到数千个甚至上万个,通过不同的优化方法,划分模型可以数十倍地提升计算性能。
GPU是由多个流处理器构成的,流处理器以块(Block)为基本调度单元,因此,对于流处理器较多的GPU,它一次可以处理的块(Block)更多,从而运算速度更快,时间更短。反之对于流处理器较少的GPU,其运算速度便会较慢[14]。这一原理可以通过图1看出来:
图1 GPU流处理器处理性能的比较
对于浮点数操作能力,CPU与GPU的能力相差在GPU更适用于计算强度高、流程简单的并行计算中。GPU拥有更多晶体管,而不是像CPU一样具有数据Cache和流控制器。这样的设计是因为并行计算的时候每个数据单元执行相同程序,不需要那么繁琐的流程控制,而更需要高计算能力,从而也不需要大Cache来缓存数据。CPU架构与GPU架构的对比如图2所示[15]。
图2 CPU架构与GPU架构的对比
随着“三金”工程尤其是金卡工程的启动,DES算法在POS、ATM、智能卡、收费站等领域被广泛使用,例如持卡人的PIN码保护加密算法、智能卡与设备交互之间的双向认证、数据包的MAC效验等均使用到了DES算法。DES算法的入口有3个参数:key、data、mode。其中key是一个8字节64位;data也是8字节64位;而key实际上的每一字节是7bit位,最后一位是奇偶效验位。DES的算法框图如图3所示[16]。
图3 DES算法框图
本文开发的基于CPU和GPU枚举DES密钥的Bitslice程序所使用的开发环境为Microsoft Windows 10 Professional 64位,Intel Core i5-2300@2.8GHz,8GB内 存,GPU为 Geforce GTX Titan Z,GPU详细的引擎规格如图4所示。由于Geforce GTX Titan Z是GPU×2。为了对比研究仅使用一个GPU芯片单元,显存大小也为原始的1/2即6GB。驱动程序版本号为344.11。CUDA开发环境为V6.5,SDK版本为4.2.9,CPU和GPU所使用的开发编译环境为Visual Studio2013。
图5为Geforce GTX Titan Z显卡的CUDA详细参数,每个处理器的最大线程数为2048个,每个block中的线程数为1024个。还包括Grid和不同维度上block块的大小(维度上不对称)。最重要的是memory pitch数据的大小为2147483647 Bytes,这个参数与复制的数组的容量有关。
图4 GPU详细的引擎规格
图5 CUDA详细参数
2 CUDA并行化的实现
GPU并行化程序的调试,首先实现在CPU上的单线程过程,通过单线程程序再改写成GPU多线程并行化程序,这样做的好处是能方便的计算出GPU相对于CPU的性能提升比。由于在支持超线程技术的单CPU上目前线程数一般不会超过48个线程。因此使用CPU进行多线程解决大数据问题并不具有可操作性。
2.1 CPU单线程Bitslice程序设计
单CPU端的单线程的Bitslice程序的流程如下:输入两个8字节的明文iv_p和pt_p这两个数字按位做异或运算,运算后得到64位的明文并存在在ptext[64]中,每一个字节的最后一位与前七位进行与操作,判断是否为非零,如果非零则输出非零的无符号长整型,否则则输出零。将这个数作为明文存在p[64]中。同样的道理将密文和预输入的key也存在c[64]和k[64]中。在32位系统中32bit为25,取位长的log值,使key的后五位与1或运算,如果这5位bit数与密文的相对应的bit位一致则进行另5bit位的比较,一次比较相当于进行了32次密钥搜索,Bitslice程序相当于自身并行运行提高破解效率。CPU单线程Bitslice程序不对密钥数据分组,会从高位一直向低位遍历,每一字节为0~255的数,通过Bitslice程序和最后一位奇偶效验位的简化,0~255被分成1,3,5,7.....255的128个数,相当于每个字节搜索128个数。如表1所示,程序将key[56]的一维数组输入到deseval函数中进行计算,如果对比发现为0,进行increment_key函数改变key的5个bit位之后的一位置为无符号的1,再进行deseval函数计算,以此类推直到所有的bit位全部对上才算找到key。
表1 CPU遍历key的查询函数
2.2 GPU多线程Bitslice程序设计
GPU多线程Bitslice程序设计的主要思想则是为了多线程的调用而优化的,它是将密钥数据进行分组到每个block中,每个block中的线程又执行不同分组的数据密钥,众多的线程累加起来完成全部的密钥空间。将密钥分块为若干个区域,每个区域作为一个流要执行的部分。再次使用与流数相关的第二个区域作为第二个流的数据区域,以此类推。
而对于每个数据分块,直到线程为止。定义每个块中含有256(ThreadN)个线程,含有BlockN个块。因此,ThreadN×BlockN就是一共含有的线程数。将大数据分成ThreadN×BlockN个小模块中处理,假设单步处理总时间为T,理想情况下所需时间应为T/ThreadN×BlockN,但是实际情况下显卡处理器性能会对其产生限制。举例说,显卡处理器为W位,在ThreadN×BlockN>W的情况下,数据处理仍是位一组,提高时速应在W倍附近。数组的列即枚举key的bit长度——56位长度。
图6 数据分组结构示意图
数据分组结构的示意图如图6所示,假设形成标准的XYZ三轴坐标轴体系,那么X轴就为BlockN表示块数。Z轴为Rows表示每块数据的一个数据的首地址。Y轴为Cols表示为经过Bitslice转换的56位密钥的数据。由CPU端到GPU端涉及到线程关联问题,即将现有问题或数据分成多块,并将这样的数据多块与每一个线程关联,得到一个与线程ID有关的函数。这样每个线程ID与没有数据块进行关联后进行计算可以得到每个线程计算的结果。如表2,GPU遍历Key的查询函数,但是数据块ID与线程ID的关联模型并不是惟一,不同关联方式对性能的影响也不一样。
表2 GPU遍历key的查询函数
在程序中Dsize的大小很重要,这个参数用于控制枚举的数量。在GPU并行多线程中,有些线程会重复操作,因此以密钥的一个字节为例:枚举{255,0,0,0,0,0,0,0}时,Dsize的数量最小为 Threads×3+1。当枚举 {255,255,0,0,0,0,0,0}时,Dsize的数量最小为 Threads×511+1。枚举 {255,255,255,0,0,0,0,0}时,Dsize的数量最小为Threads×65535+1。
我们设计的程序分块的思想为前3个字节都为0,让GPU去累加,而从第4个字节开始使用Bitslice思想让其分为128段(1,3,5,7...255)。以线程数的累加增加分块的数量同时增加密钥的累积长度,实现在同等时间内搜索密钥的长度更长。Dsize与搜索时间的关系图如图7所示。
图7 时间随Dsize变化的趋势图
由于单粒GPU的并行量和计算速度都不可能独立破解DES算法。因此,为了能验证GPU并行化所带来的优势,在一定密钥的长度下同时考验GPU和CPU所花费的时间来计算GPU加速倍数,且这个密钥的长度不能选择太长,由于CPU计算时间过长而失去验证的意义。选取密钥长度为{255,255,255,255,255,1,1,1},即密钥长度为2^36次方。使用GPU多线程计算设计线程为16384个线程,Dsize大小为:1073741824,穷举出密钥的时间为402秒,而通过CPU进行计算的Bitslice程序在计算同样密钥长度的时间为:15343秒。因此可以看出使用GPU数据分块,并行计算的性能提升是使用CPU的38.2倍。由于在一个GPU中的一个Grid中最大的线程数只能是65536个。因此使用GPU单流的最大BlockN为:256个。无论如何优化这个程序在单核单网格的情况下穷举DES的全部密钥成为天方夜谭。必然会想到使用多网格计算或多核计算继续优化并且扩大并行化操作。在使用多核计算之前,本文先介绍一种流技术的多网格优化方法。下面我们首先对使用CPU和GPU进行计算的性能差异进行简单的分析,进而论述多网格优化的Stream技术。
3 性能分析及优化
3.1 SSE2 vs CUDA
采取Microsoft Windows 10 Professional 64bit操作系统,Intel(R) Core(R)i5-2300 @2.8GHz CPU8GB 内存,软件则使用SSE2加速算法。当前的x86处理器采用乱序执行和缓存上下文读取,使用CPU的软SSE2指令集计算同样的程序,事实是CPU的执行效率和速度明显的优于GPU的单线程。使用不同的配置进行测试,SSE2优化代码(基于位计数的查表法)通常会快很多,而使用矩阵换位变换的则速度完全相同。这一点在英特尔凌动处理器上的SSE2代码更为突出,因此英特尔凌动处理器Atom的单指令多数据流(Single Instruction Multiple Data,SIMD)单元明显是经过定制优化。Nehalem微架构实现了硬件POPCNT,所以应该会更快。但不幸的是,有证据表明它并没有不执行POPCNT真正的SIMD128位版本,所以SSE+ POPCNT组合成为最后的选择[17-21]。
数据组作为GPU不变参量传入,常量内存交换速度快,但速度也就相当于CPU的Cache缓存速率。因此设定传入GPU的常量数组或数据为share类型。但是通常share类型的数据存储容量非常有限,一般的share类型的容量仅有16K。实现线程ID与数据结构ID的关联时,则必须找到起始线程和寄存器之间复制关系代码,且由于CUDA使用的是地址复制,一般在一维数组的复制上很容易编程,而在二维、三维数组的复制时,必须找到代码与硬件调用之间的良好平衡。这也就是为什么CUDA的例子里很少有使用二维以上的数组。如表3所示为主机端到设备端的二维数组复制的方法。如果使用Stream流技术进行优化的话,需增加一维为流ID,则主机端(Host)到设备端(Device)的数据复制为三维数组复制。
表3 Host-to-Device多维数组复制
cudaMemcpyHostToDevice);hKo = (unsigned int**)malloc(ROWS*sizeof(unsigned int*));ho = (unsigned int*)malloc(ROWS*COLS*sizeof(unsigned int));cudaMalloc((void**)(&dKo),ROWS*sizeof(unsigned int*));cudaMalloc((void**)(&dco),ROWS*COLS*sizeof (unsigned int));hRes = (unsigned int*)malloc(8 * sizeof(unsigned int));cudaMalloc((void**)&dRes,sizeof(unsigned int) * 8);hbuf = (unsigned int**)malloc(ROWS * sizeof(unsigned int*));for (int i = 0;i < ROWS;i++){cudaMalloc((void **)&hbuf[i],ROWS * sizeof(unsigned int));}cudaMalloc((void**)(&dKK),ROWS*sizeof(unsigned int*));for (int i = 0;i < ROWS;i++){cudaMemcpyAsync(hbuf[i],KK[i],sizeof(unsigned int)*ROWS,cudaMemcpyHostToDevice);}cudaMemcpyAsync((void*)dKK,(void*)hbuf,ROWS *sizeof(unsigned int*),cudaMemcpyHostToDevice);hKo = (unsigned int**)malloc(ROWS*sizeof(unsigned int*));for (int r = 0;r < ROWS;r++){hKo[r] = dco + r*COLS;}cudaMemcpyAsync((void*)(dKo),(void*)(hKo),ROWS*sizeof(unsigned int*),cudaMemcpyHostToDevice);
频繁进行host到device之间的上下文内存传输时,CUDA内核视调用内存上下文访问为高延迟操作。这对于相对简单的和高速整数运算尤为重要,这能提高计算程序调用硬件的效率。而对于需要频繁进行内存操作交换数据的程序则不能通过GPU进行加速。显卡没有操作系统或者受控于自身的复杂任务规划或调度,硬件的直接访问可以很好地对任务进行扩展。允许将程序分割成不均匀子任务。但是,不能并行太多不同的任务异步。由于着色器可预见的复杂和脆弱的存在。CPU-GPU中间件复杂限于结合并行在GPU上执行的任务次数。
总之,CPU编程中使用SIMD矢量编程(SSE)是需要一定编程技巧的,这里的主要问题是当前一代SSE并不是完整的图灵架构。从Michael Abrash[21]似乎看出LEBni(Larrabee的矢量指令集)将更好,Larabee有512位矢量寄存器。另外,Nvidia公司似乎并没有重点发展异构计算架构,其麦克斯韦架构也没有慢上下文切换。上下文切换计算可以使得GPU像CPU一样,利用上下文切换来管理多任务,每一个任务都可以通过分时的方法利用计算机资源。随着AMD公司的GCN异构计算架构(计算单元是基于新指令集的架构,抛弃了以往的VLIW(甚长指令字),而且每个计算单元都能同时从多个内核那里执行指令,单位周期单位面积的指令数也有所增加。这种架构相比以往利用率和吞吐量更高,多线程多任务并行执行的能力也大大增强)的发展,GCN架构的缓存体系也经过了完全重新设计,规模相当庞大而复杂。每四个计算单元共享16KB指令缓存和32KB标量数据缓存,并与二级缓存相连;每个计算单元都有自己的寄存器和本地数据共享,搭配16KB可读写一级缓存,每时钟周期带宽为64字节;二级缓存总容量768KB,可读写,对应每个显存控制器分成六组,每组容量128KB,每时钟周期带宽也是64字节;全局数据共享则用于不同计算单元之间的同步辅助。因此通用计算平台会成为GCN vs CUDA的对立格局。
3.2 Stream流技术优化
一些计算能力2.x或更高的设备可以同时并行执行多个内核函数。应用程序可以检查设备属性中的concurrentKernels项来确定设备是否支持这一功能,值为1代表支持。运算能力3.5的设备在同一时刻能够并行执行的最大内核函数数量为32,运算能力小于3.5的硬件则最多支持同时启动16个内核函数的执行。同时需要注意的是,在一个CUDA上下文中的内核函数不能与另一个CUDA上下文中的内核函数同时执行。使用很多纹理内存或者大量本地内存的内核函数也很可能无法与其他内核函数并行执行。为了并行化数据拷贝和内核执行,主机端内存必须分配为锁页(page-locked)内存其模块结构如图8所示[22-26]。
图8 CUDA流技术加速原理示意图
流的作用是任务之间的并行化,把要计算的数据分块,让不同的流执行不同的数据块,这样Stream1执行数据块1的时候,Stream2把数据块2从内存拷贝到显存,这样做不用等到所有数据拷贝完之后再计算从而提高并行力度。通过使用Stream流技术加速并扩大并行运算密钥Key的长度后,可以实现240次方的密钥长度计算时间为7.46小时。这个长度的密钥在CPU上的实现是非常困难的。但是在GPU多线程并行加速的架构下,只需执行几个小时即可完成。但是这个长度的密钥离完全穷举完全部密钥还差很远,以这样的速度计算下去全部穷举完的时间是无法接受的。
4 结束语
计算速率竞争的时代已经结束,现代的计算全面进入到并行时代。这意味着算法和系统的自适应性,并行架构的组织以满足数以万计的并行线程的现实,其中数据需要分块、重组,减少冲突和适应并行访问将成为并行计算的核心计算。GP-GPU计算的使用必须使用大量的计算资源,结合不断增长的数据挖掘集群服务器才能提供高效的计算服务。结合目前发展的快速随机访问固态硬盘SSD、高速RAM驱动器、DDR4代内存以节省带宽减缓可以提供更高效的上下文访问机制。并行计算作为一种科学手段,随着科研需求而迅猛发展,帮助学者们解决了很多复杂、冗余的问题。同时也给密码计算带来了前所未有的挑战,传统的CPU密码计算逐步的转向了分布式GPU阵列密码破译网络。通过使用大量并行多线程技术,能有效的降低密码在单线程任务上的计算容量,大量的线程任务同时工作提升了计算效率。虽然相比单CPU的Bitslice程序而言已有38.2倍的性能提升,但是其单个GPU的计算性能仍无法在可接受的时间内破解DES算法,因此,下一步的工作就是通过使用MPI多机技术将多块GPU组建成集群多机计算阵列。从上述研究可以表明DES在GPU上的运算并不消耗过多的计算资源,因此,每一个GPU的性能并不需太强,可以降低集群计算阵列的整体成本,同时也实现解决问题的目的。