Keepass OpenCL format.

src/keepass_common.h

@@ -13,7 +13,7 @@
 #define FORMAT_TAG_LEN       (sizeof(FORMAT_TAG)-1)
 #define BENCHMARK_COMMENT	""
 #define BENCHMARK_LENGTH	-1
-#define PLAINTEXT_LENGTH	125
+#define PLAINTEXT_LENGTH	124
 #define BINARY_SIZE		0
 #define BINARY_ALIGN		MEM_ALIGN_NONE
 #define SALT_SIZE		sizeof(keepass_salt_t)
@@ -29,23 +29,25 @@
 
 extern struct fmt_tests keepass_tests[];
 
+/* This format should be dyna salt instead! */
+#define MAX_CONT_SIZE 0x1000
+
 typedef struct {
 	long long offset;
 	int version;
 	int isinline;
 	int keyfilesize;
 	int have_keyfile;
 	int contentsize;
-//	unsigned char contents[LINE_BUFFER_SIZE];
-	unsigned char contents[0x30000];	// We need to fix this in some other way, now that LINE_BUFFER_SIZE has been dropped so heavily!
+	uint32_t key_transf_rounds;
+	int algorithm; // 1 for Twofish
 	unsigned char final_randomseed[32];
 	unsigned char enc_iv[16];
 	unsigned char keyfile[32];
 	unsigned char contents_hash[32];
 	unsigned char transf_randomseed[32];
 	unsigned char expected_bytes[32];
-	uint32_t key_transf_rounds;
-	int algorithm; // 1 for Twofish
+	unsigned char contents[MAX_CONT_SIZE];
 } keepass_salt_t;
 
 extern char (*keepass_key)[PLAINTEXT_LENGTH + 1];

src/keepass_common_plug.c

@@ -49,12 +49,14 @@ struct fmt_tests keepass_tests[] = {
 };
 
 char (*keepass_key)[PLAINTEXT_LENGTH + 1];
