Skip to content

Commit bb37248

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

File tree

5 files changed

+636
-461
lines changed

5 files changed

+636
-461
lines changed

Diff for: ‎run/opencl/o5logon_kernel.cl

+203-74
Original file line numberDiff line numberDiff line change
@@ -1,89 +1,218 @@
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+
* 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 (c) 2012-2025, magnum
6+
* Copyright (c) 2014 Harrison Neal
7+
* Copyright (c) 2011 by Samuele Giovanni Tonon, samu at linuxasylum dot net
8+
*
9+
* This program comes with ABSOLUTELY NO WARRANTY; express or
10+
* implied .
11+
* This is free software, and you are welcome to redistribute it
12+
* under certain conditions; as expressed here
13+
* http://www.gnu.org/licenses/gpl-2.0.html
1814
*/
1915

2016
#include "opencl_device_info.h"
2117
#include "opencl_misc.h"
2218
#include "opencl_sha1.h"
23-
#define AES_SRC_TYPE __constant
19+
#include "opencl_md5_ctx.h"
20+
#include "opencl_mask.h"
21+
22+
// Workaround for UHD Graphics 630 version "1.2(Jan 10 2025 21:16:54)"
23+
#if (__OS_X__ && gpu_intel(DEVICE_INFO))
24+
#define AES_BITSLICE 1
25+
#endif
2426
#include "opencl_aes.h"
2527

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

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

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)