1 /*
2 * Format for cracking blockchain.info "My Wallet" format wallets. Hacked
3 * together during June of 2013 by Dhiru Kholia <dhiru at openwall.com>.
4 *
5 * See https://blockchain.info/wallet/wallet-format
6
7 * This software is Copyright (c) 2012 Lukas Odzioba <ukasz at openwall.net>
8 * and Copyright (c) 2013 Dhiru Kholia <dhiru at openwall.com>, and it is
9 * hereby released to the general public under the following terms:
10 *
11 * Redistribution and use in source and binary forms, with or without
12 * modification, are permitted.
13 *
14 * Improved detection, added iteration count and handle v2 hashes, Feb, 2015, JimF.
15 */
16
17 #ifdef HAVE_OPENCL
18
19 #if FMT_EXTERNS_H
20 extern struct fmt_main fmt_opencl_blockchain;
21 #elif FMT_REGISTERS_H
22 john_register_one(&fmt_opencl_blockchain);
23 #else
24
25 #include <stdint.h>
26 #include <string.h>
27
28 #include "arch.h"
29 #include "formats.h"
30 #include "common.h"
31 #include "jumbo.h"
32 #include "opencl_common.h"
33 #include "options.h"
34 #include "blockchain_common.h"
35
36 #define FORMAT_LABEL "blockchain-opencl"
37 #define FORMAT_NAME "blockchain My Wallet"
38 #define ALGORITHM_NAME "PBKDF2-SHA1 AES OpenCL"
39 #define BENCHMARK_COMMENT " (v2 x5000)"
40 /*
41 * We'd need to be benchmarking for Many vs. Only one salt if we move the v1
42 * test vectors back to the start of the tests array.
43 */
44 #define BENCHMARK_LENGTH 0x107
45 #define MIN_KEYS_PER_CRYPT 1
46 #define MAX_KEYS_PER_CRYPT 1
47 #define BINARY_SIZE 0
48 #define PLAINTEXT_LENGTH 64
49 #define SALT_SIZE sizeof(struct custom_salt)
50 #define BINARY_ALIGN MEM_ALIGN_WORD
51 #define SALT_ALIGN 4
52
53 /* PBKDF2 parameters */
54 #define KEYLEN PLAINTEXT_LENGTH
55 #define OUTLEN 32
56 #define SALTLEN 64
57
58 typedef struct {
59 uint32_t length;
60 uint8_t v[KEYLEN];
61 } pbkdf2_password;
62
63 typedef struct {
64 uint32_t v[(OUTLEN+3)/4];
65 } pbkdf2_out;
66
67 typedef struct {
68 uint32_t iterations;
69 uint32_t outlen;
70 uint32_t skip_bytes;
71 uint8_t length;
72 uint8_t salt[SALTLEN];
73 } pbkdf2_salt;
74
75 typedef struct {
76 uint32_t cracked;
77 } blockchain_out;
78
79 typedef struct {
80 pbkdf2_salt pbkdf2;
81 uint8_t data[SAFETY_FACTOR];
82 uint32_t length;
83 } blockchain_salt;
84
85 static struct custom_salt *cur_salt;
86
87 static cl_int cl_error;
88 static pbkdf2_password *inbuffer;
89 static blockchain_out *output;
90 static blockchain_salt currentsalt;
91 static cl_mem mem_in, mem_dk, mem_salt, mem_out;
92 static struct fmt_main *self;
93
94 static int new_keys;
95
96 static size_t insize, dksize, saltsize, outsize;
97
98 #define STEP 0
99 #define SEED 256
100
101 // This file contains auto-tuning routine(s). Has to be included after formats definitions.
102 #include "opencl_autotune.h"
103
104 static const char * warn[] = {
105 "xfer: ", ", crypt: ", ", xfer: "
106 };
107
108 /* ------- Helper functions ------- */
get_task_max_work_group_size()109 static size_t get_task_max_work_group_size()
110 {
111 return autotune_get_task_max_work_group_size(FALSE, 0, crypt_kernel);
112 }
113
create_clobj(size_t gws,struct fmt_main * self)114 static void create_clobj(size_t gws, struct fmt_main *self)
115 {
116 insize = sizeof(pbkdf2_password) * gws;
117 dksize = sizeof(pbkdf2_out) * gws;
118 saltsize = sizeof(blockchain_salt);
119 outsize = sizeof(blockchain_out) * gws;
120
121 inbuffer = mem_calloc(1, insize);
122 output = mem_alloc(outsize);
123
124 // Allocate memory
125 mem_in =
126 clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, insize, NULL,
127 &cl_error);
128 HANDLE_CLERROR(cl_error, "Error allocating mem in");
129 mem_salt =
130 clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, saltsize, NULL,
131 &cl_error);
132 HANDLE_CLERROR(cl_error, "Error allocating mem salt");
133 mem_dk =
134 clCreateBuffer(context[gpu_id], CL_MEM_READ_WRITE, dksize, NULL,
135 &cl_error);
136 HANDLE_CLERROR(cl_error, "Error allocating pbkdf2 out");
137 mem_out =
138 clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY, outsize, NULL,
139 &cl_error);
140 HANDLE_CLERROR(cl_error, "Error allocating mem out");
141
142 HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 0, sizeof(mem_in),
143 &mem_in), "Error while setting mem_in kernel argument");
144 HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1, sizeof(mem_dk),
145 &mem_dk), "Error while setting mem_dk kernel argument");
146 HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 2, sizeof(mem_salt),
147 &mem_salt), "Error while setting mem_salt kernel argument");
148 HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 3, sizeof(mem_out),
149 &mem_out), "Error while setting mem_out kernel argument");
150 }
151
release_clobj(void)152 static void release_clobj(void)
153 {
154 if (output) {
155 HANDLE_CLERROR(clReleaseMemObject(mem_in), "Release mem in");
156 HANDLE_CLERROR(clReleaseMemObject(mem_dk), "Release mem dk");
157 HANDLE_CLERROR(clReleaseMemObject(mem_salt), "Release mem salt");
158 HANDLE_CLERROR(clReleaseMemObject(mem_out), "Release mem out");
159
160 MEM_FREE(inbuffer);
161 MEM_FREE(output);
162 }
163 }
164
done(void)165 static void done(void)
166 {
167 if (autotuned) {
168 release_clobj();
169
170 HANDLE_CLERROR(clReleaseKernel(crypt_kernel), "Release kernel");
171 HANDLE_CLERROR(clReleaseProgram(program[gpu_id]), "Release Program");
172
173 autotuned--;
174 }
175 }
176
init(struct fmt_main * _self)177 static void init(struct fmt_main *_self)
178 {
179 self = _self;
180 opencl_prepare_dev(gpu_id);
181 }
182
reset(struct db_main * db)183 static void reset(struct db_main *db)
184 {
185 if (!autotuned) {
186 char build_opts[128];
187
188 snprintf(build_opts, sizeof(build_opts),
189 "-DKEYLEN=%d -DSALTLEN=%d -DOUTLEN=%d -DSAFETY_FACTOR=%d",
190 PLAINTEXT_LENGTH, SALTLEN, OUTLEN, SAFETY_FACTOR);
191
192 opencl_init("$JOHN/kernels/blockchain_kernel.cl", gpu_id, build_opts);
193
194 crypt_kernel = clCreateKernel(program[gpu_id], "blockchain", &cl_error);
195 HANDLE_CLERROR(cl_error, "Error creating kernel");
196
197 // Initialize openCL tuning (library) for this format.
198 opencl_init_auto_setup(SEED, 0, NULL, warn, 1, self,
199 create_clobj, release_clobj,
200 sizeof(pbkdf2_password), 0, db);
201
202 // Auto tune execution from shared/included code.
203 autotune_run(self, 1, 0, 1000);
204 }
205 }
206
set_salt(void * salt)207 static void set_salt(void *salt)
208 {
209 cur_salt = (struct custom_salt *)salt;
210
211 memcpy((char*)currentsalt.pbkdf2.salt, cur_salt->data, 16);
212 currentsalt.pbkdf2.length = 16;
213 currentsalt.pbkdf2.iterations = cur_salt->iter;
214 currentsalt.pbkdf2.outlen = 32;
215 currentsalt.pbkdf2.skip_bytes = 0;
216
217 memcpy((char*)currentsalt.data, cur_salt->data, SAFETY_FACTOR);
218 HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_salt,
219 CL_FALSE, 0, saltsize, ¤tsalt, 0, NULL, NULL),
220 "Copy salt to gpu");
221 }
222
set_key(char * key,int index)223 static void set_key(char *key, int index)
224 {
225 uint8_t length = strlen(key);
226
227 inbuffer[index].length = length;
228 memcpy(inbuffer[index].v, key, length);
229
230 new_keys = 1;
231 }
232
get_key(int index)233 static char *get_key(int index)
234 {
235 static char ret[PLAINTEXT_LENGTH + 1];
236 uint8_t length = inbuffer[index].length;
237
238 memcpy(ret, inbuffer[index].v, length);
239 ret[length] = '\0';
240 return ret;
241 }
242
crypt_all(int * pcount,struct db_salt * salt)243 static int crypt_all(int *pcount, struct db_salt *salt)
244 {
245 const int count = *pcount;
246 size_t *lws = local_work_size ? &local_work_size : NULL;
247
248 global_work_size = GET_NEXT_MULTIPLE(count, local_work_size);
249
250 // Copy data to gpu
251 if (ocl_autotune_running || new_keys) {
252 BENCH_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_in, CL_FALSE, 0,
253 insize, inbuffer, 0, NULL, multi_profilingEvent[0]),
254 "Copy data to gpu");
255 new_keys = 0;
256 }
257
258 // Run kernel
259 BENCH_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1,
260 NULL, &global_work_size, lws, 0, NULL,
261 multi_profilingEvent[1]), "Run kernel");
262
263 // Read the result back
264 BENCH_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_out, CL_TRUE, 0,
265 outsize, output, 0, NULL, multi_profilingEvent[2]), "Copy result back");
266
267 return count;
268 }
269
cmp_all(void * binary,int count)270 static int cmp_all(void *binary, int count)
271 {
272 int index;
273
274 for (index = 0; index < count; index++)
275 if (output[index].cracked)
276 return 1;
277 return 0;
278 }
279
cmp_one(void * binary,int index)280 static int cmp_one(void *binary, int index)
281 {
282 return output[index].cracked;
283 }
284
cmp_exact(char * source,int index)285 static int cmp_exact(char *source, int index)
286 {
287 return 1;
288 }
289
290 struct fmt_main fmt_opencl_blockchain = {
291 {
292 FORMAT_LABEL,
293 FORMAT_NAME,
294 ALGORITHM_NAME,
295 BENCHMARK_COMMENT,
296 BENCHMARK_LENGTH,
297 0,
298 PLAINTEXT_LENGTH,
299 BINARY_SIZE,
300 BINARY_ALIGN,
301 SALT_SIZE,
302 SALT_ALIGN,
303 MIN_KEYS_PER_CRYPT,
304 MAX_KEYS_PER_CRYPT,
305 FMT_CASE | FMT_8_BIT | FMT_HUGE_INPUT,
306 /* FIXME: Should report iteration count as a tunable cost */
307 { NULL },
308 { FORMAT_TAG },
309 blockchain_tests
310 }, {
311 init,
312 done,
313 reset,
314 fmt_default_prepare,
315 blockchain_common_valid,
316 fmt_default_split,
317 fmt_default_binary,
318 blockchain_common_get_salt,
319 { NULL },
320 fmt_default_source,
321 {
322 fmt_default_binary_hash
323 },
324 fmt_default_salt_hash,
325 NULL,
326 set_salt,
327 set_key,
328 get_key,
329 fmt_default_clear_keys,
330 crypt_all,
331 {
332 fmt_default_get_hash
333 },
334 cmp_all,
335 cmp_one,
336 cmp_exact
337 }
338 };
339
340 #endif /* plugin stanza */
341
342 #endif /* HAVE_OPENCL */
343