+keepass_salt_t *keepass_salt;
 
 int keepass_valid(char *ciphertext, struct fmt_main *self)
 {
 	char *ctcopy;
 	char *keeptr;
 	char *p;
+	long long algo;
 	int version, res, contentsize, extra;
 
 	if (strncmp(ciphertext, FORMAT_TAG, FORMAT_TAG_LEN))
@@ -75,6 +77,13 @@ int keepass_valid(char *ciphertext, struct fmt_main *self)
 		goto err;
 	if (!isdec(p))  /* this field contained file offsets in the past, this is hard to validate */
 		goto err;
+	algo = atoll(p);
+	if (algo < 0 || algo > 2)
+		algo = 0;
+	if (version == 1 && algo > 1) /* Unsupported combo */
+		goto err;
+	if (version == 2 && algo == 1) /* TODO, v2 w/ Twofish */
+		goto err;
 	if ((p = strtokm(NULL, "*")) == NULL)	/* final random seed */
 		goto err;
 	res = hexlenl(p, &extra);
@@ -102,6 +111,18 @@ int keepass_valid(char *ciphertext, struct fmt_main *self)
 		if ((p = strtokm(NULL, "*")) == NULL)	/* content size */
 			goto err;
 		contentsize = atoi(p);
+		if (contentsize > MAX_CONT_SIZE) {
+			static int warned;
+
+			if (!ldr_in_pot && warned < contentsize) {
+				fprintf(stderr,
+				        "%s: Input rejected due to larger size than compile-time limit.\n"
+				        "Bump MAX_CONT_SIZE in keepass_common.h to >= 0x%x, and rebuild\n",
+				        self->params.label, contentsize);
+				warned = contentsize;
+			}
+			goto err;
+		}
 		if ((p = strtokm(NULL, "*")) == NULL)	/* content */
 			goto err;
 		if (!contentsize || hexlenl(p, &extra) / 2 != contentsize || extra)

src/opencl/bitlocker_kernel.cl

@@ -18,7 +18,7 @@
 #include "opencl_misc.h"
 #include "opencl_bitlocker.h"
 
-__constant unsigned int TS0[256] = {
+__const_a8 unsigned int TS0[256] = {
 	0xC66363A5U, 0xF87C7C84U, 0xEE777799U, 0xF67B7B8DU, 0xFFF2F20DU,
 	0xD66B6BBDU, 0xDE6F6FB1U, 0x91C5C554U,
 	0x60303050U, 0x02010103U, 0xCE6767A9U, 0x562B2B7DU, 0xE7FEFE19U,
@@ -85,7 +85,7 @@ __constant unsigned int TS0[256] = {
 	0xA85454FCU, 0x6DBBBBD6U, 0x2C16163AU
 };
 
-__constant unsigned int TS1[256] = {
+__const_a8 unsigned int TS1[256] = {
 	0xA5C66363U, 0x84F87C7CU, 0x99EE7777U, 0x8DF67B7BU, 0x0DFFF2F2U,
 	0xBDD66B6BU, 0xB1DE6F6FU, 0x5491C5C5U,
 	0x50603030U, 0x03020101U, 0xA9CE6767U, 0x7D562B2BU, 0x19E7FEFEU,
@@ -152,7 +152,7 @@ __constant unsigned int TS1[256] = {
 	0xFCA85454U, 0xD66DBBBBU, 0x3A2C1616U
 };
 
-__constant unsigned int TS2[256] = {
+__const_a8 unsigned int TS2[256] = {
 	0x63A5C663U, 0x7C84F87CU, 0x7799EE77U, 0x7B8DF67BU, 0xF20DFFF2U,
 	0x6BBDD66BU, 0x6FB1DE6FU, 0xC55491C5U,
 	0x30506030U, 0x01030201U, 0x67A9CE67U, 0x2B7D562BU, 0xFE19E7FEU,
@@ -219,7 +219,7 @@ __constant unsigned int TS2[256] = {
 	0x54FCA854U, 0xBBD66DBBU, 0x163A2C16U
 };
 
-__constant unsigned int TS3[256] = {
+__const_a8 unsigned int TS3[256] = {
 	0x6363A5C6U, 0x7C7C84F8U, 0x777799EEU, 0x7B7B8DF6U, 0xF2F20DFFU,
 	0x6B6BBDD6U, 0x6F6FB1DEU, 0xC5C55491U,
 	0x30305060U, 0x01010302U, 0x6767A9CEU, 0x2B2B7D56U, 0xFEFE19E7U,

src/opencl/cryptmd5_kernel.cl

@@ -127,7 +127,7 @@ typedef struct {
 
 __constant uchar cl_md5_salt_prefix[] = "$1$";
 __constant uchar cl_apr1_salt_prefix[] = "$apr1$";
-__constant uchar g[] =
+__const_a8 uchar g[] =
     { 0, 7, 3, 5, 3, 7, 1, 6, 3, 5, 3, 7, 1, 7, 2, 5, 3, 7, 1, 7, 3, 4, 3, 7,
 	    1, 7, 3, 5, 2, 7, 1, 7, 3, 5, 3, 6, 1, 7, 3, 5, 3, 7 };
 

src/opencl/keepass_kernel.cl

@@ -0,0 +1,199 @@
+/*
+ * This software is Copyright (c) 2018 magnum,
+ * 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 "opencl_misc.h"
+#define AES_KEY_TYPE __private
+#define AES_SRC_TYPE __private
+#define AES_DST_TYPE __private
+#define OCL_AES_ENCRYPT 1
+#define OCL_AES_CBC_DECRYPT 1
+#include "opencl_aes.h"
+#include "opencl_sha2_ctx.h"
+#include "opencl_chacha.h"
+#include "opencl_twofish.h"
+
+typedef struct {
+	uint32_t length;
+	uint8_t v[PLAINTEXT_LENGTH];
+} keepass_password;
+
+typedef struct {
+	uint32_t cracked;
+} keepass_result;
+
+typedef struct {
+	uint64_t offset;
+	int version;
+	int isinline;
+	int keyfilesize;
+	int have_keyfile;
+	int contentsize;
+	uint32_t key_transf_rounds;
+	int algorithm;
+	uchar final_randomseed[32];
+	uchar enc_iv[16];
+	uchar keyfile[32];
+	uchar contents_hash[32];
+	uchar transf_randomseed[32];
+	uchar expected_bytes[32];
+	uchar contents[MAX_CONT_SIZE];
+} keepass_salt_t;
+
+typedef struct {
+	uint iterations;
+	uchar hash[32];
+	AES_KEY akey;
+} keepass_state;
+
+__kernel void keepass_init(__global const keepass_password *masterkey,
+                           __constant keepass_salt_t *salt,
+                           __global keepass_state *state)
+{
+	uint gid = get_global_id(0);
+	uchar hash[32];
+	uint pwlen = masterkey[gid].length;
+	uchar pbuf[PLAINTEXT_LENGTH];
+	SHA256_CTX ctx;
+	AES_KEY akey;
+
+	// We can afford some safety because only the loop kernel is significant
+	if (pwlen > PLAINTEXT_LENGTH)
+		pwlen = 0;
+
+	memcpy_macro(pbuf, masterkey[gid].v, pwlen);
+
+	// First, hash the masterkey
+	SHA256_Init(&ctx);
+	SHA256_Update(&ctx, pbuf, pwlen);
+	SHA256_Final(hash, &ctx);
+
+	if (salt->have_keyfile) {
+		memcpy_macro(pbuf, salt->keyfile, 32);
+		SHA256_Init(&ctx);
+		SHA256_Update(&ctx, hash, 32);
+		SHA256_Update(&ctx, pbuf, 32);
+		SHA256_Final(hash, &ctx);
+	} else if (salt->version == 2) {
+		SHA256_Init(&ctx);
+		SHA256_Update(&ctx, hash, 32);
+		SHA256_Final(hash, &ctx);
+	}
+
+	// Next, encrypt the hash using the random seed as key
+	memcpy_macro(pbuf, salt->transf_randomseed, 32);
+	AES_set_encrypt_key(pbuf, 256, &akey);
+
+	// Save state for loop kernel.
+	state[gid].iterations = salt->key_transf_rounds;
+	memcpy_macro(state[gid].hash, hash, 32);
+	memcpy_pg(&state[gid].akey, &akey, sizeof(AES_KEY));
+}
+
+// Here's the heavy part. NOTHING else is significant for performance!
+__kernel void keepass_loop(__global keepass_state *state)
+{
+	uint gid = get_global_id(0);
+	AES_KEY akey;
+	uint i;
+	uchar hash[32];
+
+	i = MIN(state[gid].iterations, HASH_LOOPS);
+	state[gid].iterations -= i;
+	memcpy_macro(hash, state[gid].hash, 32);
+	memcpy_gp(&akey, &state[gid].akey, sizeof(AES_KEY));
+
+	while (i--) {
+		AES_encrypt(hash, hash, &akey);
+		AES_encrypt(hash + 16, hash + 16, &akey);
+	}
+
+	memcpy_macro(state[gid].hash, hash, 32);
+}
+
+__kernel void keepass_final(__global keepass_state *state,
+                            __constant keepass_salt_t *salt,
+                            __global keepass_result *result)
+{
+	uint gid = get_global_id(0);
+	SHA256_CTX ctx;
+	AES_KEY akey;
+	uchar pbuf[32];
+	uchar hash[32];
+	uchar iv[16];
+
+	memcpy_macro(hash, state[gid].hash, 32);
+
+	// Finally, hash it again...
+	SHA256_Init(&ctx);
+	SHA256_Update(&ctx, hash, 32);
+	SHA256_Final(hash, &ctx);
+
+	// ...and hash the result together with the random seed
+	SHA256_Init(&ctx);
+	if (salt->version == 1) {
+		memcpy_macro(pbuf, salt->final_randomseed, 16);
+		SHA256_Update(&ctx, pbuf, 16);
+	} else {
+		memcpy_macro(pbuf, salt->final_randomseed, 32);
+		SHA256_Update(&ctx, pbuf, 32);
+	}
+	SHA256_Update(&ctx, hash, 32);
+	SHA256_Final(hash, &ctx);
+
+	memcpy_macro(iv, salt->enc_iv, 16);
+
+	if (salt->version == 1) {
+		uchar content[MAX_CONT_SIZE];
+		int contentsize = salt->contentsize;
+		int datasize;
+
+		if (contentsize < 16 || contentsize > MAX_CONT_SIZE)
+			contentsize = 16;
+
+		memcpy_macro(content, salt->contents, contentsize);
+
+		if (salt->algorithm == 0) {
+			uint pad_byte;
+
+			AES_set_decrypt_key(hash, 256, &akey);
+			AES_cbc_decrypt(content, content, contentsize, &akey, iv);
+			pad_byte = content[contentsize - 1];
+			datasize = contentsize - pad_byte;
+		} else /* if (salt->algorithm == 1) */ {
+			Twofish_key tkey;
+
+			Twofish_prepare_key(hash, 32, &tkey);
+			datasize = Twofish_Decrypt(&tkey, content, content,
+			                           contentsize, iv);
+
+			if (datasize < 0 || datasize > contentsize)
+				datasize = 0;
+		}
+		SHA256_Init(&ctx);
+		SHA256_Update(&ctx, content, datasize);
+		SHA256_Final(hash, &ctx);
+		result[gid].cracked = !memcmp_pc(hash, salt->contents_hash, 32);
+	}
+	else if (salt->version == 2) {
+		uchar content[32];
+
+		memcpy_macro(content, salt->contents, 32);
+
+		if (salt->algorithm == 0) {
+			AES_set_decrypt_key(hash, 256, &akey);
+			AES_cbc_decrypt(content, hash, 32, &akey, iv);
+		} else /* if (salt->algorithm == 2) */ {
+			chacha_ctx ckey;
+
+			chacha_keysetup(&ckey, hash, 256);
+			chacha_ivsetup(&ckey, iv, 0, 12);
+			chacha_decrypt_bytes(&ckey, content, hash, 32);
+		}
+		result[gid].cracked = !memcmp_pc(hash, salt->expected_bytes, 32);
+	}
+	else
+		result[gid].cracked = 0; // We should never end up here
+}

src/opencl/lotus5_kernel.cl

@@ -16,7 +16,7 @@
 #define MAYBE_CONSTANT __local const
 #endif
 
-static __constant uint magic_table[256] = {
+__const_a8 uint magic_table[256] = {
   0xbd, 0x56, 0xea, 0xf2, 0xa2, 0xf1, 0xac, 0x2a,
   0xb0, 0x93, 0xd1, 0x9c, 0x1b, 0x33, 0xfd, 0xd0,
   0x30, 0x04, 0xb6, 0xdc, 0x7d, 0xdf, 0x32, 0x4b,

src/opencl/office_kernel.cl

@@ -224,8 +224,8 @@ inline void Decrypt(__constant ms_office_salt *salt,
 	AES_cbc_decrypt(encryptedVerifier, decryptedVerifier, length, &akey, iv);
 }
 
-__constant uint InputBlockKeyInt[] = { 0xfea7d276, 0x3b4b9e79 };
-__constant uint ValueBlockKeyInt[] = { 0xd7aa0f6d, 0x3061344e };
+__const_a8 uint InputBlockKeyInt[] = { 0xfea7d276, 0x3b4b9e79 };
+__const_a8 uint ValueBlockKeyInt[] = { 0xd7aa0f6d, 0x3061344e };
 
 __kernel
 void Generate2010key(__global ms_office_state *state,

src/opencl/wpapsk_kernel.cl

@@ -188,7 +188,7 @@ void wpapsk_pass2(MAYBE_CONSTANT wpapsk_salt *salt,
 
 //__constant uchar *text = "Pairwise key expansion\0";
 //__constant uint text[6] = { 0x72696150, 0x65736977, 0x79656b20, 0x70786520, 0x69736e61, 0x00006e6f };
-__constant uint text[6] = { 0x50616972, 0x77697365, 0x206b6579, 0x20657870, 0x616e7369, 0x6f6e0000 };
+__const_a8 uint text[6] = { 0x50616972, 0x77697365, 0x206b6579, 0x20657870, 0x616e7369, 0x6f6e0000 };
 
 inline void prf_512(const MAYBE_VECTOR_UINT *key,
                     MAYBE_CONSTANT uint *data,