基于NVIDIA GPU的MD5加速穷举(CUDA)
声明：本文仅限于技术分享，请勿将代码用于非法用途，利用本文代码所产生的各种法律问题，与本文作者无关。
1. 摘要：
MD5为非常普遍使用的消息摘要算法，很多应用系统采用该算法加密密码，在计算文件摘要值以验证文件是否被篡改方面也普遍使用，
MD5服务安全方面很多年，随着计算机技术的发展，该算法已经很不安全，穷举遍历的代价也变得没那么高，笔者建议至少采用（SHA1+盐值）
方法加密新建设的应用系统，由于目前很多网站大量的用户名密码泄露，个人的信息安全也越来越重要，目前很多系统采用的加密算法有：
1>自定义算法
2>MD5(PWD)
3>MD5(PWD+盐)
4>MD5(MD5(PWD))
5>SHA1(PWD)
6>SHA1(PWD+盐)
7>3DES+安全存储的密钥
8>明文 -- 这类系统建设者安全意识缺失严重
本文介绍使用GPU穷举MD5方式正向匹配，也称作暴力，本文技术分享只限用于密码安全研究，加强企业安全体系的建设，请勿用于非法途径！！！
单项加密算法解密的一般步骤：字典匹配->社工->彩虹表->暴力。本文只介绍暴力穷举方式，采用单机版本，可自行加入分布式模块分解搜索空间。
2.理解本文至少所需要的知识：c/c++, CUDA, MD5算法
需要准备的工作:
1>VS 2010/2012
2>CUDA 5.x or higher（本文采用5.5）
3>一台携带支持CUDA开发NVIDIA 显卡的电脑，可使用GPU-Z查看
项目结构如下：
3.源码详解
首先对各个文件作用说明：
--------------------------------------头文件--------------------------------------------------------
1. common_def.h 通用定义头，用户声明一些基本类型，如typedef unsigned int uint;
2. deviceMemoryDef.h 用于定义一些破解所需要的参数，比如搜索空间
3. findMessage.h 用于定义查找的入口
4. gpuMD5.h 用于初始化，把需要的参数从内存拷贝到显存，然后执行穷举
5. stdafx.h 通用包含头
6. utility.h 工具包
-----------------------------------------CU文件(相当于.c或者.cpp文件)----------------------------
7. findMessage.cu 对应findMessage.h的函数实现
8. gupMD5.cu 对应gpuMD5.h的实现，核心算法在这文件中
9. main.cu 应用入口
首选从main函数入口：
/** **主程序入口 **/ #include "stdafx.h" #include "gpuMD5.h" #include "findMessage.h" /** *主函数 */ int main() { //123456 = e10adc3949ba59abbe56e057f20f883e // 999999 = 52c69e3a57331081823331c4e69d3f2e //adc12d4 = 32f3db39aa85fac25c19c0c8b555dc83 string target = "32f3db39aa85fac25c19c0c8b555dc83"; //0123456789abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ string searchScope = "0123456789abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ"; //搜索空间,可自行修改 initGPU(target, searchScope); //初始化显存参数 size_t startNum = 5, endNum = 8; pair<bool, string> result = findMessage(startNum, endNum, searchScope); if(result.first){ cout<<"找到明文："<<result.second<<endl; }else{ cout<<"未搜索到相应明文."<<endl; } return 0; }
上面函数所需的包含头内容如下：
stdafx.h：
// stdafx.h : 标准系统包含文件的包含文件， // 或是经常使用但不常更改的 // 特定于项目的包含文件 // #pragma once #include <iostream> #include <string> #include <tchar.h> #include <time.h> #include <cuda_runtime.h> using namespace std; #include "common_def.h" #include "utility.h"
gpuMD5.h：
//gupMD5.H /** *GPU匹配MD5算法 */ #include "stdafx.h" #include "deviceMemoryDef.h" /** GPU初始化 参数： targetDigest 密文 searchScope 搜索的字符集 */ void initGPU(string targetDigest, string searchScope); /* GPU运算搜索 参数： d_startSearchWord 存放每个线程的起始搜索空间 useThreadNum 实际使用的线程数量 charsetLength 搜索空间长度 size 当前搜索长度 d_isFound 搜索结果 d_message 明文 */ __global__ void searchMD5(float*, float, size_t, size_t, size_t*, uchar*);
findMessage.h：
/**findMessage.h */ #include "stdafx.h" /** *搜索明文 参数： min 最小长度 max 最大长度 searchScope 搜索空间 */ pair<bool, string> findMessage(size_t min, size_t max, string searchScope) ;
当然还缺少不了通用的定义头如下：
common_def.h：
/** *公用文件：包含通用定义 */ #ifndef __COMMON_H__ #define __COMMON_H__ typedef unsigned int uint; typedef unsigned char uchar; #define MAX_PLAIN_TEXT_LENGTH 16 ////破译密码最大长度为16，对于大于16，则失去破译意义 #endif // !__COMMON_H__
deviceMemoryDef.h：
/** *设备公用显存变量定义 */ #include "stdafx.h" #ifndef __DEVICE_MEMORY_DEF_H__ #define __DEVICE_MEMORY_DEF_H__ // 比对的目标数组(a b c d)，只能由GPU设备调用 #define NUM_DIGEST_SIZE 4 __device__ __constant__ uint d_targetDigest[NUM_DIGEST_SIZE]; // 搜索字符数组 包含 a-z A-Z 0-9 ~!@#$%^&*()_+-=[]\|;:"'<,>.?/ #define NUM_POWER_SYMBOLS 96 __device__ __constant__ uchar d_powerSymbols[NUM_POWER_SYMBOLS]; //搜索长度的组合数量 #define NUM_POWER_VALUES 16 __constant__ float d_powerValues[NUM_POWER_VALUES]; #endif // !__DEVICE_MEMORY_DEF_H__
utility.h：
//utility.h /** *实用工具类 */ #ifndef __UTILITY_H__ #define __UTILITY_H__ #include "stdafx.h" #define CUDA_MALLOC_ERROR 1 //CUDA内存分配错误 #define CUDA_MEM_CPY_ERROR 2 //CUDA内存拷贝错误 /* *打印CUDA错误信息 *参数: error 错误码 msg 错误信息 errorType 错误类型 fileName 出错的文件名 line 错误在文件中的行数 */ void printCudaError(cudaError_t error, string msg,string fileName, int line); #endif // !__UTILITY_H__
上面列出了本应用所有的头文件，下面讲解实现文件：
main函数中，初始化GPU所需的initGPU函数在gpuMD5.h中定义，对应的实现文件也就是gpuMD5.cu，初始化方法中主要做了四件事情：
1>分解目标加密串，得到四个整数数组，也就是匹配这四个整数的数组（后面成为目标数组）即可，节省计算时间与内存空间。
2>为显存分配空间，将目标数组从内存拷贝至显存，将所需要的搜索空间拷贝至显存。
3>在searchMD5函数中对显卡的各个计算模块分解任务，当有个计算模块得到目标值后则改变显卡的GLOBAL参数，通知其他计算模块。
4>退出计算，回收显存与内存，打印计算结果
下面为gupMD5.cu的函数实现，需要具备一定的算法知识与MD5的理解。
//gupMD5.cu #include "gpuMD5.h" /** MD5散列函数 **/ #define F(x, y, z) (((x) & (y)) | ((~x) & (z))) #define G(x, y, z) (((x) & (z)) | ((y) & (~z))) #define H(x, y, z) ((x) ^ (y) ^ (z)) #define I(x, y, z) ((y) ^ ((x) | (~z))) #define ROTATE_LEFT(x, n) (((x) << (n)) | ((x) >> (32-(n)))) #define FF(a, b, c, d, x, s, ac) \ {(a) += F ((b), (c), (d)) + (x) + (uint)(ac); \ (a) = ROTATE_LEFT ((a), (s)); \ (a) += (b); \ } #define GG(a, b, c, d, x, s, ac) \ {(a) += G ((b), (c), (d)) + (x) + (uint)(ac); \ (a) = ROTATE_LEFT ((a), (s)); \ (a) += (b); \ } #define HH(a, b, c, d, x, s, ac) \ {(a) += H ((b), (c), (d)) + (x) + (uint)(ac); \ (a) = ROTATE_LEFT ((a), (s)); \ (a) += (b); \ } #define II(a, b, c, d, x, s, ac) \ {(a) += I ((b), (c), (d)) + (x) + (uint)(ac); \ (a) = ROTATE_LEFT ((a), (s)); \ (a) += (b); \ } // char 转化为 uchar uchar c2c (char c){ return (uchar)((c > '9') ? (c - 'a' + 10) : (c - '0')); } void initGPU(string targetDigest, string searchScope) { uint h_targetDigest[4]; //内存中的比对目标 for(int c=0;c<targetDigest.size();c+=8) { uint x = c2c(targetDigest[c]) <<4 | c2c(targetDigest[c+1]); uint y = c2c(targetDigest[c+2]) << 4 | c2c(targetDigest[c+3]); uint z = c2c(targetDigest[c+4]) << 4 | c2c(targetDigest[c+5]); uint w = c2c(targetDigest[c+6]) << 4 | c2c(targetDigest[c+7]); h_targetDigest[c/8] = w << 24 | z << 16 | y << 8 | x; } cout<<"h_targetDigest[0]="<<h_targetDigest[0]<<endl; cout<<"h_targetDigest[1]="<<h_targetDigest[1]<<endl; cout<<"h_targetDigest[2]="<<h_targetDigest[2]<<endl; cout<<"h_targetDigest[3]="<<h_targetDigest[3]<<endl; float charsetLen = searchScope.length(); cudaError_t error; //将目标散列数组 由主机拷贝到设备常量存储器 error = cudaMemcpyToSymbol(d_targetDigest, h_targetDigest, NUM_DIGEST_SIZE * sizeof(uint)); if (error != cudaSuccess){ printCudaError(error,"拷贝（目标散列数组）到（设备常量存储器）出错", __FILE__, __LINE__); } uchar h_powerSymbols[NUM_POWER_SYMBOLS]; for (size_t i = 0; i != charsetLen; ++i) { h_powerSymbols[i] = searchScope[i]; } // 拷贝搜索空间字符数组到 设备常量存储器 error = cudaMemcpyToSymbol(d_powerSymbols, h_powerSymbols, NUM_POWER_SYMBOLS * sizeof(uchar)); if (error != cudaSuccess){ printCudaError(error,"拷贝（搜索空间字符数组）到（设备常量存储器出错）", __FILE__, __LINE__); } float h_powerValues[NUM_POWER_VALUES]; for (size_t i = 0; i != NUM_POWER_VALUES; ++i) h_powerValues[i] = pow(charsetLen, (float)(NUM_POWER_VALUES - i - 1)); cudaMemcpyToSymbol(d_powerValues, h_powerValues, NUM_POWER_VALUES * sizeof(float)); } __global__ void searchMD5(float* d_startNumbers, float nIterations, size_t charsetLength, size_t size, size_t* d_isFound, uchar* message){ size_t idx = blockIdx.x * blockDim.x + threadIdx.x; float maxValue = powf(__uint2float_rz(charsetLength), __uint2float_rz(size));//最大组合数 uint in[17]; for (size_t i = 0; i != 17; ++i){ in[i] = 0x00000000; } in[14] = size << 3; uchar* toHashAsChar = (uchar*)in; for (size_t i = 0; i != size; ++i){ toHashAsChar[i] = d_powerSymbols[0]; } toHashAsChar[size] = 0x80; float numberToConvert = d_startNumbers[idx];//获取起始匹配地址 size_t toHashAsCharIndices[17];//记录当前线程需要处理的字符数在搜索空间里面的位置 if(numberToConvert < maxValue) { //得到该线程的起始搜索地址 for(size_t i = 0; i != size; ++i) { //得到该线程起始地址在当前搜索范围中的比率，然后取整 toHashAsCharIndices[i] = __float2uint_rz(floorf(numberToConvert / d_powerValues[NUM_POWER_VALUES - size + i])); //得到多出来的位数 numberToConvert = floorf(fmodf(numberToConvert, d_powerValues[NUM_POWER_VALUES - size + i])); } /*printf("线程%d的起始搜索地址：",idx); for (size_t i = 0; i != size; ++i){ toHashAsChar[i] = d_powerSymbols[toHashAsCharIndices[i]]; printf("%c",toHashAsChar[i]); } printf("\n");*/ #pragma unroll 5 for(float iterationsDone = 0; iterationsDone != nIterations; ++iterationsDone){ if (*d_isFound == 1) break; for (size_t i = 0; i != size; ++i){ toHashAsChar[i] = d_powerSymbols[toHashAsCharIndices[i]];//根据字符位置取出字符 } //MD5 HASH uint h0 = 0x67452301; uint h1 = 0xEFCDAB89; uint h2 = 0x98BADCFE; uint h3 = 0x10325476; uint a = h0; uint b = h1; uint c = h2; uint d = h3; /* Round 1 */ #define S11 7 #define S12 12 #define S13 17 #define S14 22 FF ( a, b, c, d, in[ 0], S11, 3614090360); /* 1 */ FF ( d, a, b, c, in[ 1], S12, 3905402710); /* 2 */ FF ( c, d, a, b, in[ 2], S13, 606105819); /* 3 */ FF ( b, c, d, a, in[ 3], S14, 3250441966); /* 4 */ FF ( a, b, c, d, in[ 4], S11, 4118548399); /* 5 */ FF ( d, a, b, c, in[ 5], S12, 1200080426); /* 6 */ FF ( c, d, a, b, in[ 6], S13, 2821735955); /* 7 */ FF ( b, c, d, a, in[ 7], S14, 4249261313); /* 8 */ FF ( a, b, c, d, in[ 8], S11, 1770035416); /* 9 */ FF ( d, a, b, c, in[ 9], S12, 2336552879); /* 10 */ FF ( c, d, a, b, in[10], S13, 4294925233); /* 11 */ FF ( b, c, d, a, in[11], S14, 2304563134); /* 12 */ FF ( a, b, c, d, in[12], S11, 1804603682); /* 13 */ FF ( d, a, b, c, in[13], S12, 4254626195); /* 14 */ FF ( c, d, a, b, in[14], S13, 2792965006); /* 15 */ FF ( b, c, d, a, in[15], S14, 1236535329); /* 16 */ /* Round 2 */ #define S21 5 #define S22 9 #define S23 14 #define S24 20 GG ( a, b, c, d, in[ 1], S21, 4129170786); /* 17 */ GG ( d, a, b, c, in[ 6], S22, 3225465664); /* 18 */ GG ( c, d, a, b, in[11], S23, 643717713); /* 19 */ GG ( b, c, d, a, in[ 0], S24, 3921069994); /* 20 */ GG ( a, b, c, d, in[ 5], S21, 3593408605); /* 21 */ GG ( d, a, b, c, in[10], S22, 38016083); /* 22 */ GG ( c, d, a, b, in[15], S23, 3634488961); /* 23 */ GG ( b, c, d, a, in[ 4], S24, 3889429448); /* 24 */ GG ( a, b, c, d, in[ 9], S21, 568446438); /* 25 */ GG ( d, a, b, c, in[14], S22, 3275163606); /* 26 */ GG ( c, d, a, b, in[ 3], S23, 4107603335); /* 27 */ GG ( b, c, d, a, in[ 8], S24, 1163531501); /* 28 */ GG ( a, b, c, d, in[13], S21, 2850285829); /* 29 */ GG ( d, a, b, c, in[ 2], S22, 4243563512); /* 30 */ GG ( c, d, a, b, in[ 7], S23, 1735328473); /* 31 */ GG ( b, c, d, a, in[12], S24, 2368359562); /* 32 */ /* Round 3 */ #define S31 4 #define S32 11 #define S33 16 #define S34 23 HH ( a, b, c, d, in[ 5], S31, 4294588738); /* 33 */ HH ( d, a, b, c, in[ 8], S32, 2272392833); /* 34 */ HH ( c, d, a, b, in[11], S33, 1839030562); /* 35 */ HH ( b, c, d, a, in[14], S34, 4259657740); /* 36 */ HH ( a, b, c, d, in[ 1], S31, 2763975236); /* 37 */ HH ( d, a, b, c, in[ 4], S32, 1272893353); /* 38 */ HH ( c, d, a, b, in[ 7], S33, 4139469664); /* 39 */ HH ( b, c, d, a, in[10], S34, 3200236656); /* 40 */ HH ( a, b, c, d, in[13], S31, 681279174); /* 41 */ HH ( d, a, b, c, in[ 0], S32, 3936430074); /* 42 */ HH ( c, d, a, b, in[ 3], S33, 3572445317); /* 43 */ HH ( b, c, d, a, in[ 6], S34, 76029189); /* 44 */ HH ( a, b, c, d, in[ 9], S31, 3654602809); /* 45 */ HH ( d, a, b, c, in[12], S32, 3873151461); /* 46 */ HH ( c, d, a, b, in[15], S33, 530742520); /* 47 */ HH ( b, c, d, a, in[ 2], S34, 3299628645); /* 48 */ /* Round 4 */ #define S41 6 #define S42 10 #define S43 15 #define S44 21 II ( a, b, c, d, in[ 0], S41, 4096336452); /* 49 */ II ( d, a, b, c, in[ 7], S42, 1126891415); /* 50 */ II ( c, d, a, b, in[14], S43, 2878612391); /* 51 */ II ( b, c, d, a, in[ 5], S44, 4237533241); /* 52 */ II ( a, b, c, d, in[12], S41, 1700485571); /* 53 */ II ( d, a, b, c, in[ 3], S42, 2399980690); /* 54 */ II ( c, d, a, b, in[10], S43, 4293915773); /* 55 */ II ( b, c, d, a, in[ 1], S44, 2240044497); /* 56 */ II ( a, b, c, d, in[ 8], S41, 1873313359); /* 57 */ II ( d, a, b, c, in[15], S42, 4264355552); /* 58 */ II ( c, d, a, b, in[ 6], S43, 2734768916); /* 59 */ II ( b, c, d, a, in[13], S44, 1309151649); /* 60 */ II ( a, b, c, d, in[ 4], S41, 4149444226); /* 61 */ II ( d, a, b, c, in[11], S42, 3174756917); /* 62 */ II ( c, d, a, b, in[ 2], S43, 718787259); /* 63 */ II ( b, c, d, a, in[ 9], S44, 3951481745); /* 64 */ a += h0; b += h1; c += h2; d += h3; //检查散列值是否匹配 if (a == d_targetDigest[0] && b == d_targetDigest[1] && c == d_targetDigest[2] && d == d_targetDigest[3]){ *d_isFound = 1; for (size_t i = 0; i != size; ++i){//取出结果 message[i] = toHashAsChar[i]; } }else { size_t i = size - 1; bool incrementNext = true;//是否递增，若无法递增则进位 while (incrementNext){//若后面无法进位，则把指针移到前一位进位，如[115]->[121] if (toHashAsCharIndices[i] < (charsetLength - 1)) { ++toHashAsCharIndices[i]; incrementNext = false; } else { if (toHashAsCharIndices[i] >= charsetLength) { *d_isFound = 3; } toHashAsCharIndices[i] = 0; if (i == 0) { incrementNext = false; } else { --i; } } } } } } }
初始化函数后到重要的findMessage函数，findMessage.cu的实现步骤：
1>设置显卡运算线程数，不同显卡具有的计算核心不一样，设置也需要相应的改变。
2>定义搜索目标的密码长度，比如从6位搜索至9位等等，最长16位，当密码长度达到>=12的组合密码时，计算代价非常大。
代码如下：
//findMessage.cpp #include "stdafx.h" #include "deviceMemoryDef.h" #include "gpuMD5.h" /** 搜索明文 min：明文最小长度 max：明文最大长度 searchScope：搜索空间 */ pair<bool, string> findMessage(size_t min, size_t max, string searchScope) { bool isFound = false; size_t h_isFound = -1; size_t * d_isFound; //搜索结果标识 uchar* d_message; uchar h_message[16]; //明文，最大支持长度为16 string message = ""; //GoForce GT650M 比较优秀的设置：1024*1024 int nBlocks = 1024; int nThreadsPerBlock = 1024; size_t nTotalThreads = nBlocks * nThreadsPerBlock; // 总线程数 size_t charsetLength = searchScope.length(); //搜索空间字符数长度 cudaError_t error; error = cudaMalloc((void**)&d_isFound, sizeof(size_t)); if (error != cudaSuccess){ printCudaError(error,"分配（搜索结果标识）显存出错", __FILE__, __LINE__); } error = cudaMemcpy(d_isFound, &h_isFound, sizeof(size_t), cudaMemcpyHostToDevice); if (error != cudaSuccess){ printCudaError(error,"拷贝（搜索结果标识）至显存出错", __FILE__, __LINE__); } error = cudaMalloc((void**)&d_message, 16 * sizeof(uchar)); if (error != cudaSuccess){ printCudaError(error,"分配搜索结果（明文）显存出错", __FILE__, __LINE__); } //分配每个线程的搜索起始地址 float* h_startNumbers = new float[nTotalThreads]; float* d_startNumbers; error = cudaMalloc((void**)&d_startNumbers, nTotalThreads * sizeof(float)); if (error != cudaSuccess){ printCudaError(error,"分配线程的搜索起始地址出错", __FILE__, __LINE__); } for (size_t size = min; size <= max; ++size) { cout<<"当前搜索长度："<<size<<endl; float maxValue = pow((float)charsetLength, (float)size); //最大匹配数 float nIterations = ceil(maxValue / (nBlocks * nThreadsPerBlock));//每个线程分配的任务数,即每个线程需要遍历的个数 for (size_t i = 0; i != nTotalThreads; ++i) { h_startNumbers[i] = i * nIterations; } error = cudaMemcpy(d_startNumbers, h_startNumbers, nTotalThreads * sizeof(float), cudaMemcpyHostToDevice); if (error != cudaSuccess){ printCudaError(error,"拷贝 线程的搜索起始地址 到显存出错", __FILE__, __LINE__); } clock_t start = clock(); //开始运算 searchMD5<<< nBlocks, nThreadsPerBlock >>>(d_startNumbers, nIterations, charsetLength, size, d_isFound, d_message); cudaThreadSynchronize(); cout<<"耗时："<<(clock()-start)/CLK_TCK<<endl; printf("%s\n", cudaGetErrorString(cudaGetLastError())); cudaMemcpy(&h_isFound, d_isFound, sizeof(int), cudaMemcpyDeviceToHost); printf("####################### h_isFound = %d\n", h_isFound); if (h_isFound != -1) { printf("h_isFound=%d\n", h_isFound); cudaMemcpy(h_message, d_message, 16 * sizeof(uchar), cudaMemcpyDeviceToHost); for (size_t i = 0; i != size; ++i){ message.push_back(h_message[i]); } isFound = true; cout << message << endl; break; } } //释放内存和显存 cudaFree(d_targetDigest); cudaFree(d_powerSymbols); cudaFree(d_powerValues); cudaFree(d_isFound); cudaFree(d_message); cudaFree(d_startNumbers); delete(h_startNumbers); cout<<"释放内存完毕..."<<endl; return make_pair(isFound, message); }
最后剩下一个工具类utility.cu的实现：
//utility.cu #include "utility.h" void printCudaError(cudaError_t error, string msg, string fileName, int line) { cout<<msg<<"，错误码："<<error<<"，文件("<<fileName<<")，行数("<<line<<")."<<endl; exit(EXIT_FAILURE); }
总结：本文仅用于学习，为单机版本，需要另外实现TCP通信与任务分解模块得以分布式破译。
Git：https://git.oschina.net/redcode/md5GPUCrack.git
