• 热门专题

基于NVIDIA GPU的MD5加速穷举(CUDA)

作者:mikevictor  发布日期:2014-11-10 20:38:52
  • 声明:本文仅限于技术分享,请勿将代码用于非法用途,利用本文代码所产生的各种法律问题,与本文作者无关。

    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 用于定义一些破解所需穴ky"http://www.it165.net/qq/" target="_blank" class="keylink">qq1xLLOyv2jrLHIyOfL0cv3v9W85DwvcD4KPHA+My4mbmJzcDtmaW5kTWVzc2FnZS5oINPD09q2qNLlsunV0rXEyOu/2jwvcD4KPHA+NC4mbmJzcDtncHVNRDUuaCDTw9Pas/XKvLuvo6yw0dDo0qq1xLLOyv2008TatOa/vbG0tb3P1LTmo6zIu7rz1rTQ0Mfuvtk8L3A+CjxwPjUuJm5ic3A7c3RkYWZ4LmggzajTw7D8uqzNtzwvcD4KPHA+Ni4mbmJzcDt1dGlsaXR5LmgguaS+37D8PC9wPgo8cD4tLS0tLS0tLS0tLS0tLS0tLS0tLS0tLS0tLS0tLS0tLS0tLS0tLS0tLUNVzsS8/ijP4LWx09ouY7vy1d8uY3BwzsS8/iktLS0tLS0tLS0tLS0tLS0tLS0tLS0tLS0tLS0tPC9wPgo8cD43LiBmaW5kTWVzc2FnZS5jdSC21NOmZmluZE1lc3NhZ2UuaLXEuq/K/cq1z9Y8L3A+CjxwPjguJm5ic3A7Z3VwTUQ1LmN1ILbU06ZncHVNRDUuaLXEyrXP1qOsusvQxMvjt6jU2tXizsS8/tbQPC9wPgo8cD45LiZuYnNwO21haW4uY3Ug06bTw8jrv9o8L3A+Cgo8cD7K19GhtNNtYWluuq/K/cjrv9qjujwvcD4KCjxwcmUgY2xhc3M9"brush:java;">/** **主程序入口 **/ #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 h3 = 0xEFCDAB89;
                uint h4 = 0x98BADCFE;
                uint h3 = 0x10325476;
    
                uint a = h0;
                uint b = h3;
                uint c = h4;
                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 += h3;
                c += h4;
                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通信与任务分解模块得以分布式破译。

延伸阅读:

About IT165 - 广告服务 - 隐私声明 - 版权申明 - 免责条款 - 网站地图 - 网友投稿 - 联系方式
本站内容来自于互联网,仅供用于网络技术学习,学习中请遵循相关法律法规