STRLCPY/btcrecover

add GPU debugging output for #104
Christopher Gurnee committed 7 years ago

e0b170b5

1 parent 4fab004f

Revision indexing in progress... (symbol navigation in revisions will be accurate after indexed)

■ ■ ■ ■ ■ ■

btcrecover/btcrpass.py

		skipped 28 lines
29	29		# (all optional futures for 2.7)
30	30		from __future__ import print_function, absolute_import, division, unicode_literals
31	31
32		-	__version__ = "0.17.5"
	32	+	__version__ = "0.17.5-debuggpu-0"
33	33		__ordering_version__ = b"0.6.4" # must be updated whenever password ordering changes
34	34
35	35		import sys, argparse, itertools, string, re, multiprocessing, signal, os, cPickle, gc, \
		skipped 811 lines
847	847		# many iterations should be done at a time based on iter_count and the requested int_rate,
848	848		# rounding up to maximize the number of iterations done in the last set to optimize performance
849	849		assert hasattr(self, "_iter_count") and self._iter_count, "WalletBitcoinCore.init_opencl_kernel: bitcoin core wallet or mkey has been loaded"
	850	+	self._iter_count = 3
850	851		self._iter_count_chunksize = self._iter_count // int_rate or 1
851	852		if self._iter_count_chunksize % int_rate != 0: # if not evenly divisible,
852	853		self._iter_count_chunksize += 1 # then round up
		skipped 4489 lines
5342	5343
5343	5344		# Measure the performance of the verification function
5344	5345		# (for CPU, run for about 0.5s; for GPU, run for one global-worksize chunk)
5345		-	if args.performance and args.enable_gpu: # skip this time-consuming & unnecessary measurement in this case
	5346	+	if True or args.performance and args.enable_gpu: # skip this time-consuming & unnecessary measurement in this case
5346	5347		est_secs_per_password = 0.01 # set this to something relatively big, it doesn't matter exactly what
5347	5348		else:
5348	5349		if args.enable_gpu:
		skipped 222 lines

■ ■ ■ ■ ■ ■

btcrecover/sha512-bc-kernel.cl

		skipped 179 lines
180	180		};
181	181
182	182
	183	+	#define PFOUT8 "%16lx %16lx %16lx %16lx %16lx %16lx %16lx %16lx"
	184	+	#define PFLETTERS a, b, c, d, e, f, g, h
	185	+
	186	+	#define PFOUT16 PFOUT8 " " PFOUT8
	187	+	#define PFWS w[0], w[1], w[2], w[3], w[4], w[5], w[6], w[7], w[8], w[9], w[10], w[11], w[12], w[13], w[14], w[15]
	188	+
183	189		// From sha512_kernel.cl
184	190
185	191		__kernel
		skipped 6 lines
192	198
193	199		// Get location of the hash to work on for this kernel
194	200		hashes_buffer += (get_global_id(0) << 3);
	201	+
	202	+	printf("AA: " PFOUT8 "\n", hashes_buffer[0], hashes_buffer[1], hashes_buffer[2], hashes_buffer[3], hashes_buffer[4], hashes_buffer[5], hashes_buffer[6], hashes_buffer[7]);
195	203
196	204		// Copy initial hash into local input variable and convert endianness
197	205		#pragma unroll
		skipped 19 lines
217	225		g = H6;
218	226		h = H7;
219	227
	228	+	printf("BB: " PFOUT16 "\n", PFWS);
	229	+	printf("BB: " PFOUT8 "\n", PFLETTERS);
	230	+
220	231		#pragma unroll
221	232		for (int i = 0; i < 16; i++) {
222	233		t = k[i] + w[i] + h + Sigma1(e) + Ch(e, f, g);
		skipped 7 lines
230	241		c = b;
231	242		b = a;
232	243		a = t;
	244	+
	245	+	printf("%2d: %16lx\n", i, t);
	246	+	printf("%2d: " PFOUT8 "\n", i, PFLETTERS);
233	247		}
234	248
235	249		#pragma unroll
		skipped 10 lines
246	260		c = b;
247	261		b = a;
248	262		a = t;
	263	+
	264	+	printf("%2d: %16lx\n", i, t);
	265	+	printf("%2d: " PFOUT16 "\n", i, PFWS);
	266	+	printf("%2d: " PFOUT8 "\n", i, PFLETTERS);
249	267		}
250	268
251	269		// Copy resulting SHA512 hash back into the local input variable
		skipped 12 lines
264	282		for (int i = 9; i < 15; i++)
265	283		w[i] = 0;
266	284		w[15] = 512; // The length in bits
	285	+
	286	+	printf("ZZ: " PFOUT16 "\n", PFWS);
267	287		}
268	288
269	289		// Copy iterated SHA512 hash into the I/O buffer and convert endianness
		skipped 5 lines

add GPU debugging output for #104