基於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