diff --git a/CHANGELOG.md b/CHANGELOG.md index 3cd3f2d87..198df360e 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -3,6 +3,7 @@ All notable changes to this project will be documented in this file. This project uses the changelog in accordance with [keepchangelog](http://keepachangelog.com/). Please use this to write notable changes, which is not the same as git commit log... ## [unreleased][unreleased] +- Added CUDA version of tools/mfulc_des_brute (@C2Pwn) - Added `hf mfu desbrute` command: native client support for ULC key recovery (@C2Pwn) - Added `hf mf sen` command: native client support for FM11RF08S SEN recovery (@C2Pwn) - Added Makefile `PLATFORM_FILE` variable to specify another Makefile.platform file (@doegox) diff --git a/tools/mfulc_des_brute_cuda/Makefile b/tools/mfulc_des_brute_cuda/Makefile new file mode 100644 index 000000000..464e4ac40 --- /dev/null +++ b/tools/mfulc_des_brute_cuda/Makefile @@ -0,0 +1,15 @@ +NVCC ?= nvcc +TARGET ?= mfulc_des_brute_cuda +SOURCES := main.cu desbrute.cu + +NVCCFLAGS ?= -O3 -std=c++14 + +.PHONY: all clean + +all: $(TARGET) + +$(TARGET): $(SOURCES) desbrute.h des_host.h + $(NVCC) $(NVCCFLAGS) -o $@ $(SOURCES) + +clean: + rm -f $(TARGET) \ No newline at end of file diff --git a/tools/mfulc_des_brute_cuda/des_host.h b/tools/mfulc_des_brute_cuda/des_host.h new file mode 100644 index 000000000..d13725ef6 --- /dev/null +++ b/tools/mfulc_des_brute_cuda/des_host.h @@ -0,0 +1,222 @@ +/* + * des_host.h — Self-contained CPU DES implementation (no OpenSSL, no CUDA). + * Used for LFSR detection before GPU launch. + * + * All uint64_t values are "big-endian packed": byte 0 of the DES block occupies + * bits 63-56, byte 7 occupies bits 7-0. This matches our GPU DES output format + * exactly, so the same LFSR / rotation checks work on both sides without bswap. + */ +#pragma once +#include +#include + +/* -------------------------------------------------------------------------- */ +/* DES constant tables (1-indexed bit positions, MSB = bit 1 = our bit 0) */ +/* -------------------------------------------------------------------------- */ + +static const uint8_t H_PC1[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 +}; + +static const uint8_t H_PC2[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 +}; + +static const uint8_t H_SHIFTS[16] = { + 1,1,2,2,2,2,2,2,1,2,2,2,2,2,2,1 +}; + +static const uint8_t H_IP[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 +}; + +static const uint8_t H_FP[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 +}; + +static const uint8_t H_E[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 +}; + +static const uint8_t H_P[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 +}; + +static const uint8_t H_SBOX[8][4][16] = { + { {14,4,13,1,2,15,11,8,3,10,6,12,5,9,0,7}, /* S1 */ + {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}, /* S2 */ + {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}, /* S3 */ + {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}, /* S4 */ + {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}, /* S5 */ + {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}, /* S6 */ + {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}, /* S7 */ + {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}, /* S8 */ + {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} } +}; + +/* -------------------------------------------------------------------------- */ +/* Utility */ +/* -------------------------------------------------------------------------- */ + +/* Convert 8-byte array (big-endian) → packed uint64 (byte 0 in bits 63-56) */ +static inline uint64_t bytes_to_u64(const uint8_t *b) +{ + return ((uint64_t)b[0] << 56) | ((uint64_t)b[1] << 48) | + ((uint64_t)b[2] << 40) | ((uint64_t)b[3] << 32) | + ((uint64_t)b[4] << 24) | ((uint64_t)b[5] << 16) | + ((uint64_t)b[6] << 8) | (uint64_t)b[7]; +} + +/* -------------------------------------------------------------------------- */ +/* CPU DES primitives */ +/* -------------------------------------------------------------------------- */ + +static inline uint64_t h_perm(uint64_t src, int src_bits, const uint8_t *tbl, int n) +{ + uint64_t dst = 0; + int i; + for (i = 0; i < n; i++) { + int sb = tbl[i] - 1; + dst |= ((src >> (src_bits - 1 - sb)) & 1ULL) << (n - 1 - i); + } + return dst; +} + +static inline void h_des_keyschedule(uint64_t key64, uint64_t sk[16]) +{ + uint64_t key56 = h_perm(key64, 64, H_PC1, 56); + uint32_t C = (uint32_t)(key56 >> 28) & 0x0FFFFFFFU; + uint32_t D = (uint32_t)(key56) & 0x0FFFFFFFU; + int i; + for (i = 0; i < 16; i++) { + int s = H_SHIFTS[i]; + C = ((C << s) | (C >> (28 - s))) & 0x0FFFFFFFU; + D = ((D << s) | (D >> (28 - s))) & 0x0FFFFFFFU; + sk[i] = h_perm(((uint64_t)C << 28) | D, 56, H_PC2, 48); + } +} + +static inline uint64_t h_des_crypt(uint64_t block, const uint64_t sk[16]) +{ + uint64_t ip = h_perm(block, 64, H_IP, 64); + uint32_t L = (uint32_t)(ip >> 32); + uint32_t R = (uint32_t)(ip); + int i; + for (i = 0; i < 16; i++) { + uint64_t eR = h_perm((uint64_t)R, 32, H_E, 48); + uint64_t x = eR ^ sk[i]; + uint32_t sout = 0; + int s; + for (s = 0; s < 8; s++) { + uint8_t b6 = (uint8_t)((x >> (42 - s * 6)) & 0x3F); + int row = ((b6 >> 5) & 1) * 2 + (b6 & 1); + int col = (b6 >> 1) & 0xF; + sout |= (uint32_t)H_SBOX[s][row][col] << (28 - s * 4); + } + uint32_t f = (uint32_t)h_perm((uint64_t)sout, 32, H_P, 32); + uint32_t nR = L ^ f; + L = R; R = nR; + } + return h_perm(((uint64_t)R << 32) | L, 64, H_FP, 64); +} + +static inline uint64_t h_des_decrypt(uint64_t block, const uint64_t sk[16]) +{ + uint64_t rsk[16]; + int i; + for (i = 0; i < 16; i++) rsk[i] = sk[15 - i]; + return h_des_crypt(block, rsk); +} + +/* -------------------------------------------------------------------------- */ +/* LFSR validators (input: BE uint64, byte 0 in bits 63-56) */ +/* -------------------------------------------------------------------------- */ + +static inline int h_valid_lfsr_ulcg(uint64_t x64) +{ + /* x64 is already B0=MSB; original code bswap'd from OpenSSL LE, we skip that */ + uint16_t x16 = (uint16_t)(x64 >> 48); + x16 = (uint16_t)(x16 << 15 | ((x16 >> 1) ^ ((x16 >> 3 ^ x16 >> 4 ^ x16 >> 6) & 1))); + if (x16 != (uint16_t)((x64 >> 32) & 0xFFFF)) return 0; + x16 = (uint16_t)(x16 << 15 | ((x16 >> 1) ^ ((x16 >> 3 ^ x16 >> 4 ^ x16 >> 6) & 1))); + if (x16 != (uint16_t)((x64 >> 16) & 0xFFFF)) return 0; + x16 = (uint16_t)(x16 << 15 | ((x16 >> 1) ^ ((x16 >> 3 ^ x16 >> 4 ^ x16 >> 6) & 1))); + if (x16 != (uint16_t)(x64 & 0xFFFF)) return 0; + return 1; +} + +static inline int h_valid_lfsr_mfc(uint64_t x64) +{ + uint16_t x16 = (uint16_t)(x64 & 0xFFFF); + int i; + for (i = 0; i < 16; i++) + x16 = (uint16_t)(x16 >> 1 | (x16 ^ x16 >> 2 ^ x16 >> 3 ^ x16 >> 5) << 15); + if (x16 != (uint16_t)((x64 >> 16) & 0xFFFF)) return 0; + for (i = 0; i < 16; i++) + x16 = (uint16_t)(x16 >> 1 | (x16 ^ x16 >> 2 ^ x16 >> 3 ^ x16 >> 5) << 15); + if (x16 != (uint16_t)((x64 >> 32) & 0xFFFF)) return 0; + for (i = 0; i < 16; i++) + x16 = (uint16_t)(x16 >> 1 | (x16 ^ x16 >> 2 ^ x16 >> 3 ^ x16 >> 5) << 15); + if (x16 != (uint16_t)((x64 >> 48) & 0xFFFF)) return 0; + return 1; +} + +/* + * Detect LFSR type from init_ciphertext (8 bytes). + * Decrypts with the all-zero DES key and checks LFSR structure. + * Returns 1=ULCG, 2=MFC, 0=unknown. + */ +static inline int h_detect_lfsr(const uint8_t init_ciphertext[8]) +{ + uint8_t zero_key[8]; + uint64_t sk[16]; + uint64_t ct, out; + memset(zero_key, 0, 8); + h_des_keyschedule(bytes_to_u64(zero_key), sk); + ct = bytes_to_u64(init_ciphertext); + out = h_des_decrypt(ct, sk); + if (h_valid_lfsr_ulcg(out)) return 1; + if (h_valid_lfsr_mfc(out)) return 2; + return 0; +} diff --git a/tools/mfulc_des_brute_cuda/desbrute.cu b/tools/mfulc_des_brute_cuda/desbrute.cu new file mode 100644 index 000000000..745b03aaa --- /dev/null +++ b/tools/mfulc_des_brute_cuda/desbrute.cu @@ -0,0 +1,835 @@ +/* + * desbrute.cu -- CUDA brute-force engine for 2TDEA key recovery. + */ + +#include +#include +#include +#ifdef _WIN32 +#include +#else +#include +#endif +#include +#include "desbrute.h" + +#ifdef _WIN32 +typedef LARGE_INTEGER wall_time_t; + +static void wall_time_now(wall_time_t *t) +{ + QueryPerformanceCounter(t); +} + +static double wall_time_seconds(const wall_time_t *start, const wall_time_t *now) +{ + LARGE_INTEGER qpf; + QueryPerformanceFrequency(&qpf); + return (double)(now->QuadPart - start->QuadPart) / (double)qpf.QuadPart; +} + +static void sleep_ms(unsigned int ms) +{ + Sleep(ms); +} +#else +typedef struct timespec wall_time_t; + +static void wall_time_now(wall_time_t *t) +{ + clock_gettime(CLOCK_MONOTONIC, t); +} + +static double wall_time_seconds(const wall_time_t *start, const wall_time_t *now) +{ + time_t sec = now->tv_sec - start->tv_sec; + long nsec = now->tv_nsec - start->tv_nsec; + return (double)sec + (double)nsec / 1e9; +} + +static void sleep_ms(unsigned int ms) +{ + struct timespec req; + req.tv_sec = ms / 1000; + req.tv_nsec = (long)(ms % 1000) * 1000000L; + nanosleep(&req, NULL); +} +#endif + +/* ========================================================================== */ +/* CUDA error-checking */ +/* ========================================================================== */ + +/* + * CUDA_CHECK -- check a CUDA runtime call and jump to cleanup on failure. + * + * Requires the enclosing function to have: + * • A label cleanup: near the end. + * • All device handles/pointers initialised to NULL/nullptr before first use. + */ +#define CUDA_CHECK(call) do { \ + cudaError_t _ce = (call); \ + if (_ce != cudaSuccess) { \ + fprintf(stderr, "[cuda] %s:%d %s\n", \ + __FILE__, __LINE__, cudaGetErrorString(_ce)); \ + goto cleanup; \ + } \ +} while (0) + +/* ========================================================================== */ +/* DES constant tables */ +/* ========================================================================== */ + +__constant__ static uint8_t d_IP[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__ static uint8_t d_FP[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 +}; + +/* + * Combined S+P table: SP[s*64 + b6] is the contribution of S-box s with 6-bit + * input b6 to the 32-bit P-permutation output. Size: 8x64x4 = 2 KB. + * Precomputed on CPU in desbrute_launch() and uploaded via cudaMemcpyToSymbol. + */ +__constant__ static uint32_t d_SP[512]; + +/* Precomputed DES encryption subkeys for the fixed key half */ +__constant__ static uint64_t d_fixed_sk[16]; + +/* + * Key-schedule XOR tables for the variable (candidate) key half. + * + * The 4 variable bytes contribute 28 effective bits (7 bits each, parity + * discarded) to the 16x48-bit subkeys. Because the DES key schedule is a + * sequence of linear bit operations (PC1 permutation, half-rotations, PC2 + * permutation), the contribution of each input bit to every subkey bit is + * fixed and can be precomputed. + * + * d_cand_sk_base[r] = subkeys when all variable bytes are 0 + * d_cand_sk_contrib[b*16+r] = XOR contribution of effective bit b (0..27) + * to subkey round r (0..15) + * + * In the kernel, per candidate: + * cand_sk[r] = d_cand_sk_base[r] + * for b in 0..27 (unrolled, branchless): + * mask = -(uint64_t)((idx >> b) & 1) // 0 or 0xFFFFFFFFFFFFFFFF + * cand_sk[r] ^= d_cand_sk_contrib[b*16+r] & mask + * + * This replaces DPERM(PC1,56) + 16xDPERM(PC2,48) with 28x16 ALU ops. + * Size: (16 + 28*16) * 8 = 3712 bytes. + */ +__constant__ static uint64_t d_cand_sk_base[16]; +__constant__ static uint64_t d_cand_sk_contrib[28 * 16]; + +/* ========================================================================== */ +/* Bit permutation macro (used only for IP/FP in device code) */ +/* ========================================================================== */ + +#define DPERM(src, src_bits, tbl, N, dst) do { \ + int _i; (dst) = 0; \ + _Pragma("unroll") \ + for (_i = 0; _i < (N); _i++) { \ + int _sb = (tbl)[_i] - 1; \ + (dst) |= (((src) >> ((src_bits)-1-_sb)) & 1ULL) \ + << ((N)-1-_i); \ + } \ +} while(0) + +/* ========================================================================== */ +/* Device DES primitives */ +/* ========================================================================== */ + +/* + * Feistel f function using the shared-memory SP table. + * + * E expansion is inlined as direct bit extractions from R — eliminates the + * 48-iteration DPERM(E) loop entirely. Called 3x16 = 48 times per candidate. + * + * E table maps R bits (DES 1-indexed) to 6-bit S-box inputs (0-indexed from + * MSB of 48-bit eR): + * S1: E[0..5] = DES bits 32,1,2,3,4,5 → R bits 0,31,30,29,28,27 + * S2: E[6..11] = DES bits 4,5,6,7,8,9 → R bits 28,27,26,25,24,23 + * S3: = DES bits 8..13 → R bits 24..19 + * S4: = DES bits 12..17 → R bits 20..15 + * S5: = DES bits 16..21 → R bits 16..11 + * S6: = DES bits 20..25 → R bits 12..7 + * S7: = DES bits 24..29 → R bits 8..3 + * S8: = DES bits 28,29,30,31,32,1→ R bits 4,3,2,1,0,31 + * + * In our uint32_t R: DES bit i (1-based) = bit (32-i) from LSB. + */ +__device__ static __forceinline__ +uint32_t d_des_f(uint32_t R, uint64_t K, const uint32_t * __restrict__ sp) +{ + /* Extract 6-bit S-box inputs directly from R */ + uint32_t r0 = ((R & 1u) << 5) | ((R >> 27) & 0x1Fu); /* S1 */ + uint32_t r1 = (R >> 23) & 0x3Fu; /* S2 */ + uint32_t r2 = (R >> 19) & 0x3Fu; /* S3 */ + uint32_t r3 = (R >> 15) & 0x3Fu; /* S4 */ + uint32_t r4 = (R >> 11) & 0x3Fu; /* S5 */ + uint32_t r5 = (R >> 7) & 0x3Fu; /* S6 */ + uint32_t r6 = (R >> 3) & 0x3Fu; /* S7 */ + uint32_t r7 = ((R & 0x1Fu) << 1) | ((R >> 31) & 1u); /* S8 */ + + /* XOR with subkey 6-bit chunks (K is 48 bits in bits 47..0) */ + r0 ^= (uint32_t)((K >> 42) & 0x3Fu); + r1 ^= (uint32_t)((K >> 36) & 0x3Fu); + r2 ^= (uint32_t)((K >> 30) & 0x3Fu); + r3 ^= (uint32_t)((K >> 24) & 0x3Fu); + r4 ^= (uint32_t)((K >> 18) & 0x3Fu); + r5 ^= (uint32_t)((K >> 12) & 0x3Fu); + r6 ^= (uint32_t)((K >> 6) & 0x3Fu); + r7 ^= (uint32_t)(K & 0x3Fu); + + return sp[r0] ^ sp[64+r1] ^ sp[128+r2] ^ sp[192+r3] + ^ sp[256+r4] ^ sp[320+r5] ^ sp[384+r6] ^ sp[448+r7]; +} + +/* + * 16 Feistel rounds with NO Initial or Final Permutation. + * + * rev=0 encrypt (subkeys sk[0]..sk[15]) + * rev=1 decrypt (subkeys sk[15]..sk[0]) + * + * Input `block` must already have IP applied. + * Output is the pre-FP word (R16 || L16). + */ +__device__ static __forceinline__ +uint64_t d_des_rounds(uint64_t block, const uint64_t sk[16], int rev, + const uint32_t * __restrict__ sp) +{ + uint32_t L = (uint32_t)(block >> 32); + uint32_t R = (uint32_t)(block); + #pragma unroll 16 + for (int i = 0; i < 16; i++) { + uint64_t K = rev ? sk[15 - i] : sk[i]; + uint32_t nR = L ^ d_des_f(R, K, sp); + L = R; + R = nR; + } + return ((uint64_t)R << 32) | L; /* R16 || L16 (pre-FP) */ +} + +/* + * Optimised 2TDEA decrypt. + * + * One IP at entry + three rounds-only stages + one FP at exit. + * Saves 4 permutations vs. the naive DES_full x 3 approach. + * + * mode=0 K1=candidate, K2=fixed: D_cand -> E_fixed -> D_cand + * mode=1 K2=candidate, K1=fixed: D_fixed -> E_cand -> D_fixed + */ +__device__ static __forceinline__ +uint64_t d_tdea2_dec(uint64_t cipher, const uint64_t cand_sk[16], int mode, + const uint32_t * __restrict__ sp) +{ + uint64_t ip, t1, t2, t3, out; + DPERM(cipher, 64, d_IP, 64, ip); + if (mode == 0) { + t1 = d_des_rounds(ip, cand_sk, 1, sp); + t2 = d_des_rounds(t1, d_fixed_sk, 0, sp); + t3 = d_des_rounds(t2, cand_sk, 1, sp); + } else { + t1 = d_des_rounds(ip, d_fixed_sk, 1, sp); + t2 = d_des_rounds(t1, cand_sk, 0, sp); + t3 = d_des_rounds(t2, d_fixed_sk, 1, sp); + } + DPERM(t3, 64, d_FP, 64, out); + return out; +} + +/* ========================================================================== */ +/* LFSR validators (counterfeit mode only) */ +/* ========================================================================== */ + +__device__ static __forceinline__ int d_valid_lfsr_ulcg(uint64_t x64) +{ + uint16_t x = (uint16_t)(x64 >> 48); + x = (uint16_t)(x << 15 | ((x >> 1) ^ ((x >> 3 ^ x >> 4 ^ x >> 6) & 1))); + if (x != (uint16_t)((x64 >> 32) & 0xFFFF)) return 0; + x = (uint16_t)(x << 15 | ((x >> 1) ^ ((x >> 3 ^ x >> 4 ^ x >> 6) & 1))); + if (x != (uint16_t)((x64 >> 16) & 0xFFFF)) return 0; + x = (uint16_t)(x << 15 | ((x >> 1) ^ ((x >> 3 ^ x >> 4 ^ x >> 6) & 1))); + return (x == (uint16_t)(x64 & 0xFFFF)); +} + +__device__ static __forceinline__ int d_valid_lfsr_mfc(uint64_t x64) +{ + int i; + uint16_t x = (uint16_t)(x64 & 0xFFFF); + #pragma unroll 16 + for (i = 0; i < 16; i++) + x = (uint16_t)(x >> 1 | (x ^ x >> 2 ^ x >> 3 ^ x >> 5) << 15); + if (x != (uint16_t)((x64 >> 16) & 0xFFFF)) return 0; + #pragma unroll 16 + for (i = 0; i < 16; i++) + x = (uint16_t)(x >> 1 | (x ^ x >> 2 ^ x >> 3 ^ x >> 5) << 15); + if (x != (uint16_t)((x64 >> 32) & 0xFFFF)) return 0; + #pragma unroll 16 + for (i = 0; i < 16; i++) + x = (uint16_t)(x >> 1 | (x ^ x >> 2 ^ x >> 3 ^ x >> 5) << 15); + return (x == (uint16_t)((x64 >> 48) & 0xFFFF)); +} + +/* ========================================================================== */ +/* Brute-force kernel */ +/* ========================================================================== */ + +/* + * PROGRESS_STRIDE: candidates each thread processes between progress flushes. + * Lower = finer granularity, slightly more atomic overhead. + */ +#define PROGRESS_STRIDE 4096u + +__global__ void desbrute_kernel( + uint64_t ciphertext, + uint64_t init_ciphertext, + uint64_t prev_ciphertext, + int cand_mode, + int lfsr_type, + int is_reader_mode, + int *d_found, + uint32_t *d_result_idx, + unsigned long long *d_progress +) +{ + /* + * Step 1: load SP table from __constant__ into shared memory. + * + * Using shared memory avoids constant-cache serialisation: in a 32-thread + * warp each thread accesses a different S-box entry (divergent), so constant + * memory would serialize those 32 reads. Shared memory handles them in + * parallel with ~4-cycle latency. + */ + __shared__ uint32_t s_SP[512]; + for (int i = (int)threadIdx.x; i < 512; i += (int)blockDim.x) + s_SP[i] = d_SP[i]; + __syncthreads(); /* Barrier 1: SP table visible to all threads in block */ + + /* + * Step 2: initialise per-block early-exit flag. + * Separate barrier so the intent of each sync is clear. + */ + __shared__ int s_found; + if (threadIdx.x == 0) + s_found = *d_found; + __syncthreads(); /* Barrier 2: s_found visible before loop starts */ + + const uint32_t total = 1u << 28; + const uint32_t stride = (uint32_t)(gridDim.x * blockDim.x); + uint32_t idx = (uint32_t)(blockIdx.x * blockDim.x + threadIdx.x); + + /* + * candidates_tested semantics + * ---------------------------- + * no-match : sum of per-thread counters == total (2^28) + * match : approximate work done before early-exit; not necessarily + * idx+1 because threads run in parallel and some overshoot. + */ + unsigned long long local_count = 0; + + for (; idx < total; idx += stride) { + + /* + * Periodic progress flush and early-exit check. + * + * __syncthreads() is intentionally absent here. If it were present, + * threads that finish fewer iterations than PROGRESS_STRIDE could exit + * the loop while others hit the barrier -> deadlock. + * + * Only thread 0 flushes local_count and re-reads d_found. + * A __threadfence_block() after the write makes the updated s_found + * visible to other threads in the same block within a few cycles. + */ + if ((local_count & (PROGRESS_STRIDE - 1)) == 0) { + if (s_found) goto done; + if (threadIdx.x == 0) { + atomicAdd(d_progress, local_count); + local_count = 0; + s_found = *d_found; + __threadfence_block(); /* push s_found write to shared mem */ + } + } + local_count++; + + uint64_t cand_sk[16]; + #pragma unroll 16 + for (int r = 0; r < 16; r++) + cand_sk[r] = d_cand_sk_base[r]; + + #pragma unroll 28 + for (int b = 0; b < 28; b++) { + uint64_t mask = (uint64_t)(-(int64_t)((idx >> b) & 1u)); + #pragma unroll 16 + for (int r = 0; r < 16; r++) + cand_sk[r] ^= d_cand_sk_contrib[b * 16 + r] & mask; + } + + /* Optimised 2TDEA decrypt (single IP/FP for the whole chain) */ + uint64_t out = d_tdea2_dec(ciphertext, cand_sk, cand_mode, s_SP); + + /* Match check */ + int match = 0; + if (is_reader_mode) { + uint64_t init_out = d_tdea2_dec(init_ciphertext, cand_sk, + cand_mode, s_SP); + out ^= prev_ciphertext; + uint64_t rotated = (init_out << 8) | (init_out >> 56); + match = (out == rotated); + } else { + if (lfsr_type == LFSR_ULCG) match = d_valid_lfsr_ulcg(out); + else if (lfsr_type == LFSR_MFC) match = d_valid_lfsr_mfc(out); + } + + if (match) { + if (atomicCAS(d_found, 0, 1) == 0) + *d_result_idx = idx; + s_found = 1; + /* Do NOT increment local_count again here -- it was already + * counted above. An extra ++ would make candidates_tested + * off by one. */ + goto done; + } + } + +done: + /* Flush remaining local count to the global progress counter */ + if (local_count > 0) + atomicAdd(d_progress, local_count); +} + +/* ========================================================================== */ +/* CPU helpers */ +/* ========================================================================== */ + +static void cpu_des_keyschedule(uint64_t key64, uint64_t sk[16]) +{ + static const uint8_t PC1[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 + }; + static const uint8_t PC2[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 + }; + static const uint8_t SHF[16] = {1,1,2,2,2,2,2,2,1,2,2,2,2,2,2,1}; + int i, j; + uint64_t key56 = 0; + for (i = 0; i < 56; i++) { + int sb = PC1[i] - 1; + key56 |= ((key64 >> (63 - sb)) & 1ULL) << (55 - i); + } + uint32_t C = (uint32_t)(key56 >> 28) & 0x0FFFFFFFu; + uint32_t D = (uint32_t)(key56) & 0x0FFFFFFFu; + for (i = 0; i < 16; i++) { + int s = SHF[i]; + C = ((C << s) | (C >> (28 - s))) & 0x0FFFFFFFu; + D = ((D << s) | (D >> (28 - s))) & 0x0FFFFFFFu; + uint64_t CD = ((uint64_t)C << 28) | D; + uint64_t sub = 0; + for (j = 0; j < 48; j++) { + int sb2 = PC2[j] - 1; + sub |= ((CD >> (55 - sb2)) & 1ULL) << (47 - j); + } + sk[i] = sub; + } +} + +static inline uint64_t cpu_bytes_to_u64(const uint8_t *b) +{ + return ((uint64_t)b[0] << 56) | ((uint64_t)b[1] << 48) | + ((uint64_t)b[2] << 40) | ((uint64_t)b[3] << 32) | + ((uint64_t)b[4] << 24) | ((uint64_t)b[5] << 16) | + ((uint64_t)b[6] << 8) | (uint64_t)b[7]; +} + +/* + * Precompute the combined S+P table on the CPU. + * + * For S-box s (0..7) and 6-bit input b6 (0..63): + * sval = SBOX[s][row][col] (4-bit output, 0..15) + * sval occupies sout bits [28-s*4 .. 31-s*4] (bit 0 = integer LSB) + * P permutation: pf[31-i] = sout[32 - P[i]] + * SP[s*64+b6] accumulates only pf bits from S-box s. + */ +static void cpu_compute_sp(uint32_t sp_out[512]) +{ + static const uint8_t SBOX[512] = { + /* S1 */ + 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, + /* S2 */ + 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, + /* S3 */ + 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, + /* S4 */ + 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, + /* S5 */ + 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, + /* S6 */ + 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, + /* S7 */ + 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, + /* S8 */ + 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 + }; + static const uint8_t P[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 + }; + + int s, b6, i; + for (s = 0; s < 8; s++) { + for (b6 = 0; b6 < 64; b6++) { + int row = ((b6 >> 5) & 1) * 2 + (b6 & 1); + int col = (b6 >> 1) & 0xF; + int sval = SBOX[s * 64 + row * 16 + col]; /* 4-bit, 0..15 */ + + /* Accumulate pf contribution from this S-box only */ + uint32_t sp = 0; + for (i = 0; i < 32; i++) { + int sout_bit = 32 - P[i]; /* integer bit index (0=LSB) */ + int sbox_lo = 28 - s * 4; + int sbox_hi = 31 - s * 4; + if (sout_bit >= sbox_lo && sout_bit <= sbox_hi) { + int sval_bit = sout_bit - sbox_lo; /* 0..3 */ + if ((sval >> sval_bit) & 1) + sp |= (1u << (31 - i)); + } + } + sp_out[s * 64 + b6] = sp; + } + } +} + +/* + * Precompute the per-bit XOR contribution tables for the candidate key half. + * + * The 4 variable bytes (at byte positions var_offset..var_offset+3 within + * the 8-byte half-key) contribute 28 effective bits (7 each) to the 16 + * DES subkeys. Because PC1 and PC2 are linear bit selections, each input + * bit's contribution to every subkey bit is fixed and can be tabulated. + * + * For each of the 28 effective bit positions b (0..27): + * byte_in_half = b / 7 which variable byte (0..3) + * bit_in_byte = b % 7 which bit within that byte (0..6) + * actual bit in key byte = bit (bit_in_byte + 1) (parity at bit 0 is 0) + * + * We compute the key schedule for a half-key with only that one bit set, + * giving the XOR contribution of that bit to each subkey round. + * + * Outputs: + * sk_base[16] -- subkeys for candidate half with variable bytes=0 + * sk_contrib[28 * 16] -- contrib[b*16+r] = subkey[r] XOR contrib of bit b + */ +static void cpu_compute_sk_tables( + uint64_t base_half, + int var_offset, + uint64_t sk_base[16], + uint64_t sk_contrib[28 * 16]) +{ + /* Base: key schedule with variable bytes already zeroed (base_half has them 0) */ + cpu_des_keyschedule(base_half, sk_base); + + /* Per-bit contributions */ + int b, r; + for (b = 0; b < 28; b++) { + int byte_in_half = b / 7; + int bit_in_byte = b % 7; + + /* Build half-key with only this one effective bit set */ + uint8_t tmp[8]; + memset(tmp, 0, sizeof(tmp)); + tmp[var_offset + byte_in_half] = (uint8_t)(1u << (bit_in_byte + 1)); + + uint64_t bit_key = cpu_bytes_to_u64(tmp); + uint64_t sk_for_bit[16]; + cpu_des_keyschedule(bit_key, sk_for_bit); + + /* XOR contribution: because key schedule is linear over GF(2), the + * contribution of this bit equals the key schedule result for the + * half-key with only this bit set (base has variable bytes = 0, + * which means their contribution is 0 already). */ + for (r = 0; r < 16; r++) + sk_contrib[b * 16 + r] = sk_for_bit[r]; + } +} + +/* ========================================================================== */ +/* Public launch function */ +/* ========================================================================== */ + +void desbrute_launch( + uint64_t ciphertext_be, + uint64_t init_ciphertext_be, + uint64_t prev_ciphertext_be, + const uint8_t base_key[KEY_SIZE], + int key_mode, + int lfsr_type, + int is_reader_mode, + desbrute_result_t *result +) +{ + /* -- API-level validation ----------------------------------------------- */ + if (!result) + return; + + memset(result, 0, sizeof(*result)); + result->idx = UINT32_MAX; /* sentinel: no match yet */ + + if (!base_key) { return; } + if (key_mode < 0 || key_mode > 3) { return; } + if (is_reader_mode != 0 && is_reader_mode != 1) { return; } + if (!is_reader_mode && lfsr_type == LFSR_UNDEF) { return; } + + /* -- Declare all locals up-front so no goto bypasses an initialisation -- */ + cudaStream_t stream = NULL; + cudaStream_t progress_stream = NULL; + cudaEvent_t ev_start = NULL; + cudaEvent_t ev_stop = NULL; + int *d_found = NULL; + uint32_t *d_result_idx = NULL; + unsigned long long *d_progress = NULL; + int candidate_in_k1; + int cand_mode; + int var_offset; + uint64_t fixed_u64; + uint64_t base_half; + int block_size; + int min_grid; + int dev; + int dev_mp; + int grid_size; + uint64_t total_candidates; + wall_time_t qpc_start, qpc_now; + uint64_t h_fixed_sk[16]; + uint32_t h_SP[512]; + uint8_t cand_bytes[8]; + uint64_t h_cand_sk_base[16]; + uint64_t h_cand_sk_contrib[28 * 16]; + + /* -- Candidate / fixed half selection ----------------------------------- */ + candidate_in_k1 = (key_mode < 2) ? 1 : 0; + cand_mode = candidate_in_k1 ? 0 : 1; + if (candidate_in_k1) + var_offset = (key_mode % 2) * 4; + else + var_offset = ((key_mode - 2) % 2) * 4; + + /* -- Fixed-half subkeys ------------------------------------------------- */ + fixed_u64 = candidate_in_k1 + ? cpu_bytes_to_u64(base_key + 8) + : cpu_bytes_to_u64(base_key); + cpu_des_keyschedule(fixed_u64, h_fixed_sk); + + CUDA_CHECK(cudaMemcpyToSymbol(d_fixed_sk, h_fixed_sk, 16 * sizeof(uint64_t))); + + /* -- Precompute and upload SP table ------------------------------------- */ + cpu_compute_sp(h_SP); + CUDA_CHECK(cudaMemcpyToSymbol(d_SP, h_SP, 512 * sizeof(uint32_t))); + + /* -- Candidate half: zero out variable bytes ---------------------------- */ + if (candidate_in_k1) memcpy(cand_bytes, base_key, 8); + else memcpy(cand_bytes, base_key + 8, 8); + cand_bytes[var_offset] = 0; + cand_bytes[var_offset + 1] = 0; + cand_bytes[var_offset + 2] = 0; + cand_bytes[var_offset + 3] = 0; + base_half = cpu_bytes_to_u64(cand_bytes); + + /* -- Key-schedule tables (main throughput optimisation) ----------------- */ + cpu_compute_sk_tables(base_half, var_offset, h_cand_sk_base, h_cand_sk_contrib); + + CUDA_CHECK(cudaMemcpyToSymbol(d_cand_sk_base, h_cand_sk_base, + 16 * sizeof(uint64_t))); + CUDA_CHECK(cudaMemcpyToSymbol(d_cand_sk_contrib, h_cand_sk_contrib, + 28 * 16 * sizeof(uint64_t))); + + /* -- Device buffers ----------------------------------------------------- */ + CUDA_CHECK(cudaMalloc((void **)&d_found, sizeof(int))); + CUDA_CHECK(cudaMalloc((void **)&d_result_idx, sizeof(uint32_t))); + CUDA_CHECK(cudaMalloc((void **)&d_progress, sizeof(unsigned long long))); + + CUDA_CHECK(cudaMemset(d_found, 0, sizeof(int))); + CUDA_CHECK(cudaMemset(d_result_idx, 0xFF, sizeof(uint32_t))); /* UINT32_MAX */ + CUDA_CHECK(cudaMemset(d_progress, 0, sizeof(unsigned long long))); + + /* -- Launch configuration ----------------------------------------------- */ + block_size = 256; + min_grid = 0; + CUDA_CHECK(cudaOccupancyMaxPotentialBlockSize(&min_grid, &block_size, + desbrute_kernel, 0, 0)); + if (block_size > 256) block_size = 256; + + /* + * Use the CURRENT device (not hardcoded 0) so the grid is correctly sized + * when the user passes a non-zero [cuda_dev] argument. + */ + dev = 0; + dev_mp = 1; + CUDA_CHECK(cudaGetDevice(&dev)); + CUDA_CHECK(cudaDeviceGetAttribute(&dev_mp, cudaDevAttrMultiProcessorCount, dev)); + + grid_size = dev_mp * 64; + if (grid_size > 65535) grid_size = 65535; + total_candidates = 1ULL << 28; + + printf("\n[kernel] Grid: %d blocks x %d threads (64 waves/SM on dev %d)\n", + grid_size, block_size, dev); + printf("[kernel] Keyspace: 2^28 = %llu candidates\n", + (unsigned long long)total_candidates); + fflush(stdout); + + /* -- Streams and events ------------------------------------------------- + * Two streams: + * stream -- the kernel runs here + * progress_stream -- independent stream for live progress copies + * + * Because progress_stream has no dependency on the kernel stream, + * cudaMemcpyAsync on it reads d_progress while the kernel is still running. + * Atomic updates in the kernel commit to L2/DRAM so they are visible. + */ + CUDA_CHECK(cudaStreamCreate(&stream)); + CUDA_CHECK(cudaStreamCreate(&progress_stream)); + CUDA_CHECK(cudaEventCreate(&ev_start)); + CUDA_CHECK(cudaEventCreate(&ev_stop)); + + /* -- Wall-clock start --------------------------------------------------- */ + wall_time_now(&qpc_start); + + /* -- Launch ------------------------------------------------------------- */ + CUDA_CHECK(cudaEventRecord(ev_start, stream)); + + desbrute_kernel<<>>( + ciphertext_be, init_ciphertext_be, prev_ciphertext_be, + cand_mode, lfsr_type, is_reader_mode, + d_found, d_result_idx, d_progress + ); + + /* Check for launch errors immediately */ + CUDA_CHECK(cudaGetLastError()); + +#if defined(DEBUG) || defined(_DEBUG) + /* In debug builds, synchronise now to surface async execution errors close + * to the launch site rather than at a later sync point. */ + CUDA_CHECK(cudaStreamSynchronize(stream)); +#endif + + CUDA_CHECK(cudaEventRecord(ev_stop, stream)); + + /* -- Progress polling loop ---------------------------------------------- */ + printf("[search] Running\n"); + fflush(stdout); + + { + cudaError_t q; + while ((q = cudaStreamQuery(stream)) == cudaErrorNotReady) { + sleep_ms(150); + + /* Read progress via the independent progress_stream */ + unsigned long long prog = 0; + cudaMemcpyAsync(&prog, d_progress, sizeof(unsigned long long), + cudaMemcpyDeviceToHost, progress_stream); + cudaStreamSynchronize(progress_stream); + + wall_time_now(&qpc_now); + double wall_s = wall_time_seconds(&qpc_start, &qpc_now); + + double pct = (prog >= total_candidates) + ? 100.0 + : 100.0 * (double)prog / (double)total_candidates; + double kps = (wall_s > 0.001) ? (double)prog / wall_s : 0.0; + + int hf = 0; + cudaMemcpyAsync(&hf, d_found, sizeof(int), + cudaMemcpyDeviceToHost, progress_stream); + cudaStreamSynchronize(progress_stream); + + printf("\r[search] %6.2f%% | %8.2f M keys/s | %.2f s elapsed%s", + pct, kps / 1e6, wall_s, hf ? " [FOUND!]" : " "); + fflush(stdout); + } + + /* Any status other than cudaSuccess means an async error occurred */ + if (q != cudaSuccess) { + fprintf(stderr, "\n[cuda] stream error: %s\n", cudaGetErrorString(q)); + goto cleanup; + } + } + + CUDA_CHECK(cudaStreamSynchronize(stream)); + printf("\n"); + + /* -- Gather results ----------------------------------------------------- */ + wall_time_now(&qpc_now); + result->wall_ms = 1000.0 * wall_time_seconds(&qpc_start, &qpc_now); + + { + float gpu_ms = 0.0f; + CUDA_CHECK(cudaEventElapsedTime(&gpu_ms, ev_start, ev_stop)); + result->gpu_ms = gpu_ms; + + unsigned long long final_prog = 0; + CUDA_CHECK(cudaMemcpy(&final_prog, d_progress, + sizeof(unsigned long long), cudaMemcpyDeviceToHost)); + result->candidates_tested = (uint64_t)final_prog; + + if (gpu_ms > 0.0f) + result->keys_per_sec = (double)final_prog / (gpu_ms / 1000.0); + + int h_found = 0; + uint32_t h_result_idx = UINT32_MAX; + CUDA_CHECK(cudaMemcpy(&h_found, d_found, + sizeof(int), cudaMemcpyDeviceToHost)); + CUDA_CHECK(cudaMemcpy(&h_result_idx, d_result_idx, + sizeof(uint32_t), cudaMemcpyDeviceToHost)); + result->found = h_found; + result->idx = h_result_idx; + } + + /* -- Cleanup (also reached via CUDA_CHECK goto on error) ---------------- */ +cleanup: + if (ev_stop) cudaEventDestroy(ev_stop); + if (ev_start) cudaEventDestroy(ev_start); + if (progress_stream) cudaStreamDestroy(progress_stream); + if (stream) cudaStreamDestroy(stream); + if (d_progress) cudaFree(d_progress); + if (d_result_idx) cudaFree(d_result_idx); + if (d_found) cudaFree(d_found); +} diff --git a/tools/mfulc_des_brute_cuda/desbrute.h b/tools/mfulc_des_brute_cuda/desbrute.h new file mode 100644 index 000000000..b7d7e5487 --- /dev/null +++ b/tools/mfulc_des_brute_cuda/desbrute.h @@ -0,0 +1,40 @@ +/* + * desbrute.h — Shared types and GPU launch interface. + * No external dependencies beyond CUDA runtime. + */ +#pragma once +#include + +#define BLOCK_SIZE 8 +#define KEY_SIZE 16 + +#define LFSR_UNDEF 0 +#define LFSR_ULCG 1 +#define LFSR_MFC 2 + +/* Metrics returned by desbrute_launch */ +typedef struct { + int found; /* 1 if a key was found */ + uint32_t idx; /* winning candidate index (if found) */ + double wall_ms; /* wall-clock time in milliseconds */ + float gpu_ms; /* GPU kernel time (CUDA events) */ + uint64_t candidates_tested; /* total candidates processed before exit */ + double keys_per_sec; /* throughput: candidates / gpu_sec */ +} desbrute_result_t; + +/* + * desbrute_launch — brute-force 2^28 candidate key indices on the GPU. + * + * All uint64_t block values are big-endian packed (byte 0 in bits 63-56). + * The caller must have already called cudaSetDevice() if needed. + */ +void desbrute_launch( + uint64_t ciphertext_be, + uint64_t init_ciphertext_be, + uint64_t prev_ciphertext_be, + const uint8_t base_key[KEY_SIZE], + int key_mode, + int lfsr_type, + int is_reader_mode, + desbrute_result_t *result +); diff --git a/tools/mfulc_des_brute_cuda/main.cu b/tools/mfulc_des_brute_cuda/main.cu new file mode 100644 index 000000000..eaaa3845b --- /dev/null +++ b/tools/mfulc_des_brute_cuda/main.cu @@ -0,0 +1,332 @@ +/* + * main.cu - CUDA-accelerated 2TDEA key recovery. + * Written by C2Pwn + * + * Usage + * ----- + * Counterfeit mode (-c): + * desbrute -c [device] + * + * Reader nonce mode (-r): + * desbrute -r [device] + * + * [device] is optional. When omitted the best available GPU is chosen + * automatically (highest SM count). Pass an integer to force a specific + * CUDA device index. + */ + +#include +#include +#include +#include +#ifdef _WIN32 +#include +#endif +#include + +#include "des_host.h" +#include "desbrute.h" + +/* ========================================================================== */ +/* Helpers */ +/* ========================================================================== */ + +static int hex_to_bytes(const char *hex, uint8_t *buf, size_t len) +{ + size_t i; + if (strlen(hex) != len * 2) return 0; + for (i = 0; i < len; i++) { + unsigned int byte; + if (sscanf(hex + 2 * i, "%2x", &byte) != 1) return 0; + buf[i] = (uint8_t)byte; + } + return 1; +} + +static void print_hex(const uint8_t *buf, size_t len) +{ + size_t i; + for (i = 0; i < len; i++) printf("%02X", buf[i]); +} + +static void print_separator(void) +{ + printf("---------------------------------------------------------------\n"); +} + +static void print_usage(const char *name) +{ + fprintf(stderr, + "\nUsage:\n" + " Counterfeit: %s -c " + " [device]\n" + " Reader nonce: %s -r " + " [device]\n" + "\n" + " seg : 1-4 (which 4-byte block of the 16-byte key to brute-force)\n" + " device : optional CUDA device index (default: auto-select best GPU)\n" + "\n", + name, name); + exit(1); +} + +/* ========================================================================== */ +/* Auto device selection */ +/* ========================================================================== */ + +/* + * Pick the CUDA device with the most streaming multiprocessors. + * On a single-GPU system this always returns 0. + * On multi-GPU systems it chooses the most powerful card. + * Prints a one-line summary of all available devices. + */ +static int pick_best_device(void) +{ + int count = 0; + if (cudaGetDeviceCount(&count) != cudaSuccess || count == 0) { + fprintf(stderr, "[cuda] No CUDA devices found.\n"); + exit(1); + } + + int best_dev = 0, best_mp = 0; + int d; + for (d = 0; d < count; d++) { + struct cudaDeviceProp p; + if (cudaGetDeviceProperties(&p, d) != cudaSuccess) continue; + printf(" [dev %d] %s SM %d.%d %d MPs\n", + d, p.name, p.major, p.minor, p.multiProcessorCount); + if (p.multiProcessorCount > best_mp) { + best_mp = p.multiProcessorCount; + best_dev = d; + } + } + return best_dev; +} + +/* ========================================================================== */ +/* GPU info banner */ +/* ========================================================================== */ + +static void print_gpu_info(int dev) +{ + struct cudaDeviceProp p; + if (cudaGetDeviceProperties(&p, dev) != cudaSuccess) return; + + double bw_gbps = 2.0 * p.memoryClockRate * 1e3 + * p.memoryBusWidth / 8.0 / 1e9; + + /* Approximate CUDA core count — varies by architecture */ + int cores_per_mp = + (p.major == 9) ? 128 : + (p.major == 8 && p.minor == 9) ? 128 : + (p.major == 8 && p.minor == 6) ? 128 : + (p.major == 8 && p.minor == 0) ? 64 : + (p.major == 7 && p.minor == 5) ? 64 : + (p.major == 7) ? 64 : + (p.major == 6 && p.minor == 1) ? 128 : + (p.major == 6 && p.minor == 0) ? 64 : + (p.major == 5) ? 128 : 64; + + print_separator(); + printf(" GPU : %s [device %d]\n", p.name, dev); + printf(" SM : %d.%d | MPs: %d | CUDA cores: ~%d\n", + p.major, p.minor, p.multiProcessorCount, + p.multiProcessorCount * cores_per_mp); + printf(" Clock: %d MHz (mem: %d MHz, bus: %d-bit)\n", + p.clockRate / 1000, p.memoryClockRate / 1000, p.memoryBusWidth); + printf(" VRAM : %.0f MB | Peak BW: %.0f GB/s\n", + (double)p.totalGlobalMem / (1024.0 * 1024.0), bw_gbps); + printf(" L2 : %d KB | Warp size: %d\n", + p.l2CacheSize / 1024, p.warpSize); + print_separator(); +} + +/* ========================================================================== */ +/* main */ +/* ========================================================================== */ + +int main(int argc, char **argv) +{ + int is_reader_mode, seg, key_mode, cuda_dev, lfsr_type; + uint8_t init_ciphertext[BLOCK_SIZE]; + uint8_t tmp_blocks[2 * BLOCK_SIZE]; + uint8_t ciphertext[BLOCK_SIZE]; + uint8_t prev_ciphertext[BLOCK_SIZE]; + uint8_t base_key[KEY_SIZE]; + desbrute_result_t result; + + /* --- Argument parsing ------------------------------------------------- */ + if (argc < 2) print_usage(argv[0]); + + if (strcmp(argv[1], "-c") == 0) { + is_reader_mode = 0; + if (argc < 6 || argc > 7) { + fprintf(stderr, "Error: -c requires exactly 4 arguments" + " plus an optional device index.\n"); + print_usage(argv[0]); + } + } else if (strcmp(argv[1], "-r") == 0) { + is_reader_mode = 1; + if (argc < 6 || argc > 7) { + fprintf(stderr, "Error: -r requires exactly 4 arguments" + " plus an optional device index.\n"); + print_usage(argv[0]); + } + } else { + fprintf(stderr, "Error: first argument must be -c or -r\n"); + print_usage(argv[0]); + } + + /* --- Parse fixed positional args ------------------------------------- */ + if (is_reader_mode) { + if (!hex_to_bytes(argv[2], init_ciphertext, BLOCK_SIZE)) { + fprintf(stderr, "Error: ERndB must be exactly 16 hex chars (8 bytes)\n"); + return 1; + } + if (!hex_to_bytes(argv[3], tmp_blocks, 2 * BLOCK_SIZE)) { + fprintf(stderr, "Error: ERndA||RndB' must be exactly 32 hex chars (16 bytes)\n"); + return 1; + } + memcpy(prev_ciphertext, tmp_blocks, BLOCK_SIZE); + memcpy(ciphertext, tmp_blocks + BLOCK_SIZE, BLOCK_SIZE); + } else { + if (!hex_to_bytes(argv[2], init_ciphertext, BLOCK_SIZE)) { + fprintf(stderr, "Error: null-key ERndB must be exactly 16 hex chars (8 bytes)\n"); + return 1; + } + if (!hex_to_bytes(argv[3], ciphertext, BLOCK_SIZE)) { + fprintf(stderr, "Error: target ERndB must be exactly 16 hex chars (8 bytes)\n"); + return 1; + } + memset(prev_ciphertext, 0, BLOCK_SIZE); + } + + if (!hex_to_bytes(argv[4], base_key, KEY_SIZE)) { + fprintf(stderr, "Error: base key must be exactly 32 hex chars (16 bytes)\n"); + return 1; + } + + seg = atoi(argv[5]); + if (seg < 1 || seg > 4) { + fprintf(stderr, "Error: segment must be 1, 2, 3, or 4\n"); + return 1; + } + key_mode = seg - 1; + + /* --- CUDA device selection -------------------------------------------- */ + print_separator(); + printf(" DESBRUTE - CUDA 2TDEA key recovery engine\n"); + print_separator(); + + if (argc == 7) { + /* Explicit device requested */ + cuda_dev = atoi(argv[6]); + int count = 0; + cudaGetDeviceCount(&count); + if (cuda_dev < 0 || cuda_dev >= count) { + fprintf(stderr, + "[cuda] Device %d does not exist (%d device(s) available).\n" + " Re-run without the [device] argument to auto-select.\n", + cuda_dev, count); + return 1; + } + printf(" Device: %d (user-specified)\n", cuda_dev); + } else { + /* Auto-select best device */ + printf(" Auto-selecting best CUDA device...\n"); + cuda_dev = pick_best_device(); + printf(" Selected: device %d\n", cuda_dev); + } + + if (cudaSetDevice(cuda_dev) != cudaSuccess) { + fprintf(stderr, "[cuda] Failed to set device %d\n", cuda_dev); + return 1; + } + + /* --- LFSR detection (counterfeit mode only, CPU) ---------------------- */ + lfsr_type = LFSR_UNDEF; + if (!is_reader_mode) { + lfsr_type = h_detect_lfsr(init_ciphertext); + switch (lfsr_type) { + case LFSR_ULCG: + printf("[lfsr] Detected: ULCG\n"); break; + case LFSR_MFC: + printf("[lfsr] Detected: MFC (USCUID-UL / FJ8010)\n"); break; + default: + fprintf(stderr, "[lfsr] ERROR: unrecognised LFSR - cannot proceed\n"); + return 1; + } + } + + /* --- Banner ----------------------------------------------------------- */ + print_gpu_info(cuda_dev); + printf(" Mode : %s\n", is_reader_mode ? "Reader nonce (-r)" : "Counterfeit (-c)"); + printf(" Segment : %d (key_mode %d, bytes %d-%d of full key)\n", + seg, key_mode, key_mode * 4, key_mode * 4 + 3); + printf(" Base key : "); print_hex(base_key, KEY_SIZE); printf("\n"); + printf(" Ciphertext : "); print_hex(ciphertext, BLOCK_SIZE); printf("\n"); + printf(" Init CT : "); print_hex(init_ciphertext, BLOCK_SIZE); printf("\n"); + if (is_reader_mode) { + printf(" Prev CT : "); print_hex(prev_ciphertext, BLOCK_SIZE); printf("\n"); + } + print_separator(); + fflush(stdout); + + /* --- Launch GPU search ------------------------------------------------ */ + uint64_t ct_be = bytes_to_u64(ciphertext); + uint64_t init_be = bytes_to_u64(init_ciphertext); + uint64_t prev_be = bytes_to_u64(prev_ciphertext); + + desbrute_launch(ct_be, init_be, prev_be, + base_key, key_mode, lfsr_type, is_reader_mode, + &result); + + /* --- Metrics ---------------------------------------------------------- */ + print_separator(); + printf(" PERFORMANCE METRICS\n"); + print_separator(); + printf(" GPU kernel time : %.3f ms\n", result.gpu_ms); + printf(" Wall-clock time : %.3f ms\n", result.wall_ms); + printf(" Candidates : %llu / %llu\n", + (unsigned long long)result.candidates_tested, + (unsigned long long)(1ULL << 28)); + printf(" Throughput : %.2f M keys/s\n", result.keys_per_sec / 1e6); + { + double overhead_pct = (result.wall_ms > 0) + ? 100.0 * (result.wall_ms - result.gpu_ms) / result.wall_ms + : 0.0; + printf(" Host overhead : %.1f%%\n", overhead_pct); + } + print_separator(); + + /* --- Result ----------------------------------------------------------- */ + if (!result.found) { + printf(" RESULT: No matching key found.\n"); + print_separator(); + return 1; + } + + /* Reconstruct full key from winning index */ + uint8_t b0 = (uint8_t)(( result.idx & 0x7Fu) << 1); + uint8_t b1 = (uint8_t)(((result.idx >> 7) & 0x7Fu) << 1); + uint8_t b2 = (uint8_t)(((result.idx >> 14) & 0x7Fu) << 1); + uint8_t b3 = (uint8_t)(((result.idx >> 21) & 0x7Fu) << 1); + + uint8_t full_key[KEY_SIZE]; + memcpy(full_key, base_key, KEY_SIZE); + { + int seg_offset = key_mode * 4; + full_key[seg_offset] = b0; + full_key[seg_offset + 1] = b1; + full_key[seg_offset + 2] = b2; + full_key[seg_offset + 3] = b3; + } + + printf(" RESULT: KEY FOUND\n"); + printf(" Index : %u (0x%08X)\n", result.idx, result.idx); + printf(" Full key: "); print_hex(full_key, KEY_SIZE); printf("\n"); + print_separator(); + + cudaDeviceReset(); + return 0; +}