声明： 本文仅限于技术分享，请勿将代码用于非法用途，利用本文代码所产生的各种法律问题，与本文作者无关。

1. 摘要：

MD5为非常普遍使用的消息摘要算法，很多应用系统采用该算法加密密码，在计算文件摘要值以验证文件是否被篡改方面也普遍使用，

MD5服务安全方面很多年，随着计算机技术的发展，该算法已经很不安全，穷举遍历的代价也变得没那么高，笔者建议至少采用（SHA1+盐值）

方法加密新建设的应用系统，由于目前很多网站大量的用户名密码泄露，个人的信息安全也越来越重要，目前很多系统采用的加密算法有：

1>自定义算法

2>MD5(PWD)

3>MD5(PWD+盐)

4>MD5(MD5(PWD))

5>SHA1(PWD)

6>SHA1(PWD+盐)

7>3DES+安全存储的密钥

8>明文   — 这类系统建设者安全意识缺失严重

单项加密算法解密的一般步骤：字典匹配->社工->彩虹表->暴力。本文只介绍暴力穷举方式，采用单机版本，可自行加入分布式模块分解搜索空间。

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通信与任务分解模块得以分布式破译。