From 81da98f41aaa5cbbea199b71ac786285c82597a7 Mon Sep 17 00:00:00 2001 From: Dhiru Kholia Date: Mon, 23 Jul 2012 20:26:02 +0530 Subject: [PATCH] OpenCL implementation of ZIP AES format. --- src/Makefile | 2 +- src/john.c | 3 + src/opencl/zip_kernel.cl | 519 +++++++++++++++++++++++++++++++++++++++++++++++ src/opencl_zip_fmt.c | 319 +++++++++++++++++++++++++++++ 4 files changed, 842 insertions(+), 1 deletion(-) create mode 100644 src/opencl/zip_kernel.cl create mode 100644 src/opencl_zip_fmt.c diff --git a/src/Makefile b/src/Makefile index 9db559c..1c84b48 100644 --- a/src/Makefile +++ b/src/Makefile @@ -134,7 +134,7 @@ 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_keychain_fmt.o opencl_agilekeychain_fmt.o opencl_zip_fmt.o CUDA_OBJS = \ cuda_common.o \ diff --git a/src/john.c b/src/john.c index 7eb2fa1..50c4b15 100644 --- a/src/john.c +++ b/src/john.c @@ -132,6 +132,7 @@ extern struct fmt_main fmt_opencl_mscash2; 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_xsha512; extern struct fmt_main fmt_opencl_rawsha512; extern struct fmt_main fmt_opencl_bf; @@ -158,6 +159,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" @@ -269,6 +271,7 @@ static void john_register_all(void) john_register_one(&fmt_opencl_wpapsk); john_register_one(&fmt_opencl_keychain); john_register_one(&fmt_opencl_agilekeychain); + john_register_one(&fmt_opencl_zip); john_register_one(&fmt_opencl_xsha512); john_register_one(&fmt_opencl_rawsha512); john_register_one(&fmt_opencl_bf); diff --git a/src/opencl/zip_kernel.cl b/src/opencl/zip_kernel.cl new file mode 100644 index 0000000..9a3f2d9 --- /dev/null +++ b/src/opencl/zip_kernel.cl @@ -0,0 +1,519 @@ +/* + * Modified by Dhiru Kholia for ZIP AES 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]; +} zip_password; + +typedef struct { + uint32_t v[17]; // 16*4=64 +} zip_hash; + +typedef struct { + uint8_t length; + uint8_t salt[64]; +} zip_salt; + + +# define SWAP(n) \ + (((n) << 24) | (((n) & 0xff00) << 8) | (((n) >> 8) & 0xff00) | ((n) >> 24)) + +#define ITERATIONS 1000 + +#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[16]; + + for (i = 0; i < keylen; i++) + ipad[i] = var1 ^ key[i]; + for (i = keylen; i < 16; i++) + ipad[i] = var1; + + for (i = 0; i < 4; i++) + GET_WORD_32_BE(W[i], ipad, i * 4); + + for (i = 4; 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); + + // memcpy(out2, tmp_out, 20); + 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); + // memcpy(out2, tmp_out, 6); + for(i = 0; i < 6; i++) + dst[i] = src[i]; +} + +__kernel void zip(__global const zip_password * inbuffer, + __global zip_hash * outbuffer, __global const zip_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_zip_fmt.c b/src/opencl_zip_fmt.c new file mode 100644 index 0000000..7143689 --- /dev/null +++ b/src/opencl_zip_fmt.c @@ -0,0 +1,319 @@ +/* + * 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 +#include "common-opencl.h" +#ifdef _OPENMP +#include +#endif + +#define FORMAT_LABEL "zip-opencl" +#define FORMAT_NAME "ZIP-AES PBKDF2-HMAC-SHA-1" +#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 SWAP(n) \ + (((n) << 24) | (((n) & 0xff00) << 8) | (((n) >> 8) & 0xff00) | ((n) >> 24)) + +#define BINARY_SIZE 16 +#define PLAINTEXT_LENGTH 15 +#define SALT_SIZE sizeof(zip_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]; +} zip_password; + +typedef struct { + uint32_t v[17]; +} zip_hash; + +typedef struct { + uint8_t length; + uint8_t salt[64]; +} zip_salt; + +static int *cracked; + +typedef struct { + uint8_t length; + uint8_t salt[20]; + int type; /* type of zip file */ + int mode; + unsigned char passverify[2]; +} zip_cpu_salt; + +zip_cpu_salt *cur_salt; + +static struct fmt_tests zip_tests[] = { + {"$zip$*0*1*8005b1b7d077708d*dee4", "testpassword#"}, + {NULL} +}; + +static zip_password *inbuffer; +static zip_hash *outbuffer; +static zip_salt currentsalt; +static cl_mem mem_in, mem_out, mem_setting; +static size_t insize = sizeof(zip_password) * KEYS_PER_CRYPT; +static size_t outsize = sizeof(zip_hash) * KEYS_PER_CRYPT; +static size_t settingsize = sizeof(zip_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 = + (zip_password *) malloc(sizeof(zip_password) * + MAX_KEYS_PER_CRYPT); + outbuffer = + (zip_hash *) malloc(sizeof(zip_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/zip_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], "zip", &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, "$zip$", 5); +} + +static void *get_salt(char *ciphertext) +{ + char *ctcopy = strdup(ciphertext); + char *keeptr = ctcopy; + int i; + char *p; + char *encoded_salt; + static zip_cpu_salt cs; + int strength, n; + + ctcopy += 6; /* skip over "$zip$*" */ + cs.type = atoi(strtok(ctcopy, "*")); + strength = atoi(strtok(NULL, "*")); + cs.mode = strength; + switch (strength) { + case 1: + n = 8; + break; + case 2: + n = 12; + break; + case 3: + n = 16; + break; + default: + fprintf(stderr, "ZIP: Unsupported strength %d\n", strength); + error(); + n = 0; /* Not reached */ + } + cs.length = n; + encoded_salt = strtok(NULL, "*"); + for (i = 0; i < n; i++) + cs.salt[i] = atoi16[ARCH_INDEX(encoded_salt[i * 2])] * 16 + + atoi16[ARCH_INDEX(encoded_salt[i * 2 + 1])]; + p = strtok(NULL, "*"); + for (i = 0; i < 2; i++) + cs.passverify[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 = (zip_cpu_salt*)salt; + memcpy((char*)currentsalt.salt, cur_salt->salt, cur_salt->length); + currentsalt.length = cur_salt->length; +} + +#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 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++) + { + unsigned char pwd_ver[2] = { 0 }; + unsigned char *p; + p = (unsigned char*)outbuffer[index].v; +#ifdef DEBUG + print_hex(p, 34); +#endif + memcpy(pwd_ver, &p[32], 2); + if(!memcmp(pwd_ver, cur_salt->passverify, 2)) + cracked[index] = 1; + else + cracked[index] = 0; + } + +} + +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_zip = { + { + 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, + zip_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.2