Help us understand the problem. What is going on with this article?

「数字6桁パスワードのMD5ハッシュ値の総当たり」CUDAを使うと0.20秒まではやくなったよ

More than 3 years have passed since last update.

http://tech-sketch.jp/2014/03/md5-openmp.html
↑を見つけてその後
「数字6桁パスワードのMD5ハッシュ値の総当たり」OpenMPを使うと0.70秒までキタにたどり着いたので、これらを参考にせっかくC言語で書いてあるのでCUDA化してみました。

環境

Windows8.1 Pro
CUDA : 7.5
CPU : corei7-4770K @ 3.50GHz
GPU : GeForce GTX 760

今回使用したGPUは1152コアあります。最新のものは5000コア以上あるものもあるのでスペック的にはまあまあといった感じのものです。

ちなみにGPUのスペックはCUDAを入れていればdeviceQueryを実行すれば調べられます。
私の場合は

C:\ProgramData\NVIDIA Corporation\CUDA Samples\v7.5\1_Utilities\deviceQuery

に入っていました。

結果

C言語版のループを抜けるまでの時間を計測しました。

C言語:4.91秒
CUDA:0.20秒

高速化できましたね^ ^。

ただ、個人的にはもう少し早くなるかと思っていました。今回はCUDAに移植しただけであまり最適化的なことはやっていないので、そこらへんもやればもう少し早くなるのではないかと思います。

CUDA化手順

かなり雑です。
一応コードはこれです。
そのうちきれいにしたい...

CUDAは基本的にC言語の拡張のようなものなのです。なので、
「数字6桁パスワードのMD5ハッシュ値の総当たり」OpenMPを使うと0.70秒までキタの「Cで実装してみた」とhttp://openwall.info/wiki/people/solar/software/public-domain-source-code/md5
で公開されているC実装をもとにCUDA化していきます。

残念ながらCUDAではsprintf_sとかが使えない(はず?)のでそこら辺を書き換えが面倒でした。

まずはmain部分です。

#include <stdio.h>
#include <string.h>
#include <time.h>
#include "md5.cuh"

int main(int argc, char* argv[])
{
    if (argc != 3) return 1;

    char* SALT = argv[1];
    char* PW = argv[2];
    char result[7];

    /*-----並列化された処理-----*/
    char *dev_SALT, *dev_PW, *dev_result;

    cudaMalloc((void**)&dev_SALT, strlen(SALT) *sizeof(char));
    cudaMalloc((void**)&dev_PW, strlen(PW) *sizeof(char));
    cudaMalloc((void**)&dev_result, 7 * sizeof(char));

    cudaMemcpy(dev_SALT, SALT, strlen(SALT) * sizeof(char), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_PW, PW, strlen(PW) * sizeof(char), cudaMemcpyHostToDevice);

    md5_brute_force<<<1000, 1000 >>>(dev_SALT, dev_PW, dev_result);

    cudaMemcpy(result, dev_result, 7 * sizeof(char), cudaMemcpyDeviceToHost);
    /*---------------------*/
    printf("%s\n", result);

    return 0;
}

さらにhttp://openwall.info/wiki/people/solar/software/public-domain-source-code/md5 のmd5.c md5.hの関数の前にGPU側で呼ばれる関数であることを示す__device__をつけます。

そしてmd5.cに

// http://libc.blog47.fc2.com/blog-entry-30.html
//strlenのCUDA用
__device__ __host__ size_t strlen(const char *s)
{
    size_t n;
    for (n = 0; *s != '\0'; s++, n++);
    return n;
}

// http://www.geeksforgeeks.org/implement-itoa/
// itoaもCUDAに移すために上のリンクの説明を参考にCUDA化します
/* A utility function to reverse a string  */
__device__ void reverse(char str[], int length)
{
    int start = 0;
    int end = length - 1;
    while (start < end)
    {
        char temp = *(str + start);
        *(str + start) = *(str + end);
        *(str + end) = temp;

        //swap(*(str + start), *(str + end));
        start++;
        end--;
    }
}

