1 /*
2 * This software is Copyright (c) 2014 Sayantan Datta <std2048 at gmail dot com>
3 * and Copyright (c) 2014-2016 magnum
4 * and it is hereby released to the general public under the following terms:
5 * Redistribution and use in source and binary forms, with or without modification, are permitted.
6 * Based on CPU version by Jeff Fay, bartavelle and Solar Designer.
7 */
8
9 #ifdef HAVE_OPENCL
10
11 #if FMT_EXTERNS_H
12 extern struct fmt_main fmt_opencl_1otus5;
13 #elif FMT_REGISTERS_H
14 john_register_one(&fmt_opencl_1otus5);
15 #else
16
17 #include <stdio.h>
18 #include <string.h>
19 #include <assert.h>
20
21 #include "misc.h"
22 #include "formats.h"
23 #include "common.h"
24 #include "opencl_common.h"
25 #include "opencl_lotus5_fmt.h"
26 #include "options.h"
27
28 /*preprocessor constants that John The Ripper likes*/
29 #define FORMAT_LABEL "lotus5-opencl"
30 #define FORMAT_NAME "Lotus Notes/Domino 5"
31 #define ALGORITHM_NAME "OpenCL"
32 #define BENCHMARK_COMMENT ""
33 #define BENCHMARK_LENGTH 0x107
34 #define CIPHERTEXT_LENGTH 32
35 #define SALT_SIZE 0
36 #define BINARY_ALIGN MEM_ALIGN_WORD
37 #define SALT_ALIGN 1
38 #define MIN_KEYS_PER_CRYPT 1
39 #define MAX_KEYS_PER_CRYPT 1
40 #define KEY_SIZE_IN_BYTES sizeof(lotus5_key)
41
42 /*A struct used for JTR's benchmarks*/
43 static struct fmt_tests tests[] = {
44 {"06E0A50B579AD2CD5FFDC48564627EE7", "secret"},
45 {"355E98E7C7B59BD810ED845AD0FD2FC4", "password"},
46 {"CD2D90E8E00D8A2A63A81F531EA8A9A3", "lotus"},
47 {"69D90B46B1AC0912E5CCF858094BBBFC", "dirtydog"},
48 {NULL}
49 };
50
51 /*Some more JTR variables*/
52 static cl_uint *crypt_key;
53 static lotus5_key *saved_key;
54 static struct fmt_main *self;
55
56 static cl_mem cl_tx_keys, cl_tx_binary;
57
58 #define STEP 0
59 #define SEED 256
60
61 // This file contains auto-tuning routine(s). Has to be included after formats definitions.
62 #include "opencl_autotune.h"
63
64 static const char *warn[] = {
65 "xfer: ", ", crypt: ", ", xfer: "
66 };
67
68 /* ------- Helper functions ------- */
get_task_max_work_group_size()69 static size_t get_task_max_work_group_size()
70 {
71 size_t max_lws = get_kernel_max_lws(gpu_id, crypt_kernel);
72
73 if (cpu(device_info[gpu_id]))
74 return get_platform_vendor_id(platform_id) == DEV_INTEL ?
75 max_lws : 1;
76 return max_lws;
77 }
78
create_clobj(size_t gws,struct fmt_main * self)79 static void create_clobj(size_t gws, struct fmt_main *self)
80 {
81 size_t mem_alloc_sz;
82
83 mem_alloc_sz = KEY_SIZE_IN_BYTES * gws;
84 cl_tx_keys = clCreateBuffer(context[gpu_id],
85 CL_MEM_READ_ONLY,
86 mem_alloc_sz, NULL, &ret_code);
87 HANDLE_CLERROR(ret_code, "Failed to create buffer cl_tx_keys.");
88
89 mem_alloc_sz = BINARY_SIZE * gws;
90 cl_tx_binary = clCreateBuffer(context[gpu_id],
91 CL_MEM_WRITE_ONLY,
92 mem_alloc_sz, NULL, &ret_code);
93 HANDLE_CLERROR(ret_code, "Failed to create buffer cl_tx_binary.");
94
95 HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 0,
96 sizeof(cl_mem), &cl_tx_keys),
97 "Failed to set kernel argument 0, cl_tx_keys.");
98 HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1,
99 sizeof(cl_mem), &cl_tx_binary),
100 "Failed to set kernel argument 1, cl_tx_binary.");
101
102 crypt_key = mem_calloc(gws, BINARY_SIZE);
103 saved_key = mem_calloc(gws, KEY_SIZE_IN_BYTES);
104 }
105
release_clobj(void)106 static void release_clobj(void)
107 {
108 if (crypt_key) {
109 HANDLE_CLERROR(clReleaseMemObject(cl_tx_keys),
110 "Failed to release buffer cl_tx_keys.");
111 HANDLE_CLERROR(clReleaseMemObject(cl_tx_binary),
112 "Failed to release buffer cl_tx_binary.");
113
114 MEM_FREE(saved_key);
115 MEM_FREE(crypt_key);
116 }
117 }
118
init(struct fmt_main * _self)119 static void init(struct fmt_main *_self)
120 {
121 self = _self;
122 opencl_prepare_dev(gpu_id);
123 }
124
reset(struct db_main * db)125 static void reset(struct db_main *db)
126 {
127 if (!autotuned) {
128 size_t gws_limit;
129
130 opencl_init("$JOHN/kernels/lotus5_kernel.cl", gpu_id, NULL);
131
132 crypt_kernel = clCreateKernel(program[gpu_id], "lotus5", &ret_code);
133 HANDLE_CLERROR(ret_code, "Failed to create kernel lotus5.");
134
135 gws_limit = get_max_mem_alloc_size(gpu_id) / KEY_SIZE_IN_BYTES;
136
137 if (gws_limit & (gws_limit - 1)) {
138 get_power_of_two(gws_limit);
139 gws_limit >>= 1;
140 }
141
142 // Initialize openCL tuning (library) for this format.
143 opencl_init_auto_setup(SEED, 0, NULL, warn, 1, self,
144 create_clobj, release_clobj,
145 KEY_SIZE_IN_BYTES, gws_limit, db);
146
147 // Auto tune execution from shared/included code.
148 autotune_run_extra(self, 1, gws_limit, 200, CL_TRUE);
149 }
150 }
151
done(void)152 static void done(void)
153 {
154 if (autotuned) {
155 release_clobj();
156 HANDLE_CLERROR(clReleaseKernel(crypt_kernel),
157 "Release kernel lotus5.");
158 HANDLE_CLERROR(clReleaseProgram(program[gpu_id]),
159 "Release Program");
160
161 autotuned--;
162 }
163 }
164
165 /*Utility function to convert hex to bin */
get_binary(char * ciphertext)166 static void *get_binary(char *ciphertext)
167 {
168 static char realcipher[BINARY_SIZE];
169 int i;
170 for (i = 0; i < BINARY_SIZE; i++)
171 realcipher[i] = atoi16[ARCH_INDEX(ciphertext[i*2])]*16 + atoi16[ARCH_INDEX(ciphertext[i*2+1])];
172 return ((void *) realcipher);
173 }
174
175 /*Another function required by JTR: decides whether we have a valid
176 * ciphertext */
valid(char * ciphertext,struct fmt_main * self)177 static int valid (char *ciphertext, struct fmt_main *self)
178 {
179 int i;
180
181 for (i = 0; i < CIPHERTEXT_LENGTH; i++)
182 if (!(((ciphertext[i] >= '0') && (ciphertext[i] <= '9'))
183 || ((ciphertext[i] >= 'A') && (ciphertext[i] <= 'F'))))
184 {
185 return 0;
186 }
187 return !ciphertext[i];
188 }
189
190 /*sets the value of saved_key so we can play with it*/
set_key(char * key,int index)191 static void set_key (char *key, int index)
192 {
193 int len = strlen(key);
194
195 memset(saved_key[index].v.c, 0, PLAINTEXT_LENGTH);
196 memcpy(saved_key[index].v.c, key, len);
197 saved_key[index].l = len;
198 }
199
200 /*retrieves the saved key; used by JTR*/
get_key(int index)201 static char *get_key (int index)
202 {
203 static char out[PLAINTEXT_LENGTH + 1];
204 int len = saved_key[index].l;
205
206 memcpy(out, saved_key[index].v.c, len);
207 out[len] = 0;
208
209 return out;
210 }
211
cmp_all(void * binary,int count)212 static int cmp_all (void *binary, int count)
213 {
214 int index;
215 for (index = 0; index < count; index++)
216 if (!memcmp(binary, crypt_key + index * BINARY_SIZE_IN_uint32_t, BINARY_SIZE))
217 return 1;
218 return 0;
219 }
220
cmp_one(void * binary,int index)221 static int cmp_one (void *binary, int index)
222 {
223 return !memcmp(binary, crypt_key + index * BINARY_SIZE_IN_uint32_t, BINARY_SIZE);
224 }
225
cmp_exact(char * source,int index)226 static int cmp_exact (char *source, int index)
227 {
228 return 1;
229 }
230
231 /*the last public function; generates ciphertext*/
crypt_all(int * pcount,struct db_salt * salt)232 static int crypt_all(int *pcount, struct db_salt *salt)
233 {
234 const int count = *pcount;
235 size_t mem_cpy_sz;
236 size_t N, *M;
237
238 mem_cpy_sz = count * KEY_SIZE_IN_BYTES;
239 BENCH_CLERROR(clEnqueueWriteBuffer(queue[gpu_id],
240 cl_tx_keys, CL_FALSE, 0,
241 mem_cpy_sz, saved_key,
242 0, NULL, multi_profilingEvent[0]),
243 "Failed to write buffer cl_tx_keys.");
244
245 M = local_work_size ? &local_work_size : NULL;
246 N = GET_NEXT_MULTIPLE(count, local_work_size);
247
248 BENCH_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id],
249 crypt_kernel, 1,
250 NULL, &N, M, 0, NULL, multi_profilingEvent[1]),
251 "Failed to enqueue kernel lotus5.");
252
253 mem_cpy_sz = count * BINARY_SIZE;
254 BENCH_CLERROR(clEnqueueReadBuffer(queue[gpu_id],
255 cl_tx_binary, CL_TRUE, 0,
256 mem_cpy_sz, crypt_key, 0,
257 NULL, multi_profilingEvent[2]),
258 "Failed to read buffer cl_tx_binary.");
259
260 return count;
261 }
262
get_hash_0(int index)263 static int get_hash_0(int index) { return crypt_key[index * BINARY_SIZE_IN_uint32_t] & PH_MASK_0; }
get_hash_1(int index)264 static int get_hash_1(int index) { return crypt_key[index * BINARY_SIZE_IN_uint32_t] & PH_MASK_1; }
get_hash_2(int index)265 static int get_hash_2(int index) { return crypt_key[index * BINARY_SIZE_IN_uint32_t] & PH_MASK_2; }
get_hash_3(int index)266 static int get_hash_3(int index) { return crypt_key[index * BINARY_SIZE_IN_uint32_t] & PH_MASK_3; }
get_hash_4(int index)267 static int get_hash_4(int index) { return crypt_key[index * BINARY_SIZE_IN_uint32_t] & PH_MASK_4; }
get_hash_5(int index)268 static int get_hash_5(int index) { return crypt_key[index * BINARY_SIZE_IN_uint32_t] & PH_MASK_5; }
get_hash_6(int index)269 static int get_hash_6(int index) { return crypt_key[index * BINARY_SIZE_IN_uint32_t] & PH_MASK_6; }
270
271 /* C's version of a class specifier */
272 struct fmt_main fmt_opencl_1otus5 = {
273 {
274 FORMAT_LABEL,
275 FORMAT_NAME,
276 ALGORITHM_NAME,
277 BENCHMARK_COMMENT,
278 BENCHMARK_LENGTH,
279 0,
280 PLAINTEXT_LENGTH,
281 BINARY_SIZE,
282 BINARY_ALIGN,
283 SALT_SIZE,
284 SALT_ALIGN,
285 MIN_KEYS_PER_CRYPT,
286 MAX_KEYS_PER_CRYPT,
287 FMT_CASE | FMT_8_BIT,
288 { NULL },
289 { NULL },
290 tests
291 }, {
292 init,
293 done,
294 reset,
295 fmt_default_prepare,
296 valid,
297 fmt_default_split,
298 get_binary,
299 fmt_default_salt,
300 { NULL },
301 fmt_default_source,
302 {
303 fmt_default_binary_hash_0,
304 fmt_default_binary_hash_1,
305 fmt_default_binary_hash_2,
306 fmt_default_binary_hash_3,
307 fmt_default_binary_hash_4,
308 fmt_default_binary_hash_5,
309 fmt_default_binary_hash_6
310 },
311 fmt_default_salt_hash,
312 NULL,
313 fmt_default_set_salt,
314 set_key,
315 get_key,
316 fmt_default_clear_keys,
317 crypt_all,
318 {
319 get_hash_0,
320 get_hash_1,
321 get_hash_2,
322 get_hash_3,
323 get_hash_4,
324 get_hash_5,
325 get_hash_6
326 },
327 cmp_all,
328 cmp_one,
329 cmp_exact
330 }
331 };
332
333 #endif /* plugin stanza */
334
335 #endif /* HAVE_OPENCL */
336