From 163a2e2ee14758644162dcdeedbc14f0657a8d93 Mon Sep 17 00:00:00 2001 From: Dhiru Kholia Date: Tue, 24 Jul 2012 07:51:15 +0530 Subject: [PATCH] OpenCL implementation of ODF format --- src/Makefile | 9 +- src/john.c | 3 + src/opencl/odf_kernel.cl | 516 +++++++++++++++++++++++++++++++++++++++++++++++ src/opencl_odf_fmt.c | 400 ++++++++++++++++++++++++++++++++++++ 4 files changed, 924 insertions(+), 4 deletions(-) create mode 100644 src/opencl/odf_kernel.cl create mode 100644 src/opencl_odf_fmt.c diff --git a/src/Makefile b/src/Makefile index 1c84b48..efade1b 100644 --- a/src/Makefile +++ b/src/Makefile @@ -25,7 +25,7 @@ CPPFLAGS = -E OMPFLAGS = # gcc with OpenMP -#OMPFLAGS = -fopenmp +OMPFLAGS = -fopenmp #OMPFLAGS = -fopenmp -msse2 # Sun Studio with OpenMP (set the OMP_NUM_THREADS env var at runtime) #OMPFLAGS = -xopenmp @@ -57,12 +57,12 @@ ifdef HAVE_NSS NSS_LDFLAGS = `pkg-config --libs nss` endif -CFLAGS = -c -Wall -O2 -fomit-frame-pointer -Wdeclaration-after-statement -I/usr/local/include $(HAVE_NSS) $(OMPFLAGS) $(JOHN_CFLAGS) $(AMDAPP) +CFLAGS = -c -Wall -ggdb -fomit-frame-pointer -Wdeclaration-after-statement -I/usr/local/include $(HAVE_NSS) $(OMPFLAGS) $(JOHN_CFLAGS) $(AMDAPP) # -DHAVE_SKEY # CFLAGS for use on the main john.c file only CFLAGS_MAIN = $(CFLAGS) ASFLAGS = -c $(JOHN_CFLAGS) $(OMPFLAGS) -LDFLAGS = -s -L/usr/local/lib -L/usr/local/ssl/lib -lssl -lcrypto -lm -lz $(JOHN_CFLAGS) $(OMPFLAGS) $(NSS_LDFLAGS) +LDFLAGS = -L/usr/local/lib -L/usr/local/ssl/lib -lssl -lcrypto -lm -lz $(JOHN_CFLAGS) $(OMPFLAGS) $(NSS_LDFLAGS) # -lskey LDFLAGS_SOLARIS = -lrt -lnsl -lsocket -lm -lz -lcrypto -lssl LDFLAGS_MKV = -s -lm @@ -134,7 +134,8 @@ OCL_OBJS = \ opencl_cryptsha512_fmt.o opencl_mscash2_fmt.o opencl_wpapsk_fmt.o \ 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_keychain_fmt.o opencl_agilekeychain_fmt.o opencl_zip_fmt.o \ + opencl_odf_fmt.o CUDA_OBJS = \ cuda_common.o \ diff --git a/src/john.c b/src/john.c index 762f9d1..4948b29 100644 --- a/src/john.c +++ b/src/john.c @@ -133,6 +133,7 @@ extern struct fmt_main fmt_opencl_wpapsk; extern struct fmt_main fmt_opencl_keychain; extern struct fmt_main fmt_opencl_agilekeychain; extern struct fmt_main fmt_opencl_zip; +extern struct fmt_main fmt_opencl_odf; extern struct fmt_main fmt_opencl_xsha512; extern struct fmt_main fmt_opencl_rawsha512; extern struct fmt_main fmt_opencl_bf; @@ -159,6 +160,7 @@ extern struct fmt_main fmt_pdf; extern struct fmt_main rar_fmt; extern struct fmt_main zip_fmt; extern struct fmt_main fmt_wpapsk; +extern struct fmt_main fmt_wpapsk; #include "fmt_externs.h" @@ -271,6 +273,7 @@ static void john_register_all(void) john_register_one(&fmt_opencl_keychain); john_register_one(&fmt_opencl_agilekeychain); john_register_one(&fmt_opencl_zip); + john_register_one(&fmt_opencl_odf); john_register_one(&fmt_opencl_xsha512); john_register_one(&fmt_opencl_rawsha512); john_register_one(&fmt_opencl_bf); diff --git a/src/opencl/odf_kernel.cl b/src/opencl/odf_kernel.cl new file mode 100644 index 0000000..9e2d579 --- /dev/null +++ b/src/opencl/odf_kernel.cl @@ -0,0 +1,516 @@ +/* + * Modified by Dhiru Kholia for ODF 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. */ + +#define uint8_t unsigned char +#define uint16_t unsigned short +#define uint32_t unsigned int + +typedef struct { + uint8_t length; + uint8_t v[24]; +} odf_password; + +typedef struct { + uint32_t v[17]; // 16*4=64 +} odf_hash; + +typedef struct { + uint8_t length; + uint8_t salt[64]; +} odf_salt; + + +# define SWAP(n) \ + (((n) << 24) | (((n) & 0xff00) << 8) | (((n) >> 8) & 0xff00) | ((n) >> 24)) + +#define ITERATIONS 1024 + +#define INIT_A 0x67452301 +#define INIT_B 0xefcdab89 +#define INIT_C 0x98badcfe +#define INIT_D 0x10325476 +#define INIT_E 0xc3d2e1f0 + +#define SQRT_2 0x5a827999 +#define SQRT_3 0x6ed9eba1 + +#define SHA1_DIGEST_LENGTH 20 + +#define K1 0x5a827999 +#define K2 0x6ed9eba1 +#define K3 0x8f1bbcdc +#define K4 0xca62c1d6 + +#define F1(x,y,z) (z ^ (x & (y ^ z))) +#define F2(x,y,z) (x ^ y ^ z) +#define F3(x,y,z) ((x & y) | (z & (x | y))) +#define F4(x,y,z) (x ^ y ^ z) + +#ifndef GET_WORD_32_BE +#define GET_WORD_32_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_WORD_32_BE +#define PUT_WORD_32_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 + +#define S(x,n) ((x << n) | ((x) >> (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 R2(t) \ +( \ + S((W[(t - 3) & 0x0F] ^ W[(t - 8) & 0x0F] ^ \ + W[(t - 14) & 0x0F] ^ W[ t & 0x0F]),1) \ +) + +#define P1(a,b,c,d,e,x) \ +{ \ + e += S(a,5) + F1(b,c,d) + K1 + x; b = S(b,30); \ +} + +#define P2(a,b,c,d,e,x) \ +{ \ + e += S(a,5) + F2(b,c,d) + K2 + x; b = S(b,30); \ +} + +#define P3(a,b,c,d,e,x) \ +{ \ + e += S(a,5) + F3(b,c,d) + K3 + x; b = S(b,30); \ +} + +#define P4(a,b,c,d,e,x) \ +{ \ + e += S(a,5) + F4(b,c,d) + K4 + x; b = S(b,30); \ +} + +#define PZ(a,b,c,d,e) \ +{ \ + e += S(a,5) + F1(b,c,d) + K1 ; b = S(b,30); \ +} + +#define SHA1(A,B,C,D,E,W) \ + P1(A, B, C, D, E, W[0] );\ + P1(E, A, B, C, D, W[1] );\ + P1(D, E, A, B, C, W[2] );\ + P1(C, D, E, A, B, W[3] );\ + P1(B, C, D, E, A, W[4] );\ + P1(A, B, C, D, E, W[5] );\ + P1(E, A, B, C, D, W[6] );\ + P1(D, E, A, B, C, W[7] );\ + P1(C, D, E, A, B, W[8] );\ + P1(B, C, D, E, A, W[9] );\ + P1(A, B, C, D, E, W[10]);\ + P1(E, A, B, C, D, W[11]);\ + P1(D, E, A, B, C, W[12]);\ + P1(C, D, E, A, B, W[13]);\ + P1(B, C, D, E, A, W[14]);\ + P1(A, B, C, D, E, W[15]);\ + P1(E, A, B, C, D, R(16));\ + P1(D, E, A, B, C, R(17));\ + P1(C, D, E, A, B, R(18));\ + P1(B, C, D, E, A, R(19));\ + P2(A, B, C, D, E, R(20));\ + P2(E, A, B, C, D, R(21));\ + P2(D, E, A, B, C, R(22));\ + P2(C, D, E, A, B, R(23));\ + P2(B, C, D, E, A, R(24));\ + P2(A, B, C, D, E, R(25));\ + P2(E, A, B, C, D, R(26));\ + P2(D, E, A, B, C, R(27));\ + P2(C, D, E, A, B, R(28));\ + P2(B, C, D, E, A, R(29));\ + P2(A, B, C, D, E, R(30));\ + P2(E, A, B, C, D, R(31));\ + P2(D, E, A, B, C, R(32));\ + P2(C, D, E, A, B, R(33));\ + P2(B, C, D, E, A, R(34));\ + P2(A, B, C, D, E, R(35));\ + P2(E, A, B, C, D, R(36));\ + P2(D, E, A, B, C, R(37));\ + P2(C, D, E, A, B, R(38));\ + P2(B, C, D, E, A, R(39));\ + P3(A, B, C, D, E, R(40));\ + P3(E, A, B, C, D, R(41));\ + P3(D, E, A, B, C, R(42));\ + P3(C, D, E, A, B, R(43));\ + P3(B, C, D, E, A, R(44));\ + P3(A, B, C, D, E, R(45));\ + P3(E, A, B, C, D, R(46));\ + P3(D, E, A, B, C, R(47));\ + P3(C, D, E, A, B, R(48));\ + P3(B, C, D, E, A, R(49));\ + P3(A, B, C, D, E, R(50));\ + P3(E, A, B, C, D, R(51));\ + P3(D, E, A, B, C, R(52));\ + P3(C, D, E, A, B, R(53));\ + P3(B, C, D, E, A, R(54));\ + P3(A, B, C, D, E, R(55));\ + P3(E, A, B, C, D, R(56));\ + P3(D, E, A, B, C, R(57));\ + P3(C, D, E, A, B, R(58));\ + P3(B, C, D, E, A, R(59));\ + P4(A, B, C, D, E, R(60));\ + P4(E, A, B, C, D, R(61));\ + P4(D, E, A, B, C, R(62));\ + P4(C, D, E, A, B, R(63));\ + P4(B, C, D, E, A, R(64));\ + P4(A, B, C, D, E, R(65));\ + P4(E, A, B, C, D, R(66));\ + P4(D, E, A, B, C, R(67));\ + P4(C, D, E, A, B, R(68));\ + P4(B, C, D, E, A, R(69));\ + P4(A, B, C, D, E, R(70));\ + P4(E, A, B, C, D, R(71));\ + P4(D, E, A, B, C, R(72));\ + P4(C, D, E, A, B, R(73));\ + P4(B, C, D, E, A, R(74));\ + P4(A, B, C, D, E, R(75));\ + P4(E, A, B, C, D, R(76));\ + P4(D, E, A, B, C, R(77));\ + P4(C, D, E, A, B, R(78));\ + P4(B, C, D, E, A, R(79)); + +#define SHA2BEG(A,B,C,D,E,W) \ + P1(A, B, C, D, E, W[0]);\ + P1(E, A, B, C, D, W[1]);\ + P1(D, E, A, B, C, W[2]);\ + P1(C, D, E, A, B, W[3]);\ + P1(B, C, D, E, A, W[4]);\ + P1(A, B, C, D, E, W[5]);\ + PZ(E, A, B, C, D);\ + PZ(D, E, A, B, C);\ + PZ(C, D, E, A, B);\ + PZ(B, C, D, E, A);\ + PZ(A, B, C, D, E);\ + PZ(E, A, B, C, D);\ + PZ(D, E, A, B, C);\ + PZ(C, D, E, A, B);\ + PZ(B, C, D, E, A);\ + P1(A, B, C, D, E, W[15]);\ + +#define Q16 (W[0] = S((W[2] ^ W[0]),1)) +#define Q17 (W[1] = S((W[3] ^ W[1]),1)) +#define Q18 (W[2] = S((W[15] ^ W[4] ^ W[2]),1)) +#define Q19 (W[3] = S((W[0] ^ W[5] ^ W[3]),1)) +#define Q20 (W[4] = S((W[1] ^ W[4]),1)) +#define Q21 (W[5] = S((W[2] ^ W[5]),1)) +#define Q22 (W[6] = S(W[3],1)) +#define Q23 (W[7] = S((W[4] ^ W[15]),1)) +#define Q24 (W[8] = S((W[5] ^ W[0]),1)) +#define Q25 (W[9] = S((W[6] ^ W[1]),1)) +#define Q26 (W[10] = S((W[7] ^ W[2]),1)) +#define Q27 (W[11] = S((W[8] ^ W[3]),1)) +#define Q28 (W[12] = S((W[9] ^ W[4]),1)) +#define Q29 (W[13] = S((W[10] ^ W[5] ^ W[15]),1)) +#define Q30 (W[14] = S((W[11] ^ W[6] ^ W[0]),1)) +#define SHA2END(A,B,C,D,E,W)\ + P1(E, A, B, C, D, Q16);\ + P1(D, E, A, B, C, Q17);\ + P1(C, D, E, A, B, Q18);\ + P1(B, C, D, E, A, Q19);\ + P2(A, B, C, D, E, Q20);\ + P2(E, A, B, C, D, Q21);\ + P2(D, E, A, B, C, Q22);\ + P2(C, D, E, A, B, Q23);\ + P2(B, C, D, E, A, Q24);\ + P2(A, B, C, D, E, Q25);\ + P2(E, A, B, C, D, Q26);\ + P2(D, E, A, B, C, Q27);\ + P2(C, D, E, A, B, Q28);\ + P2(B, C, D, E, A, Q29);\ + P2(A, B, C, D, E, Q30);\ + P2(E, A, B, C, D, R(31));\ + P2(D, E, A, B, C, R(32));\ + P2(C, D, E, A, B, R(33));\ + P2(B, C, D, E, A, R(34));\ + P2(A, B, C, D, E, R(35));\ + P2(E, A, B, C, D, R(36));\ + P2(D, E, A, B, C, R(37));\ + P2(C, D, E, A, B, R(38));\ + P2(B, C, D, E, A, R(39));\ + P3(A, B, C, D, E, R(40));\ + P3(E, A, B, C, D, R(41));\ + P3(D, E, A, B, C, R(42));\ + P3(C, D, E, A, B, R(43));\ + P3(B, C, D, E, A, R(44));\ + P3(A, B, C, D, E, R(45));\ + P3(E, A, B, C, D, R(46));\ + P3(D, E, A, B, C, R(47));\ + P3(C, D, E, A, B, R(48));\ + P3(B, C, D, E, A, R(49));\ + P3(A, B, C, D, E, R(50));\ + P3(E, A, B, C, D, R(51));\ + P3(D, E, A, B, C, R(52));\ + P3(C, D, E, A, B, R(53));\ + P3(B, C, D, E, A, R(54));\ + P3(A, B, C, D, E, R(55));\ + P3(E, A, B, C, D, R(56));\ + P3(D, E, A, B, C, R(57));\ + P3(C, D, E, A, B, R(58));\ + P3(B, C, D, E, A, R(59));\ + P4(A, B, C, D, E, R(60));\ + P4(E, A, B, C, D, R(61));\ + P4(D, E, A, B, C, R(62));\ + P4(C, D, E, A, B, R(63));\ + P4(B, C, D, E, A, R(64));\ + P4(A, B, C, D, E, R(65));\ + P4(E, A, B, C, D, R(66));\ + P4(D, E, A, B, C, R(67));\ + P4(C, D, E, A, B, R(68));\ + P4(B, C, D, E, A, R(69));\ + P4(A, B, C, D, E, R(70));\ + P4(E, A, B, C, D, R(71));\ + P4(D, E, A, B, C, R(72));\ + P4(C, D, E, A, B, R(73));\ + P4(B, C, D, E, A, R(74));\ + P4(A, B, C, D, E, R(75));\ + P4(E, A, B, C, D, R(76));\ + P4(D, E, A, B, C, R2(77));\ + P4(C, D, E, A, B, R2(78));\ + P4(B, C, D, E, A, R2(79)); + +#define SHA2(A,B,C,D,E,W) SHA2BEG(A,B,C,D,E,W) SHA2END(A,B,C,D,E,W) + + +inline void preproc(__global const uint8_t * key, uint32_t keylen, + __private uint32_t * state, uint8_t var1, uint32_t var4) +{ + int i; + uint32_t W[16], temp; + uint8_t ipad[20]; + + for (i = 0; i < keylen; i++) + ipad[i] = var1 ^ key[i]; + for (i = keylen; i < 20; i++) + ipad[i] = var1; + + for (i = 0; i < 5; i++) + GET_WORD_32_BE(W[i], ipad, i * 4); + + for (i = 5; i < 16; i++) + W[i] = var4; + + uint32_t A = INIT_A; + uint32_t B = INIT_B; + uint32_t C = INIT_C; + uint32_t D = INIT_D; + uint32_t E = INIT_E; + + SHA1(A, B, C, D, E, W); + + state[0] = A + INIT_A; + state[1] = B + INIT_B; + state[2] = C + INIT_C; + state[3] = D + INIT_D; + state[4] = E + INIT_E; + +} + +inline void hmac_sha1_(__private uint32_t * output, + __private uint32_t * ipad_state, + __private uint32_t * opad_state, + __global const uint8_t * salt, int saltlen, uint8_t add) +{ + int i; + uint32_t temp, W[16]; + uint32_t A, B, C, D, E; + uint8_t buf[64]; + uint32_t *src = (uint32_t *) buf; + i = 64 / 4; + while (i--) + *src++ = 0; + //memcpy(buf, salt, saltlen); + for (i = 0; i < saltlen; i++) + buf[i] = salt[i]; + + buf[saltlen + 4] = 0x80; + buf[saltlen + 3] = add; + PUT_WORD_32_BE((64 + saltlen + 4) << 3, buf, 60); + + A = ipad_state[0]; + B = ipad_state[1]; + C = ipad_state[2]; + D = ipad_state[3]; + E = ipad_state[4]; + + for (i = 0; i < 16; i++) + GET_WORD_32_BE(W[i], buf, i * 4); + + SHA1(A, B, C, D, E, W); + + A += ipad_state[0]; + B += ipad_state[1]; + C += ipad_state[2]; + D += ipad_state[3]; + E += ipad_state[4]; + + PUT_WORD_32_BE(A, buf, 0); + PUT_WORD_32_BE(B, buf, 4); + PUT_WORD_32_BE(C, buf, 8); + PUT_WORD_32_BE(D, buf, 12); + PUT_WORD_32_BE(E, buf, 16); + PUT_WORD_32_BE(0, buf, 20); + PUT_WORD_32_BE(0, buf, 24); + + + buf[20] = 0x80; + PUT_WORD_32_BE(0x2A0, buf, 60); + + A = opad_state[0]; + B = opad_state[1]; + C = opad_state[2]; + D = opad_state[3]; + E = opad_state[4]; + + for (i = 0; i < 16; i++) + GET_WORD_32_BE(W[i], buf, i * 4); + + SHA1(A, B, C, D, E, W); + + A += opad_state[0]; + B += opad_state[1]; + C += opad_state[2]; + D += opad_state[3]; + E += opad_state[4]; + + output[0] = A; + output[1] = B; + output[2] = C; + output[3] = D; + output[4] = E; +} + + + +inline void big_hmac_sha1(__private uint32_t * input, uint32_t inputlen, + __private uint32_t * ipad_state, + __private uint32_t * opad_state, __private uint32_t * tmp_out) +{ + int i, lo; + uint32_t temp, W[16]; + uint32_t A, B, C, D, E; + + for (i = 0; i < 5; i++) + W[i] = input[i]; + + for (lo = 1; lo < ITERATIONS; lo++) { + + A = ipad_state[0]; + B = ipad_state[1]; + C = ipad_state[2]; + D = ipad_state[3]; + E = ipad_state[4]; + + W[5] = 0x80000000; + W[15] = 0x2A0; + + SHA2(A, B, C, D, E, W); + + A += ipad_state[0]; + B += ipad_state[1]; + C += ipad_state[2]; + D += ipad_state[3]; + E += ipad_state[4]; + + W[0] = A; + W[1] = B; + W[2] = C; + W[3] = D; + W[4] = E; + W[5] = 0x80000000; + W[15] = 0x2A0; + + A = opad_state[0]; + B = opad_state[1]; + C = opad_state[2]; + D = opad_state[3]; + E = opad_state[4]; + + SHA2(A, B, C, D, E, W); + + A += opad_state[0]; + B += opad_state[1]; + C += opad_state[2]; + D += opad_state[3]; + E += opad_state[4]; + + W[0] = A; + W[1] = B; + W[2] = C; + W[3] = D; + W[4] = E; + + tmp_out[0] ^= A; + tmp_out[1] ^= B; + tmp_out[2] ^= C; + tmp_out[3] ^= D; + tmp_out[4] ^= E; + } + + for (i = 0; i < 5; i++) + tmp_out[i] = SWAP(tmp_out[i]); +} + +inline void pbkdf2(__global const uint8_t * pass, int passlen, + __global const uint8_t * salt, int saltlen, __global uint32_t * out) +{ + uint32_t ipad_state[5]; + uint32_t opad_state[5]; + uint32_t tmp_out[5]; + int i; + + preproc(pass, passlen, ipad_state, 0x36, 0x36363636); + preproc(pass, passlen, opad_state, 0x5c, 0x5c5c5c5c); + + uint8_t rnd = 0x01; + __global unsigned char *dst = (__global unsigned char*)out; + unsigned char *src; + for (; rnd < 0x04;) { + hmac_sha1_(tmp_out, ipad_state, opad_state, salt, saltlen, + rnd++); + + big_hmac_sha1(tmp_out, SHA1_DIGEST_LENGTH, ipad_state, + opad_state, tmp_out); + src = (unsigned char*)tmp_out; + for(i = 0; i < 20; i++) + dst[i] = src[i]; + dst+=(5*4); + } + hmac_sha1_(tmp_out, ipad_state, opad_state, salt, saltlen, 0x04); + big_hmac_sha1(tmp_out, SHA1_DIGEST_LENGTH, ipad_state, opad_state, + tmp_out); + for(i = 0; i < 6; i++) + dst[i] = src[i]; +} + +__kernel void odf(__global const odf_password * inbuffer, + __global odf_hash * outbuffer, __global const odf_salt * salt) +{ + uint32_t idx = get_global_id(0); + + pbkdf2(inbuffer[idx].v, inbuffer[idx].length, + salt->salt, salt->length, outbuffer[idx].v); +} diff --git a/src/opencl_odf_fmt.c b/src/opencl_odf_fmt.c new file mode 100644 index 0000000..124c845 --- /dev/null +++ b/src/opencl_odf_fmt.c @@ -0,0 +1,400 @@ +/* + * Modified by Dhiru Kholia for Keychain 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 "common.h" +#include "formats.h" +#include "common-opencl.h" +#include +#include +#include +#ifdef _OPENMP +#include +#endif + +#define FORMAT_LABEL "odf-opencl" +#define FORMAT_NAME "ODF SHA-1 Blowfish / SHA-256 AES" +#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 20 +#define PLAINTEXT_LENGTH 15 +#define SALT_SIZE sizeof(odf_cpu_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]; +} odf_password; + +typedef struct { + uint32_t v[17]; +} odf_hash; + +typedef struct { + uint8_t length; + uint8_t salt[64]; +} odf_salt; + +static char (*saved_key)[PLAINTEXT_LENGTH + 1]; +static ARCH_WORD_32 (*crypt_out)[32 / sizeof(ARCH_WORD_32)]; + +typedef struct { + int cipher_type; + int checksum_type; + int iterations; + int key_size; + int iv_length; + int salt_length; + unsigned char iv[16]; + unsigned char salt[32]; + unsigned char content[1024]; +} odf_cpu_salt; + +odf_cpu_salt *cur_salt; + +static struct fmt_tests odf_tests[] = { + {"$odf$*0*0*1024*16*df6c10f64d191a841812af53874b636d014ce3fe*8*07e28aff39d2660e*16*b124be9f3346fb77e0ebcc3bb80028f8*0*2276a1077f6a2a027bd565ce89824d6a20086e378876be05c4b8e3796a460e828c9803a692caf7a53492c220d1d7ecbf4e2d336c7abf5a7672acc804ca267318252cbc13676616d1fde38820f9fbeef1360067d9de096ba8c1032ae947bde1d0fedaf37b6020663d49faf36b7c095c5b9aae11c8fc2be74148f008edbdbb180b44028ad8259f1215b483542bf3027f56dee5f962448333b30f88e6ae4790b60d24abb286edff9adee831a4b3351fc47259043f0d683d7a25be7e47aff3aedca140005d866e218c8efcca32093c19bbece50bd96656d0f94a712d3c60d1e5342db86482fc73f05faf513ca0b137378126597b95986c372b412c953e97011259aab0839fe453c756559497a28ba88dce009e1e7980436131029d38e56a34f608e6471970d9959068808c898608024db9eb394c4feae7a364ea9272ec4ea2315a9f0407a4b27d5e49a8ab1e3ddce5c84927d5aecd7e68e4437a820ea8743c6b5b4e2abbb47b0001e2f77ceac4603e8774e4ccbc1adde794428c11ae4a7492727b620334302e63f72b0c06c1cf83800366916ee8295176819272d557863a831ee0a576841191482959aad69095831fa1d64e3e0e6f6c6a751bcdadf0fbaa27a17458709f708c04587cb208984c9525da6786e0e5aabefe30ad1dbbef66e85ce9d6dbe456fd85e4135de5cf16d9455976d7ca8de7b1b530661c74c0fae90c0fff1a2b5fcdfab19fcff75fadcec445ed8af6ab5babf1463e08458918be8045083de6db988c37e4be582cfac5cdf741d1f0322fb2902665c7ff347813348109e5d442e91fcb010c28f042da481e807084fcb4759b40ccf2cae77bad00cdfbfba4acf36aa1f74c30a315e3d7f1ca522b6306e8903352aafa51dc523d582d418934398d5eb88120e3656bfb640a239db507b285302a86855ea850ddc9af72fc62dc79336c9bc29ee8314c65adb0574e9c701d73d7fa977edd1d52a1ff2da5b8b94e1a0fdd01ffcc6583758f0a1f51750e45f12b58c6d38b140e5676cf3474224520ef7c52ca5e634f85456651f3d6f43d016ed7cc5da54ea640a3bc50c2b9d3dea8f93c0340d66ccd06efc5ae002108c33cf3a470c4a50f6a6ca2f11b8ad15511688c282b94ba6f1c332e239d10946dc46f763f08d12cb9edc1e79c0e07f7151f548e6d7d20ec13b52d911bf980cac60694e192651403c9a69abea045190e847be093fc9ba43fec55b32f77f5796ddca25b441f259d5c51e06df6c6588c6414899481ba9e06bcebec58f82ff3021b09c6beae13a5d22bc94870f72ab813d0c0be01d91f3d075192e7a5de765599d72244757d09539529a8347e077a36678166e5ed9f73a5aad2e147d8154095c397e3e5e4ba1987ca64c1301a0c6c3e438097ede9b701a105ec38fcb54abb31b367c7740cd9ac459e561094a34f01acee555e60267157e6", "test"}, +// {"$odf$*1*1*1024*32*61802eba18eab842de1d053809ba40927fd40b26c69ddeca6a8a652ed9c16a28*16*c5c0815b931f313627100d592a9c972f*16*e9a48b7daff738deaabe442007fb2ec4*0*be3b65ea09642c2b4fdc23e553e1f5304bc5df222b624c6373d53e674f5df01fdb8873cdab7a5a685fa45ad5441a9d8869401b7fa076c488ad53fd9971e97244ecc9416484450d4fb2ee4ec08af4044d7def937e6545dea2ce36bd5c57b1f46b11b9cf90c8fb3accff149ce2d54820b181b9124db9aac131f6436d77cf716423f04d42438eed6f9ca14bd24b9b17d3478176addd5fa0254bf986fccd879e326485790e28b94ad5306868734b5ac1b1ddb3f876382dee6e9428e8230e84bf11b7e85ccbae8b4b424cd73160c380f874b37fbe3c7e88c13ef4bde74b56507d17095c2c32bb8bcded0637e4403107bb33252f72f5886a91b7720fe32a8659a09c217717e4c74a7c2e09fc40b46aa288309a36e86b9f1856e1bce176bc9690555431e05c7b67ff95df64f8f40053079bfc9dda021ab2714fecf74398b867ebef675958f29eaa15eb631845e358a0c5caff0b824a2a69a6eabee069d3d6236d77709fd60438c9e3ad9e42b26810375e1e587eff105ac295327ef8bf66f6462388b7727ec32d6abde2f8d6126b185124bb437753663f6ab1f321ddfdb36d9f1f528729492e0b1bb8d3b9eda3c86c1997c92b902f5160f77587c37e45b5c133b5d9709fea910a2e9b54c0960b0ebc870cdbb858aabe07ed27cba86d29a7e64c6e3863131859314a14e64c1168d4a2d5ca0697853fb1fe969ba968e31359881d51edce287eff415de8e60cec2068bb82157fbcf0cf9a95e92cb23f32e6156daced4bee6ba8c8b41174d01fcd7662911bcc10d5b4478f8209ce3b91075d10529780be4f17e841a1f1833d432c3dc854908643e58b03c8860dfbc710a29f79f75ea262cfcef9cd67fb67d73f55b300d42f4577445af2b9f224620204cfb88de2cbf57931ac0e0f8d98259a41d744cad6a58abc7761c266f4e93aca19356b07073c09ae9d1976f4f2e1a76c350cc7764c27ae257eb69ba4213dd0a7794fa83d220439a398efd988b6dbf0de4c08bc3e4830c9e482b9e0fd1679f14e6f132cf06bae1d763dde7ce6f525ff9a0ebad28aeca16496194f2a6263a20e7afeb43d83c8c936130d6508f2bf68b5ca50375948424193a7fb1106fdf63ff72896e1b2633907f01a693218e3303436542bcf2af24cc4a41621c36768ce9a84d32cc9f3c2b108bfc78c25b1c2ea94e6e0d65406f78bdb8bc33c94a9550e5cc3e995cfbd31da03afb929418acdc89b099415f9bdb7dab7a75d44a696e14b031d601ad8d907e14a28044706c0c2955df2cb34ffea82af367e487b6cc928dc87a33fc7555173e7faa5cfd1af6d3d6f496f23a9579db22dd4a2c16e950fdc90696d95a81183765a4fbddb42c488d40ac1de28483cf1cdddf821d3f859c57b13cb7f21a916bd0d89438a17634c68637f23e2544589e8ae5ee5bced91680c087cb3105cd74a09e88d3aae17d75e", "test"}, + {NULL} +}; + +static odf_password *inbuffer; +static odf_hash *outbuffer; +static odf_salt currentsalt; +static cl_mem mem_in, mem_out, mem_setting; +static size_t insize = sizeof(odf_password) * KEYS_PER_CRYPT; +static size_t outsize = sizeof(odf_hash) * KEYS_PER_CRYPT; +static size_t settingsize = sizeof(odf_salt); + +#define DEBUG + +#ifdef DEBUG +static void print_hex(unsigned char *str, int len) +{ + int i; + for (i = 0; i < len; ++i) + printf("%02x", str[i]); + printf("\n"); +} +#endif + +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 *pFmt) +{ + cl_int cl_error; + + global_work_size = MAX_KEYS_PER_CRYPT; + + inbuffer = + (odf_password *) malloc(sizeof(odf_password) * + MAX_KEYS_PER_CRYPT); + outbuffer = + (odf_hash *) malloc(sizeof(odf_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; + } + saved_key = mem_calloc_tiny(sizeof(*saved_key) * + pFmt->params.max_keys_per_crypt, MEM_ALIGN_NONE); + + crypt_out = mem_calloc_tiny(sizeof(*crypt_out) * pFmt->params.max_keys_per_crypt, MEM_ALIGN_WORD); + + //listOpenCLdevices(); + opencl_init("$JOHN/odf_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], "odf", &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(pFmt); + + atexit(release_all); +} + +static int valid(char *ciphertext, struct fmt_main *pFmt) +{ + return !strncmp(ciphertext, "$odf$", 5); +} + +static void *get_salt(char *ciphertext) +{ + char *ctcopy = strdup(ciphertext); + char *keeptr = ctcopy; + int i; + char *p; + static odf_cpu_salt cs; + ctcopy += 6; /* skip over "$odf$*" */ + p = strtok(ctcopy, "*"); + cs.cipher_type = atoi(p); + p = strtok(NULL, "*"); + cs.checksum_type = atoi(p); + p = strtok(NULL, "*"); + cs.iterations = atoi(p); + p = strtok(NULL, "*"); + cs.key_size = atoi(p); + p = strtok(NULL, "*"); + /* skip checksum field */ + p = strtok(NULL, "*"); + cs.iv_length = atoi(p); + p = strtok(NULL, "*"); + for (i = 0; i < cs.iv_length; i++) + cs.iv[i] = atoi16[ARCH_INDEX(p[i * 2])] * 16 + + atoi16[ARCH_INDEX(p[i * 2 + 1])]; + p = strtok(NULL, "*"); + cs.salt_length = atoi(p); + p = strtok(NULL, "*"); + for (i = 0; i < cs.salt_length; i++) + cs.salt[i] = atoi16[ARCH_INDEX(p[i * 2])] * 16 + + atoi16[ARCH_INDEX(p[i * 2 + 1])]; + p = strtok(NULL, "*"); + p = strtok(NULL, "*"); + for (i = 0; i < 1024; i++) + cs.content[i] = atoi16[ARCH_INDEX(p[i * 2])] * 16 + + atoi16[ARCH_INDEX(p[i * 2 + 1])]; + free(keeptr); + return (void *)&cs; +} + +static void *get_binary(char *ciphertext) +{ + static union { + unsigned char c[BINARY_SIZE+1]; + ARCH_WORD dummy; + } buf; + unsigned char *out = buf.c; + char *p; + int i; + char *ctcopy = strdup(ciphertext); + char *keeptr = ctcopy; + ctcopy += 6; /* skip over "$odf$*" */ + p = strtok(ctcopy, "*"); + p = strtok(NULL, "*"); + p = strtok(NULL, "*"); + p = strtok(NULL, "*"); + p = strtok(NULL, "*"); + for (i = 0; i < BINARY_SIZE; i++) { + out[i] = + (atoi16[ARCH_INDEX(*p)] << 4) | + atoi16[ARCH_INDEX(p[1])]; + p += 2; + } + free(keeptr); + return out; +} + +static void set_salt(void *salt) +{ + cur_salt = (odf_cpu_salt*)salt; + memcpy((char*)currentsalt.salt, cur_salt->salt, cur_salt->salt_length); + currentsalt.length = cur_salt->salt_length; +} + +static int binary_hash_0(void *binary) { return *(ARCH_WORD_32 *)binary & 0xf; } +static int binary_hash_1(void *binary) { return *(ARCH_WORD_32 *)binary & 0xff; } +static int binary_hash_2(void *binary) { return *(ARCH_WORD_32 *)binary & 0xfff; } +static int binary_hash_3(void *binary) { return *(ARCH_WORD_32 *)binary & 0xffff; } +static int binary_hash_4(void *binary) { return *(ARCH_WORD_32 *)binary & 0xfffff; } +static int binary_hash_5(void *binary) { return *(ARCH_WORD_32 *)binary & 0xffffff; } +static int binary_hash_6(void *binary) { return *(ARCH_WORD_32 *)binary & 0x7ffffff; } + +static int get_hash_0(int index) { return crypt_out[index][0] & 0xf; } +static int get_hash_1(int index) { return crypt_out[index][0] & 0xff; } +static int get_hash_2(int index) { return crypt_out[index][0] & 0xfff; } +static int get_hash_3(int index) { return crypt_out[index][0] & 0xffff; } +static int get_hash_4(int index) { return crypt_out[index][0] & 0xfffff; } +static int get_hash_5(int index) { return crypt_out[index][0] & 0xffffff; } +static int get_hash_6(int index) { return crypt_out[index][0] & 0x7ffffff; } + +#undef set_key +static void set_key(char *key, int index) +{ + int saved_key_length = strlen(key); + if (saved_key_length > PLAINTEXT_LENGTH) + saved_key_length = PLAINTEXT_LENGTH; + memcpy(saved_key[index], key, saved_key_length); + saved_key[index][saved_key_length] = 0; + + +} + +static char *get_key(int index) +{ + return saved_key[index]; +} + +static void crypt_all(int count) +{ + int index; +#ifdef _OPENMP +#pragma omp parallel for + for(index = 0; index < count; index++) +#else + for(index = 0; index < count; index++) +#endif + { + unsigned char hash[32]; + SHA_CTX ctx; + SHA1_Init(&ctx); + SHA1_Update(&ctx, (unsigned char *)saved_key[index], strlen(saved_key[index])); + SHA1_Final((unsigned char *)hash, &ctx); + inbuffer[index].length = 20; + memcpy(inbuffer[index].v, hash, 20); + } + + /// 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"); + +#ifdef _OPENMP +#pragma omp parallel for + for(index = 0; index < count; index++) +#else + for(index = 0; index < count; index++) +#endif + { + BF_KEY bf_key; + SHA_CTX ctx; + int bf_ivec_pos; + unsigned char ivec[8]; + unsigned char output[1024]; + bf_ivec_pos = 0; + memcpy(ivec, cur_salt->iv, 8); + BF_set_key(&bf_key, cur_salt->key_size, outbuffer[index].v); + BF_cfb64_encrypt(cur_salt->content, output, 1024, &bf_key, ivec, &bf_ivec_pos, 0); + SHA1_Init(&ctx); + SHA1_Update(&ctx, output, 1024); + SHA1_Final((unsigned char*)crypt_out[index], &ctx); + } +} + +static int cmp_all(void *binary, int count) +{ + int index = 0; +#ifdef _OPENMP + for (; index < count; index++) +#endif + if (!memcmp(binary, crypt_out[index], BINARY_SIZE)) + return 1; + return 0; +} + +static int cmp_one(void *binary, int index) +{ + return !memcmp(binary, crypt_out[index], BINARY_SIZE); +} + +static int cmp_exact(char *source, int index) +{ + return 1; +} + +struct fmt_main fmt_opencl_odf = { + { + 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, + odf_tests + }, { + init, + fmt_default_prepare, + valid, + fmt_default_split, + get_binary, + get_salt, + { + binary_hash_0, + binary_hash_1, + binary_hash_2, + binary_hash_3, + binary_hash_4, + binary_hash_5, + binary_hash_6 + }, + fmt_default_salt_hash, + set_salt, + set_key, + get_key, + fmt_default_clear_keys, + crypt_all, + { + get_hash_0, + get_hash_1, + get_hash_2, + get_hash_3, + get_hash_4, + get_hash_5, + get_hash_6 + }, + cmp_all, + cmp_one, + cmp_exact + } +}; -- 1.7.11.2