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);

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

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

環境

結果

CUDA化手順