Skip to content

Commit eeb7ba4

Browse files
magnumrippermagnumripper
committed
o5logon-opencl: Support Oracle 12
Also implements internal mask, for a 3-4x boost. Closes openwall#5648
1 parent 8b9089a commit eeb7ba4

File tree

5 files changed

+620
-461
lines changed

5 files changed

+620
-461
lines changed

Diff for: ‎run/opencl/o5logon_kernel.cl

+198-74
Original file line numberDiff line numberDiff line change
@@ -1,89 +1,213 @@
11
/*
2-
This code was largely inspired by
3-
pyrit opencl kernel sha1 routines, royger's sha1 sample,
4-
and md5_opencl_kernel.cl inside jtr.
5-
Copyright 2011 by Samuele Giovanni Tonon
6-
samu at linuxasylum dot net
7-
and Copyright (c) 2012-2017, magnum
8-
This program comes with ABSOLUTELY NO WARRANTY; express or
9-
implied .
10-
This is free software, and you are welcome to redistribute it
11-
under certain conditions; as expressed here
12-
http://www.gnu.org/licenses/gpl-2.0.html
13-
*/
14-
15-
/*
16-
* Modifications (c) 2014 Harrison Neal.
17-
* Licensed GPLv2
2+
* Copyright (c) 2012-2025, magnum
3+
* Copyright (c) 2014 Harrison Neal
4+
* Copyright (c) 2011 by Samuele Giovanni Tonon, samu at linuxasylum dot net
5+
*
6+
* This program comes with ABSOLUTELY NO WARRANTY; express or
7+
* implied .
8+
* This is free software, and you are welcome to redistribute it
9+
* under certain conditions; as expressed here
10+
* http://www.gnu.org/licenses/gpl-2.0.html
1811
*/
1912

2013
#include "opencl_device_info.h"
2114
#include "opencl_misc.h"
2215
#include "opencl_sha1.h"
23-
#define AES_SRC_TYPE __constant
16+
#include "opencl_md5_ctx.h"
17+
#include "opencl_mask.h"
18+
19+
// Workaround for UHD Graphics 630 version "1.2(Jan 10 2025 21:16:54)"
20+
#if (__OS_X__ && gpu_intel(DEVICE_INFO))
21+
#define AES_BITSLICE 1
22+
#endif
2423
#include "opencl_aes.h"
2524

2625
typedef struct {
27-
uint salt[3]; // ((SALT_LENGTH + 3)/4)
28-
uchar ct[48]; // CIPHERTEXT_LENGTH
29-
} salt_t;
26+
uint pw_len; /* AUTH_PASSWORD length (blocks) */
27+
uint salt[(SALT_LENGTH + 1 + 3) / 4]; /* AUTH_VFR_DATA */
28+
uchar ct[CIPHERTEXT_LENGTH]; /* Server's AUTH_SESSKEY */
29+
uchar csk[CIPHERTEXT_LENGTH]; /* Client's AUTH_SESSKEY */
30+
uchar pw[PLAINTEXT_LENGTH + 16]; /* Client's AUTH_PASSWORD, padded */
31+
} o5logon_salt;
32+
33+
#define SECRET_LEN (CIPHERTEXT_LENGTH - 16)
3034

