diff --git a/decoded/bruteforce/results.json b/decoded/bruteforce/results.json index 23f216d..674ba54 100644 --- a/decoded/bruteforce/results.json +++ b/decoded/bruteforce/results.json @@ -1,5 +1,5 @@ { - "elapsed_sec": 261.712, + "elapsed_sec": 58.691, "pad_matches": [], "confirmed": [] } diff --git a/ios/DESCracker.xcodeproj/project.pbxproj b/ios/DESCracker.xcodeproj/project.pbxproj index d22ee71..6cffc2e 100644 --- a/ios/DESCracker.xcodeproj/project.pbxproj +++ b/ios/DESCracker.xcodeproj/project.pbxproj @@ -34,7 +34,7 @@ A20000000000000000000008 /* CpuEngine.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = CpuEngine.swift; sourceTree = ""; }; A20000000000000000000009 /* des.c */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.c; name = des.c; path = ../native/common/des.c; sourceTree = SOURCE_ROOT; }; A20000000000000000000010 /* des.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; name = des.h; path = ../native/common/des.h; sourceTree = SOURCE_ROOT; }; - A20000000000000000000011 /* des_bruteforce.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; name = des_bruteforce.metal; path = ../native/metal/des_bruteforce.metal; sourceTree = SOURCE_ROOT; }; + A20000000000000000000011 /* des_bruteforce.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = des_bruteforce.metal; sourceTree = ""; }; A20000000000000000000012 /* Info.plist */ = {isa = PBXFileReference; lastKnownFileType = text.plist.xml; path = Info.plist; sourceTree = ""; }; A20000000000000000000013 /* Assets.xcassets */ = {isa = PBXFileReference; lastKnownFileType = folder.assetcatalog; path = Assets.xcassets; sourceTree = ""; }; A20000000000000000000014 /* Metal.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = Metal.framework; path = System/Library/Frameworks/Metal.framework; sourceTree = SDKROOT; }; @@ -103,6 +103,7 @@ A20000000000000000000006 /* BruteEngine.swift */, A20000000000000000000007 /* MetalEngine.swift */, A20000000000000000000008 /* CpuEngine.swift */, + A20000000000000000000011 /* des_bruteforce.metal */, ); path = Compute; sourceTree = ""; @@ -112,7 +113,6 @@ children = ( A20000000000000000000009 /* des.c */, A20000000000000000000010 /* des.h */, - A20000000000000000000011 /* des_bruteforce.metal */, ); name = Native; sourceTree = ""; @@ -263,7 +263,6 @@ GCC_WARN_UNUSED_FUNCTION = YES; GCC_WARN_UNUSED_VARIABLE = YES; IPHONEOS_DEPLOYMENT_TARGET = 16.0; - MTL_COMPILER_FLAGS = "-fno-unroll-loops"; MTL_ENABLE_DEBUG_INFO = NO; MTL_FAST_MATH = YES; ONLY_ACTIVE_ARCH = YES; @@ -312,7 +311,6 @@ GCC_WARN_UNUSED_FUNCTION = YES; GCC_WARN_UNUSED_VARIABLE = YES; IPHONEOS_DEPLOYMENT_TARGET = 16.0; - MTL_COMPILER_FLAGS = "-fno-unroll-loops"; MTL_ENABLE_DEBUG_INFO = NO; MTL_FAST_MATH = YES; SDKROOT = iphoneos; diff --git a/ios/DESCracker/Compute/des_bruteforce.metal b/ios/DESCracker/Compute/des_bruteforce.metal new file mode 100644 index 0000000..2f5818c --- /dev/null +++ b/ios/DESCracker/Compute/des_bruteforce.metal @@ -0,0 +1,178 @@ +#include +using namespace metal; + +// DES bit numbering: bit 1 is the MSB of a 64-bit word. + +constant uchar IP_TBL[64] = { + 58,50,42,34,26,18,10, 2, 60,52,44,36,28,20,12, 4, + 62,54,46,38,30,22,14, 6, 64,56,48,40,32,24,16, 8, + 57,49,41,33,25,17, 9, 1, 59,51,43,35,27,19,11, 3, + 61,53,45,37,29,21,13, 5, 63,55,47,39,31,23,15, 7 +}; +constant uchar FP_TBL[64] = { + 40, 8,48,16,56,24,64,32, 39, 7,47,15,55,23,63,31, + 38, 6,46,14,54,22,62,30, 37, 5,45,13,53,21,61,29, + 36, 4,44,12,52,20,60,28, 35, 3,43,11,51,19,59,27, + 34, 2,42,10,50,18,58,26, 33, 1,41, 9,49,17,57,25 +}; +constant uchar E_TBL[48] = { + 32, 1, 2, 3, 4, 5, 4, 5, 6, 7, 8, 9, + 8, 9,10,11,12,13, 12,13,14,15,16,17, + 16,17,18,19,20,21, 20,21,22,23,24,25, + 24,25,26,27,28,29, 28,29,30,31,32, 1 +}; +constant uchar P_TBL[32] = { + 16, 7,20,21,29,12,28,17, 1,15,23,26, 5,18,31,10, + 2, 8,24,14,32,27, 3, 9, 19,13,30, 6,22,11, 4,25 +}; +constant uchar PC1_TBL[56] = { + 57,49,41,33,25,17, 9, 1,58,50,42,34,26,18, + 10, 2,59,51,43,35,27, 19,11, 3,60,52,44,36, + 63,55,47,39,31,23,15, 7,62,54,46,38,30,22, + 14, 6,61,53,45,37,29, 21,13, 5,28,20,12, 4 +}; +constant uchar PC2_TBL[48] = { + 14,17,11,24, 1, 5, 3,28,15, 6,21,10, + 23,19,12, 4,26, 8, 16, 7,27,20,13, 2, + 41,52,31,37,47,55, 30,40,51,45,33,48, + 44,49,39,56,34,53, 46,42,50,36,29,32 +}; +constant uchar SHIFTS[16] = {1,1,2,2,2,2,2,2,1,2,2,2,2,2,2,1}; + +constant uchar SBOX[8][64] = { + {14,4,13,1,2,15,11,8,3,10,6,12,5,9,0,7, 0,15,7,4,14,2,13,1,10,6,12,11,9,5,3,8, 4,1,14,8,13,6,2,11,15,12,9,7,3,10,5,0, 15,12,8,2,4,9,1,7,5,11,3,14,10,0,6,13}, + {15,1,8,14,6,11,3,4,9,7,2,13,12,0,5,10, 3,13,4,7,15,2,8,14,12,0,1,10,6,9,11,5, 0,14,7,11,10,4,13,1,5,8,12,6,9,3,2,15, 13,8,10,1,3,15,4,2,11,6,7,12,0,5,14,9}, + {10,0,9,14,6,3,15,5,1,13,12,7,11,4,2,8, 13,7,0,9,3,4,6,10,2,8,5,14,12,11,15,1, 13,6,4,9,8,15,3,0,11,1,2,12,5,10,14,7, 1,10,13,0,6,9,8,7,4,15,14,3,11,5,2,12}, + {7,13,14,3,0,6,9,10,1,2,8,5,11,12,4,15, 13,8,11,5,6,15,0,3,4,7,2,12,1,10,14,9, 10,6,9,0,12,11,7,13,15,1,3,14,5,2,8,4, 3,15,0,6,10,1,13,8,9,4,5,11,12,7,2,14}, + {2,12,4,1,7,10,11,6,8,5,3,15,13,0,14,9, 14,11,2,12,4,7,13,1,5,0,15,10,3,9,8,6, 4,2,1,11,10,13,7,8,15,9,12,5,6,3,0,14, 11,8,12,7,1,14,2,13,6,15,0,9,10,4,5,3}, + {12,1,10,15,9,2,6,8,0,13,3,4,14,7,5,11, 10,15,4,2,7,12,9,5,6,1,13,14,0,11,3,8, 9,14,15,5,2,8,12,3,7,0,4,10,1,13,11,6, 4,3,2,12,9,5,15,10,11,14,1,7,6,0,8,13}, + {4,11,2,14,15,0,8,13,3,12,9,7,5,10,6,1, 13,0,11,7,4,9,1,10,14,3,5,12,2,15,8,6, 1,4,11,13,12,3,7,14,10,15,6,8,0,5,9,2, 6,11,13,8,1,4,10,7,9,5,0,15,14,2,3,12}, + {13,2,8,4,6,15,11,1,10,9,3,14,5,0,12,7, 1,15,13,8,10,3,7,4,12,5,6,11,0,14,9,2, 7,11,4,1,9,12,14,2,0,6,10,13,15,3,5,8, 2,1,14,7,4,10,8,13,15,12,9,0,3,5,6,11} +}; + +inline ulong perm(ulong src, constant uchar *tbl, uint nout, uint srcbits) { + ulong outv = 0; + for (uint i = 0; i < nout; i++) { + ulong bit = (src >> (srcbits - tbl[i])) & 1UL; + outv = (outv << 1) | bit; + } + return outv; +} + +inline uint rotl28(uint v, uint s) { + return ((v << s) | (v >> (28 - s))) & 0x0FFFFFFFu; +} + +inline void des_key_schedule(ulong key, thread ulong sk[16]) { + ulong cd = perm(key, PC1_TBL, 56, 64); + uint c = uint(cd >> 28); + uint d = uint(cd & 0x0FFFFFFFUL); + for (uint r = 0; r < 16; r++) { + c = rotl28(c, SHIFTS[r]); + d = rotl28(d, SHIFTS[r]); + ulong cd2 = (ulong(c) << 28) | ulong(d); + sk[r] = perm(cd2, PC2_TBL, 48, 56); + } +} + +inline uint feistel(uint r, ulong subkey) { + ulong er = perm(ulong(r), E_TBL, 48, 32) ^ subkey; + uint s = 0; + for (uint i = 0; i < 8; i++) { + uint chunk = uint((er >> (42 - 6 * i)) & 0x3F); + uint row = ((chunk & 0x20) >> 4) | (chunk & 1); + uint col = (chunk >> 1) & 0xF; + s = (s << 4) | uint(SBOX[i][row * 16 + col]); + } + return uint(perm(ulong(s), P_TBL, 32, 32)); +} + +inline ulong des_crypt(ulong block, thread ulong sk[16], bool decrypt) { + ulong ip = perm(block, IP_TBL, 64, 64); + uint l = uint(ip >> 32); + uint r = uint(ip & 0xFFFFFFFFUL); + for (uint i = 0; i < 16; i++) { + uint rnd = decrypt ? (15 - i) : i; + uint n = l ^ feistel(r, sk[rnd]); + l = r; + r = n; + } + ulong pre = (ulong(r) << 32) | ulong(l); + return perm(pre, FP_TBL, 64, 64); +} + +struct Params { + ulong start; + uint key_len; + uint charset_len; + uint pad_byte; + uint fill_count; + uint batch_count; + uint _pad; + ulong target; + ulong fills[8]; +}; + +struct Hit { + ulong index; + ulong key; + ulong plain; +}; + +inline ulong make_key(ulong index, constant Params &p, const device uchar *charset) { + uchar bytes[8]; + for (uint i = 0; i < 8; i++) { + bytes[i] = uchar(p.pad_byte); + } + ulong n = index; + uint clen = p.charset_len; + for (int pos = int(p.key_len) - 1; pos >= 0; pos--) { + bytes[pos] = charset[n % clen]; + n /= clen; + } + ulong key = 0; + for (uint i = 0; i < 8; i++) { + key = (key << 8) | ulong(bytes[i]); + } + return key; +} + +inline bool is_fill(ulong pt, constant Params &p) { + for (uint i = 0; i < p.fill_count; i++) { + if (pt == p.fills[i]) { + return true; + } + } + return false; +} + +kernel void des_known_answer(device ulong *out [[buffer(0)]], + uint gid [[thread_position_in_grid]]) { + if (gid != 0) return; + ulong key = 0x133457799BBCDFF1UL; + ulong pt = 0x0123456789ABCDEFUL; + ulong sk[16]; + des_key_schedule(key, sk); + out[0] = des_crypt(pt, sk, false); + out[1] = des_crypt(out[0], sk, true); +} + +kernel void des_brute(constant Params ¶ms [[buffer(0)]], + const device uchar *charset [[buffer(1)]], + device Hit *hits [[buffer(2)]], + device atomic_uint *hit_count [[buffer(3)]], + uint gid [[thread_position_in_grid]]) { + if (gid >= params.batch_count) return; + ulong index = params.start + ulong(gid); + ulong key = make_key(index, params, charset); + ulong sk[16]; + des_key_schedule(key, sk); + ulong pt = des_crypt(params.target, sk, true); + if (!is_fill(pt, params)) return; + uint slot = atomic_fetch_add_explicit(hit_count, 1u, memory_order_relaxed); + if (slot < 256) { + hits[slot].index = index; + hits[slot].key = key; + hits[slot].plain = pt; + } +} diff --git a/native/metal/des_bruteforce.metal b/native/metal/des_bruteforce.metal index 520de91..2f5818c 100644 --- a/native/metal/des_bruteforce.metal +++ b/native/metal/des_bruteforce.metal @@ -39,22 +39,20 @@ constant uchar PC2_TBL[48] = { }; constant uchar SHIFTS[16] = {1,1,2,2,2,2,2,2,1,2,2,2,2,2,2,1}; -constant uchar SBOX[512] = { - 14,4,13,1,2,15,11,8,3,10,6,12,5,9,0,7, 0,15,7,4,14,2,13,1,10,6,12,11,9,5,3,8, 4,1,14,8,13,6,2,11,15,12,9,7,3,10,5,0, 15,12,8,2,4,9,1,7,5,11,3,14,10,0,6,13, - 15,1,8,14,6,11,3,4,9,7,2,13,12,0,5,10, 3,13,4,7,15,2,8,14,12,0,1,10,6,9,11,5, 0,14,7,11,10,4,13,1,5,8,12,6,9,3,2,15, 13,8,10,1,3,15,4,2,11,6,7,12,0,5,14,9, - 10,0,9,14,6,3,15,5,1,13,12,7,11,4,2,8, 13,7,0,9,3,4,6,10,2,8,5,14,12,11,15,1, 13,6,4,9,8,15,3,0,11,1,2,12,5,10,14,7, 1,10,13,0,6,9,8,7,4,15,14,3,11,5,2,12, - 7,13,14,3,0,6,9,10,1,2,8,5,11,12,4,15, 13,8,11,5,6,15,0,3,4,7,2,12,1,10,14,9, 10,6,9,0,12,11,7,13,15,1,3,14,5,2,8,4, 3,15,0,6,10,1,13,8,9,4,5,11,12,7,2,14, - 2,12,4,1,7,10,11,6,8,5,3,15,13,0,14,9, 14,11,2,12,4,7,13,1,5,0,15,10,3,9,8,6, 4,2,1,11,10,13,7,8,15,9,12,5,6,3,0,14, 11,8,12,7,1,14,2,13,6,15,0,9,10,4,5,3, - 12,1,10,15,9,2,6,8,0,13,3,4,14,7,5,11, 10,15,4,2,7,12,9,5,6,1,13,14,0,11,3,8, 9,14,15,5,2,8,12,3,7,0,4,10,1,13,11,6, 4,3,2,12,9,5,15,10,11,14,1,7,6,0,8,13, - 4,11,2,14,15,0,8,13,3,12,9,7,5,10,6,1, 13,0,11,7,4,9,1,10,14,3,5,12,2,15,8,6, 1,4,11,13,12,3,7,14,10,15,6,8,0,5,9,2, 6,11,13,8,1,4,10,7,9,5,0,15,14,2,3,12, - 13,2,8,4,6,15,11,1,10,9,3,14,5,0,12,7, 1,15,13,8,10,3,7,4,12,5,6,11,0,14,9,2, 7,11,4,1,9,12,14,2,0,6,10,13,15,3,5,8, 2,1,14,7,4,10,8,13,15,12,9,0,3,5,6,11 +constant uchar SBOX[8][64] = { + {14,4,13,1,2,15,11,8,3,10,6,12,5,9,0,7, 0,15,7,4,14,2,13,1,10,6,12,11,9,5,3,8, 4,1,14,8,13,6,2,11,15,12,9,7,3,10,5,0, 15,12,8,2,4,9,1,7,5,11,3,14,10,0,6,13}, + {15,1,8,14,6,11,3,4,9,7,2,13,12,0,5,10, 3,13,4,7,15,2,8,14,12,0,1,10,6,9,11,5, 0,14,7,11,10,4,13,1,5,8,12,6,9,3,2,15, 13,8,10,1,3,15,4,2,11,6,7,12,0,5,14,9}, + {10,0,9,14,6,3,15,5,1,13,12,7,11,4,2,8, 13,7,0,9,3,4,6,10,2,8,5,14,12,11,15,1, 13,6,4,9,8,15,3,0,11,1,2,12,5,10,14,7, 1,10,13,0,6,9,8,7,4,15,14,3,11,5,2,12}, + {7,13,14,3,0,6,9,10,1,2,8,5,11,12,4,15, 13,8,11,5,6,15,0,3,4,7,2,12,1,10,14,9, 10,6,9,0,12,11,7,13,15,1,3,14,5,2,8,4, 3,15,0,6,10,1,13,8,9,4,5,11,12,7,2,14}, + {2,12,4,1,7,10,11,6,8,5,3,15,13,0,14,9, 14,11,2,12,4,7,13,1,5,0,15,10,3,9,8,6, 4,2,1,11,10,13,7,8,15,9,12,5,6,3,0,14, 11,8,12,7,1,14,2,13,6,15,0,9,10,4,5,3}, + {12,1,10,15,9,2,6,8,0,13,3,4,14,7,5,11, 10,15,4,2,7,12,9,5,6,1,13,14,0,11,3,8, 9,14,15,5,2,8,12,3,7,0,4,10,1,13,11,6, 4,3,2,12,9,5,15,10,11,14,1,7,6,0,8,13}, + {4,11,2,14,15,0,8,13,3,12,9,7,5,10,6,1, 13,0,11,7,4,9,1,10,14,3,5,12,2,15,8,6, 1,4,11,13,12,3,7,14,10,15,6,8,0,5,9,2, 6,11,13,8,1,4,10,7,9,5,0,15,14,2,3,12}, + {13,2,8,4,6,15,11,1,10,9,3,14,5,0,12,7, 1,15,13,8,10,3,7,4,12,5,6,11,0,14,9,2, 7,11,4,1,9,12,14,2,0,6,10,13,15,3,5,8, 2,1,14,7,4,10,8,13,15,12,9,0,3,5,6,11} }; -template -inline ulong perm_n(ulong src, constant uchar *tbl, uint srcbits) { +inline ulong perm(ulong src, constant uchar *tbl, uint nout, uint srcbits) { ulong outv = 0; -#pragma clang loop unroll(disable) - for (uint i = 0; i < Nout; i++) { + for (uint i = 0; i < nout; i++) { ulong bit = (src >> (srcbits - tbl[i])) & 1UL; outv = (outv << 1) | bit; } @@ -66,36 +64,33 @@ inline uint rotl28(uint v, uint s) { } inline void des_key_schedule(ulong key, thread ulong sk[16]) { - ulong cd = perm_n<56>(key, PC1_TBL, 64); + ulong cd = perm(key, PC1_TBL, 56, 64); uint c = uint(cd >> 28); uint d = uint(cd & 0x0FFFFFFFUL); -#pragma clang loop unroll(disable) for (uint r = 0; r < 16; r++) { c = rotl28(c, SHIFTS[r]); d = rotl28(d, SHIFTS[r]); ulong cd2 = (ulong(c) << 28) | ulong(d); - sk[r] = perm_n<48>(cd2, PC2_TBL, 56); + sk[r] = perm(cd2, PC2_TBL, 48, 56); } } inline uint feistel(uint r, ulong subkey) { - ulong er = perm_n<48>(ulong(r), E_TBL, 32) ^ subkey; + ulong er = perm(ulong(r), E_TBL, 48, 32) ^ subkey; uint s = 0; -#pragma clang loop unroll(disable) for (uint i = 0; i < 8; i++) { uint chunk = uint((er >> (42 - 6 * i)) & 0x3F); uint row = ((chunk & 0x20) >> 4) | (chunk & 1); uint col = (chunk >> 1) & 0xF; - s = (s << 4) | uint(SBOX[i * 64 + row * 16 + col]); + s = (s << 4) | uint(SBOX[i][row * 16 + col]); } - return uint(perm_n<32>(ulong(s), P_TBL, 32)); + return uint(perm(ulong(s), P_TBL, 32, 32)); } inline ulong des_crypt(ulong block, thread ulong sk[16], bool decrypt) { - ulong ip = perm_n<64>(block, IP_TBL, 64); + ulong ip = perm(block, IP_TBL, 64, 64); uint l = uint(ip >> 32); uint r = uint(ip & 0xFFFFFFFFUL); -#pragma clang loop unroll(disable) for (uint i = 0; i < 16; i++) { uint rnd = decrypt ? (15 - i) : i; uint n = l ^ feistel(r, sk[rnd]); @@ -103,7 +98,7 @@ inline ulong des_crypt(ulong block, thread ulong sk[16], bool decrypt) { r = n; } ulong pre = (ulong(r) << 32) | ulong(l); - return perm_n<64>(pre, FP_TBL, 64); + return perm(pre, FP_TBL, 64, 64); } struct Params { @@ -126,19 +121,16 @@ struct Hit { inline ulong make_key(ulong index, constant Params &p, const device uchar *charset) { uchar bytes[8]; -#pragma clang loop unroll(disable) for (uint i = 0; i < 8; i++) { bytes[i] = uchar(p.pad_byte); } ulong n = index; uint clen = p.charset_len; -#pragma clang loop unroll(disable) for (int pos = int(p.key_len) - 1; pos >= 0; pos--) { bytes[pos] = charset[n % clen]; n /= clen; } ulong key = 0; -#pragma clang loop unroll(disable) for (uint i = 0; i < 8; i++) { key = (key << 8) | ulong(bytes[i]); } @@ -146,7 +138,6 @@ inline ulong make_key(ulong index, constant Params &p, const device uchar *chars } inline bool is_fill(ulong pt, constant Params &p) { -#pragma clang loop unroll(disable) for (uint i = 0; i < p.fill_count; i++) { if (pt == p.fills[i]) { return true;