3DS LFCS brute force

bfcl.c

@@ -69,16 +69,18 @@ int main(int argc, const char *argv[]) {
 			puts(invalid_parameters);
 			ret = -1;
 		}
-	} else if(argc == 4){
+	} else if(argc == 4 && !strcmp(argv[1], "msky")){
 		uint32_t msky[4], ver[4];
 		hex2bytes((unsigned char*)msky, 16, argv[2], 1);
 		hex2bytes((unsigned char*)ver, 16, argv[3], 1);
-		if (!strcmp(argv[1], "msky")) {
+		ret = ocl_brute_msky(msky, ver);
-			ret = ocl_brute_msky(msky, ver);
+	} else if(argc == 5 && !strcmp(argv[1], "lfcs")){
-		} else {
+		uint32_t lfcs, ver[2];
-			puts(invalid_parameters);
+		uint16_t newflag;
-			ret = -1;
+		hex2bytes((unsigned char*)&lfcs, 4, argv[2], 1);
-		}
+		hex2bytes((unsigned char*)&newflag, 2, argv[3], 1);
+		hex2bytes((unsigned char*)ver, 8, argv[4], 1);
+		ret = ocl_brute_lfcs(lfcs, newflag, ver);
 	} else {
 		printf(invalid_parameters);
 		ret = -1;

cl/common.h

@@ -2,6 +2,7 @@
 typedef unsigned int uint32_t;
 
 typedef unsigned char u8;
+typedef unsigned short u16;
 typedef unsigned int u32;
 typedef unsigned long u64;
 

cl/kernel_lfcs.cl

@@ -0,0 +1,26 @@
+
+__kernel void test_lfcs(
+	u32 lfcs, u16 newflag,
+	u32 v0, u32 v1,
+	__global u32 *out)
+{
+	if (*out) {
+		return;
+	}
+
+	u32 gid = (u32)get_global_id(0);
+	// u32 gid = 0x5128;
+
+	u32 io[3];
+	io[0] = lfcs + (gid >> 16); // lfcs += gid_h
+	io[1] = ((u32)newflag) | (gid << 16); // rand = gid_l
+	io[2] = 0;
+
+	sha256_12((u8*)io);
+
+	// *out = io[1] - v1 + 1;
+	if (io[0] == v0 && io[1] == v1){
+		*out = gid;
+	}
+}
+

cl/sha256_16.cl

@@ -1,5 +1,5 @@
-/* sha256_16
+/* sha256_16/12
- * again specialized to only take 16 bytes input and spit out the first 16 bytes
+ * again specialized to only take 16/12 bytes input and spit out the first 16/last 8 bytes
  * again code dug out from mbed TLS
  * https://github.com/ARMmbed/mbedtls/blob/development/library/sha256.c
  */
@@ -49,7 +49,11 @@ __constant const uint32_t K[] =
     d += temp1; h = temp1 + temp2;              \
 }
 
+#ifdef SHA256_16
 void sha256_16(unsigned char *io)
+#elif defined SHA256_12
+void sha256_12(unsigned char *io)
+#endif
 {
 	uint32_t temp1, temp2, W[64];
 	uint32_t A[8] = {
@@ -64,14 +68,25 @@ void sha256_16(unsigned char *io)
 	};
 	unsigned int i;
 
-	// padding and msglen identical to sha1_16
+	// padding and msglen identical/similar to sha1_16
 	GET_UINT32_BE(W[0], io, 0);
 	GET_UINT32_BE(W[1], io, 4);
 	GET_UINT32_BE(W[2], io, 8);
+#ifdef SHA256_16
 	GET_UINT32_BE(W[3], io, 12);
-	W[4] = 0x80000000u; W[5] = 0; W[6] = 0; W[7] = 0;
+	W[4] = 0x80000000u;
+#elif defined SHA256_12
+	W[3] = 0x80000000u;
+	W[4] = 0;
+#endif
+	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;
+	W[12] = 0; W[13] = 0; W[14] = 0;
+#ifdef SHA256_16
+	W[15] = 0x80u;
+#elif defined SHA256_12
+	W[15] = 0x60u;
+#endif
 
 	for (i = 0; i < 16; i += 8)
 	{
@@ -97,6 +112,7 @@ void sha256_16(unsigned char *io)
 		P(A[1], A[2], A[3], A[4], A[5], A[6], A[7], A[0], R(i + 7), K[i + 7]);
 	}
 
+#ifdef SHA256_16
 	A[0] += 0x6A09E667;
 	A[1] += 0xBB67AE85;
 	A[2] += 0x3C6EF372;
@@ -106,4 +122,11 @@ void sha256_16(unsigned char *io)
 	PUT_UINT32_BE(A[1], io, 4);
 	PUT_UINT32_BE(A[2], io, 8);
 	PUT_UINT32_BE(A[3], io, 12);
+#elif defined SHA256_12
+	A[6] += 0x1F83D9AB;
+	A[7] += 0x5BE0CD19;
+
+	PUT_UINT32_BE(A[6], io, 0);
+	PUT_UINT32_BE(A[7], io, 4);
+#endif
 }

ocl_brute.c