3135
__kernel void
32-
o5logon_kernel(__global const uint *keys, __constant salt_t *salt,
33-
__global const uint *index, __global uint *result)
36+
o5logon_kernel(__global const uchar* key_buf, __global const uint* const key_idx,
37+
__constant o5logon_salt* salt,
38+
__global volatile uint* crack_count_ret,
39+
__global uint* const out_index,
40+
__global const uint* const int_key_loc,
41+
__global const uint* const int_keys)
3442
{
3543
__local aes_local_t lt;
3644
AES_KEY akey; akey.lt = <
37-
uint W[16] = { 0 }, salt_s[3], output[5];
38-
uint gid = get_global_id(0);
39-
uint base = index[gid];
40-
uint len = base & 63;
41-
uint i;
42-
uint shift = len % 4;
43-
uint sr = 8 * shift;
44-
uint sl = 32 - sr;
45-
uint sra = (0xffffffff - (1 << sr)) + 1;
46-
uint sla = 0xffffffff - sra;
47-
union {
48-
uchar c[24];
49-
uint w[24 / 4];
50-
} aes_key;
51-
union {
52-
uchar c[16];
53-
ulong l[16 / 8];
54-
} pt;
55-
uchar iv[16];
56-
57-
keys += base >> 6;
58-
59-
for (i = 0; i < (len + 3) / 4; i++)
60-
W[i] = SWAP32(*keys++);
61-
62-
// Do the typical byte swapping...
63-
for (i = 0; i < 3; i++)
64-
salt_s[i] = SWAP32(salt->salt[i]);
65-
66-
// Shift the salt bytes into place after the given key.
67-
W[len / 4] |= (salt_s[0] & sra) >> sr;
68-
W[len / 4 + 1] = ((salt_s[0] & sla) << sl) | ((salt_s[1] & sra) >> sr);
69-
W[len / 4 + 2] = ((salt_s[1] & sla) << sl) | ((salt_s[2] & sra) >> sr);
70-
W[len / 4 + 3] = (salt_s[2] & sla) << sl;
71-
72-
// The 0x80 ending character was added to the salt before we receive it
73-
74-
W[15] = (len + 10) << 3;
75-
76-
sha1_single(uint, W, output);
77-
78-
for (i = 0; i < 5; i++)
79-
aes_key.w[i] = SWAP32(output[i]);
80-
aes_key.w[5] = 0;
81-
82-
for (i = 0; i < 16; i++)
83-
iv[i] = salt->ct[16 + i];
84-
85-
AES_set_decrypt_key(aes_key.c, 192, &akey);
86-
AES_cbc_decrypt(&salt->ct[32], pt.c, 16, &akey, iv);
87-
88-
result[gid] = (pt.l[1] == 0x0808080808080808UL);
45+
46+
const uint gid = get_global_id(0);
47+
48+
#if NUM_INT_KEYS > 1 && !IS_STATIC_GPU_MASK
49+
const uint ikl = int_key_loc[gid];
50+
const uint loc0 = ikl & 0xff;
51+
#if MASK_FMT_INT_PLHDR > 1
52+
#if LOC_1 >= 0
53+
const uint loc1 = (ikl & 0xff00) >> 8;
54+
#endif
55+
#endif
56+
#if MASK_FMT_INT_PLHDR > 2
57+
#if LOC_2 >= 0
58+
const uint loc2 = (ikl & 0xff0000) >> 16;
59+
#endif
60+
#endif
61+
#if MASK_FMT_INT_PLHDR > 3
62+
#if LOC_3 >= 0
63+
const uint loc3 = (ikl & 0xff000000) >> 24;
64+
#endif
65+
#endif
66+
#endif
67+
68+
#if !IS_STATIC_GPU_MASK
69+
#define GPU_LOC_0 loc0
70+
#define GPU_LOC_1 loc1
71+
#define GPU_LOC_2 loc2
72+
#define GPU_LOC_3 loc3
73+
#else
74+
#define GPU_LOC_0 LOC_0
75+
#define GPU_LOC_1 LOC_1
76+
#define GPU_LOC_2 LOC_2
77+
#define GPU_LOC_3 LOC_3
78+
#endif
79+
80+
const uint base = key_idx[gid];
81+
const uint len = key_idx[gid + 1] - base;
82+
key_buf += base;
83+
84+
const uint shift = len % 4;
85+
const uint sr = 8 * shift;
86+
const uint sl = 32 - sr;
87+
const uint sra = (0xffffffff - (1 << sr)) + 1;
88+
const uint sla = 0xffffffff - sra;
89+
90+
// Endian swap salt
91+
uint salt_be[sizeof(salt->salt) / 4];
92+
for (uint i = 0; i < sizeof(salt->salt) / 4; i++)
93+
salt_be[i] = SWAP32(salt->salt[i]);
94+
95+
for (uint idx = 0; idx < NUM_INT_KEYS; idx++) {
96+
const uint gidx = gid * NUM_INT_KEYS + idx;
97+
uchar password[PLAINTEXT_LENGTH] = { 0 };
98+
uint i;
99+
100+
for (i = 0; i < len; i++)
101+
password[i] = key_buf[i];
102+
103+
#if NUM_INT_KEYS > 1
104+
password[GPU_LOC_0] = (int_keys[idx] & 0xff);
105+
#if MASK_FMT_INT_PLHDR > 1
106+
#if LOC_1 >= 0
107+
password[GPU_LOC_1] = (int_keys[idx] & 0xff >> 8);
108+
#endif
109+
#endif
110+
#if MASK_FMT_INT_PLHDR > 2
111+
#if LOC_2 >= 0
112+
password[GPU_LOC_2] = (int_keys[idx] & 0xff >> 16);
113+
#endif
114+
#endif
115+
#if MASK_FMT_INT_PLHDR > 3
116+
#if LOC_3 >= 0
117+
password[GPU_LOC_3] = (int_keys[idx] & 0xff >> 24);
118+
#endif
119+
#endif
120+
#endif
121+
uint W[16] = { 0 };
122+
for (i = 0; i < PLAINTEXT_LENGTH / 4; i++)
123+
GET_UINT32BE(W[i], password, 4 * i);
124+
125+
// Shift the salt bytes into place after the given key.
126+
W[len / 4] |= (salt_be[0] & sra) >> sr;
127+
W[len / 4 + 1] = ((salt_be[0] & sla) << sl) | ((salt_be[1] & sra) >> sr);
128+
W[len / 4 + 2] = ((salt_be[1] & sla) << sl) | ((salt_be[2] & sra) >> sr);
129+
W[len / 4 + 3] = (salt_be[2] & sla) << sl;
130+
131+
// The Merkel-Damgård 0x80 ending byte was already added to the salt
132+
// on host side, here's the length.
133+
W[15] = (len + 10) << 3;
134+
135+
uint output[160 / 32];
136+
sha1_single(uint, W, output);
137+
138+
union {
139+
uchar c[192 / 8];
140+
uint w[0];
141+
} key;
142+
for (i = 0; i < 5; i++)
143+
key.w[i] = SWAP32(output[i]);
144+
key.w[5] = 0;
145+
146+
uchar iv[16];
147+
148+
AES_set_decrypt_key(key.c, 192, &akey);
149+
150+
if (salt->pw_len) {
151+
const uint blen = (len + 15) / 16;
152+
153+
// Early reject
154+
if (salt->pw_len != blen)
155+
return;
156+
157+
memcpy_cp(iv, salt->ct, 16);
158+
uchar ct[SECRET_LEN];
159+
memcpy_cp(ct, salt->ct + 16, SECRET_LEN);
160+
161+
uchar s_secret[SECRET_LEN];
162+
//AES_set_decrypt_key(key.c, 192, &akey);
163+
AES_cbc_decrypt(ct, s_secret, SECRET_LEN, &akey, iv);
164+
165+
memcpy_cp(iv, salt->csk, 16);
166+
uchar csk[SECRET_LEN];
167+
memcpy_cp(csk, salt->csk + 16, SECRET_LEN);
168+
uchar c_secret[SECRET_LEN];
169+
//AES_set_decrypt_key(key.c, 192, &akey);
170+
AES_cbc_decrypt(csk, c_secret, SECRET_LEN, &akey, iv);
171+
172+
uchar combined_sk[SECRET_LEN];
173+
for (i = 0; i < SECRET_LEN; i++)
174+
combined_sk[i] = s_secret[i] ^ c_secret[i];
175+
176+
uchar final_key[32];
177+
MD5_CTX ctx;
178+
MD5_Init(&ctx);
179+
MD5_Update(&ctx, combined_sk, 16);
180+
MD5_Final(final_key, &ctx);
181+
MD5_Init(&ctx);
182+
MD5_Update(&ctx, combined_sk + 16, 8);
183+
MD5_Final(final_key + 16, &ctx);
184+
185+
memcpy_cp(iv, salt->pw, 16);
186+
187+
uchar pw[PLAINTEXT_LENGTH];
188+
memcpy_cp(pw, salt->pw + 16, PLAINTEXT_LENGTH);
189+
190+
uchar dec_pw[PLAINTEXT_LENGTH + 16];
191+
AES_set_decrypt_key(final_key, 192, &akey);
192+
AES_cbc_decrypt(pw, dec_pw, salt->pw_len * 16, &akey, iv);
193+
194+
if (!memcmp_pp(dec_pw, password, len) &&
195+
check_pkcs_pad(dec_pw, salt->pw_len * 16, AES_BLOCK_SIZE))
196+
out_index[atomic_inc(crack_count_ret)] = gidx;
197+
} else {
198+
union {
199+
uchar c[16];
200+
ulong l[0];
201+
} pt;
202+
203+
memcpy_cp(iv, &salt->ct[16], 16);
204+
uchar ct[16];
205+
memcpy_cp(ct, &salt->ct[32], 16);
206+
//AES_set_decrypt_key(key.c, 192, &akey);
207+
AES_cbc_decrypt(ct, pt.c, 16, &akey, iv);
208+
209+
if (pt.l[1] == 0x0808080808080808UL)
210+
out_index[atomic_inc(crack_count_ret)] = gidx;
211+
}
212+
}
89213
}

