Skip to content

Commit 175d236

Browse files
Jimmy-ZJimmy-Z
committed
better EMMC CID brute suggested by dark_samus3, done
1 parent 7eb91ba commit 175d236

11 files changed

Lines changed: 189 additions & 93 deletions

File tree

‎Makefile‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@ LDFLAGS += -L$(INTELOCLSDKROOT)/lib/x64
77
all : $(PNAME)
88

99
$(PNAME) : $(OBJS)
10-
$(CC) $(LDFLAGS) -o $@ $^ -lOpenCL
10+
$(CC) $(LDFLAGS) -o $@ $^ -lOpenCL -static -lmbedcrypto
1111

1212
clean :
1313
rm $(PNAME) *.o

‎aes_128.c‎

Lines changed: 87 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,91 @@
11

2+
#include <stdio.h>
23
#include <mbedtls/config.h>
4+
#include <mbedtls/version.h>
35
#include <mbedtls/aes.h>
46
#include <mbedtls/aesni.h>
5-
#include <mbedtls/version.h>
7+
#include "crypto.h"
8+
9+
static mbedtls_aes_context ctx;
10+
11+
static int (*p_aes_crypt_ecb)(mbedtls_aes_context*, int, const unsigned char *, unsigned char *) = NULL;
12+
13+
static void (*p_aes_set_key_enc_128)(const unsigned char *key) = NULL;
14+
15+
static void (*p_aes_set_key_dec_128)(const unsigned char *key) = NULL;
16+
17+
#define AES_KEY_LEN 128
18+
#define NR 10
19+
20+
// I hope eliminating the AESNI check can make it a bit faster
21+
static void aes_set_key_enc_128_aesni(const unsigned char *key){
22+
// mbedtls_aes_setkey_enc(&ctx, key, 128);
23+
ctx.nr = NR;
24+
ctx.rk = ctx.buf;
25+
mbedtls_aesni_setkey_enc((unsigned char *)ctx.rk, key, AES_KEY_LEN);
26+
}
27+
28+
static void aes_set_key_dec_128_aesni(const unsigned char *key) {
29+
mbedtls_aes_context cty;
30+
cty.nr = NR;
31+
cty.rk = cty.buf;
32+
mbedtls_aesni_setkey_enc((unsigned char *)cty.rk, key, AES_KEY_LEN);
33+
ctx.nr = cty.nr;
34+
ctx.rk = ctx.buf;
35+
mbedtls_aesni_inverse_key((unsigned char *)ctx.rk, (const unsigned char *)cty.rk, ctx.nr);
36+
}
37+
38+
static void aes_set_key_enc_128_c(const unsigned char *key) {
39+
mbedtls_aes_setkey_enc(&ctx, key, AES_KEY_LEN);
40+
}
41+
42+
static void aes_set_key_dec_128_c(const unsigned char *key) {
43+
mbedtls_aes_setkey_dec(&ctx, key, AES_KEY_LEN);
44+
}
45+
46+
void aes_init(){
47+
fputs(MBEDTLS_VERSION_STRING_FULL, stdout);
48+
mbedtls_aes_init(&ctx);
49+
// prevent runtime checks
50+
if(mbedtls_aesni_has_support(MBEDTLS_AESNI_AES)){
51+
puts(", AES-NI supported");
52+
p_aes_crypt_ecb = mbedtls_aesni_crypt_ecb;
53+
p_aes_set_key_enc_128 = aes_set_key_enc_128_aesni;
54+
p_aes_set_key_dec_128 = aes_set_key_dec_128_aesni;
55+
}else {
56+
puts(", AES-NI not supported");
57+
p_aes_crypt_ecb = mbedtls_aes_crypt_ecb;
58+
p_aes_set_key_enc_128 = aes_set_key_enc_128_c;
59+
p_aes_set_key_dec_128 = aes_set_key_dec_128_c;
60+
}
61+
#ifndef MBEDTLS_AES_ROM_TABLES
62+
// it will error out but also get aes_gen_tables done
63+
mbedtls_aes_setkey_enc(&ctx, NULL, 0);
64+
#endif
65+
}
66+
67+
void aes_set_key_enc_128(const unsigned char *key) {
68+
p_aes_set_key_enc_128(key);
69+
}
70+
71+
void aes_set_key_dec_128(const unsigned char *key) {
72+
p_aes_set_key_dec_128(key);
73+
}
74+
75+
void aes_encrypt_128(const unsigned char *in, unsigned char *out){
76+
p_aes_crypt_ecb(&ctx, MBEDTLS_AES_ENCRYPT, in, out);
77+
}
78+
79+
void aes_decrypt_128(const unsigned char *in, unsigned char *out){
80+
p_aes_crypt_ecb(&ctx, MBEDTLS_AES_DECRYPT, in, out);
81+
}
82+
83+
void aes_encrypt_128_bulk(const unsigned char *in, unsigned char *out, unsigned len){
84+
len >>= 4;
85+
for(unsigned i = 0; i < len; ++i){
86+
p_aes_crypt_ecb(&ctx, MBEDTLS_AES_ENCRYPT, in, out);
87+
in += AES_BLOCK_SIZE;
88+
out += AES_BLOCK_SIZE;
89+
}
90+
}
91+

