1 /*
2 * This software is Copyright (c) 2018 Dhiru Kholia <kholia at kth dot se> and
3 * 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 * Based on opencl_pfx_fmt_plug.c file, and other files which are,
9 *
10 * Copyright (c) 2012 Lukas Odzioba <ukasz@openwall.net>, Copyright (c) JimF,
11 * and Copyright (c) magnum.
12 */
13
14 #ifdef HAVE_OPENCL
15
16 #if FMT_EXTERNS_H
17 extern struct fmt_main fmt_opencl_ssh;
18 #elif FMT_REGISTERS_H
19 john_register_one(&fmt_opencl_ssh);
20 #else
21
22 #include <stdint.h>
23 #include <string.h>
24
25 #include "misc.h"
26 #include "arch.h"
27 #include "params.h"
28 #include "common.h"
29 #include "formats.h"
30 #include "opencl_common.h"
31 #include "options.h"
32 #include "ssh_common.h"
33 #include "ssh_variable_code.h"
34
35 #define FORMAT_LABEL "ssh-opencl"
36 #define FORMAT_NAME ""
37 #define ALGORITHM_NAME "RSA/DSA/EC (SSH private keys) OpenCL"
38 #define BENCHMARK_COMMENT ""
39 #define BENCHMARK_LENGTH 0x107
40 #define BINARY_SIZE 0
41 #define BINARY_ALIGN sizeof(uint32_t)
42 #define SALT_SIZE sizeof(*cur_salt)
43 #define SALT_ALIGN sizeof(int)
44 #define PLAINTEXT_LENGTH 32
45 #define MIN_KEYS_PER_CRYPT 1
46 #define MAX_KEYS_PER_CRYPT 1
47
48 // input
49 typedef struct {
50 uint32_t length;
51 unsigned char v[PLAINTEXT_LENGTH];
52 } ssh_password;
53
54 typedef struct {
55 uint32_t cracked;
56 } ssh_out;
57
58 // input
59 typedef struct {
60 unsigned char salt[16];
61 unsigned char ct[N];
62 int cipher;
63 int ctl;
64 int sl;
65 int rounds;
66 int ciphertext_begin_offset;
67 } ssh_salt;
68
69 static ssh_out *output;
70 static struct custom_salt *cur_salt;
71 static cl_int cl_error;
72 static ssh_password *inbuffer;
73 static ssh_salt currentsalt;
74 static cl_mem mem_in, mem_out, mem_setting;
75 static struct fmt_main *self;
76
77 static size_t insize, outsize, settingsize;
78
79 #define STEP 0
80 #define SEED 256
81
82 // This file contains auto-tuning routine(s). Has to be included after formats definitions.
83 #include "opencl_autotune.h"
84
85 static const char *warn[] = {
86 "xfer: ", ", crypt: ", ", xfer: "
87 };
88
89 /* ------- Helper functions ------- */
get_task_max_work_group_size()90 static size_t get_task_max_work_group_size()
91 {
92 return autotune_get_task_max_work_group_size(FALSE, 0, crypt_kernel);
93 }
94
create_clobj(size_t gws,struct fmt_main * self)95 static void create_clobj(size_t gws, struct fmt_main *self)
96 {
97 insize = sizeof(ssh_password) * gws;
98 outsize = sizeof(ssh_out) * gws;
99 settingsize = sizeof(ssh_salt);
100
101 inbuffer = mem_calloc(1, insize);
102 output = mem_alloc(outsize);
103
104 // Allocate memory
105 mem_in =
106 clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, insize, NULL,
107 &cl_error);
108 HANDLE_CLERROR(cl_error, "Error allocating mem in");
109 mem_out =
110 clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY, outsize, NULL,
111 &cl_error);
112 HANDLE_CLERROR(cl_error, "Error allocating mem out");
113 mem_setting =
114 clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, settingsize,
115 NULL, &cl_error);
116 HANDLE_CLERROR(cl_error, "Error allocating mem setting");
117
118 HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 0, sizeof(mem_in),
119 &mem_in), "Error while setting mem_in kernel argument");
120 HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1, sizeof(mem_out),
121 &mem_out), "Error while setting mem_out kernel argument");
122 HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 2, sizeof(mem_setting),
123 &mem_setting), "Error while setting mem_salt kernel argument");
124 }
125
release_clobj(void)126 static void release_clobj(void)
127 {
128 if (output) {
129 HANDLE_CLERROR(clReleaseMemObject(mem_in), "Release mem in");
130 HANDLE_CLERROR(clReleaseMemObject(mem_setting), "Release mem setting");
131 HANDLE_CLERROR(clReleaseMemObject(mem_out), "Release mem out");
132
133 MEM_FREE(inbuffer);
134 MEM_FREE(output);
135 }
136 }
137
init(struct fmt_main * _self)138 static void init(struct fmt_main *_self)
139 {
140 self = _self;
141 opencl_prepare_dev(gpu_id);
142 }
143
reset(struct db_main * db)144 static void reset(struct db_main *db)
145 {
146 if (!autotuned) {
147 char build_opts[64];
148
149 snprintf(build_opts, sizeof(build_opts),
150 "-DPLAINTEXT_LENGTH=%d -DCTLEN=%d -DSAFETY_FACTOR=%d",
151 PLAINTEXT_LENGTH, N, 16);
152 opencl_init("$JOHN/kernels/ssh_kernel.cl",
153 gpu_id, build_opts);
154
155 crypt_kernel = clCreateKernel(program[gpu_id], "ssh", &cl_error);
156 HANDLE_CLERROR(cl_error, "Error creating kernel");
157
158 // Initialize openCL tuning (library) for this format.
159 opencl_init_auto_setup(SEED, 0, NULL, warn, 1, self,
160 create_clobj, release_clobj,
161 sizeof(ssh_password), 0, db);
162
163 // Auto tune execution from shared/included code.
164 autotune_run(self, 1, 0, 500);
165 }
166 }
167
done(void)168 static void done(void)
169 {
170 if (autotuned) {
171 release_clobj();
172
173 HANDLE_CLERROR(clReleaseKernel(crypt_kernel), "Release kernel");
174 HANDLE_CLERROR(clReleaseProgram(program[gpu_id]), "Release Program");
175
176 autotuned--;
177 }
178 }
179
set_salt(void * salt)180 static void set_salt(void *salt)
181 {
182 cur_salt = (struct custom_salt*)salt;
183
184 currentsalt.rounds = cur_salt->rounds;
185 currentsalt.cipher = cur_salt->cipher;
186 currentsalt.sl = cur_salt->sl;
187 currentsalt.ctl = cur_salt->ctl;
188 currentsalt.ciphertext_begin_offset = cur_salt->ciphertext_begin_offset;
189
190 memcpy((char*)currentsalt.salt, cur_salt->salt, currentsalt.sl);
191 memcpy((char*)currentsalt.ct, cur_salt->ct, currentsalt.ctl);
192
193 HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_setting,
194 CL_FALSE, 0, settingsize, ¤tsalt, 0, NULL, NULL),
195 "Copy setting to gpu");
196 }
197
ssh_set_key(char * key,int index)198 static void ssh_set_key(char *key, int index)
199 {
200 uint32_t length = strlen(key);
201
202 if (length > PLAINTEXT_LENGTH)
203 length = PLAINTEXT_LENGTH;
204 inbuffer[index].length = length;
205 memcpy(inbuffer[index].v, key, length);
206 }
207
get_key(int index)208 static char *get_key(int index)
209 {
210 static char ret[PLAINTEXT_LENGTH + 1];
211 uint32_t length = inbuffer[index].length;
212
213 memcpy(ret, inbuffer[index].v, length);
214 ret[length] = '\0';
215
216 return ret;
217 }
218
crypt_all(int * pcount,struct db_salt * salt)219 static int crypt_all(int *pcount, struct db_salt *salt)
220 {
221 const int count = *pcount;
222 size_t *lws = local_work_size ? &local_work_size : NULL;
223 size_t gws = GET_NEXT_MULTIPLE(count, local_work_size);
224
225 // Copy data to gpu
226 BENCH_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_in, CL_FALSE, 0,
227 insize, inbuffer, 0, NULL, multi_profilingEvent[0]),
228 "Copy data to gpu");
229
230 // Run kernel
231 BENCH_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1,
232 NULL, &gws, lws, 0, NULL,
233 multi_profilingEvent[1]),
234 "Run kernel");
235
236 // Read the result back
237 BENCH_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_out, CL_TRUE, 0, outsize, output, 0, NULL, multi_profilingEvent[5]), "Copy result back");
238
239 return count;
240 }
241
cmp_all(void * binary,int count)242 static int cmp_all(void *binary, int count)
243 {
244 int index;
245
246 for (index = 0; index < count; index++)
247 if (output[index].cracked)
248 return 1;
249 return 0;
250 }
251
cmp_one(void * binary,int index)252 static int cmp_one(void *binary, int index)
253 {
254 return output[index].cracked;
255 }
256
cmp_exact(char * source,int index)257 static int cmp_exact(char *source, int index)
258 {
259 return 1;
260 }
261
262 struct fmt_main fmt_opencl_ssh = {
263 {
264 FORMAT_LABEL,
265 FORMAT_NAME,
266 ALGORITHM_NAME,
267 BENCHMARK_COMMENT,
268 BENCHMARK_LENGTH,
269 0,
270 PLAINTEXT_LENGTH,
271 BINARY_SIZE,
272 BINARY_ALIGN,
273 SALT_SIZE,
274 SALT_ALIGN,
275 MIN_KEYS_PER_CRYPT,
276 MAX_KEYS_PER_CRYPT,
277 FMT_CASE | FMT_8_BIT | FMT_HUGE_INPUT,
278 {
279 "KDF/cipher [0=MD5/AES 1=MD5/3DES 2=Bcrypt/AES]",
280 "iteration count",
281 },
282 { FORMAT_TAG },
283 ssh_tests
284 }, {
285 init,
286 done,
287 reset,
288 fmt_default_prepare,
289 ssh_valid,
290 fmt_default_split,
291 fmt_default_binary,
292 ssh_get_salt,
293 {
294 ssh_kdf,
295 ssh_iteration_count,
296 },
297 fmt_default_source,
298 {
299 fmt_default_binary_hash
300 },
301 fmt_default_salt_hash,
302 NULL,
303 set_salt,
304 ssh_set_key,
305 get_key,
306 fmt_default_clear_keys,
307 crypt_all,
308 {
309 fmt_default_get_hash
310 },
311 cmp_all,
312 cmp_one,
313 cmp_exact
314 }
315 };
316
317 #endif /* plugin stanza */
318
319 #endif /* HAVE_OPENCL */
320