Diff for: ‎src/o5logon_common.h

+48
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,48 @@
1+
/*
2+
* This software is Copyright (c) 2025 magnum
3+
* and it is hereby released to the general public under the following terms:
4+
*
5+
* Redistribution and use in source and binary forms, with or without
6+
* modification, are permitted.
7+
*/
8+
9+
#ifndef O5LOGON_COMMON
10+
#define O5LOGON_COMMON
11+
12+
#include <string.h>
13+
14+
#include "arch.h"
15+
#include "formats.h"
16+
#include "misc.h"
17+
#include "common.h"
18+
#include "params.h"
19+
#include "options.h"
20+
21+
#define FORMAT_NAME "Oracle O5LOGON protocol"
22+
#define FORMAT_TAG "$o5logon$"
23+
#define FORMAT_TAG_LEN (sizeof(FORMAT_TAG)-1)
24+
25+
#define BENCHMARK_COMMENT ""
26+
#define BENCHMARK_LENGTH 7
27+
#define PLAINTEXT_LENGTH 32 /* Can't be bumped for OpenCL */
28+
#define CIPHERTEXT_LENGTH 48
29+
#define SALT_LENGTH 10
30+
#define BINARY_SIZE 0
31+
#define BINARY_ALIGN 1
32+
#define SALT_SIZE sizeof(o5logon_salt)
33+
#define SALT_ALIGN sizeof(int32_t)
34+
35+
typedef struct {
36+
unsigned int pw_len; /* AUTH_PASSWORD length (blocks) */
37+
unsigned char salt[(SALT_LENGTH + 1 + 3) / 4 * 4]; /* AUTH_VFR_DATA */
38+
unsigned char ct[CIPHERTEXT_LENGTH]; /* Server's AUTH_SESSKEY */
39+
unsigned char csk[CIPHERTEXT_LENGTH]; /* Client's AUTH_SESSKEY */
40+
unsigned char pw[PLAINTEXT_LENGTH + 16]; /* Client's AUTH_PASSWORD */
41+
} o5logon_salt;
42+
43+
extern struct fmt_tests o5logon_tests[];
44+
45+
extern int o5logon_valid(char *ciphertext, struct fmt_main *self);
46+
extern void *o5logon_get_salt(char *ciphertext);
47+
48+
#endif

0 commit comments

Comments
 (0)