‎bfcl.c‎

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -14,6 +14,8 @@ static inline cl_ushort u16be(const unsigned char *in){
1414
return out;
1515
}
1616

17+
const char invalid_parameters[] = "invalid parameters\n";
18+
1719
int main(int argc, const char *argv[]) {
1820
int ret = 0;
1921
if (argc == 1) {
@@ -43,7 +45,7 @@ int main(int argc, const char *argv[]) {
4345
ret = ocl_brute_console_id(console_id, 0,
4446
u16be(offset0), src0, ver0, u16be(offset1), src1, ver1, CTR);
4547
} else {
46-
puts("invalid parameters\n");
48+
puts(invalid_parameters);
4749
ret = -1;
4850
}
4951
} else if (argc == 7) {
@@ -63,11 +65,11 @@ int main(int argc, const char *argv[]) {
6365
} else if (!strcmp(argv[1], "emmc_cid")) {
6466
ret = ocl_brute_emmc_cid(console_id, emmc_cid, u16be(offset), src, ver);
6567
} else {
66-
puts("invalid parameters\n");
68+
puts(invalid_parameters);
6769
ret = -1;
6870
}
6971
} else {
70-
printf("invalid parameters\n");
72+
printf(invalid_parameters);
7173
ret = -1;
7274
}
7375
#ifdef _WIN32

‎cl/common.h‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,7 @@
1+
2+
typedef unsigned int uint32_t;
3+
4+
typedef unsigned char u8;
5+
typedef unsigned int u32;
6+
typedef unsigned long u64;
7+

‎cl/kernel_emmc_cid.cl‎

Lines changed: 4 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -7,14 +7,12 @@ __kernel void test_emmc_cid(
77
if (*out) {
88
return;
99
}
10-
u8 emmc_cid[16];
11-
*(u64*)emmc_cid = emmc_cid_l;
12-
*(u64*)(emmc_cid + 8) = emmc_cid_h;
13-
*(u32*)(emmc_cid + 1) |= get_global_id(0);
10+
u64 emmc_cid[2] = { emmc_cid_l, emmc_cid_h };
11+
*(u32*)(((u8*)emmc_cid) + 1) |= get_global_id(0);
1412

15-
sha1_16((u32*)emmc_cid);
13+
sha1_16((u8*)emmc_cid);
1614

17-
if (sha1_16_l == *(u64*)emmc_cid && sha1_16_h == *(u64*)(emmc_cid + 8)) {
15+
if (sha1_16_l == emmc_cid[0] && sha1_16_h == emmc_cid[1]) {
1816
*out = get_global_id(0);
1917
}
2018
}

‎cl/sha1_16.cl‎

Lines changed: 25 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -8,12 +8,30 @@ __constant const uint32_t
88
h3 = 0x10325476,
99
h4 = 0xC3D2E1F0;
1010

11-
void sha1_16(uint32_t *io)
11+
#define GET_UINT32_BE(n,b,i) \
12+
{ \
13+
(n) = ( (uint32_t) (b)[(i) ] << 24 ) \
14+
| ( (uint32_t) (b)[(i) + 1] << 16 ) \
15+
| ( (uint32_t) (b)[(i) + 2] << 8 ) \
16+
| ( (uint32_t) (b)[(i) + 3] ); \
17+
}
18+
19+
#define PUT_UINT32_BE(n,b,i) \
20+
{ \
21+
(b)[(i) ] = (unsigned char) ( (n) >> 24 ); \
22+
(b)[(i) + 1] = (unsigned char) ( (n) >> 16 ); \
23+
(b)[(i) + 2] = (unsigned char) ( (n) >> 8 ); \
24+
(b)[(i) + 3] = (unsigned char) ( (n) ); \
25+
}
26+
27+
void sha1_16(unsigned char *io)
1228
{
1329
uint32_t temp, W[16],
1430
A = h0, B = h1, C = h2, D = h3, E = h4;
15-
16-
W[0] = io[0]; W[1] = io[1]; W[2] = io[2]; W[3] = io[3];
31+
GET_UINT32_BE(W[0], io, 0);
32+
GET_UINT32_BE(W[1], io, 4);
33+
GET_UINT32_BE(W[2], io, 8);
34+
GET_UINT32_BE(W[3], io, 12);
1735
W[4] = 0x80000000u; W[5] = 0; W[6] = 0; W[7] = 0;
1836
W[8] = 0; W[9] = 0; W[10] = 0; W[11] = 0;
1937
W[12] = 0; W[13] = 0; W[14] = 0; W[15] = 0x80u;
@@ -149,9 +167,9 @@ void sha1_16(uint32_t *io)
149167
C += h2;
150168
D += h3;
151169

152-
io[0] = A;
153-
io[1] = B;
154-
io[2] = C;
155-
io[3] = D;
170+
PUT_UINT32_BE(A, io, 0);
171+
PUT_UINT32_BE(B, io, 4);
172+
PUT_UINT32_BE(C, io, 8);
173+
PUT_UINT32_BE(D, io, 12);
156174
}
157175

‎crypto.h‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -11,10 +11,9 @@ void sha1_16(const unsigned char in[16], unsigned char out[16]);
1111
void aes_init(void);
1212

1313
void aes_set_key_enc_128(const unsigned char *key);
14-
15-
void aes_encrypt_128(const unsigned char input[16], unsigned char output[16]);
16-
1714
void aes_set_key_dec_128(const unsigned char *key);
1815

16+
void aes_encrypt_128(const unsigned char input[16], unsigned char output[16]);
1917
void aes_decrypt_128(const unsigned char input[16], unsigned char output[16]);
2018

19+
void aes_encrypt_128_bulk(const unsigned char input[16], unsigned char output[16], unsigned len);

‎dsi.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -39,7 +39,7 @@ static inline u32 u32be(const u8 *in){
3939
}
4040

4141
// CAUTION this one doesn't work in-place
42-
static inline void byte_reverse_16(u8 *out, const u8 *in){
42+
inline void byte_reverse_16(u8 *out, const u8 *in){
4343
out[0] = in[15];
4444
out[1] = in[14];
4545
out[2] = in[13];

‎ocl_brute.c‎

Lines changed: 51 additions & 59 deletions
Original file line numberDiff line numberDiff line change
@@ -34,6 +34,33 @@ int ocl_brute_console_id(const cl_uchar *console_id, const cl_uchar *emmc_cid,
3434
cl_uint offset1, const cl_uchar *src1, const cl_uchar *ver1,
3535
ocl_brute_mode mode)
3636
{
37+
// preparing args
38+
cl_ulong console_id_template = u64be(console_id);
39+
cl_ulong xor0[2] = { 0 }, xor1[2] = { 0 };
40+
dsi_make_xor((u8*)xor0, src0, ver0);
41+
if (src1 != 0) {
42+
dsi_make_xor((u8*)xor1, src1, ver1);
43+
}
44+
cl_uint ctr[4] = { 0 };
45+
if (emmc_cid != 0) {
46+
dsi_make_ctr((u8*)ctr, emmc_cid, offset0);
47+
}
48+
cl_ulong out = 0;
49+
#if DEBUG
50+
{
51+
printf("XOR : %s\n", hexdump(xor0, 16, 0));
52+
u8 aes_key[16];
53+
dsi_make_key(aes_key, u64be(console_id));
54+
printf("AES KEY : %s\n", hexdump(aes_key, 16, 0));
55+
aes_init();
56+
aes_set_key_enc_128(aes_key);
57+
printf("CTR : %s\n", hexdump(ctr, 16, 0));
58+
aes_encrypt_128((u8*)ctr, (u8*)xor0);
59+
printf("XOR TRY : %s\n", hexdump(xor0, 16, 0));
60+
// exit(1);
61+
}
62+
#endif
63+
3764
TimeHP t0, t1; long long td = 0;
3865

3966
cl_int err;
@@ -73,35 +100,6 @@ int ocl_brute_console_id(const cl_uchar *console_id, const cl_uchar *emmc_cid,
73100
OCL_ASSERT(clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL));
74101
printf("local work size: %u\n", (unsigned)local);
75102

76-
// preparing args
77-
cl_ulong console_id_template = u64be(console_id);
78-
cl_ulong xor0[2] = { 0 }, xor1[2] = { 0 };
79-
dsi_make_xor((u8*)xor0, src0, ver0);
80-
if (src1 != 0) {
81-
dsi_make_xor((u8*)xor1, src1, ver1);
82-
}
83-
cl_uint ctr[4] = { 0 };
84-
if (emmc_cid != 0) {
85-
dsi_make_ctr((u8*)ctr, emmc_cid, offset0);
86-
}
87-
cl_ulong out = 0;
88-
#if DEBUG
89-
{
90-
printf("XOR : %s\n", hexdump(xor, 16, 0));
91-
u8 aes_key[16];
92-
dsi_make_key(aes_key, u64be(console_id));
93-
printf("AES KEY : %s\n", hexdump(aes_key, 16, 0));
94-
cl_uint aes_rk[RK_LEN];
95-
aes_gen_tables();
96-
aes_set_key_enc_128(aes_rk, aes_key);
97-
printf("AES RK : %s\n", hexdump(aes_rk, 48, 0));
98-
printf("CTR : %s\n", hexdump(ctr, 16, 0));
99-
aes_encrypt_128(aes_rk, (u8*)ctr, (u8*)xor);
100-
printf("XOR TRY : %s\n", hexdump(xor, 16, 0));
101-
// exit(1);
102-
}
103-
#endif
104-
105103
// there's no option to create it zero initialized
106104
cl_mem mem_out = OCL_ASSERT2(clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_ulong), NULL, &err));
107105
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,
208206
int ocl_brute_emmc_cid(const cl_uchar *console_id, cl_uchar *emmc_cid,
209207
cl_uint offset, const cl_uchar *src, const cl_uchar *ver)
210208
{
209+
// preparing args
210+
u8 aes_key[16];
211+
dsi_make_key(aes_key, u64be(console_id));
212+
aes_init();
213+
aes_set_key_dec_128(aes_key);
214+
cl_ulong xor[2];
215+
dsi_make_xor((u8*)xor, src, ver);
216+
cl_ulong ctr[2];
217+
aes_decrypt_128((u8*)xor , (u8*)ctr);
218+
cl_ulong emmc_cid_sha1_16[2];
219+
byte_reverse_16((u8*)emmc_cid_sha1_16, (u8*)ctr);
220+
sub_128_64(emmc_cid_sha1_16, offset);
221+
cl_ulong out = 0;
222+
#ifdef DEBUG
223+
{
224+
printf("SHA1 A: %s\n", hexdump(emmc_cid_sha1_16, 16, 0));
225+
u8 sha1_verify[16];
226+
sha1_16(emmc_cid, sha1_verify);
227+
printf("SHA1 B: %s\n", hexdump(sha1_verify, 16, 0));
228+
}
229+
#endif
230+
211231
TimeHP t0, t1; long long td = 0;
212232

213233
cl_int err;
@@ -234,33 +254,6 @@ int ocl_brute_emmc_cid(const cl_uchar *console_id, cl_uchar *emmc_cid,
234254
OCL_ASSERT(clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL));
235255
printf("local work size: %u\n", (unsigned)local);
236256

237-
// preparing args
238-
u8 aes_key[16];
239-
dsi_make_key(aes_key, u64be(console_id));
240-
aes_init();
241-
aes_set_key_dec_128(aes_key);
242-
cl_ulong xor[2];
243-
dsi_make_xor((u8*)xor, src, ver);
244-
cl_ulong ctr[2];
245-
aes_decrypt_128((u8*)xor , (u8*)ctr);
246-
cl_ulong emmc_cid_sha1_16[2];
247-
byte_reverse_16((u8*)emmc_cid_sha1_16, (u8*)ctr);
248-
sub_128_64(emmc_cid_sha1_16, offset);
249-
cl_ulong out = 0;
250-
#ifdef DEBUG
251-
{
252-
printf("XOR : %s\n", hexdump(xor, 16, 0));
253-
printf("AES KEY : %s\n", hexdump(aes_key, 16, 0));
254-
printf("AES RK : %s\n", hexdump(aes_rk, 48, 0));
255-
u8 ctr[16];
256-
dsi_make_ctr(ctr, emmc_cid, u_offset);
257-
printf("CTR : %s\n", hexdump(ctr, 16, 0));
258-
aes_encrypt_128(aes_rk, ctr, (u8*)xor);
259-
printf("XOR TRY : %s\n", hexdump(xor, 16, 0));
260-
// exit(1);
261-
}
262-
#endif
263-
264257
// there's no option to create it zero initialized
265258
cl_mem mem_out = OCL_ASSERT2(clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_ulong), NULL, &err));
266259
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,
280273
puts(hexdump(emmc_cid, 16, 0));
281274
OCL_ASSERT(clSetKernelArg(kernel, 0, sizeof(cl_ulong), emmc_cid));
282275
OCL_ASSERT(clSetKernelArg(kernel, 1, sizeof(cl_ulong), emmc_cid + 8));
283-
OCL_ASSERT(clSetKernelArg(kernel, 2, sizeof(cl_ulong), &emmc_cid_sha1_16[0]));
284-
OCL_ASSERT(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &emmc_cid_sha1_16[1]));
276+
OCL_ASSERT(clSetKernelArg(kernel, 2, sizeof(cl_ulong), emmc_cid_sha1_16));
277+
OCL_ASSERT(clSetKernelArg(kernel, 3, sizeof(cl_ulong), emmc_cid_sha1_16 + 1));
285278
OCL_ASSERT(clSetKernelArg(kernel, 4, sizeof(cl_mem), &mem_out));
286279

287280
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,
308301
printf("%.2f seconds, %.2f M/s\n", td / 1000000.0, tested * 1.0 / td);
309302

310303
clReleaseKernel(kernel);
311-
clReleaseMemObject(mem_rk);
312304
clReleaseMemObject(mem_out);
313305
clReleaseProgram(program);
314306
clReleaseCommandQueue(command_queue);

0 commit comments

Comments
 (0)