From 8c843ce93cb0a0fb4607e1cc88494d66e7d070cf Mon Sep 17 00:00:00 2001 From: Dhiru Kholia Date: Tue, 14 Aug 2012 21:47:12 +0530 Subject: [PATCH] OpenCL implementation of GPG format --- src/Makefile | 2 +- src/john.c | 2 + src/opencl/gpg_kernel.cl | 433 ++++++++++++++++++++++++++++++++++++++ src/opencl_gpg_fmt.c | 529 +++++++++++++++++++++++++++++++++++++++++++++++ 4 files changed, 965 insertions(+), 1 deletion(-) create mode 100644 src/opencl/gpg_kernel.cl create mode 100644 src/opencl_gpg_fmt.c diff --git a/src/Makefile b/src/Makefile index 7f3518e..ba0d768 100644 --- a/src/Makefile +++ b/src/Makefile @@ -136,7 +136,7 @@ OCL_OBJS = \ opencl_xsha512_fmt.o opencl_rawsha512_fmt.o opencl_bf_std.o \ opencl_bf_fmt.o opencl_pwsafe_fmt.o opencl_rawmd4_fmt.o \ opencl_keychain_fmt.o opencl_agilekeychain_fmt.o opencl_zip_fmt.o \ - opencl_encfs_fmt.o opencl_odf_fmt.o opencl_sxc_fmt.o + opencl_encfs_fmt.o opencl_odf_fmt.o opencl_sxc_fmt.o opencl_gpg_fmt.o CUDA_OBJS = \ cuda_common.o \ diff --git a/src/john.c b/src/john.c index cdb02bc..e14e82b 100644 --- a/src/john.c +++ b/src/john.c @@ -136,6 +136,7 @@ extern struct fmt_main fmt_opencl_zip; extern struct fmt_main fmt_opencl_encfs; extern struct fmt_main fmt_opencl_odf; extern struct fmt_main fmt_opencl_sxc; +extern struct fmt_main fmt_opencl_gpg; extern struct fmt_main fmt_opencl_xsha512; extern struct fmt_main fmt_opencl_rawsha512; extern struct fmt_main fmt_opencl_bf; @@ -279,6 +280,7 @@ static void john_register_all(void) john_register_one(&fmt_opencl_encfs); john_register_one(&fmt_opencl_odf); john_register_one(&fmt_opencl_sxc); + john_register_one(&fmt_opencl_gpg); john_register_one(&fmt_opencl_xsha512); john_register_one(&fmt_opencl_rawsha512); john_register_one(&fmt_opencl_bf); diff --git a/src/opencl/gpg_kernel.cl b/src/opencl/gpg_kernel.cl new file mode 100644 index 0000000..81a065a --- /dev/null +++ b/src/opencl/gpg_kernel.cl @@ -0,0 +1,433 @@ +/* + * FIPS-180-1 compliant SHA-1 implementation + * + * Copyright (C) 2006-2010, Brainspark B.V. + * + * This file is part of PolarSSL (http://www.polarssl.org) + * Lead Maintainer: Paul Bakker + * + * All rights reserved. + * + * This program is free software; you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation; either version 2 of the License, or + * (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License along + * with this program; if not, write to the Free Software Foundation, Inc., + * 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA. + */ +/* + * The SHA-1 standard was published by NIST in 1993. + * + * http://www.itl.nist.gov/fipspubs/fip180-1.htm + */ + +/* + * 32-bit integer manipulation macros (big endian) + */ + +#ifndef GET_ULONG_BE +#define GET_ULONG_BE(n,b,i) \ +{ \ + (n) = ( (unsigned long) (b)[(i) ] << 24 ) \ + | ( (unsigned long) (b)[(i) + 1] << 16 ) \ + | ( (unsigned long) (b)[(i) + 2] << 8 ) \ + | ( (unsigned long) (b)[(i) + 3] ); \ +} +#endif + +#ifndef PUT_ULONG_BE +#define PUT_ULONG_BE(n,b,i) \ +{ \ + (b)[(i) ] = (unsigned char) ( (n) >> 24 ); \ + (b)[(i) + 1] = (unsigned char) ( (n) >> 16 ); \ + (b)[(i) + 2] = (unsigned char) ( (n) >> 8 ); \ + (b)[(i) + 3] = (unsigned char) ( (n) ); \ +} +#endif + +void* _memcpy(void* dest, __global const uchar *src, int count) +{ + char* dst8 = (char*)dest; + __global uchar* src8 = (__global uchar*)src; + + while (count--) { + *dst8++ = *src8++; + } + return dest; +} + +void* _memcpy_(void* dest, const uchar *src, int count) +{ + char* dst8 = (char*)dest; + uchar* src8 = (uchar*)src; + + while (count--) { + *dst8++ = *src8++; + } + return dest; +} + + + +#define uint8_t unsigned char +#define uint16_t unsigned short +#define uint32_t unsigned int + +typedef struct { + uint8_t length; + uint8_t v[24]; +} gpg_password; + +typedef struct { + uint8_t v[16]; +} gpg_hash; + +typedef struct { + uint8_t length; + int count; + uint8_t salt[8]; +} gpg_salt; + + +/* + * SHA-1 context setup + */ + +typedef struct +{ + unsigned long total[2]; /*!< number of bytes processed */ + unsigned long state[5]; /*!< intermediate digest state */ + unsigned char buffer[64]; /*!< data block being processed */ +} +sha1_context; + +void sha1_starts( sha1_context *ctx ) +{ + ctx->total[0] = 0; + ctx->total[1] = 0; + + ctx->state[0] = 0x67452301; + ctx->state[1] = 0xEFCDAB89; + ctx->state[2] = 0x98BADCFE; + ctx->state[3] = 0x10325476; + ctx->state[4] = 0xC3D2E1F0; +} + +static void sha1_process( sha1_context *ctx, const unsigned char data[64] ) +{ + unsigned long temp, W[16], A, B, C, D, E; + + GET_ULONG_BE( W[ 0], data, 0 ); + GET_ULONG_BE( W[ 1], data, 4 ); + GET_ULONG_BE( W[ 2], data, 8 ); + GET_ULONG_BE( W[ 3], data, 12 ); + GET_ULONG_BE( W[ 4], data, 16 ); + GET_ULONG_BE( W[ 5], data, 20 ); + GET_ULONG_BE( W[ 6], data, 24 ); + GET_ULONG_BE( W[ 7], data, 28 ); + GET_ULONG_BE( W[ 8], data, 32 ); + GET_ULONG_BE( W[ 9], data, 36 ); + GET_ULONG_BE( W[10], data, 40 ); + GET_ULONG_BE( W[11], data, 44 ); + GET_ULONG_BE( W[12], data, 48 ); + GET_ULONG_BE( W[13], data, 52 ); + GET_ULONG_BE( W[14], data, 56 ); + GET_ULONG_BE( W[15], data, 60 ); + +#define S(x,n) ((x << n) | ((x & 0xFFFFFFFF) >> (32 - n))) + +#define R(t) \ +( \ + temp = W[(t - 3) & 0x0F] ^ W[(t - 8) & 0x0F] ^ \ + W[(t - 14) & 0x0F] ^ W[ t & 0x0F], \ + ( W[t & 0x0F] = S(temp,1) ) \ +) + +#define P(a,b,c,d,e,x) \ +{ \ + e += S(a,5) + F(b,c,d) + K + x; b = S(b,30); \ +} + + A = ctx->state[0]; + B = ctx->state[1]; + C = ctx->state[2]; + D = ctx->state[3]; + E = ctx->state[4]; + +#define F(x,y,z) (z ^ (x & (y ^ z))) +#define K 0x5A827999 + + P( A, B, C, D, E, W[0] ); + P( E, A, B, C, D, W[1] ); + P( D, E, A, B, C, W[2] ); + P( C, D, E, A, B, W[3] ); + P( B, C, D, E, A, W[4] ); + P( A, B, C, D, E, W[5] ); + P( E, A, B, C, D, W[6] ); + P( D, E, A, B, C, W[7] ); + P( C, D, E, A, B, W[8] ); + P( B, C, D, E, A, W[9] ); + P( A, B, C, D, E, W[10] ); + P( E, A, B, C, D, W[11] ); + P( D, E, A, B, C, W[12] ); + P( C, D, E, A, B, W[13] ); + P( B, C, D, E, A, W[14] ); + P( A, B, C, D, E, W[15] ); + P( E, A, B, C, D, R(16) ); + P( D, E, A, B, C, R(17) ); + P( C, D, E, A, B, R(18) ); + P( B, C, D, E, A, R(19) ); + +#undef K +#undef F + +#define F(x,y,z) (x ^ y ^ z) +#define K 0x6ED9EBA1 + + P( A, B, C, D, E, R(20) ); + P( E, A, B, C, D, R(21) ); + P( D, E, A, B, C, R(22) ); + P( C, D, E, A, B, R(23) ); + P( B, C, D, E, A, R(24) ); + P( A, B, C, D, E, R(25) ); + P( E, A, B, C, D, R(26) ); + P( D, E, A, B, C, R(27) ); + P( C, D, E, A, B, R(28) ); + P( B, C, D, E, A, R(29) ); + P( A, B, C, D, E, R(30) ); + P( E, A, B, C, D, R(31) ); + P( D, E, A, B, C, R(32) ); + P( C, D, E, A, B, R(33) ); + P( B, C, D, E, A, R(34) ); + P( A, B, C, D, E, R(35) ); + P( E, A, B, C, D, R(36) ); + P( D, E, A, B, C, R(37) ); + P( C, D, E, A, B, R(38) ); + P( B, C, D, E, A, R(39) ); + +#undef K +#undef F + +#define F(x,y,z) ((x & y) | (z & (x | y))) +#define K 0x8F1BBCDC + + P( A, B, C, D, E, R(40) ); + P( E, A, B, C, D, R(41) ); + P( D, E, A, B, C, R(42) ); + P( C, D, E, A, B, R(43) ); + P( B, C, D, E, A, R(44) ); + P( A, B, C, D, E, R(45) ); + P( E, A, B, C, D, R(46) ); + P( D, E, A, B, C, R(47) ); + P( C, D, E, A, B, R(48) ); + P( B, C, D, E, A, R(49) ); + P( A, B, C, D, E, R(50) ); + P( E, A, B, C, D, R(51) ); + P( D, E, A, B, C, R(52) ); + P( C, D, E, A, B, R(53) ); + P( B, C, D, E, A, R(54) ); + P( A, B, C, D, E, R(55) ); + P( E, A, B, C, D, R(56) ); + P( D, E, A, B, C, R(57) ); + P( C, D, E, A, B, R(58) ); + P( B, C, D, E, A, R(59) ); + +#undef K +#undef F + +#define F(x,y,z) (x ^ y ^ z) +#define K 0xCA62C1D6 + + P( A, B, C, D, E, R(60) ); + P( E, A, B, C, D, R(61) ); + P( D, E, A, B, C, R(62) ); + P( C, D, E, A, B, R(63) ); + P( B, C, D, E, A, R(64) ); + P( A, B, C, D, E, R(65) ); + P( E, A, B, C, D, R(66) ); + P( D, E, A, B, C, R(67) ); + P( C, D, E, A, B, R(68) ); + P( B, C, D, E, A, R(69) ); + P( A, B, C, D, E, R(70) ); + P( E, A, B, C, D, R(71) ); + P( D, E, A, B, C, R(72) ); + P( C, D, E, A, B, R(73) ); + P( B, C, D, E, A, R(74) ); + P( A, B, C, D, E, R(75) ); + P( E, A, B, C, D, R(76) ); + P( D, E, A, B, C, R(77) ); + P( C, D, E, A, B, R(78) ); + P( B, C, D, E, A, R(79) ); + +#undef K +#undef F + + ctx->state[0] += A; + ctx->state[1] += B; + ctx->state[2] += C; + ctx->state[3] += D; + ctx->state[4] += E; +} + +/* + * SHA-1 process buffer + */ +void sha1_update( sha1_context *ctx, const unsigned char *input, int ilen ) +{ + int fill; + unsigned long left; + + if( ilen <= 0 ) + return; + + left = ctx->total[0] & 0x3F; + fill = 64 - left; + + ctx->total[0] += (unsigned long) ilen; + ctx->total[0] &= 0xFFFFFFFF; + + if( ctx->total[0] < (unsigned long) ilen ) + ctx->total[1]++; + + if( left && ilen >= fill ) + { + _memcpy_( (void *) (ctx->buffer + left), + input, fill ); + sha1_process( ctx, ctx->buffer ); + input += fill; + ilen -= fill; + left = 0; + } + + while( ilen >= 64 ) + { + sha1_process( ctx, input ); + input += 64; + ilen -= 64; + } + + if( ilen > 0 ) + { + _memcpy_( (void *) (ctx->buffer + left), + input, ilen ); + } +} + +static const unsigned char sha1_padding[64] = +{ + 0x80, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, + 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 +}; + +/* + * SHA-1 final digest + */ +void sha1_finish( sha1_context *ctx, unsigned char output[20] ) +{ + unsigned long last, padn; + unsigned long high, low; + unsigned char msglen[8]; + + high = ( ctx->total[0] >> 29 ) + | ( ctx->total[1] << 3 ); + low = ( ctx->total[0] << 3 ); + + PUT_ULONG_BE( high, msglen, 0 ); + PUT_ULONG_BE( low, msglen, 4 ); + + last = ctx->total[0] & 0x3F; + padn = ( last < 56 ) ? ( 56 - last ) : ( 120 - last ); + + sha1_update( ctx, (unsigned char *) sha1_padding, padn ); + sha1_update( ctx, msglen, 8 ); + + PUT_ULONG_BE( ctx->state[0], output, 0 ); + PUT_ULONG_BE( ctx->state[1], output, 4 ); + PUT_ULONG_BE( ctx->state[2], output, 8 ); + PUT_ULONG_BE( ctx->state[3], output, 12 ); + PUT_ULONG_BE( ctx->state[4], output, 16 ); +} + +/* + * output = SHA-1( input buffer ) + */ +void sha1( const unsigned char *input, int ilen, unsigned char output[20] ) +{ + sha1_context ctx; + + sha1_starts( &ctx ); + sha1_update( &ctx, input, ilen ); + sha1_finish( &ctx, output ); + + // memset( &ctx, 0, sizeof( sha1_context ) ); +} + +#define KEYBUFFER_LENGTH 8192 +#define SHA_DIGEST_LENGTH 20 + +static void S2KItSaltedSHA1Generator(__global const uchar *password, int password_length, __global const uchar *salt, int count, __global uchar *key, int length) +{ + unsigned char keybuf[KEYBUFFER_LENGTH]; + sha1_context ctx; + int i, j; + int tl; + int mul; + int bs; + unsigned char *bptr; + int n; + uchar lkey[8192]; + _memcpy(lkey, key, length); + + unsigned int numHashes = (length + SHA_DIGEST_LENGTH - 1) / SHA_DIGEST_LENGTH; + _memcpy(keybuf, salt, 8); + unsigned char wtf = '\0'; + + // TODO: This is not very efficient with multiple hashes + for (i = 0; i < numHashes; i++) { + sha1_starts(&ctx); + for (j = 0; j < i; j++) { + sha1_update(&ctx, &wtf, 1); + } + // Find multiplicator + tl = password_length + 8; + mul = 1; + while (mul < tl && ((64 * mul) % tl)) { + ++mul; + } + // Try to feed the hash function with 64-byte blocks + bs = mul * 64; + bptr = keybuf + tl; + n = bs / tl; + _memcpy(keybuf + 8, password, password_length); + while (n-- > 1) { + _memcpy_(bptr, keybuf, tl); + bptr += tl; + } + n = count / bs; + while (n-- > 0) { + sha1_update(&ctx, keybuf, bs); + } + sha1_update(&ctx, keybuf, count % bs); + sha1_finish(&ctx, lkey + (i * SHA_DIGEST_LENGTH)); + } + for(i = 0; i < length; i++) + key[i] = lkey[i]; +} + +__kernel void gpg(__global const gpg_password * inbuffer, + __global gpg_hash * outbuffer, __global const gpg_salt * salt) +{ + uint32_t idx = get_global_id(0); + + S2KItSaltedSHA1Generator(inbuffer[idx].v, inbuffer[idx].length, + salt->salt, salt->count, outbuffer[idx].v, 16); +} diff --git a/src/opencl_gpg_fmt.c b/src/opencl_gpg_fmt.c new file mode 100644 index 0000000..aeff84f --- /dev/null +++ b/src/opencl_gpg_fmt.c @@ -0,0 +1,529 @@ +/* + * Modified by Dhiru Kholia for GPG format. + * + * This software is Copyright (c) 2012 Lukas Odzioba + * and it is hereby released to the general public under the following terms: + * Redistribution and use in source and binary forms, with or without + * modification, are permitted. */ + +#include +#include "arch.h" +#include "formats.h" +#include "common.h" +#include "misc.h" +#include +#include "common-opencl.h" +#include +#include +#include +#include +#include +#include "sha2.h" +#ifdef _OPENMP +#include +#endif + +#define FORMAT_LABEL "gpg-opencl" +#define FORMAT_NAME "OpenPGP / GnuPG Secret Key" +#define ALGORITHM_NAME "OpenCL" +#define BENCHMARK_COMMENT "" +#define BENCHMARK_LENGTH -1 +#define KEYS_PER_CRYPT 1024*9 +#define MIN_KEYS_PER_CRYPT KEYS_PER_CRYPT +#define MAX_KEYS_PER_CRYPT KEYS_PER_CRYPT +#define BINARY_SIZE 16 +#define PLAINTEXT_LENGTH 15 +#define SALT_SIZE sizeof(struct custom_salt) + +#define uint8_t unsigned char +#define uint16_t unsigned short +#define uint32_t unsigned int + +typedef struct { + uint8_t length; + uint8_t v[24]; +} gpg_password; + +typedef struct { + uint8_t v[16]; +} gpg_hash; + +typedef struct { + uint8_t length; + int count; + uint8_t salt[8]; +} gpg_salt; + +static int *cracked; + + +#define KEYBUFFER_LENGTH 8192 +#ifndef MD5_DIGEST_LENGTH +#define MD5_DIGEST_LENGTH 16 +#endif + +// Minimum number of bits when checking the first BN +#define MIN_BN_BITS 64 + +static int *cracked; + +enum { + SPEC_SIMPLE = 0, + SPEC_SALTED = 1, + SPEC_ITERATED_SALTED = 3 +}; + + +enum { + PKA_UNKOWN = 0, + PKA_RSA_ENCSIGN = 1, + PKA_DSA = 17 +}; + +enum { + CIPHER_UNKOWN = -1, + CIPHER_CAST5 = 3, + CIPHER_BLOWFISH = 4, + CIPHER_AES128 = 7, + CIPHER_AES192 = 8, + CIPHER_AES256 = 9 +}; + +enum { + HASH_UNKOWN = -1, + HASH_MD5 = 1, + HASH_SHA1 = 2, + HASH_RIPEMD160 = 3, + HASH_SHA256 = 8, + HASH_SHA384 = 9, + HASH_SHA512 = 10, + HASH_SHA224 = 11 +}; + +static struct custom_salt { + int datalen; + unsigned char data[4096]; + char spec; + char pk_algorithm; + char hash_algorithm; + char cipher_algorithm; + int usage; + int bits; + unsigned char salt[8]; + unsigned char iv[16]; + int ivlen; + int count; + void (*s2kfun)(char *, unsigned char*, int); +} *cur_salt; + +static struct fmt_tests gpg_tests[] = { + {"$gpg$*1*667*2048*387de4c9e2c1018aed84af75922ecaa92d1bc68d48042144c77dfe168de1fd654e4db77bfbc60ec68f283483382413cbfddddcfad714922b2d558f8729f705fbf973ab1839e756c26207a4bc8796eeb567bf9817f73a2a81728d3e4bc0894f62ad96e04e60752d84ebc01316703b0fd0f618f6120289373347027924606712610c583b25be57c8a130bc4dd796964f3f03188baa057d6b8b1fd36675af94d45847eeefe7fff63b755a32e8abe26b7f3f58bb091e5c7b9250afe2180b3d0abdd2c1db3d4fffe25e17d5b7d5b79367d98c523a6c280aafef5c1975a42fd97242ba86ced73c5e1a9bcab82adadd11ef2b64c3aad23bc930e62fc8def6b1d362e954795d87fa789e5bc2807bfdc69bba7e66065e3e3c2df0c25eab0fde39fbe54f32b26f07d88f8b05202e55874a1fa37d540a5af541e28370f27fe094ca8758cd7ff7b28df1cbc475713d7604b1af22fd758ebb3a83876ed83f003285bc8fdc7a5470f7c5a9e8a93929941692a9ff9f1bc146dcc02aab47e2679297d894f28b62da16c8baa95cd393d838fa63efc9d3f88de93dc970c67022d5dc88dce25decec8848f8e6f263d7c2c0238d36aa0013d7edefd43dac1299a54eb460d9b82cb53cf86fcb7c8d5dba95795a1adeb729a705b47b8317594ac3906424b2c0e425343eca019e53d927e6bc32688bd9e87ee808fb1d8eeee8ab938855131b839776c7da79a33a6d66e57eadb430ef04809009794e32a03a7e030b8792be5d53ceaf480ffd98633d1993c43f536a90bdbec8b9a827d0e0a49155450389beb53af5c214c4ec09712d83b175671358d8e9d54da7a8187f72aaaca5203372841af9b89a07b8aadecafc0f2901b8aec13a5382c6f94712d629333b301afdf52bdfa62534de2b10078cd4d0e781c88efdfe4e5252e39a236af449d4d62081cee630ab*3*254*2*3*8*b1fdf3772bb57e1f*65536*2127ccd55e721ba0", "polished"}, + {NULL} +}; + +static gpg_password *inbuffer; +static gpg_hash *outbuffer; +static gpg_salt currentsalt; +static cl_mem mem_in, mem_out, mem_setting; +static size_t insize = sizeof(gpg_password) * KEYS_PER_CRYPT; +static size_t outsize = sizeof(gpg_hash) * KEYS_PER_CRYPT; +static size_t settingsize = sizeof(gpg_salt); + +// Returns the block size (in bytes) of a given cipher +static uint32_t blockSize(char algorithm) +{ + switch (algorithm) { + case CIPHER_CAST5: + return CAST_BLOCK; + case CIPHER_BLOWFISH: + return BF_BLOCK; + case CIPHER_AES128: + case CIPHER_AES192: + case CIPHER_AES256: + return AES_BLOCK_SIZE; + default: break; + } + return 0; +} + +// Returns the key size (in bytes) of a given cipher +static uint32_t keySize(char algorithm) +{ + switch (algorithm) { + case CIPHER_CAST5: + return CAST_KEY_LENGTH; + case CIPHER_BLOWFISH: + return 16; + case CIPHER_AES128: + return 16; + case CIPHER_AES192: + return 24; + case CIPHER_AES256: + return 32; + default: break; + } + return 0; +} + +// Returns the digest size (in bytes) of a given hash algorithm +static uint32_t digestSize(char algorithm) +{ + switch (algorithm) { + case HASH_MD5: + return 16; + case HASH_SHA1: + return 20; + case HASH_SHA512: + return 64; + case HASH_SHA256: + return 32; + case HASH_RIPEMD160: + return 20; + default: break; + } + return 0; +} + +static void release_all(void) +{ + HANDLE_CLERROR(clReleaseKernel(crypt_kernel), "Release Kernel"); + HANDLE_CLERROR(clReleaseMemObject(mem_in), "Release mem in"); + HANDLE_CLERROR(clReleaseMemObject(mem_setting), "Release mem setting"); + HANDLE_CLERROR(clReleaseMemObject(mem_out), "Release mem out"); + HANDLE_CLERROR(clReleaseCommandQueue(queue[gpu_id]), "Release Queue"); +} +static void init(struct fmt_main *self) +{ + cl_int cl_error; + + global_work_size = MAX_KEYS_PER_CRYPT; + + inbuffer = + (gpg_password *) malloc(sizeof(gpg_password) * + MAX_KEYS_PER_CRYPT); + outbuffer = + (gpg_hash *) malloc(sizeof(gpg_hash) * MAX_KEYS_PER_CRYPT); + + /* Zeroize the lengths in case crypt_all() is called with some keys still + * not set. This may happen during self-tests. */ + { + int i; + for (i = 0; i < MAX_KEYS_PER_CRYPT; i++) + inbuffer[i].length = 0; + } + + cracked = mem_calloc_tiny(sizeof(*cracked) * + KEYS_PER_CRYPT, MEM_ALIGN_WORD); + + //listOpenCLdevices(); + opencl_init("$JOHN/gpg_kernel.cl", gpu_id, platform_id); + /// Alocate memory + mem_in = + clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, insize, NULL, + &cl_error); + HANDLE_CLERROR(cl_error, "Error alocating mem in"); + mem_setting = + clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, settingsize, + NULL, &cl_error); + HANDLE_CLERROR(cl_error, "Error alocating mem setting"); + mem_out = + clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY, outsize, NULL, + &cl_error); + HANDLE_CLERROR(cl_error, "Error alocating mem out"); + + crypt_kernel = clCreateKernel(program[gpu_id], "gpg", &cl_error); + HANDLE_CLERROR(cl_error, "Error creating kernel"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 0, sizeof(mem_in), + &mem_in), "Error while setting mem_in kernel argument"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1, sizeof(mem_out), + &mem_out), "Error while setting mem_out kernel argument"); + HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 2, sizeof(mem_setting), + &mem_setting), "Error while setting mem_salt kernel argument"); + opencl_find_best_workgroup(self); + + atexit(release_all); +} + +static int valid(char *ciphertext, struct fmt_main *self) +{ + return !strncmp(ciphertext, "$gpg$", 5); +} + +static void *get_salt(char *ciphertext) +{ + char *ctcopy = strdup(ciphertext); + char *keeptr = ctcopy; + int i; + char *p; + static struct custom_salt cs; + ctcopy += 5; /* skip over "$gpg$" marker */ + p = strtok(ctcopy, "*"); + cs.pk_algorithm = atoi(p); + p = strtok(NULL, "*"); + cs.datalen = atoi(p); + p = strtok(NULL, "*"); + cs.bits = atoi(p); + p = strtok(NULL, "*"); + for (i = 0; i < cs.datalen; i++) + cs.data[i] = + atoi16[ARCH_INDEX(p[i * 2])] * 16 + + atoi16[ARCH_INDEX(p[i * 2 + 1])]; + p = strtok(NULL, "*"); + cs.spec = atoi(p); + p = strtok(NULL, "*"); + cs.usage = atoi(p); + p = strtok(NULL, "*"); + cs.hash_algorithm = atoi(p); + p = strtok(NULL, "*"); + cs.cipher_algorithm = atoi(p); + p = strtok(NULL, "*"); + cs.ivlen = atoi(p); + p = strtok(NULL, "*"); + for (i = 0; i < cs.ivlen; i++) + cs.iv[i] = + atoi16[ARCH_INDEX(p[i * 2])] * 16 + + atoi16[ARCH_INDEX(p[i * 2 + 1])]; + p = strtok(NULL, "*"); + cs.count = atoi(p); + p = strtok(NULL, "*"); + for (i = 0; i < 8; i++) + cs.salt[i] = + atoi16[ARCH_INDEX(p[i * 2])] * 16 + + atoi16[ARCH_INDEX(p[i * 2 + 1])]; + free(keeptr); + return (void *)&cs; +} + + +static void set_salt(void *salt) +{ + cur_salt = (struct custom_salt *)salt; + memcpy((char*)currentsalt.salt, cur_salt->salt, 8); + currentsalt.length = 8;; + currentsalt.count = cur_salt->count; +} + +#undef set_key +static void set_key(char *key, int index) +{ + uint8_t length = strlen(key); + if (length > PLAINTEXT_LENGTH) + length = PLAINTEXT_LENGTH; + inbuffer[index].length = length; + memcpy(inbuffer[index].v, key, length); +} + +static char *get_key(int index) +{ + static char ret[PLAINTEXT_LENGTH + 1]; + uint8_t length = inbuffer[index].length; + memcpy(ret, inbuffer[index].v, length); + ret[length] = '\0'; + return ret; +} + + +static int check(unsigned char *keydata, int ks) +{ + // Decrypt first data block in order to check the first two bits of + // the MPI. If they are correct, there's a good chance that the + // password is correct, too. + unsigned char ivec[32]; + unsigned char out[4096]; + int tmp = 0; + uint32_t num_bits; + int checksumOk; + int i; + + // Quick Hack + memcpy(ivec, cur_salt->iv, blockSize(cur_salt->cipher_algorithm)); + switch (cur_salt->cipher_algorithm) { + case CIPHER_CAST5: { + CAST_KEY ck; + CAST_set_key(&ck, ks, keydata); + CAST_cfb64_encrypt(cur_salt->data, out, CAST_BLOCK, &ck, ivec, &tmp, CAST_DECRYPT); + } + break; + case CIPHER_BLOWFISH: { + BF_KEY ck; + BF_set_key(&ck, ks, keydata); + BF_cfb64_encrypt(cur_salt->data, out, BF_BLOCK, &ck, ivec, &tmp, BF_DECRYPT); + } + break; + case CIPHER_AES128: + case CIPHER_AES192: + case CIPHER_AES256: { + AES_KEY ck; + AES_set_encrypt_key(keydata, ks * 8, &ck); + AES_cfb128_encrypt(cur_salt->data, out, AES_BLOCK_SIZE, &ck, ivec, &tmp, AES_DECRYPT); + } + break; + default: + break; + } + num_bits = ((out[0] << 8) | out[1]); + if (num_bits < MIN_BN_BITS || num_bits > cur_salt->bits) { + return 0; + } + // Decrypt all data + memcpy(ivec, cur_salt->iv, blockSize(cur_salt->cipher_algorithm)); + tmp = 0; + switch (cur_salt->cipher_algorithm) { + case CIPHER_CAST5: { + CAST_KEY ck; + CAST_set_key(&ck, ks, keydata); + CAST_cfb64_encrypt(cur_salt->data, out, cur_salt->datalen, &ck, ivec, &tmp, CAST_DECRYPT); + } + break; + case CIPHER_BLOWFISH: { + BF_KEY ck; + BF_set_key(&ck, ks, keydata); + BF_cfb64_encrypt(cur_salt->data, out, cur_salt->datalen, &ck, ivec, &tmp, BF_DECRYPT); + } + break; + case CIPHER_AES128: + case CIPHER_AES192: + case CIPHER_AES256: { + AES_KEY ck; + AES_set_encrypt_key(keydata, ks * 8, &ck); + AES_cfb128_encrypt(cur_salt->data, out, cur_salt->datalen, &ck, ivec, &tmp, AES_DECRYPT); + } + break; + default: + break; + } + + // Verify + checksumOk = 0; + switch (cur_salt->usage) { + case 254: { + uint8_t checksum[SHA_DIGEST_LENGTH]; + SHA_CTX ctx; + SHA1_Init(&ctx); + SHA1_Update(&ctx, out, cur_salt->datalen - SHA_DIGEST_LENGTH); + SHA1_Final(checksum, &ctx); + if (memcmp(checksum, out + cur_salt->datalen - SHA_DIGEST_LENGTH, SHA_DIGEST_LENGTH) == 0) { + checksumOk = 1; + } + } break; + case 0: + case 255: { + uint16_t sum = 0; + for (i = 0; i < cur_salt->datalen - 2; i++) { + sum += out[i]; + } + if (sum == ((out[cur_salt->datalen - 2] << 8) | out[cur_salt->datalen - 1])) { + checksumOk = 1; + } + } break; + default: + break; + } + // If the checksum is ok, try to parse the first MPI of the private key + if (checksumOk) { + BIGNUM *b = NULL; + uint32_t blen = (num_bits + 7) / 8; + if (blen < cur_salt->datalen && ((b = BN_bin2bn(out + 2, blen, NULL)) != NULL)) { + BN_free(b); + return 1; + } + } + return 0; +} + +static void crypt_all(int count) +{ + int index; + /// Copy data to gpu + HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_in, CL_FALSE, 0, + insize, inbuffer, 0, NULL, NULL), "Copy data to gpu"); + HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_setting, + CL_FALSE, 0, settingsize, ¤tsalt, 0, NULL, NULL), + "Copy setting to gpu"); + + /// Run kernel + HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1, + NULL, &global_work_size, &local_work_size, 0, NULL, &profilingEvent), + "Run kernel"); + HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish"); + + /// Read the result back + HANDLE_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_out, CL_FALSE, 0, + outsize, outbuffer, 0, NULL, NULL), "Copy result back"); + + /// Await completion of all the above + HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish"); + + for (index = 0; index < count; index++) + { + // allocate string2key buffer + int res; + int ks = keySize(cur_salt->cipher_algorithm); + int ds = digestSize(cur_salt->hash_algorithm); + unsigned char keydata[ds * ((ks + ds- 1) / ds)]; + memcpy(keydata, outbuffer[index].v, ks); + res = check(keydata, ks); + if(res) { + cracked[index] = 1; + } + } + +} + +static int cmp_all(void *binary, int count) +{ + int index; + for (index = 0; index < count; index++) + if (cracked[index]) + return 1; + return 0; +} + +static int cmp_one(void *binary, int index) +{ + return cracked[index]; +} + +static int cmp_exact(char *source, int index) +{ + return 1; +} + +struct fmt_main fmt_opencl_gpg = { + { + FORMAT_LABEL, + FORMAT_NAME, + ALGORITHM_NAME, + BENCHMARK_COMMENT, + BENCHMARK_LENGTH, + PLAINTEXT_LENGTH, + BINARY_SIZE, + SALT_SIZE, + MIN_KEYS_PER_CRYPT, + MAX_KEYS_PER_CRYPT, + FMT_CASE | FMT_8_BIT | FMT_OMP | FMT_NOT_EXACT, + gpg_tests + }, { + init, + fmt_default_prepare, + valid, + fmt_default_split, + fmt_default_binary, + get_salt, + { + fmt_default_binary_hash + }, + fmt_default_salt_hash, + set_salt, + set_key, + get_key, + fmt_default_clear_keys, + crypt_all, + { + fmt_default_get_hash + }, + cmp_all, + cmp_one, + cmp_exact + } +}; -- 1.7.11.4