// Implementation of itoa()
__device__ char* itoa(int num, char* str, int base)
{
    int i = 0;
    bool isNegative = false;

    /* Handle 0 explicitely, otherwise empty string is printed for 0 */
    if (num == 0)
    {
        str[i++] = '0';
        str[i] = '\0';
        return str;
    }

    // In standard itoa(), negative numbers are handled only with 
    // base 10. Otherwise numbers are considered unsigned.
    if (num < 0 && base == 10)
    {
        isNegative = true;
        num = -num;
    }

    // Process individual digits
    while (num != 0)
    {
        int rem = num % base;
        str[i++] = (rem > 9) ? (rem - 10) + 'a' : rem + '0';
        num = num / base;
    }

    // If number is negative, append '-'
    if (isNegative)
        str[i++] = '-';

    str[i] = '\0'; // Append string terminator

    // Reverse the string
    reverse(str, i);

    return str;
}


// CUDAではsprintf_sが使えない(んだった気がする)ので下のように地道に置き換えます。
__device__ void six_digit_decimal(char num_c[7], int num_i){
    char num[7] = "";
    int n = num_i;
    itoa(n, num, 10);
    int dig = 1;
    n /= 10;
    while (n != 0){
        dig++;
        n /= 10;
    }

    for (int j = 0; j < dig; j++){
        num_c[6 - dig + j] = num[j];
    }
}

__device__ void two_digit_hexadecimal(char num_c[3], int num_i){
    if (num_i < 16) {
        itoa(num_i, num_c, 16);
        num_c[1] = num_c[0];
        num_c[0] = '0';
    }
    else itoa(num_i, num_c, 16);
}

// http://stackoverflow.com/questions/20201335/add-char-arrays-in-cuda
// C言語のstrcatをCUDAで使うため上のURLから借りてきます。
__device__ char * my_strcpy(char *dest, const char *src){
    int i = 0;
    do {
        dest[i] = src[i];
    } while (src[i++] != 0);
    return dest;
}

__device__ char * my_strcat(char *dest, const char *src){
    int i = 0;
    while (dest[i] != 0) i++;
    my_strcpy(dest + i, src);
    return dest;
}

// 上のを参考にstrcmpもCUDAで実行できるようにします。
__device__ bool my_strcmp(char *dest, char *src){
    int i = 0;
    do {
        if (dest[i] != src[i]) return false;
    } while (src[i++] != 0);
    return true;
}

__global__ void md5_brute_force(char *SALT, char *PW, char *ret_result){
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    MD5_CTX ctx;
    unsigned char hash[MD5_DIGEST_LENGTH];
    char result[MD5_DIGEST_LENGTH * 2 + 1];
    result[0] = '\0';
    char num_c[7] = "000000";
    six_digit_decimal(num_c, i);
    char origin[256] = "";
    my_strcat(origin, SALT);
    my_strcat(origin, "$");
    my_strcat(origin, num_c);

    MD5_Init(&ctx);
    MD5_Update(&ctx, origin, (unsigned long)strlen(origin));
    MD5_Final(hash, &ctx);
    for (int j = 0; j < MD5_DIGEST_LENGTH; j++) {
        char hex[3] = "";
        two_digit_hexadecimal(hex, hash[j]);
        my_strcat(result, hex);
    }
    if (my_strcmp(result, PW)) {
        six_digit_decimal(ret_result, i);
    }
}

を追加します。
md5.hには

extern __device__ __host__ size_t strlen(const char *s);
extern __device__ void reverse(char str[], int length);
extern __device__ char* itoa(int num, char* str, int base);
extern __device__ void six_digit_decimal(char num_c[7], int num_i);
extern __device__ void two_digit_hexadecimal(char num_c[3], int num_i);
extern __device__ char * my_strcpy(char *dest, const char *src);
extern __device__ char * my_strcat(char *dest, const char *src);
extern __device__ bool my_strcmp(char *dest, char *src);
extern __global__ void md5_brute_force(char *SALT, char *PW, char *ret_result);

を追加します。
これで実行できるはず。

Why not register and get more from Qiita?
  1. We will deliver articles that match you
    By following users and tags, you can catch up information on technical fields that you are interested in as a whole
  2. you can read useful information later efficiently
    By "stocking" the articles you like, you can search right away
Comments
No comments
Sign up for free and join this conversation.
If you already have a Qiita account
Why do not you register as a user and use Qiita more conveniently?
You need to log in to use this function. Qiita can be used more conveniently after logging in.
You seem to be reading articles frequently this month. Qiita can be used more conveniently after logging in.
  1. We will deliver articles that match you
    By following users and tags, you can catch up information on technical fields that you are interested in as a whole
  2. you can read useful information later efficiently
    By "stocking" the articles you like, you can search right away
ユーザーは見つかりませんでした