@@ -333,7 +333,7 @@ int ocl_brute_msky(const cl_uint *msky, const cl_uint *ver)
 		"cl/sha256_16.cl",
 		"cl/kernel_msky.cl" };
 	cl_program program = ocl_build_from_sources(sizeof(source_names) / sizeof(char *),
-		source_names, context, device_id, "-w -Werror");
+		source_names, context, device_id, "-w -Werror -DSHA256_16");
 
 	cl_kernel kernel = OCL_ASSERT2(clCreateKernel(program, "test_msky", &err));
 
@@ -405,3 +405,92 @@ int ocl_brute_msky(const cl_uint *msky, const cl_uint *ver)
 	return !out;
 }
 
+// LFCS brute force, https://gist.github.com/zoogie/4046726878dba89eddfa1fc07c8a27da
+int ocl_brute_lfcs(cl_uint lfcs_template, cl_ushort newflag, const cl_uint *ver)
+{
+	TimeHP t0, t1; long long td = 0;
+
+	cl_int err;
+	cl_platform_id platform_id;
+	cl_device_id device_id;
+	ocl_get_device(&platform_id, &device_id);
+	if (platform_id == NULL || device_id == NULL) {
+		return -1;
+	}
+
+	cl_context context = OCL_ASSERT2(clCreateContext(0, 1, &device_id, NULL, NULL, &err));
+	cl_command_queue command_queue = OCL_ASSERT2(clCreateCommandQueue(context, device_id, 0, &err));
+
+	const char *source_names[] = {
+		"cl/common.h",
+		"cl/sha256_16.cl",
+		"cl/kernel_lfcs.cl" };
+	cl_program program = ocl_build_from_sources(sizeof(source_names) / sizeof(char *),
+		source_names, context, device_id, "-w -Werror -DSHA256_12");
+
+	cl_kernel kernel = OCL_ASSERT2(clCreateKernel(program, "test_lfcs", &err));
+
+	size_t local;
+	OCL_ASSERT(clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL));
+	printf("local work size: %u\n", (unsigned)local);
+
+	// there's no option to create it zero initialized
+	cl_uint out = 0;
+	cl_mem mem_out = OCL_ASSERT2(clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_uint), NULL, &err));
+	OCL_ASSERT(clEnqueueWriteBuffer(command_queue, mem_out, CL_TRUE, 0, sizeof(cl_uint), &out, 0, NULL, NULL));
+
+	unsigned brute_bits = 32;
+	unsigned group_bits = 28;
+	unsigned loop_bits = brute_bits - group_bits;
+	unsigned loops = 1ull << loop_bits;
+	size_t num_items = 1ull << group_bits;
+	// it needs to be aligned, a little overhead hurts nobody
+	if (num_items % local) {
+		num_items = (num_items / local + 1) * local;
+	}
+	OCL_ASSERT(clSetKernelArg(kernel, 1, sizeof(cl_ushort), &newflag));
+	OCL_ASSERT(clSetKernelArg(kernel, 2, sizeof(cl_uint), &ver[0]));
+	OCL_ASSERT(clSetKernelArg(kernel, 3, sizeof(cl_uint), &ver[1]));
+	OCL_ASSERT(clSetKernelArg(kernel, 4, sizeof(cl_mem), &mem_out));
+	get_hp_time(&t0);
+	int fan_range = 0x10000; // "fan out" full 16 bits
+	unsigned i, j;
+	for (j = 0; j < fan_range; ++j) {
+		int fan = (j & 1 ? 1 : -1) * ((j + 1) >> 1);
+		printf("%d\r", fan);
+		for (i = 0; i < loops; ++i) {
+			cl_uint lfcs = lfcs_template + fan * 0x10000 + (i << (group_bits - 16));
+			OCL_ASSERT(clSetKernelArg(kernel, 0, sizeof(cl_uint), &lfcs));
+
+			OCL_ASSERT(clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &num_items, &local, 0, NULL, NULL));
+			clFinish(command_queue);
+
+			OCL_ASSERT(clEnqueueReadBuffer(command_queue, mem_out, CL_TRUE, 0, sizeof(cl_uint), &out, 0, NULL, NULL));
+			if (out) {
+				get_hp_time(&t1); td = hp_time_diff(&t0, &t1);
+				lfcs += out >> 16;
+				printf("got a hit: %s\n", hexdump(&lfcs, 4, 0));
+				break;
+			}
+		}
+		if (out) {
+			break;
+		}
+	}
+	u64 tested = 0;
+	if (!out) {
+		tested = (1ull << brute_bits) * fan_range;
+		get_hp_time(&t1); td = hp_time_diff(&t0, &t1);
+	} else {
+		tested = out + (1ull << brute_bits) * j;
+	}
+	printf("%.2f seconds, %.2f M/s\n", td / 1000000.0, tested * 1.0 / td);
+
+	clReleaseKernel(kernel);
+	clReleaseMemObject(mem_out);
+	clReleaseProgram(program);
+	clReleaseCommandQueue(command_queue);
+	clReleaseContext(context);
+	return !out;
+}
+

ocl_brute.h

@@ -15,3 +15,6 @@ 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);
 
 int ocl_brute_msky(const cl_uint *msky, const cl_uint *ver);
+
+int ocl_brute_lfcs(cl_uint lfcs_template, cl_ushort newflag, const cl_uint *ver);
+