void MFNHashFileSaltedProtobuf_SaltedHash::MergeFrom(const MFNHashFileSaltedProtobuf_SaltedHash& from) {
GOOGLE_CHECK_NE(&from, this);
if (from._has_bits_[0 / 32] & (0xffu << (0 % 32))) {
if (from.has_hash()) {
set_hash(from.hash());
}
if (from.has_salt()) {
set_salt(from.salt());
}
if (from.has_iteration_count()) {
set_iteration_count(from.iteration_count());
}
if (from.has_other_data_1()) {
set_other_data_1(from.other_data_1());
}
if (from.has_other_data_2()) {
set_other_data_2(from.other_data_2());
}
if (from.has_other_data_3()) {
set_other_data_3(from.other_data_3());
}
if (from.has_other_data_4()) {
set_other_data_4(from.other_data_4());
}
if (from.has_other_data_5()) {
set_other_data_5(from.other_data_5());
}
}
mutable_unknown_fields()->MergeFrom(from.unknown_fields());
}
static cl_ulong gws_test(int gws, int do_benchmark, struct fmt_main *self)
{
cl_ulong startTime, endTime;
cl_command_queue queue_prof;
cl_event Event[6];
cl_int ret_code;
int i;
size_t scalar_gws = VF * gws;
create_clobj(gws, self);
queue_prof = clCreateCommandQueue(context[ocl_gpu_id], devices[ocl_gpu_id], CL_QUEUE_PROFILING_ENABLE, &ret_code);
for (i = 0; i < scalar_gws; i++)
set_key(tests[0].plaintext, i);
set_salt(get_salt(tests[0].ciphertext));
HANDLE_CLERROR(clEnqueueWriteBuffer(queue_prof, cl_saved_key, CL_TRUE, 0, UNICODE_LENGTH * scalar_gws, saved_key, 0, NULL, &Event[0]), "Failed transferring keys");
HANDLE_CLERROR(clEnqueueWriteBuffer(queue_prof, cl_saved_len, CL_TRUE, 0, sizeof(int) * scalar_gws, saved_len, 0, NULL, &Event[1]), "Failed transferring lengths");
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue_prof, GenerateSHA1pwhash, 1, NULL, &scalar_gws, &local_work_size, 0, NULL, &Event[2]), "running kernel");
for (i = 0; i < 50000 / HASH_LOOPS - 1; i++)
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue_prof, crypt_kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL), "running kernel");
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue_prof, crypt_kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, &Event[3]), "running kernel");
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue_prof, Generate2007key, 1, NULL, &global_work_size, &local_work_size, 0, NULL, &Event[4]), "running kernel");
HANDLE_CLERROR(clEnqueueReadBuffer(queue_prof, cl_key, CL_TRUE, 0, 16 * scalar_gws, key, 0, NULL, &Event[5]), "failed in reading key back");
#if 0
HANDLE_CLERROR(clGetEventProfilingInfo(Event[2],
CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &startTime,
NULL), "Failed to get profiling info");
HANDLE_CLERROR(clGetEventProfilingInfo(Event[2],
CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &endTime,
NULL), "Failed to get profiling info");
fprintf(stderr, "GenerateSHA1pwhash kernel duration: %llu us, ", (endTime-startTime)/1000ULL);
#endif
HANDLE_CLERROR(clGetEventProfilingInfo(Event[3],
CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &startTime,
NULL), "Failed to get profiling info");
HANDLE_CLERROR(clGetEventProfilingInfo(Event[3],
CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &endTime,
NULL), "Failed to get profiling info");
if (do_benchmark)
fprintf(stderr, "%.2f ms x %u = %.2f s\t", (float)((endTime - startTime)/1000000.), 50000/HASH_LOOPS, (float)(50000/HASH_LOOPS) * (endTime - startTime) / 1000000000.);
/* 200 ms duration limit for GCN to avoid ASIC hangs */
if (amd_gcn(device_info[ocl_gpu_id]) && endTime - startTime > 200000000) {
if (do_benchmark)
fprintf(stderr, "- exceeds 200 ms\n");
clReleaseCommandQueue(queue_prof);
release_clobj();
return 0;
}
#if 0
HANDLE_CLERROR(clGetEventProfilingInfo(Event[4],
CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &startTime,
NULL), "Failed to get profiling info");
HANDLE_CLERROR(clGetEventProfilingInfo(Event[4],
CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &endTime,
NULL), "Failed to get profiling info");
fprintf(stderr, "Generate2007key kernel duration: %llu us\n", (endTime-startTime)/1000ULL);
#endif
HANDLE_CLERROR(clGetEventProfilingInfo(Event[0],
CL_PROFILING_COMMAND_SUBMIT, sizeof(cl_ulong), &startTime,
NULL), "Failed to get profiling info");
HANDLE_CLERROR(clGetEventProfilingInfo(Event[5],
CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &endTime,
NULL), "Failed to get profiling info");
clReleaseCommandQueue(queue_prof);
release_clobj();
return (endTime - startTime);
}
//Do the proper test using different sizes.
static cl_ulong gws_test(size_t num, struct fmt_main * self) {
cl_event myEvent;
cl_int ret_code;
cl_uint *tmpbuffer;
cl_ulong startTime, endTime, runtime;
int i, loops;
//Prepare buffers.
create_clobj(num, self);
tmpbuffer = mem_alloc(sizeof(sha512_hash) * num);
if (tmpbuffer == NULL) {
fprintf(stderr, "Malloc failure in find_best_gws\n");
exit(EXIT_FAILURE);
}
queue_prof = clCreateCommandQueue(context[ocl_gpu_id], devices[ocl_gpu_id],
CL_QUEUE_PROFILING_ENABLE, &ret_code);
HANDLE_CLERROR(ret_code, "Failed in clCreateCommandQueue");
// Set salt.
set_salt(get_salt("$6$saltstring$"));
salt->initial = salt->rounds - get_multiple(salt->rounds, HASH_LOOPS);
// Set keys
for (i = 0; i < num; i++) {
set_key("aaabaabaaa", i);
}
//** Get execution time **//
HANDLE_CLERROR(clEnqueueWriteBuffer(queue_prof, salt_buffer, CL_FALSE, 0,
sizeof(sha512_salt), salt, 0, NULL, &myEvent),
"Failed in clEnqueueWriteBuffer");
HANDLE_CLERROR(clFinish(queue_prof), "Failed in clFinish");
HANDLE_CLERROR(clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_SUBMIT,
sizeof(cl_ulong), &startTime, NULL),
"Failed in clGetEventProfilingInfo I");
HANDLE_CLERROR(clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_END,
sizeof(cl_ulong), &endTime, NULL),
"Failed in clGetEventProfilingInfo II");
HANDLE_CLERROR(clReleaseEvent(myEvent), "Failed in clReleaseEvent");
runtime = endTime - startTime;
//** Get execution time **//
HANDLE_CLERROR(clEnqueueWriteBuffer(queue_prof, pass_buffer, CL_FALSE, 0,
sizeof(sha512_password) * num, plaintext, 0, NULL, &myEvent),
"Failed in clEnqueueWriteBuffer");
HANDLE_CLERROR(clFinish(queue_prof), "Failed in clFinish");
HANDLE_CLERROR(clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_SUBMIT,
sizeof(cl_ulong), &startTime, NULL),
"Failed in clGetEventProfilingInfo I");
HANDLE_CLERROR(clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_END,
sizeof(cl_ulong), &endTime, NULL),
"Failed in clGetEventProfilingInfo II");
HANDLE_CLERROR(clReleaseEvent(myEvent), "Failed in clReleaseEvent");
runtime += endTime - startTime;
//** Get execution time **//
if (gpu(source_in_use) || use_local(source_in_use)) {
ret_code = clEnqueueNDRangeKernel(queue_prof, prepare_kernel,
1, NULL, &num, &local_work_size, 0, NULL, &myEvent);
HANDLE_CLERROR(clFinish(queue_prof), "Failed in clFinish");
HANDLE_CLERROR(clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_SUBMIT,
sizeof(cl_ulong), &startTime, NULL),
"Failed in clGetEventProfilingInfo I");
HANDLE_CLERROR(clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_END,
sizeof(cl_ulong), &endTime, NULL),
"Failed in clGetEventProfilingInfo II");
HANDLE_CLERROR(clReleaseEvent(myEvent), "Failed in clReleaseEvent");
runtime += endTime - startTime;
}
loops = gpu(source_in_use) || use_local(source_in_use) ? (salt->rounds / HASH_LOOPS) : 1;
//** Get execution time **//
for (i = 0; i < loops; i++)
{
ret_code = clEnqueueNDRangeKernel(queue_prof, crypt_kernel,
1, NULL, &num, &local_work_size, 0, NULL, &myEvent);
HANDLE_CLERROR(clFinish(queue_prof), "Failed in clFinish");
HANDLE_CLERROR(clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_SUBMIT,
sizeof(cl_ulong), &startTime, NULL),
"Failed in clGetEventProfilingInfo I");
HANDLE_CLERROR(clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_END,
sizeof(cl_ulong), &endTime, NULL),
"Failed in clGetEventProfilingInfo II");
HANDLE_CLERROR(clReleaseEvent(myEvent), "Failed in clReleaseEvent");
runtime += endTime - startTime;
}
//** Get execution time **//
HANDLE_CLERROR(clEnqueueReadBuffer(queue_prof, hash_buffer, CL_FALSE, 0,
sizeof(sha512_hash) * num, tmpbuffer, 0, NULL, &myEvent),
"Failed in clEnqueueReadBuffer");
HANDLE_CLERROR(clFinish(queue_prof), "Failed in clFinish");
HANDLE_CLERROR(clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_SUBMIT,
sizeof(cl_ulong), &startTime, NULL),
"Failed in clGetEventProfilingInfo I");
HANDLE_CLERROR(clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_END,
sizeof(cl_ulong), &endTime, NULL),
"Failed in clGetEventProfilingInfo II");
HANDLE_CLERROR(clReleaseEvent(myEvent), "Failed in clReleaseEvent");
runtime += endTime - startTime;
MEM_FREE(tmpbuffer);
HANDLE_CLERROR(clReleaseCommandQueue(queue_prof),
"Failed in clReleaseCommandQueue");
release_clobj();
if (ret_code != CL_SUCCESS) {
if (ret_code != CL_INVALID_WORK_GROUP_SIZE)
fprintf(stderr, "Error %d\n", ret_code);
return 0;
}
return runtime;
}
