OpenCL AES: Correction for supporting two or more AES contexts at once

This currently applies to AES-XTS, used by diskcryptor and truecrypt.

Also separate Truecrypt kernel from the (potentially) shared PBKDF2-RIPE-
MD160 code for good measure.

Nvidia 550.144.03 bugs out on this for truecrypt so we revert that format
to bitsliced AES (it doesn't matter for speed).

Author: magnumripper

run/opencl/diskcryptor_aes_kernel.cl

@@ -9,18 +9,7 @@
 #define AES_SRC_TYPE MAYBE_CONSTANT
 
 #include "pbkdf2_hmac_sha512_kernel.cl"
-
-/*
- * AES_256_XTS uses two AES keys at once so need double the
- * shared memory.
- */
-#define AES_SHARED_THREADS_DECREASED  1
-#if gpu_amd(DEVICE_INFO)
-#define AES_SHARED_THREADS            (WARP_SIZE >> 2)
-#else
-#define AES_SHARED_THREADS            (WARP_SIZE >> 1)
-#endif
-#include "opencl_aes.h"
+#include "opencl_aes_xts.h"
 
 typedef struct {
 	salt_t pbkdf2;

run/opencl/opencl_aes.h

@@ -313,74 +313,7 @@ INLINE void AES_cfb_decrypt(AES_SRC_TYPE void *_in,
 	}
 }
 
-INLINE void AES_256_XTS_first_sector(AES_SRC_TYPE uint *in, AES_DST_TYPE uint *out,
-                                     AES_KEY_TYPE uchar *double_key,
-                                     __local aes_local_t *lt1, __local aes_local_t *lt2)
-{
-	uint tweak[4] = { 0 };
-	uint buf[4];
-	int i;
-	AES_KEY akey1, akey2; akey1.lt = lt1; akey2.lt = lt2;
-
-	AES_set_decrypt_key(double_key, 256, &akey1);
-	AES_set_encrypt_key(double_key + 32, 256, &akey2);
-
-	AES_encrypt((uchar*)tweak, (uchar*)tweak, &akey2);
-
-	for (i = 0; i < 4; i++)
-		buf[i] = in[i] ^ tweak[i];
-
-	AES_decrypt((uchar*)buf, (uchar*)buf, &akey1);
-
-	for (i = 0; i < 4; i++)
-		out[i] = buf[i] ^ tweak[i];
-}
-
-INLINE void AES_256_XTS_DiskCryptor(AES_SRC_TYPE uchar *data, AES_DST_TYPE uchar *output,
-                                    AES_KEY_TYPE uchar *double_key, int len,
-                                    __local aes_local_t *lt1, __local aes_local_t *lt2)
-{
-	uchar buf[16];
-	int i, j, cnt;
-	AES_KEY key1; key1.lt = lt1;
-	AES_KEY key2; key2.lt = lt2;
-	int bits = 256;
-	uchar buffer[96];
-	uchar *out = buffer;
-	unsigned char tweak[16] = { 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 };
-
-	AES_set_decrypt_key(double_key, bits, &key1);
-	AES_set_encrypt_key(&double_key[bits / 8], bits, &key2);
-
-	// first aes tweak, we do it right over tweak
-	AES_encrypt(tweak, tweak, &key2);
-
-	cnt = len / 16;
-	for (j = 0;;) {
-		for (i = 0; i < 16; ++i) buf[i] = data[i]^tweak[i];
-		AES_decrypt(buf, out, &key1);
-		for (i = 0; i < 16; ++i) out[i] ^= tweak[i];
-		++j;
-		if (j == cnt)
-			break;
-		else {
-			unsigned char Cin, Cout;
-			unsigned x;
-			Cin = 0;
-			for (x = 0; x < 16; ++x) {
-				Cout = (tweak[x] >> 7) & 1;
-				tweak[x] = ((tweak[x] << 1) + Cin) & 0xFF;
-				Cin = Cout;
-			}
-			if (Cout)
-				tweak[0] ^= 135; // GF_128_FDBK;
-		}
-		data += 16;
-		out += 16;
-	}
-
-	memcpy_macro(output, buffer, 96);
-}
+/* AES-XTS has its own file opencl_aes_xts.h */
 
 #define N_WORDS (AES_BLOCK_SIZE / sizeof(unsigned long))
 

run/opencl/opencl_aes_plain.h

@@ -27,22 +27,26 @@
 #define AES_LOCAL_TABLES	1
 #endif
 
+/*
+ * Formats using two or more keys at once (AES-XTS uses two) must set this. It
+ * needs to be a power of two so is named and used as a shift.
+ */
+#ifndef AES_SIMULTANEOUS_CTX_SHIFT
+#define AES_SIMULTANEOUS_CTX_SHIFT    0
+#endif
+
 /*
  * Even with 64K LDS, an AMD device can't fit exclusive tables to every thread
- * in a wavefront, so we have to decrease the number.
- * A format can force this if it uses two or more keys at once. (diskcryptor
- * does)
+ * in a wavefront of 64 threads, so we have to decrease the number.
+ * Also, the number of simultaneous AES contexts need to be considered per above.
  */
-#ifndef AES_SHARED_THREADS_DECREASED
 #if SHARED_MEM_SIZE < (WARP_SIZE * (256*4 + 256) + 2*4 + 4)
-#define AES_SHARED_THREADS_DECREASED  1
-#define AES_SHARED_THREADS            (WARP_SIZE >> 1)
+#define AES_SHARED_THREADS            (WARP_SIZE >> (AES_SIMULTANEOUS_CTX_SHIFT + 1))
 #else
-#define AES_SHARED_THREADS            WARP_SIZE
+#define AES_SHARED_THREADS            (WARP_SIZE >> (AES_SIMULTANEOUS_CTX_SHIFT))
 #endif
-#endif	/* AES_SHARED_THREADS_DECREASED */
 
-#define AES_SHARED_THREADS_MASK	(AES_SHARED_THREADS - 1)
+#define AES_SHARED_THREADS_MASK       (AES_SHARED_THREADS - 1)
 
 #include "opencl_aes_tables.h"
 #if AES_LOCAL_TABLES

run/opencl/opencl_aes_xts.h

@@ -0,0 +1,91 @@
+/*
+ * AES OpenCL XTS functions
+ *
+ * Copyright (c) 2017-2025, magnum.
+ *
+ * This software 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.
+ */
+
+#ifndef _OPENCL_AES_XTS_H_
+#define _OPENCL_AES_XTS_H_
+
+#ifdef _OPENCL_AES_H_
+#error "opencl_aes_xts.h cannot be sourced after opencl_aes.h"
+#endif
+
+/* Tell the AES code we use two contexts simultaneously */
+#define AES_SIMULTANEOUS_CTX_SHIFT    1
+#include "opencl_aes.h"
+
+INLINE void AES_256_XTS_first_sector(AES_SRC_TYPE uint *in, AES_DST_TYPE uint *out,
+                                     AES_KEY_TYPE uchar *double_key,
+                                     __local aes_local_t *lt1, __local aes_local_t *lt2)
+{
+	uint tweak[4] = { 0 };
+	uint buf[4];
+	int i;
+	AES_KEY akey1, akey2; akey1.lt = lt1; akey2.lt = lt2;
+
+	AES_set_decrypt_key(double_key, 256, &akey1);
+	AES_set_encrypt_key(double_key + 32, 256, &akey2);
+
+	AES_encrypt((uchar*)tweak, (uchar*)tweak, &akey2);
+
+	for (i = 0; i < 4; i++)
+		buf[i] = in[i] ^ tweak[i];
+
+	AES_decrypt((uchar*)buf, (uchar*)buf, &akey1);
+
+	for (i = 0; i < 4; i++)
+		out[i] = buf[i] ^ tweak[i];
+}
+
+INLINE void AES_256_XTS_DiskCryptor(AES_SRC_TYPE uchar *data, AES_DST_TYPE uchar *output,
+                                    AES_KEY_TYPE uchar *double_key, int len,
+                                    __local aes_local_t *lt1, __local aes_local_t *lt2)
+{
+	uchar buf[16];
+	int i, j, cnt;
+	AES_KEY key1; key1.lt = lt1;
+	AES_KEY key2; key2.lt = lt2;
+	int bits = 256;
+	uchar buffer[96];
+	uchar *out = buffer;
+	unsigned char tweak[16] = { 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 };
+
+	AES_set_decrypt_key(double_key, bits, &key1);
+	AES_set_encrypt_key(&double_key[bits / 8], bits, &key2);
+
+	// first aes tweak, we do it right over tweak
+	AES_encrypt(tweak, tweak, &key2);
+
+	cnt = len / 16;
+	for (j = 0;;) {
+		for (i = 0; i < 16; ++i) buf[i] = data[i]^tweak[i];
+		AES_decrypt(buf, out, &key1);
+		for (i = 0; i < 16; ++i) out[i] ^= tweak[i];
+		++j;
+		if (j == cnt)
+			break;
+		else {
+			unsigned char Cin, Cout;
+			unsigned x;
+			Cin = 0;
+			for (x = 0; x < 16; ++x) {
+				Cout = (tweak[x] >> 7) & 1;
+				tweak[x] = ((tweak[x] << 1) + Cin) & 0xFF;
+				Cin = Cout;
+			}
+			if (Cout)
+				tweak[0] ^= 135; // GF_128_FDBK;
+		}
+		data += 16;
+		out += 16;
+	}
+
+	memcpy_macro(output, buffer, 96);
+}
+
+#endif	/* _OPENCL_AES_XTS_H_ */

run/opencl/pbkdf2_ripemd160_kernel.cl renamed to run/opencl/pbkdf2_ripemd160.h

@@ -1,5 +1,5 @@
 /*
- * RIPEMD-160 implementation. Copyright (c) 2015, magnum
+ * RIPEMD-160 implementation. Copyright (c) 2015-2025, magnum
  * This software 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.
@@ -11,19 +11,6 @@
 
 #include "opencl_misc.h"
 #include "opencl_ripemd.h"
-#define AES_SRC_TYPE __constant
-#define AES_DST_TYPE __global
-/*
- * AES_256_XTS uses two AES keys at once so need double the
- * shared memory. These two are only needed if we re-test with table AES
- */
-#define AES_SHARED_THREADS_DECREASED  1
-#if gpu_amd(DEVICE_INFO)
-#define AES_SHARED_THREADS            (WARP_SIZE >> 2)
-#else
-#define AES_SHARED_THREADS            (WARP_SIZE >> 1)
-#endif
-#include "opencl_aes.h"
 
 #define ITERATIONS 2000
 
@@ -32,15 +19,6 @@ typedef struct {
 	uchar v[KEYLEN];
 } pbkdf2_password;
 
-typedef struct {
-	uint v[16 / 4];
-} tc_hash;
-
-typedef struct {
-	uint salt[SALTLEN / 4];
-	uint bin[(512 - 64) / 4];
-} tc_salt;
-
 INLINE void preproc(__global const uchar *key, uint keylen, uint *state,
                     uint padding)
 {
@@ -152,20 +130,3 @@ INLINE void pbkdf2(__global const uchar *pass, uint passlen,
 			PUTCHAR(out, t, ((uchar*)tmp_out)[i]);
 	}
 }
-
-__kernel void tc_ripemd_aesxts(__global const pbkdf2_password *inbuffer,
-                               __global tc_hash *outbuffer,
-                               __constant tc_salt *salt)
-{
-	__local aes_local_t lt1;
-	__local aes_local_t lt2;
-	uint idx = get_global_id(0);
-	union {
-		uint u32[64 / 4];
-		uchar uc[64];
-	} key;
-
-	pbkdf2(inbuffer[idx].v, inbuffer[idx].length, salt->salt, key.u32);
-
-	AES_256_XTS_first_sector(salt->bin, outbuffer[idx].v, key.uc, &lt1, &lt2);
-}

run/opencl/truecrypt_kernel.cl

@@ -0,0 +1,38 @@
+/*
+ * Truecrypt implementation. Copyright (c) 2015-2025, magnum
+ * This software 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_SRC_TYPE __constant
+#define AES_DST_TYPE __global
+#include "opencl_aes_xts.h"
+#include "pbkdf2_ripemd160.h"
+
+typedef struct {
+	uint v[16 / 4];
+} tc_hash;
+
+typedef struct {
+	uint salt[SALTLEN / 4];
+	uint bin[(512 - 64) / 4];
+} tc_salt;
+
+__kernel void tc_ripemd_aesxts(__global const pbkdf2_password *inbuffer,
+                               __global tc_hash *outbuffer,
+                               __constant tc_salt *salt)
+{
+	__local aes_local_t lt1;
+	__local aes_local_t lt2;
+	uint idx = get_global_id(0);
+	union {
+		uint u32[64 / 4];
+		uchar uc[64];
+	} key;
+
+	pbkdf2(inbuffer[idx].v, inbuffer[idx].length, salt->salt, key.u32);
+
+	AES_256_XTS_first_sector(salt->bin, outbuffer[idx].v, key.uc, &lt1, &lt2);
+}

src/opencl_tc_fmt_plug.c

@@ -194,7 +194,7 @@ static void reset(struct db_main *db)
 		         (int)sizeof(inbuffer->v),
 		         (int)sizeof(currentsalt.salt),
 		         OUTLEN);
-		opencl_init("$JOHN/opencl/pbkdf2_ripemd160_kernel.cl",
+		opencl_init("$JOHN/opencl/truecrypt_kernel.cl",
 		            gpu_id, build_opts);
 
 		crypt_kernel = clCreateKernel(program[gpu_id], "tc_ripemd_aesxts",