OpenCL AES formats: Adapt to new shared code

Enable local memory for table-based AES. Closes #5594 Bitlocker format is not affected as it has it's own implementation, but AES performance is insignificant for it anyway.
openwall · Dec 17, 2024 · 4c5de1b · 4c5de1b
1 parent c22161a
commit 4c5de1b
Show file tree

Hide file tree

Showing 38 changed files with 151 additions and 108 deletions.
diff --git a/run/opencl/7z_kernel.cl b/run/opencl/7z_kernel.cl
@@ -153,6 +153,7 @@ __kernel void sevenzip_final(__global const sevenzip_password *inbuffer,
 __kernel void sevenzip_aes(__constant sevenzip_salt *salt,
                            __global sevenzip_hash *outbuffer)
 {
+	__local aes_local_t lt;
 	uint gid = get_global_id(0);
 	uint i;
 	uint pad;
@@ -162,7 +163,7 @@ __kernel void sevenzip_aes(__constant sevenzip_salt *salt,
 	/* Early rejection if possible (only decrypt last 16 bytes) */
 	if (pad > 0 && salt->length >= 32) {
 		uint8_t buf[16];
-		AES_KEY akey;
+		AES_KEY akey; akey.lt = &lt;
 		unsigned char iv[16];
 
 		for (i = 0; i < 16; i++)

diff --git a/run/opencl/agile_kernel.cl b/run/opencl/agile_kernel.cl
@@ -30,8 +30,9 @@ __kernel void dk_decrypt(__global pbkdf2_password *password,
                          __global agile_out *agile_out,
                          __constant agile_salt *salt)
 {
+	__local aes_local_t lt;
+	AES_KEY akey; akey.lt = &lt;
 	uint idx = get_global_id(0);
-	AES_KEY akey;
 	uchar iv[16];
 	uchar plaintext[16];
 	uint i;

diff --git a/run/opencl/androidbackup_kernel.cl b/run/opencl/androidbackup_kernel.cl
@@ -21,12 +21,12 @@ typedef struct {
 	uchar masterkey_blob[MAX_MASTERKEYBLOB_LEN];
 } ab_salt;
 
-inline int ab_decrypt(__global uchar *key, MAYBE_CONSTANT ab_salt *salt)
+inline int ab_decrypt(__global uchar *key, MAYBE_CONSTANT ab_salt *salt, __local aes_local_t *lt)
 {
 	uchar out[MAX_MASTERKEYBLOB_LEN];
 	const int length = salt->masterkey_blob_length;
 	uchar aiv[16];
-	AES_KEY akey;
+	AES_KEY akey; akey.lt = lt;
 	int pad_byte;
 
 	memcpy_macro(aiv, salt->iv, 16);
@@ -51,7 +51,8 @@ void ab_final(MAYBE_CONSTANT ab_salt *salt,
                __global pbkdf2_out *pbkdf2,
                __global ab_out *out)
 {
+	__local aes_local_t lt;
 	uint gid = get_global_id(0);
 
-	out[gid].cracked = ab_decrypt((__global uchar*)pbkdf2[gid].dk, salt);
+	out[gid].cracked = ab_decrypt((__global uchar*)pbkdf2[gid].dk, salt, &lt);
 }
diff --git a/run/opencl/axcrypt2_kernel.cl b/run/opencl/axcrypt2_kernel.cl
@@ -35,6 +35,8 @@ __kernel void axcrypt2_final(__global crack_t *pbkdf2,
                              __constant axcrypt2_salt_t *salt,
                              __global out_t *out)
 {
+	__local aes_local_t lt;
+	AES_KEY akey; akey.lt = &lt;
 	uint gid = get_global_id(0);
 
 	int i, k, j, nb_iterations = salt->key_wrapping_rounds;
@@ -49,7 +51,6 @@ __kernel void axcrypt2_final(__global crack_t *pbkdf2,
 		key.u[i] = SWAP64(pbkdf2[gid].hash[i]);
 
 	uchar KEK[32];
-	AES_KEY akey;
 	int halfblocklen = 16 / 2;
 	int wrappedkeylen = 56 - halfblocklen;
 	union {

diff --git a/run/opencl/axcrypt_kernel.cl b/run/opencl/axcrypt_kernel.cl
@@ -43,7 +43,7 @@ typedef struct {
                 (cp)[2] ^= (uchar)((value) >> 16), \
                 (cp)[3] ^= (uchar)((value) >> 24 ) )
 
-inline int axcrypt_decrypt(__global const axcrypt_password *inbuffer, uint gid, __constant axcrypt_salt *cur_salt, __global axcrypt_out *output)
+inline int axcrypt_decrypt(__global const axcrypt_password *inbuffer, uint gid, __constant axcrypt_salt *cur_salt, __global axcrypt_out *output, __local aes_local_t *lt)
 {
 	uchar password[PLAINTEXT_LENGTH];
 	uchar keyfile[4096];
@@ -59,7 +59,7 @@ inline int axcrypt_decrypt(__global const axcrypt_password *inbuffer, uint gid,
 		uint32_t w[4];
 		uint64_t l[2];
 	} cipher;
-	AES_KEY akey;
+	AES_KEY akey; akey.lt = lt;
 	SHA_CTX ctx;
 	uint i;
 	int j, nb_iterations = cur_salt->key_wrapping_rounds;
@@ -117,7 +117,8 @@ void axcrypt(__global const axcrypt_password *inbuffer,
                 __global axcrypt_out *out,
                 __constant axcrypt_salt *salt)
 {
+	__local aes_local_t lt;
 	uint idx = get_global_id(0);
 
-	out[idx].cracked = axcrypt_decrypt(inbuffer, idx, salt, out);
+	out[idx].cracked = axcrypt_decrypt(inbuffer, idx, salt, out, &lt);
 }
diff --git a/run/opencl/bitcoin_kernel.cl b/run/opencl/bitcoin_kernel.cl
@@ -84,6 +84,8 @@ __kernel void loop_sha512(__global hash512_t *state, uint count)
 
 __kernel void bitcoin_final(__constant salt_t *salt, __global hash512_t *state, __global uint *cracked)
 {
+	__local aes_local_t lt;
+	AES_KEY aes_key; aes_key.lt = &lt;
 	uint gid = get_global_id(0);
 	uchar iv[16];  // updated IV for the final block
 	memcpy_cp(iv, salt->cry_master + salt->cry_master_length - 32, 16);
@@ -93,7 +95,6 @@ __kernel void bitcoin_final(__constant salt_t *salt, __global hash512_t *state,
 		state[gid].W[i] = SWAP64(state[gid].W[i]);
 
 	uchar output[16];
-	AES_KEY aes_key;
 	AES_set_decrypt_key(state[gid].b, 256, &aes_key);
 	AES_cbc_decrypt(salt->cry_master + salt->cry_master_length - 16, output, 16, &aes_key, iv);
 

diff --git a/run/opencl/bitwarden_kernel.cl b/run/opencl/bitwarden_kernel.cl
@@ -31,9 +31,10 @@ __kernel void bitwarden_decrypt(MAYBE_CONSTANT bitwarden_salt_t *salt,
                                 __global crack_t *out,
                                 __global uint32_t *cracked)
 {
+	__local aes_local_t lt;
+	AES_KEY akey; akey.lt = &lt;
 	uint32_t gid = get_global_id(0);
 	int32_t i;
-	AES_KEY akey;
 	union {
 		uchar c[32];
 		uint  w[32 / 4];

diff --git a/run/opencl/blockchain_kernel.cl b/run/opencl/blockchain_kernel.cl
@@ -34,10 +34,10 @@ typedef struct {
 } blockchain_salt;
 
 inline int blockchain_decrypt(__global uchar *derived_key,
-                              __constant uchar *data)
+                              __constant uchar *data, __local aes_local_t *lt)
 {
+	AES_KEY akey; akey.lt = lt;
 	uchar out[SAFETY_FACTOR];
-	AES_KEY akey;
 	uchar iv[16];
 
 	AES_set_decrypt_key(derived_key, 256, &akey);
@@ -68,11 +68,12 @@ __kernel void blockchain(__global const pbkdf2_password *inbuffer,
                          __constant blockchain_salt *salt,
                          __global blockchain_out *out)
 {
+	__local aes_local_t lt;
 	uint idx = get_global_id(0);
 
 	pbkdf2(inbuffer[idx].v, inbuffer[idx].length,
 	       salt->pbkdf2.salt, salt->pbkdf2.length, salt->pbkdf2.iterations,
 	       dk[idx].v, salt->pbkdf2.outlen, salt->pbkdf2.skip_bytes);
 
-	out[idx].cracked = blockchain_decrypt((__global uchar*)dk[idx].v, salt->data);
+	out[idx].cracked = blockchain_decrypt((__global uchar*)dk[idx].v, salt->data, &lt);
 }
diff --git a/run/opencl/bsd_softraid_kernel.cl b/run/opencl/bsd_softraid_kernel.cl
@@ -29,11 +29,12 @@ __kernel
 void softraid_final(MAYBE_CONSTANT softraid_salt *salt,
                     __global pbkdf2_out *out)
 {
+	__local aes_local_t lt;
+	AES_KEY akey; akey.lt = &lt;
 	uint gid = get_global_id(0);
 	uint dk[OUTLEN / 4];
 	uchar unmasked_keys[MASKED_KEY_SIZE];
 	uchar hashed_mask_key[SHA1_DIGEST_LENGTH];
-	AES_KEY akey;
 	SHA_CTX ctx;
 
 	memcpy_gp(dk, out[gid].dk, OUTLEN);

diff --git a/run/opencl/cryptosafe_kernel.cl b/run/opencl/cryptosafe_kernel.cl
@@ -83,6 +83,8 @@ void cryptoSafe(__global const uchar *pwbuf,
 #define GPU_LOC_3 LOC_3
 #endif
 
+	__local aes_local_t lt;
+
 	/* Prepare password, pad to length 32 with ASCII '0's */
 	prepare(pwbuf, index, password);
 
@@ -113,7 +115,7 @@ void cryptoSafe(__global const uchar *pwbuf,
 #endif
 #endif
 
-		AES_KEY aes_decrypt_key;
+		AES_KEY aes_decrypt_key; aes_decrypt_key.lt = &lt;
 		unsigned char plain[16], iv[16] = { 0 };
 
 		AES_set_decrypt_key(password, 256, &aes_decrypt_key);

diff --git a/run/opencl/diskcryptor_aes_kernel.cl b/run/opencl/diskcryptor_aes_kernel.cl
@@ -24,6 +24,8 @@ __kernel void diskcryptor_final(__global crack_t *pbkdf2,
                          __constant diskcryptor_salt_t *salt,
                          __global out_t *out)
 {
+	__local aes_local_t lt;
+
 	uint gid = get_global_id(0);
 
 	uchar output[96];
@@ -48,7 +50,7 @@ __kernel void diskcryptor_final(__global crack_t *pbkdf2,
 	for (i = 0; i < 8; i++)
 		key.u[i] = SWAP64(pbkdf2[gid].hash[i]);
 
-	AES_256_XTS_DiskCryptor(salt->header, output, key.c, 96);
+	AES_256_XTS_DiskCryptor(salt->header, output, key.c, 96, &lt);
 	memcpy_pp(version.c, output + 72, 2);
 	memcpy_pp(algorithm.c, output + 82, 4);
 	if ((!memcmp_pc(output + 64, "DCRP", 4)) && (version.value == 2 || version.value == 1) && (algorithm.value >= 0 && algorithm.value <= 7)) {

diff --git a/run/opencl/dmg_kernel.cl b/run/opencl/dmg_kernel.cl
@@ -12,9 +12,6 @@ typedef struct {
 
 #define pbkdf2_out dmg_out
 #include "pbkdf2_hmac_sha1_kernel.cl"
-#if __OS_X__
-#define AES_NO_BITSLICE
-#endif
 #define AES_SRC_TYPE MAYBE_CONSTANT
 #include "opencl_aes.h"
 #include "opencl_hmac_sha1.h"
@@ -94,10 +91,10 @@ inline int check_v1hash(const uchar *derived_key,
 }
 
 inline int check_v2hash(const uchar *derived_key,
-                        MAYBE_CONSTANT dmg_salt *salt)
+                        MAYBE_CONSTANT dmg_salt *salt, __local aes_local_t *lt)
 {
 	des3_context ks;
-	AES_KEY aes_decrypt_key;
+	AES_KEY aes_decrypt_key; aes_decrypt_key.lt = lt;
 	uint buf[8192/4];
 	uchar *outbuf = (uchar*)buf;
 	uchar iv[20];
@@ -155,10 +152,11 @@ __kernel
 void dmg_final_v2(MAYBE_CONSTANT dmg_salt *salt,
                   __global dmg_out *out)
 {
+	__local aes_local_t lt;
 	uint gid = get_global_id(0);
 	uint dk[OUTLEN / 4];
 
 	memcpy_gp(dk, out[gid].dk, OUTLEN);
 
-	out[gid].cracked = check_v2hash((uchar*)dk, salt);
+	out[gid].cracked = check_v2hash((uchar*)dk, salt, &lt);
 }
diff --git a/run/opencl/encfs_kernel.cl b/run/opencl/encfs_kernel.cl
@@ -130,10 +130,10 @@ inline uint encfs_common_MAC_32(MAYBE_CONSTANT encfs_salt *salt, uchar *src,
 
 inline void encfs_common_streamDecode(MAYBE_CONSTANT encfs_salt *salt,
                                       uchar *buf, uint size, uint64_t iv64,
-                                      uchar *key)
+                                      uchar *key, __local aes_local_t *lt)
 {
 	uchar ivec[MAX_IVLENGTH];
-	AES_KEY akey;
+	AES_KEY akey; akey.lt = lt;
 
 	encfs_common_setIVec(salt, ivec, iv64 + 1, key);
 	AES_set_encrypt_key(key, salt->keySize * 8, &akey);
@@ -151,6 +151,7 @@ void encfs_final(MAYBE_CONSTANT encfs_salt *salt,
                  __global pbkdf2_out *pbkdf2,
                  __global encfs_out *out)
 {
+	__local aes_local_t lt;
 	uint gid = get_global_id(0);
 	uint i;
 	uchar master[MAX_KEYLENGTH + MAX_IVLENGTH];
@@ -165,7 +166,7 @@ void encfs_final(MAYBE_CONSTANT encfs_salt *salt,
 		checksum = (checksum << 8) | salt->data[i];
 
 	memcpy_mcp(tmpBuf, salt->data + KEY_CHECKSUM_BYTES, salt->keySize + salt->ivLength);
-	encfs_common_streamDecode(salt, tmpBuf, salt->keySize + salt->ivLength ,checksum, master);
+	encfs_common_streamDecode(salt, tmpBuf, salt->keySize + salt->ivLength ,checksum, master, &lt);
 	checksum2 = encfs_common_MAC_32(salt, tmpBuf, salt->keySize + salt->ivLength, master);
 
 	out[gid].cracked = (checksum2 == checksum);

diff --git a/run/opencl/enpass_kernel.cl b/run/opencl/enpass_kernel.cl
@@ -53,6 +53,8 @@ void enpass5_final(MAYBE_CONSTANT enpass_salt *salt,
                    __global enpass_out *out,
                    __global pbkdf2_state *state)
 {
+	__local aes_local_t lt;
+	AES_KEY akey; akey.lt = &lt;
 	uint gid = get_global_id(0);
 	uint i;
 	uint base = state[gid].pass++ * 5;
@@ -80,7 +82,6 @@ void enpass5_final(MAYBE_CONSTANT enpass_salt *salt,
 #endif
 	} else {
 		uchar data[16];
-		AES_KEY akey;
 		union {
 			uchar c[256/8];
 			uint  w[256/8/4];
@@ -155,10 +156,11 @@ void enpass6_final(MAYBE_CONSTANT enpass_salt *salt,
                    __global enpass_out *out,
                    __global crack_t *out512)
 {
+	__local aes_local_t lt;
+	AES_KEY akey; akey.lt = &lt;
 	uint gid = get_global_id(0);
 	uint i;
 	uchar data[16];
-	AES_KEY akey;
 	union {
 		uchar c[256/8];
 		ulong  w[256/8/8];

diff --git a/run/opencl/ethereum_kernel.cl b/run/opencl/ethereum_kernel.cl
@@ -60,8 +60,9 @@ __kernel void ethereum_presale_process(__global crack_t *pbkdf2_out,
                                        __global state_t *state,
                                        __global hash_t *out)
 {
+	__local aes_local_t lt;
+	AES_KEY akey; akey.lt = &lt;
 	uint32_t gid = get_global_id(0);
-	AES_KEY akey;
 	uchar iv[16];
 	int i;
 	uchar seed[1024 + 1];

diff --git a/run/opencl/fvde_kernel.cl b/run/opencl/fvde_kernel.cl
@@ -29,6 +29,8 @@ __kernel void fvde_decrypt(MAYBE_CONSTANT fvde_salt_t *salt,
                            __global crack_t *out,
                            __global uint32_t *cracked)
 {
+	__local aes_local_t lt;
+	AES_KEY akey; akey.lt = &lt;
 	uint32_t gid = get_global_id(0);
 	MAYBE_CONSTANT uint64_t *C = salt->blob.qword; // len(C) == 3 or 5 (AES-256)
 	int32_t n = BLOBLEN / 8 - 1;  // len(C) - 1
@@ -38,7 +40,6 @@ __kernel void fvde_decrypt(MAYBE_CONSTANT fvde_salt_t *salt,
 		uint8_t stream[16];
 	} todecrypt;
 	int32_t i, j;
-	AES_KEY akey;
 	uint64_t A = C[0];
 
 	if (salt->type == 1) {

diff --git a/run/opencl/geli_kernel.cl b/run/opencl/geli_kernel.cl
@@ -38,6 +38,8 @@ __kernel void geli_final(__global crack_t *pbkdf2,
                          __constant geli_salt_t *salt,
                          __global out_t *out)
 {
+	__local aes_local_t lt;
+	AES_KEY aes_decrypt_key; aes_decrypt_key.lt = &lt;
 	uint gid = get_global_id(0);
 	__constant uchar *mmkey;
 	const uchar nullstring[1] = { 0 };
@@ -64,7 +66,6 @@ __kernel void geli_final(__global crack_t *pbkdf2,
 	for (nkey = 0; nkey < G_ELI_MAXMKEYS; nkey++, mmkey += G_ELI_MKEYLEN) {
 		int bit = (1 << nkey);
 		uchar iv[16] = { 0 };
-		AES_KEY aes_decrypt_key;
 		uchar tmpmkey[G_ELI_MKEYLEN];
 		const uchar *odhmac; /* On-disk HMAC. */
 		uchar chmac[SHA512_MDLEN]; /* Calculated HMAC. */

diff --git a/run/opencl/iwork_kernel.cl b/run/opencl/iwork_kernel.cl
@@ -39,9 +39,10 @@ void iwork_final(MAYBE_CONSTANT iwork_salt *salt,
                  __global iwork_out *result,
                  __global pbkdf2_state *state)
 {
+	__local aes_local_t lt;
+	AES_KEY akey; akey.lt = &lt;
 	uint gid = get_global_id(0);
 	uint i;
-	AES_KEY akey;
 	int success = 1; // hash was cracked
 	union {
 		uchar c[BLOBLEN];