better EMMC CID brute suggested by dark_samus3, done

Makefile

@@ -7,7 +7,7 @@ LDFLAGS += -L$(INTELOCLSDKROOT)/lib/x64
 all : $(PNAME)
 
 $(PNAME) : $(OBJS)
-	$(CC) $(LDFLAGS) -o $@ $^ -lOpenCL
+	$(CC) $(LDFLAGS) -o $@ $^ -lOpenCL -static -lmbedcrypto
 
 clean :
 	rm $(PNAME) *.o

aes_128.c

@@ -1,5 +1,91 @@
 
+#include <stdio.h>
 #include <mbedtls/config.h>
+#include <mbedtls/version.h>
 #include <mbedtls/aes.h>
 #include <mbedtls/aesni.h>
-#include <mbedtls/version.h>
+#include "crypto.h"
+
+static mbedtls_aes_context ctx;
+
+static int (*p_aes_crypt_ecb)(mbedtls_aes_context*, int, const unsigned char *, unsigned char *) = NULL;
+
+static void (*p_aes_set_key_enc_128)(const unsigned char *key) = NULL;
+
+static void (*p_aes_set_key_dec_128)(const unsigned char *key) = NULL;
+
+#define AES_KEY_LEN 128
+#define NR 10
+
+// I hope eliminating the AESNI check can make it a bit faster
+static void aes_set_key_enc_128_aesni(const unsigned char *key){
+	// mbedtls_aes_setkey_enc(&ctx, key, 128);
+	ctx.nr = NR;
+	ctx.rk = ctx.buf;
+	mbedtls_aesni_setkey_enc((unsigned char *)ctx.rk, key, AES_KEY_LEN);
+}
+
+static void aes_set_key_dec_128_aesni(const unsigned char *key) {
+	mbedtls_aes_context cty;
+	cty.nr = NR;
+	cty.rk = cty.buf;
+	mbedtls_aesni_setkey_enc((unsigned char *)cty.rk, key, AES_KEY_LEN);
+	ctx.nr = cty.nr;
+	ctx.rk = ctx.buf;
+	mbedtls_aesni_inverse_key((unsigned char *)ctx.rk, (const unsigned char *)cty.rk, ctx.nr);
+}
+
+static void aes_set_key_enc_128_c(const unsigned char *key) {
+	mbedtls_aes_setkey_enc(&ctx, key, AES_KEY_LEN);
+}
+
+static void aes_set_key_dec_128_c(const unsigned char *key) {
+	mbedtls_aes_setkey_dec(&ctx, key, AES_KEY_LEN);
+}
+
+void aes_init(){
+	fputs(MBEDTLS_VERSION_STRING_FULL, stdout);
+	mbedtls_aes_init(&ctx);
+	// prevent runtime checks
+	if(mbedtls_aesni_has_support(MBEDTLS_AESNI_AES)){
+		puts(", AES-NI supported");
+		p_aes_crypt_ecb = mbedtls_aesni_crypt_ecb;
+		p_aes_set_key_enc_128 = aes_set_key_enc_128_aesni;
+		p_aes_set_key_dec_128 = aes_set_key_dec_128_aesni;
+	}else {
+		puts(", AES-NI not supported");
+		p_aes_crypt_ecb = mbedtls_aes_crypt_ecb;
+		p_aes_set_key_enc_128 = aes_set_key_enc_128_c;
+		p_aes_set_key_dec_128 = aes_set_key_dec_128_c;
+	}
+#ifndef MBEDTLS_AES_ROM_TABLES
+	// it will error out but also get aes_gen_tables done
+	mbedtls_aes_setkey_enc(&ctx, NULL, 0);
+#endif
+}
+
+void aes_set_key_enc_128(const unsigned char *key) {
+	p_aes_set_key_enc_128(key);
+}
+
+void aes_set_key_dec_128(const unsigned char *key) {
+	p_aes_set_key_dec_128(key);
+}
+
+void aes_encrypt_128(const unsigned char *in, unsigned char *out){
+	p_aes_crypt_ecb(&ctx, MBEDTLS_AES_ENCRYPT, in, out);
+}
+
+void aes_decrypt_128(const unsigned char *in, unsigned char *out){
+	p_aes_crypt_ecb(&ctx, MBEDTLS_AES_DECRYPT, in, out);
+}
+
+void aes_encrypt_128_bulk(const unsigned char *in, unsigned char *out, unsigned len){
+	len >>= 4;
+	for(unsigned i = 0; i < len; ++i){
+		p_aes_crypt_ecb(&ctx, MBEDTLS_AES_ENCRYPT, in, out);
+		in += AES_BLOCK_SIZE;
+		out += AES_BLOCK_SIZE;
+	}
+}
+

