码迷,mamicode.com
首页 > 其他好文 > 详细

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

时间:2014-11-09 13:48:24      阅读:154      评论:0      收藏:0      [点我收藏+]

标签:des   style   blog   http   io   color   ar   os   使用   

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

 

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查看

项目结构如下:

bubuko.com,布布扣

 

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

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

标签:des   style   blog   http   io   color   ar   os   使用   

原文地址:http://www.cnblogs.com/mikevictor07/p/4084168.html

(0)
(0)
   
举报
评论 一句话评论(0
登录后才能评论!
© 2014 mamicode.com 版权所有  联系我们:gaon5@hotmail.com
迷上了代码!