bfcl.c

@@ -14,6 +14,8 @@ static inline cl_ushort u16be(const unsigned char *in){
 	return out;
 }
 
+const char invalid_parameters[] = "invalid parameters\n";
+
 int main(int argc, const char *argv[]) {
 	int ret = 0;
 	if (argc == 1) {
@@ -43,7 +45,7 @@ int main(int argc, const char *argv[]) {
 			ret = ocl_brute_console_id(console_id, 0,
 				u16be(offset0), src0, ver0, u16be(offset1), src1, ver1, CTR);
 		} else {
-			puts("invalid parameters\n");
+			puts(invalid_parameters);
 			ret = -1;
 		}
 	} else if (argc == 7) {
@@ -63,11 +65,11 @@ int main(int argc, const char *argv[]) {
 		} else if (!strcmp(argv[1], "emmc_cid")) {
 			ret = ocl_brute_emmc_cid(console_id, emmc_cid, u16be(offset), src, ver);
 		} else {
-			puts("invalid parameters\n");
+			puts(invalid_parameters);
 			ret = -1;
 		}
 	} else {
-		printf("invalid parameters\n");
+		printf(invalid_parameters);
 		ret = -1;
 	}
 #ifdef _WIN32

cl/common.h

@@ -0,0 +1,7 @@
+
+typedef unsigned int uint32_t;
+
+typedef unsigned char u8;
+typedef unsigned int u32;
+typedef unsigned long u64;
+

cl/kernel_emmc_cid.cl

@@ -7,14 +7,12 @@ __kernel void test_emmc_cid(
 	if (*out) {
 		return;
 	}
-	u8 emmc_cid[16];
+	u64 emmc_cid[2] = { emmc_cid_l, emmc_cid_h };
-	*(u64*)emmc_cid = emmc_cid_l;
+	*(u32*)(((u8*)emmc_cid) + 1) |= get_global_id(0);
-	*(u64*)(emmc_cid + 8) = emmc_cid_h;
-	*(u32*)(emmc_cid + 1) |= get_global_id(0);
 
-	sha1_16((u32*)emmc_cid);
+	sha1_16((u8*)emmc_cid);
 
-	if (sha1_16_l == *(u64*)emmc_cid && sha1_16_h == *(u64*)(emmc_cid + 8)) {
+	if (sha1_16_l == emmc_cid[0] && sha1_16_h == emmc_cid[1]) {
 		*out = get_global_id(0);
 	}
 }

cl/sha1_16.cl

@@ -8,12 +8,30 @@ __constant const uint32_t
 	h3 = 0x10325476,
 	h4 = 0xC3D2E1F0;
 
-void sha1_16(uint32_t *io)
+#define GET_UINT32_BE(n,b,i)                            \
+{                                                       \
+	(n) = ( (uint32_t) (b)[(i)    ] << 24 )             \
+		| ( (uint32_t) (b)[(i) + 1] << 16 )             \
+		| ( (uint32_t) (b)[(i) + 2] <<  8 )             \
+		| ( (uint32_t) (b)[(i) + 3]       );            \
+}
+
+#define PUT_UINT32_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)       );       \
+}
+
+void sha1_16(unsigned char *io)
 {
 	uint32_t temp, W[16],
 		A = h0, B = h1, C = h2, D = h3, E = h4;
-
+	GET_UINT32_BE(W[0], io, 0);
-	W[0] = io[0]; W[1] = io[1]; W[2] = io[2]; W[3] = io[3];
+	GET_UINT32_BE(W[1], io, 4);
+	GET_UINT32_BE(W[2], io, 8);
+	GET_UINT32_BE(W[3], io, 12);
 	W[4] = 0x80000000u; W[5] = 0; W[6] = 0; W[7] = 0;
 	W[8] = 0; W[9] = 0; W[10] = 0; W[11] = 0;
 	W[12] = 0; W[13] = 0; W[14] = 0; W[15] = 0x80u;
@@ -149,9 +167,9 @@ void sha1_16(uint32_t *io)
 	C += h2;
 	D += h3;
 
-	io[0] = A;
+	PUT_UINT32_BE(A, io, 0);
-	io[1] = B;
+	PUT_UINT32_BE(B, io, 4);
-	io[2] = C;
+	PUT_UINT32_BE(C, io, 8);
-	io[3] = D;
+	PUT_UINT32_BE(D, io, 12);
 }
 

crypto.h

@@ -11,10 +11,9 @@ void sha1_16(const unsigned char in[16], unsigned char out[16]);
 void aes_init(void);
 
 void aes_set_key_enc_128(const unsigned char *key);
-
-void aes_encrypt_128(const unsigned char input[16], unsigned char output[16]);
-
 void aes_set_key_dec_128(const unsigned char *key);
 
+void aes_encrypt_128(const unsigned char input[16], unsigned char output[16]);
 void aes_decrypt_128(const unsigned char input[16], unsigned char output[16]);
 
+void aes_encrypt_128_bulk(const unsigned char input[16], unsigned char output[16], unsigned len);

dsi.h

@@ -39,7 +39,7 @@ static inline u32 u32be(const u8 *in){
 }
 
 // CAUTION this one doesn't work in-place
-static inline void byte_reverse_16(u8 *out, const u8 *in){
+inline void byte_reverse_16(u8 *out, const u8 *in){
 	out[0] = in[15];
 	out[1] = in[14];
 	out[2] = in[13];

ocl_brute.c

@@ -34,6 +34,33 @@ int ocl_brute_console_id(const cl_uchar *console_id, const cl_uchar *emmc_cid,
 	cl_uint offset1, const cl_uchar *src1, const cl_uchar *ver1,
 	ocl_brute_mode mode)
 {
+	// preparing args
+	cl_ulong console_id_template = u64be(console_id);
+	cl_ulong xor0[2] = { 0 }, xor1[2] = { 0 };
+	dsi_make_xor((u8*)xor0, src0, ver0);
+	if (src1 != 0) {
+		dsi_make_xor((u8*)xor1, src1, ver1);
+	}
+	cl_uint ctr[4] = { 0 };
+	if (emmc_cid != 0) {
+		dsi_make_ctr((u8*)ctr, emmc_cid, offset0);
+	}
+	cl_ulong out = 0;
+#if DEBUG
+	{
+		printf("XOR     : %s\n", hexdump(xor0, 16, 0));
+		u8 aes_key[16];
+		dsi_make_key(aes_key, u64be(console_id));
+		printf("AES KEY : %s\n", hexdump(aes_key, 16, 0));
+		aes_init();
+		aes_set_key_enc_128(aes_key);
+		printf("CTR     : %s\n", hexdump(ctr, 16, 0));
+		aes_encrypt_128((u8*)ctr, (u8*)xor0);
+		printf("XOR TRY : %s\n", hexdump(xor0, 16, 0));
+		// exit(1);
+	}
+#endif
+
 	TimeHP t0, t1; long long td = 0;
 
 	cl_int err;
@@ -73,35 +100,6 @@ int ocl_brute_console_id(const cl_uchar *console_id, const cl_uchar *emmc_cid,
 	OCL_ASSERT(clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL));
 	printf("local work size: %u\n", (unsigned)local);
 
-	// preparing args
-	cl_ulong console_id_template = u64be(console_id);
-	cl_ulong xor0[2] = { 0 }, xor1[2] = { 0 };
-	dsi_make_xor((u8*)xor0, src0, ver0);
-	if (src1 != 0) {
-		dsi_make_xor((u8*)xor1, src1, ver1);
-	}
-	cl_uint ctr[4] = { 0 };
-	if (emmc_cid != 0) {
-		dsi_make_ctr((u8*)ctr, emmc_cid, offset0);
-	}
-	cl_ulong out = 0;
-#if DEBUG
-	{
-		printf("XOR     : %s\n", hexdump(xor, 16, 0));
-		u8 aes_key[16];
-		dsi_make_key(aes_key, u64be(console_id));
-		printf("AES KEY : %s\n", hexdump(aes_key, 16, 0));
-		cl_uint aes_rk[RK_LEN];
-		aes_gen_tables();
-		aes_set_key_enc_128(aes_rk, aes_key);
-		printf("AES RK  : %s\n", hexdump(aes_rk, 48, 0));
-		printf("CTR     : %s\n", hexdump(ctr, 16, 0));
-		aes_encrypt_128(aes_rk, (u8*)ctr, (u8*)xor);
-		printf("XOR TRY : %s\n", hexdump(xor, 16, 0));
-		// exit(1);
-	}
-#endif
-
 	// there's no option to create it zero initialized
 	cl_mem mem_out = OCL_ASSERT2(clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_ulong), NULL, &err));
 	OCL_ASSERT(clEnqueueWriteBuffer(command_queue, mem_out, CL_TRUE, 0, sizeof(cl_ulong), &out, 0, NULL, NULL));
@@ -208,6 +206,28 @@ int ocl_brute_console_id(const cl_uchar *console_id, const cl_uchar *emmc_cid,
 int ocl_brute_emmc_cid(const cl_uchar *console_id, cl_uchar *emmc_cid,
 	cl_uint offset, const cl_uchar *src, const cl_uchar *ver)
 {
+	// preparing args
+	u8 aes_key[16];
+	dsi_make_key(aes_key, u64be(console_id));
+	aes_init();
+	aes_set_key_dec_128(aes_key);
+	cl_ulong xor[2];
+	dsi_make_xor((u8*)xor, src, ver);
+	cl_ulong ctr[2];
+	aes_decrypt_128((u8*)xor , (u8*)ctr);
+	cl_ulong emmc_cid_sha1_16[2];
+	byte_reverse_16((u8*)emmc_cid_sha1_16, (u8*)ctr);
+	sub_128_64(emmc_cid_sha1_16, offset);
+	cl_ulong out = 0;
+#ifdef DEBUG
+	{
+		printf("SHA1 A: %s\n", hexdump(emmc_cid_sha1_16, 16, 0));
+		u8 sha1_verify[16];
+		sha1_16(emmc_cid, sha1_verify);
+		printf("SHA1 B: %s\n", hexdump(sha1_verify, 16, 0));
+	}
+#endif
+
 	TimeHP t0, t1; long long td = 0;
 
 	cl_int err;
@@ -234,33 +254,6 @@ int ocl_brute_emmc_cid(const cl_uchar *console_id, cl_uchar *emmc_cid,
 	OCL_ASSERT(clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL));
 	printf("local work size: %u\n", (unsigned)local);
 
-	// preparing args
-	u8 aes_key[16];
-	dsi_make_key(aes_key, u64be(console_id));
-	aes_init();
-	aes_set_key_dec_128(aes_key);
-	cl_ulong xor[2];
-	dsi_make_xor((u8*)xor, src, ver);
-	cl_ulong ctr[2];
-	aes_decrypt_128((u8*)xor , (u8*)ctr);
-	cl_ulong emmc_cid_sha1_16[2];
-	byte_reverse_16((u8*)emmc_cid_sha1_16, (u8*)ctr);
-	sub_128_64(emmc_cid_sha1_16, offset);
-	cl_ulong out = 0;
-#ifdef DEBUG
-	{
-		printf("XOR     : %s\n", hexdump(xor, 16, 0));
-		printf("AES KEY : %s\n", hexdump(aes_key, 16, 0));
-		printf("AES RK  : %s\n", hexdump(aes_rk, 48, 0));
-		u8 ctr[16];
-		dsi_make_ctr(ctr, emmc_cid, u_offset);
-		printf("CTR     : %s\n", hexdump(ctr, 16, 0));
-		aes_encrypt_128(aes_rk, ctr, (u8*)xor);
-		printf("XOR TRY : %s\n", hexdump(xor, 16, 0));
-		// exit(1);
-	}
-#endif
-
 	// there's no option to create it zero initialized
 	cl_mem mem_out = OCL_ASSERT2(clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_ulong), NULL, &err));
 	OCL_ASSERT(clEnqueueWriteBuffer(command_queue, mem_out, CL_TRUE, 0, sizeof(cl_ulong), &out, 0, NULL, NULL));
@@ -280,8 +273,8 @@ int ocl_brute_emmc_cid(const cl_uchar *console_id, cl_uchar *emmc_cid,
 		puts(hexdump(emmc_cid, 16, 0));
 		OCL_ASSERT(clSetKernelArg(kernel, 0, sizeof(cl_ulong), emmc_cid));
 		OCL_ASSERT(clSetKernelArg(kernel, 1, sizeof(cl_ulong), emmc_cid + 8));
-		OCL_ASSERT(clSetKernelArg(kernel, 2, sizeof(cl_ulong), &emmc_cid_sha1_16[0]));
+		OCL_ASSERT(clSetKernelArg(kernel, 2, sizeof(cl_ulong), emmc_cid_sha1_16));
-		OCL_ASSERT(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &emmc_cid_sha1_16[1]));
+		OCL_ASSERT(clSetKernelArg(kernel, 3, sizeof(cl_ulong), emmc_cid_sha1_16 + 1));
 		OCL_ASSERT(clSetKernelArg(kernel, 4, sizeof(cl_mem), &mem_out));
 
 		OCL_ASSERT(clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &num_items, &local, 0, NULL, NULL));
@@ -308,7 +301,6 @@ int ocl_brute_emmc_cid(const cl_uchar *console_id, cl_uchar *emmc_cid,
 	printf("%.2f seconds, %.2f M/s\n", td / 1000000.0, tested * 1.0 / td);
 
 	clReleaseKernel(kernel);
-	clReleaseMemObject(mem_rk);
 	clReleaseMemObject(mem_out);
 	clReleaseProgram(program);
 	clReleaseCommandQueue(command_queue);

ocl_test.c

@@ -77,8 +77,7 @@ int ocl_test() {
 
 	srand(2501);
 	cl_uchar key[16];
-	unsigned int aes_rk[RK_LEN];
+	aes_init();
-	aes_gen_tables();
 	for (unsigned i = 0; i < 16; ++i) {
 		key[i] = rand() & 0xff;
 	}
@@ -96,17 +95,16 @@ int ocl_test() {
 		// rand() & 0xff is about ~60 MB/s @ X230
 		// it's worse than that AES single thread C, so OFB it is
 		// ~240 MB/s, even faster than RDRAND
-		unsigned int aes_rk[RK_LEN];
 		unsigned char key_iv[16 * 2];
 		for (unsigned i = 0; i < 16 * 2; ++i) {
 			key_iv[i] = rand() & 0xff;
 		}
-		aes_set_key_enc_128(aes_rk, key_iv);
+		aes_set_key_enc_128(key_iv);
-		aes_encrypt_128(aes_rk, key_iv + 16, buf_in);
+		aes_encrypt_128(key_iv + 16, buf_in);
 		unsigned char *p_in = buf_in, *p_out = buf_in + 16,
 			*p_end = buf_in + BUF_SIZE;
 		while (p_out < p_end) {
-			aes_encrypt_128(aes_rk, p_in, p_out);
+			aes_encrypt_128(p_in, p_out);
 			p_in = p_out;
 			p_out += 16;
 		}
@@ -174,8 +172,8 @@ int ocl_test() {
 	for (unsigned offset = 0; offset < BUF_SIZE; offset += BLOCK_SIZE) {
 		// setting the same key over and over is stupid
 		// we do this to make the results comparable
-		aes_set_key_enc_128(aes_rk, key);
+		aes_set_key_enc_128(key);
-		aes_encrypt_128(aes_rk, buf_in + offset, buf_verify + offset);
+		aes_encrypt_128(buf_in + offset, buf_verify + offset);
 	}
 	get_hp_time(&t1); td = hp_time_diff(&t0, &t1);
 	printf("%.3f seconds for reference C(single thread), %.2f MB/s\n", td / 1000000.0, BUF_SIZE * 1.0f / td);

sha1_16.c

@@ -11,7 +11,6 @@
  * https://github.com/Jimmy-Z/TWLbf/blob/master/sha1_16.c
  */
 
-#ifndef GET_UINT32_BE
 #define GET_UINT32_BE(n,b,i)                            \
 {                                                       \
     (n) = ( (uint32_t) (b)[(i)    ] << 24 )             \
@@ -19,9 +18,7 @@
         | ( (uint32_t) (b)[(i) + 2] <<  8 )             \
         | ( (uint32_t) (b)[(i) + 3]       );            \
 }
-#endif
 
-#ifndef PUT_UINT32_BE
 #define PUT_UINT32_BE(n,b,i)                            \
 {                                                       \
     (b)[(i)    ] = (unsigned char) ( (n) >> 24 );       \
@@ -29,7 +26,6 @@
     (b)[(i) + 2] = (unsigned char) ( (n) >>  8 );       \
     (b)[(i) + 3] = (unsigned char) ( (n)       );       \
 }
-#endif
 
 static const uint32_t
 	h0 = 0x67452301,