static void set_salt(void *salt)
{
cur_salt = (pbkdf2_salt*)salt;
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_salt, CL_FALSE, 0, sizeof(pbkdf2_salt), cur_salt, 0, NULL, NULL), "Copy salt to gpu");
#if 0
fprintf(stderr, "\n%s(%.*s) len %u iter %u\n", __FUNCTION__, cur_salt->length, cur_salt->salt, cur_salt->length, cur_salt->iterations);
dump_stuff_msg("salt", cur_salt->salt, cur_salt->length);
#endif
}
static void init(struct fmt_main *_self)
{
self = _self;
opencl_init("$JOHN/kernels/phpass_kernel.cl", gpu_id, NULL);
crypt_kernel = clCreateKernel(program[gpu_id], "phpass", &cl_error);
HANDLE_CLERROR(cl_error, "Error creating kernel");
}
static int crypt_all(int *pcount, struct db_salt *salt)
{
const int count = *pcount;
int index;
size_t *lws = local_work_size ? &local_work_size : NULL;
global_work_size = local_work_size ? (count + local_work_size - 1) / local_work_size * local_work_size : count;
if (any_cracked) {
memset(cracked, 0, cracked_size);
any_cracked = 0;
}
/// Copy data to gpu
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_in, CL_FALSE, 0,
insize, inbuffer, 0, NULL, multi_profilingEvent[0]),
"Copy data to gpu");
/// Run kernel
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1,
NULL, &global_work_size, lws, 0, NULL,
multi_profilingEvent[1]), "Run kernel");
/// Read the result back
HANDLE_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_out, CL_TRUE, 0,
outsize, outbuffer, 0, NULL, multi_profilingEvent[2]), "Copy result back");
#ifdef _OPENMP
#pragma omp parallel for
#endif
for (index = 0; index < count; index++)
if (!kcdecrypt((unsigned char*)outbuffer[index].v,
salt_struct->iv, salt_struct->ct))
{
cracked[index] = 1;
#ifdef _OPENMP
#pragma omp atomic
#endif
any_cracked |= 1;
}
return count;
}
static void crypt_all(int count)
{
cl_int code;
code = clEnqueueWriteBuffer(queue[ocl_gpu_id], buffer_keys, CL_TRUE, 0,
(PLAINTEXT_LENGTH) * global_work_size, saved_plain, 0, NULL, NULL);
HANDLE_CLERROR(code, "failed in clEnqueueWriteBuffer saved_plain");
code = clEnqueueNDRangeKernel(queue[ocl_gpu_id], crypt_kernel, 1, NULL,
&global_work_size, &local_work_size, 0, NULL, profilingEvent);
HANDLE_CLERROR(code, "failed in clEnqueueNDRangeKernel");
HANDLE_CLERROR(clFinish(queue[ocl_gpu_id]), "clFinish error");
// read back partial hashes
code = clEnqueueReadBuffer(queue[ocl_gpu_id], buffer_out, CL_TRUE, 0,
sizeof(cl_uint) * global_work_size, outbuffer, 0, NULL, NULL);
HANDLE_CLERROR(code, "failed in clEnqueueReadBuffer -reading partial hashes");
have_full_hashes = 0;
}
static void create_clobj(int kpc){
pinned_saved_keys = clCreateBuffer(context[ocl_gpu_id], CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
(PLAINTEXT_LENGTH + 1) * kpc, NULL, &ret_code);
HANDLE_CLERROR(ret_code, "Error creating page-locked memory pinned_saved_keys");
saved_plain = (char *) clEnqueueMapBuffer(queue[ocl_gpu_id], pinned_saved_keys,
CL_TRUE, CL_MAP_WRITE | CL_MAP_READ, 0,
(PLAINTEXT_LENGTH + 1) * kpc, 0, NULL, NULL, &ret_code);
HANDLE_CLERROR(ret_code, "Error mapping page-locked memory saved_plain");
res_hashes = malloc(sizeof(cl_uint) * 3 * kpc);
pinned_partial_hashes = clCreateBuffer(context[ocl_gpu_id],
CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, 4 * kpc, NULL, &ret_code);
HANDLE_CLERROR(ret_code, "Error creating page-locked memory pinned_partial_hashes");
partial_hashes = (cl_uint *) clEnqueueMapBuffer(queue[ocl_gpu_id],
pinned_partial_hashes, CL_TRUE, CL_MAP_READ, 0, 4 * kpc, 0, NULL, NULL, &ret_code);
HANDLE_CLERROR(ret_code, "Error mapping page-locked memory partial_hashes");
// create and set arguments
buffer_keys = clCreateBuffer(context[ocl_gpu_id], CL_MEM_READ_ONLY,
(PLAINTEXT_LENGTH + 1) * kpc, NULL, &ret_code);
HANDLE_CLERROR(ret_code, "Error creating buffer argument buffer_keys");
buffer_out = clCreateBuffer(context[ocl_gpu_id], CL_MEM_WRITE_ONLY,
BINARY_SIZE * kpc, NULL, &ret_code);
HANDLE_CLERROR(ret_code, "Error creating buffer argument buffer_out");
data_info = clCreateBuffer(context[ocl_gpu_id], CL_MEM_READ_ONLY, sizeof(unsigned int) * 2, NULL, &ret_code);
HANDLE_CLERROR(ret_code, "Error creating data_info out argument");
HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 0, sizeof(data_info),
(void *) &data_info), "Error setting argument 0");
HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1, sizeof(buffer_keys),
(void *) &buffer_keys), "Error setting argument 1");
HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 2, sizeof(buffer_out),
(void *) &buffer_out), "Error setting argument 2");
datai[0] = PLAINTEXT_LENGTH;
datai[1] = kpc;
global_work_size = kpc;
}
static void crypt_all(int count)
{
#ifdef DEBUGVERBOSE
int i, j;
unsigned char *p = (unsigned char *) saved_plain;
count--;
for (i = 0; i < count + 1; i++) {
fprintf(stderr, "\npassword : "******"%02x ", p[i * 64 + j]);
}
}
fprintf(stderr, "\n");
#endif
// copy keys to the device
HANDLE_CLERROR( clEnqueueWriteBuffer(queue[ocl_gpu_id], data_info, CL_TRUE, 0,
sizeof(unsigned int) * 2, datai, 0, NULL, NULL),
"failed in clEnqueueWriteBuffer data_info");
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[ocl_gpu_id], buffer_keys, CL_TRUE, 0,
(PLAINTEXT_LENGTH + 1) * max_keys_per_crypt, saved_plain, 0, NULL, NULL),
"failed in clEnqueueWriteBuffer buffer_keys");
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[ocl_gpu_id], crypt_kernel, 1, NULL,
&global_work_size, &local_work_size, 0, NULL, &profilingEvent),
"failed in clEnqueueNDRangeKernel");
HANDLE_CLERROR(clFinish(queue[ocl_gpu_id]),"failed in clFinish");
// read back partial hashes
HANDLE_CLERROR(clEnqueueReadBuffer(queue[ocl_gpu_id], buffer_out, CL_TRUE, 0,
sizeof(cl_uint) * max_keys_per_crypt, partial_hashes, 0, NULL, NULL),
"failed in reading data back");
have_full_hashes = 0;
#ifdef DEBUGVERBOSE
p = (unsigned char *) partial_hashes;
for (i = 0; i < 2; i++) {
fprintf(stderr, "\n\npartial_hashes : ");
for (j = 0; j < 16; j++)
fprintf(stderr, "%02x ", p[i * 16 + j]);
}
fprintf(stderr, "\n");;
#endif
}
static void init(struct fmt_main *_self)
{
char build_opts[128];
self = _self;
snprintf(build_opts, sizeof(build_opts),
"-DHASH_LOOPS=%u -DPLAINTEXT_LENGTH=%d -DMAX_SALT_SIZE=%d",
HASH_LOOPS, PLAINTEXT_LENGTH, MAX_SALT_SIZE);
opencl_init("$JOHN/kernels/pbkdf2_hmac_sha512_kernel.cl",
gpu_id, build_opts);
crypt_kernel = clCreateKernel(program[gpu_id], KERNEL_NAME, &cl_error);
HANDLE_CLERROR(cl_error, "Error creating kernel");
split_kernel =
clCreateKernel(program[gpu_id], SPLIT_KERNEL_NAME, &cl_error);
HANDLE_CLERROR(cl_error, "Error creating split kernel");
}
static void set_salt(void *salt)
{
cur_salt = (struct custom_salt *)salt;
memcpy((char*)currentsalt.salt, cur_salt->salt, cur_salt->SaltSize);
currentsalt.length = cur_salt->SaltSize;
currentsalt.iterations = cur_salt->NumCyclesPower;
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_salt,
CL_FALSE, 0, saltsize, ¤tsalt, 0, NULL, NULL),
"Transfer salt to gpu");
}
static void set_salt(void *salt)
{
salt_struct = (struct custom_salt *)salt;
memcpy((char*)currentsalt.salt, salt_struct->salt, 20);
currentsalt.length = 20;
currentsalt.iterations = 1000;
currentsalt.outlen = 24;
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_setting,
CL_FALSE, 0, settingsize, ¤tsalt, 0, NULL, NULL),
"Copy salt to gpu");
}
static void set_salt(void *salt)
{
cur_salt = (struct custom_salt *)salt;
memcpy((char*)currentsalt.salt, cur_salt->salt, cur_salt->saltlen[0]);
currentsalt.length = cur_salt->saltlen[0];
currentsalt.iterations = cur_salt->iterations[0];
currentsalt.outlen = 16;
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_setting,
CL_FALSE, 0, settingsize, ¤tsalt, 0, NULL, NULL),
"Copy salt to gpu");
}
static void create_clobj(size_t global_work_size, struct fmt_main *self)
{
cl_int cl_error;
inbuffer = (sevenzip_password*) mem_calloc(1, insize);
outbuffer = (sevenzip_hash*) mem_alloc(outsize);
cracked = mem_calloc(1, cracked_size);
// Allocate memory
mem_in =
clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, insize, NULL,
&cl_error);
HANDLE_CLERROR(cl_error, "Error allocating mem in");
mem_salt =
clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, saltsize,
NULL, &cl_error);
HANDLE_CLERROR(cl_error, "Error allocating mem salt");
mem_state =
clCreateBuffer(context[gpu_id], CL_MEM_READ_WRITE, statesize,
NULL, &cl_error);
HANDLE_CLERROR(cl_error, "Error allocating mem state");
mem_out =
clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY, outsize, NULL,
&cl_error);
HANDLE_CLERROR(cl_error, "Error allocating mem out");
HANDLE_CLERROR(clSetKernelArg(sevenzip_init, 0, sizeof(mem_in),
&mem_in), "Error while setting mem_in kernel argument");
HANDLE_CLERROR(clSetKernelArg(sevenzip_init, 1, sizeof(mem_salt),
&mem_salt), "Error while setting mem_salt kernel argument");
HANDLE_CLERROR(clSetKernelArg(sevenzip_init, 2, sizeof(mem_state),
&mem_state), "Error while setting mem_state kernel argument");
HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 0, sizeof(mem_state),
&mem_state), "Error while setting mem_state kernel argument");
HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1, sizeof(mem_salt),
&mem_salt), "Error while setting mem_salt kernel argument");
HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 2, sizeof(mem_out),
&mem_out), "Error while setting mem_out kernel argument");
}
static void init(struct fmt_main *self)
{
char build_opts[64];
if (pers_opts.target_enc == UTF_8) {
max_len = self->params.plaintext_length = 3 * PLAINTEXT_LENGTH;
tests[1].plaintext = "\xC3\xBC"; // German u-umlaut in UTF-8
tests[1].ciphertext = "$mskrb5$$$958db4ddb514a6cc8be1b1ccf82b0191$090408357a6f41852d17f3b4bb4634adfd388db1be64d3fe1a1d75ee4338d2a4aea387e5";
tests[2].plaintext = "\xC3\x9C\xC3\x9C"; // 2x uppercase of them
tests[2].ciphertext = "$mskrb5$$$057cd5cb706b3de18e059912b1f057e3$fe2e561bd4e42767e972835ea99f08582ba526e62a6a2b6f61364e30aca7c6631929d427";
} else {
if (CP_to_Unicode[0xfc] == 0x00fc) {
tests[1].plaintext = "\xFC"; // German u-umlaut in many ISO-8859-x
tests[1].ciphertext = "$mskrb5$$$958db4ddb514a6cc8be1b1ccf82b0191$090408357a6f41852d17f3b4bb4634adfd388db1be64d3fe1a1d75ee4338d2a4aea387e5";
}
if (CP_to_Unicode[0xdc] == 0x00dc) {
tests[2].plaintext = "\xDC\xDC"; // 2x uppercase of them
tests[2].ciphertext = "$mskrb5$$$057cd5cb706b3de18e059912b1f057e3$fe2e561bd4e42767e972835ea99f08582ba526e62a6a2b6f61364e30aca7c6631929d427";
}
}
snprintf(build_opts, sizeof(build_opts),
"-D%s -DPLAINTEXT_LENGTH=%u",
cp_id2macro(pers_opts.target_enc), PLAINTEXT_LENGTH);
opencl_init("$JOHN/kernels/krb5pa-md5_kernel.cl", gpu_id, build_opts);
/* create kernels to execute */
krb5pa_md5_nthash = clCreateKernel(program[gpu_id], "krb5pa_md5_nthash", &ret_code);
HANDLE_CLERROR(ret_code, "Error creating kernel. Double-check kernel name?");
crypt_kernel = clCreateKernel(program[gpu_id], "krb5pa_md5_final", &ret_code);
HANDLE_CLERROR(ret_code, "Error creating kernel. Double-check kernel name?");
//Initialize openCL tuning (library) for this format.
opencl_init_auto_setup(SEED, 0, NULL,
warn, 2, self, create_clobj, release_clobj,
PLAINTEXT_LENGTH, 0);
//Auto tune execution from shared/included code.
autotune_run(self, 1, 0, 200);
}
static void set_salt(void *salt)
{
cur_salt = (sxc_cpu_salt*)salt;
memcpy((char*)currentsalt.salt, cur_salt->salt, cur_salt->salt_length);
currentsalt.length = cur_salt->salt_length;
currentsalt.iterations = cur_salt->iterations;
currentsalt.outlen = cur_salt->key_size;
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_setting,
CL_FALSE, 0, settingsize, ¤tsalt, 0, NULL, NULL),
"Copy salt to gpu");
}
static int cmp_all(void *binary, int count)
{
uint32_t result;
///Copy binary to GPU memory
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[ocl_gpu_id], mem_binary, CL_FALSE,
0, sizeof(uint64_t), ((uint64_t*)binary)+3, 0, NULL, NULL), "Copy mem_binary");
///Run kernel
HANDLE_CLERROR(clEnqueueNDRangeKernel
(queue[ocl_gpu_id], cmp_kernel, 1, NULL, &global_work_size, &local_work_size,
0, NULL, NULL), "Set ND range");
/// Copy result out
HANDLE_CLERROR(clEnqueueReadBuffer(queue[ocl_gpu_id], mem_cmp, CL_FALSE, 0,
sizeof(uint32_t), &result, 0, NULL, NULL), "Copy data back");
///Await completion of all the above
HANDLE_CLERROR(clFinish(queue[ocl_gpu_id]), "clFinish error");
return result;
}
static void crypt_all(int count)
{
///Copy data to GPU memory
HANDLE_CLERROR(clEnqueueWriteBuffer
(queue[gpu_id], mem_in, CL_FALSE, 0, insize, inbuffer, 0, NULL,
NULL), "Copy memin");
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_salt, CL_FALSE,
0, saltsize, &host_salt, 0, NULL, NULL), "Copy memsalt");
///Run kernel
size_t worksize = KEYS_PER_CRYPT;
size_t localworksize = local_work_size;
HANDLE_CLERROR(clEnqueueNDRangeKernel
(queue[gpu_id], crypt_kernel, 1, NULL, &worksize, &localworksize,
0, NULL, NULL), "Set ND range");
HANDLE_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_out, CL_FALSE, 0,
outsize, outbuffer, 0, NULL, NULL), "Copy data back");
///Await completion of all the above
HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish error");
}
static void set_salt(void *salt)
{
cur_salt = (zip_cpu_salt*)salt;
memcpy((char*)currentsalt.salt, cur_salt->salt, cur_salt->length);
currentsalt.length = cur_salt->length;
currentsalt.iterations = KEYING_ITERATIONS;
currentsalt.outlen = 2 * KEY_LENGTH(cur_salt->mode) + PWD_VER_LENGTH;
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[ocl_gpu_id], mem_setting,
CL_FALSE, 0, settingsize, ¤tsalt, 0, NULL, NULL),
"Copy setting to gpu");
}
static int crypt_all(int *pcount, struct db_salt *salt)
{
int count = *pcount;
opencl_limit_gws(count);
#ifdef DEBUG
printf("crypt_all(%d)\n", count);
printf("LWS = %d, GWS = %d\n",(int)local_work_size, (int)global_work_size);
#endif
/// Copy data to gpu
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[ocl_gpu_id], mem_in, CL_FALSE, 0,
global_work_size * sizeof(pass_t), host_pass, 0, NUUL, NULL), "Copy data to gpu");
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[ocl_gpu_id], mem_salt,
CL_FALSE, 0, sizeof(salt_t), host_salt, 0, NUUL, NULL),
"Copy salt to gpu");
/// Run kernel
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[ocl_gpu_id], crypt_kernel, 1,
NUUL, &global_work_size, &local_work_size, 0, NULL,
profilingEvent), "Run kernel");
HANDLE_CLERROR(clFinish(queue[ocl_gpu_id]), "clFinish");
/// Read the result back
HANDLE_CLERROR(clEnqueueReadBuffer(queue[ocl_gpu_id], mem_out, CL_FALSE, 0,
global_work_size * sizeof(crack_t), host_crack, 0, NUUL, NULL), "Copy result back");
/// Await completion of all the above
HANDLE_CLERROR(clFinish(queue[ocl_gpu_id]), "clFinish");
return count;
}
static void crypt_all(int count)
{
int index;
size_t scalar_gws = VF * global_work_size;
if (new_keys) {
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[ocl_gpu_id], cl_saved_key, CL_FALSE, 0, UNICODE_LENGTH * VF * global_work_size, saved_key, 0, NULL, NULL), "failed in clEnqueueWriteBuffer saved_key");
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[ocl_gpu_id], cl_saved_len, CL_FALSE, 0, sizeof(int) * VF * global_work_size, saved_len, 0, NULL, NULL), "failed in clEnqueueWriteBuffer saved_len");
new_keys = 0;
}
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[ocl_gpu_id], GenerateSHA1pwhash, 1, NULL, &scalar_gws, &local_work_size, 0, NULL, firstEvent), "failed in clEnqueueNDRangeKernel");
for (index = 0; index < 50000 / HASH_LOOPS; index++)
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[ocl_gpu_id], crypt_kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL), "failed in clEnqueueNDRangeKernel");
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[ocl_gpu_id], Generate2007key, 1, NULL, &global_work_size, &local_work_size, 0, NULL, lastEvent), "failed in clEnqueueNDRangeKernel");
// read back aes key
HANDLE_CLERROR(clEnqueueReadBuffer(queue[ocl_gpu_id], cl_key, CL_TRUE, 0, 16 * VF * global_work_size, key, 0, NULL, NULL), "failed in reading key back");
#ifdef _OPENMP
#pragma omp parallel for
#endif
for (index = 0; index < count; index++)
cracked[index] = PasswordVerifier(&key[index*16]);
}
static void create_clobj(size_t kpc, struct fmt_main *self)
{
kpc *= 8;
insize = sizeof(phpass_password) * kpc;
outsize = sizeof(phpass_hash) * kpc;
settingsize = sizeof(uint8_t) * ACTUAL_SALT_SIZE + 4;
inbuffer = mem_calloc(1, insize);
outbuffer = mem_alloc(outsize);
// Allocate memory
mem_in =
clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, insize, NULL,
&cl_error);
HANDLE_CLERROR(cl_error, "Error allocating mem in");
mem_setting =
clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, settingsize,
NULL, &cl_error);
HANDLE_CLERROR(cl_error, "Error allocating mem setting");
mem_out =
clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY, outsize, NULL,
&cl_error);
HANDLE_CLERROR(cl_error, "Error allocating mem out");
HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 0, sizeof(mem_in),
&mem_in), "Error while setting mem_in kernel argument");
HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1, sizeof(mem_out),
&mem_out), "Error while setting mem_out kernel argument");
HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 2, sizeof(mem_setting),
&mem_setting), "Error while setting mem_salt kernel argument");
}
static int crypt_all(int *pcount, struct db_salt *salt)
{
const int count = *pcount;
int i;
size_t scalar_gws;
global_work_size = ((count + (v_width * local_work_size - 1)) / (v_width * local_work_size)) * local_work_size;
scalar_gws = global_work_size * v_width;
#if 0
fprintf(stderr, "%s(%d) lws "Zu" gws "Zu" sgws "Zu"\n", __FUNCTION__,
count, local_work_size, global_work_size, scalar_gws);
#endif
/// Copy data to gpu
if (new_keys) {
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_in, CL_FALSE, 0, key_buf_size, inbuffer, 0, NULL, NULL), "Copy data to gpu");
new_keys = 0;
}
/// Run kernels
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], pbkdf1_init, 1, NULL, &global_work_size, &local_work_size, 0, NULL, firstEvent), "Run initial kernel");
for (i = 0; i < LOOP_COUNT; i++) {
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], pbkdf1_loop, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL), "Run loop kernel");
HANDLE_CLERROR(clFinish(queue[gpu_id]), "Error running loop kernel");
opencl_process_event();
}
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], pbkdf1_final, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL), "Run intermediate kernel");
/// Read the result back
HANDLE_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_out, CL_TRUE, 0, sizeof(pbkdf1_out) * scalar_gws, host_crack, 0, NULL, NULL), "Copy result back");
return count;
}
static void crypt_all(int count)
{
cl_int code;
//memcpy(inbuffer,saved_key,PLAINTEXT_LENGTH*count);
code = clEnqueueWriteBuffer(queue[gpu_id], data_info, CL_TRUE, 0,
sizeof(unsigned int) * 2, datai, 0, NULL, NULL);
HANDLE_CLERROR(code, "failed in clEnqueueWriteBuffer data_info");
code = clEnqueueWriteBuffer(queue[gpu_id], mysalt, CL_TRUE, 0, SALT_SIZE,
saved_salt, 0, NULL, NULL);
HANDLE_CLERROR(code, "failed in clEnqueueWriteBuffer mysalt");
code = clEnqueueWriteBuffer(queue[gpu_id], buffer_keys, CL_TRUE, 0,
(PLAINTEXT_LENGTH) * max_keys_per_crypt, inbuffer, 0, NULL, NULL);
HANDLE_CLERROR(code, "failed in clEnqueueWriteBuffer inbuffer");
code = clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1, NULL,
&global_work_size, &local_work_size, 0, NULL, NULL);
HANDLE_CLERROR(code, "failed in clEnqueueNDRangeKernel");
HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish error");
// read back partial hashes
code = clEnqueueReadBuffer(queue[gpu_id], buffer_out, CL_TRUE, 0,
sizeof(cl_uint) * max_keys_per_crypt, outbuffer, 0, NULL, NULL);
HANDLE_CLERROR(code, "failed in clEnqueueReadBuffer -reading partial hashes");
have_full_hashes = 0;
}
static void create_clobj(size_t gws, struct fmt_main *self)
{
insize = sizeof(sxc_password) * gws;
outsize = sizeof(sxc_hash) * gws;
settingsize = sizeof(sxc_salt);
inbuffer = mem_calloc(1, insize);
outbuffer = mem_alloc(outsize);
saved_key = mem_calloc(gws, sizeof(*saved_key));
crypt_out = mem_calloc(gws, sizeof(*crypt_out));
/// Allocate memory
mem_in =
clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, insize, NULL,
&cl_error);
HANDLE_CLERROR(cl_error, "Error allocating mem in");
mem_setting =
clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, settingsize,
NULL, &cl_error);
HANDLE_CLERROR(cl_error, "Error allocating mem setting");
mem_out =
clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY, outsize, NULL,
&cl_error);
HANDLE_CLERROR(cl_error, "Error allocating mem out");
HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 0, sizeof(mem_in),
&mem_in), "Error while setting mem_in kernel argument");
HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 1, sizeof(mem_out),
&mem_out), "Error while setting mem_out kernel argument");
HANDLE_CLERROR(clSetKernelArg(crypt_kernel, 2, sizeof(mem_setting),
&mem_setting), "Error while setting mem_salt kernel argument");
}
static void crypt_all(int count)
{
/// Copy data to gpu
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_in, CL_FALSE, 0,
insize, inbuffer, 0, NULL, NULL), "Copy data to gpu");
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_setting,
CL_FALSE, 0, settingsize, ¤tsalt, 0, NULL, NULL),
"Copy setting to gpu");
/// Run kernel
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1,
NULL, &global_work_size, &local_work_size, 0, NULL, NULL),
"Run kernel");
HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish");
/// Read the result back
HANDLE_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_out, CL_FALSE, 0,
outsize, outbuffer, 0, NULL, NULL), "Copy result back");
/// Await completion of all the above
HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish");
///Make last computations on CPU
wpapsk_postprocess(KEYS_PER_CRYPT);
}
static void init(struct fmt_main *self)
{
char *temp;
if ((temp = getenv("LWS")))
local_work_size = atoi(temp);
else
local_work_size = cpu(device_info[ocl_gpu_id]) ? 1 : 64;
if ((temp = getenv("GWS")))
global_work_size = atoi(temp);
else
global_work_size = MAX_KEYS_PER_CRYPT;
opencl_init("$JOHN/kernels/sha512_kernel.cl", ocl_gpu_id);
gkey = mem_calloc(global_work_size * sizeof(sha512_key));
ghash = mem_calloc(global_work_size * sizeof(sha512_hash));
///Allocate memory on the GPU
mem_in =
clCreateBuffer(context[ocl_gpu_id], CL_MEM_READ_ONLY, insize, NULL,
&ret_code);
HANDLE_CLERROR(ret_code,"Error while allocating memory for passwords");
mem_out =
clCreateBuffer(context[ocl_gpu_id], CL_MEM_WRITE_ONLY, outsize, NULL,
&ret_code);
HANDLE_CLERROR(ret_code,"Error while allocating memory for hashes");
mem_binary =
clCreateBuffer(context[ocl_gpu_id], CL_MEM_READ_ONLY, sizeof(uint64_t), NULL,
&ret_code);
HANDLE_CLERROR(ret_code,"Error while allocating memory for binary");
mem_cmp =
clCreateBuffer(context[ocl_gpu_id], CL_MEM_WRITE_ONLY, sizeof(uint32_t), NULL,
&ret_code);
HANDLE_CLERROR(ret_code,"Error while allocating memory for cmp_all result");
///Assign crypt kernel parameters
crypt_kernel = clCreateKernel(program[ocl_gpu_id], KERNEL_NAME, &ret_code);
HANDLE_CLERROR(ret_code,"Error while creating crypt_kernel");
clSetKernelArg(crypt_kernel, 0, sizeof(mem_in), &mem_in);
clSetKernelArg(crypt_kernel, 1, sizeof(mem_out), &mem_out);
///Assign cmp kernel parameters
cmp_kernel = clCreateKernel(program[ocl_gpu_id], CMP_KERNEL_NAME, &ret_code);
HANDLE_CLERROR(ret_code,"Error while creating cmp_kernel");
clSetKernelArg(cmp_kernel, 0, sizeof(mem_binary), &mem_binary);
clSetKernelArg(cmp_kernel, 1, sizeof(mem_out), &mem_out);
clSetKernelArg(cmp_kernel, 2, sizeof(mem_cmp), &mem_cmp);
self->params.max_keys_per_crypt = global_work_size;
if (!local_work_size)
opencl_find_best_workgroup(self);
self->params.min_keys_per_crypt = local_work_size;
if (options.verbosity > 2)
fprintf(stderr, "Local worksize (LWS) %d, Global worksize (GWS) %d\n",(int)local_work_size, (int)global_work_size);
}
static void init(struct fmt_main *pFmt)
{
assert(sizeof(hccap_t) == HCCAP_SIZE);
inbuffer =
(wpapsk_password *) malloc(sizeof(wpapsk_password) *
MAX_KEYS_PER_CRYPT);
outbuffer =
(wpapsk_hash *) malloc(sizeof(wpapsk_hash) * MAX_KEYS_PER_CRYPT);
mic = (mic_t *) malloc(sizeof(mic_t) * MAX_KEYS_PER_CRYPT);
//listOpenCLdevices();
opencl_init("$JOHN/wpapsk_kernel.cl", gpu_id, platform_id);
/// Alocate memory
cl_int cl_error;
mem_in =
clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, insize, NULL,
&cl_error);
HANDLE_CLERROR(cl_error, "Error alocating mem in");
mem_setting =
clCreateBuffer(context[gpu_id], CL_MEM_READ_ONLY, settingsize,
NULL, &cl_error);
HANDLE_CLERROR(cl_error, "Error alocating mem setting");
mem_out =
clCreateBuffer(context[gpu_id], CL_MEM_WRITE_ONLY, outsize, NULL,
&cl_error);
HANDLE_CLERROR(cl_error, "Error alocating mem out");
crypt_kernel = clCreateKernel(program[gpu_id], "wpapsk", &cl_error);
HANDLE_CLERROR(cl_error, "Error creating kernel");
clSetKernelArg(crypt_kernel, 0, sizeof(mem_in), &mem_in);
clSetKernelArg(crypt_kernel, 1, sizeof(mem_out), &mem_out);
clSetKernelArg(crypt_kernel, 2, sizeof(mem_setting), &mem_setting);
find_best_workgroup();
atexit(release_all);
}
static void set_salt(void *salt)
{
saved_salt = *((my_salt**)salt);
memcpy((char*)currentsalt.salt, saved_salt->salt, SALT_LENGTH(saved_salt->v.mode));
currentsalt.length = SALT_LENGTH(saved_salt->v.mode);
currentsalt.iterations = KEYING_ITERATIONS;
currentsalt.outlen = 2 * KEY_LENGTH(saved_salt->v.mode) + PWD_VER_LENGTH;
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_setting,
CL_FALSE, 0, settingsize, ¤tsalt, 0, NULL, NULL),
"Copy setting to gpu");
}
static int crypt_all(int *pcount, struct db_salt *salt)
{
int i;
const int count = *pcount;
int loops = (host_salt->rounds + HASH_LOOPS - 1) / HASH_LOOPS;
opencl_limit_gws(count);
#if 0
printf("crypt_all(%d)\n", count);
printf("LWS = %d, GWS = %d, loops=%d\n",(int)local_work_size, (int)global_work_size, loops);
#endif
/// Copy data to gpu
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[gpu_id], mem_in, CL_FALSE, 0,
global_work_size * sizeof(pass_t), host_pass, 0, NUUL,
NULL), "Copy data to gpu");
/// Run kernel
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1,
NUUL, &global_work_size, &local_work_size, 0, NULL,
NULL), "Run kernel");
for(i = 0; i < loops; i++) {
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[gpu_id],
split_kernel,
1, NULL, &global_work_size, &local_work_size, 0, NULL,
NULL), "Run split kernel");
HANDLE_CLERROR(clFinish(queue[gpu_id]), "clFinish");
opencl_process_event();
}
/// Read the result back
HANDLE_CLERROR(clEnqueueReadBuffer(queue[gpu_id], mem_out, CL_TRUE, 0,
global_work_size * sizeof(crack_t), host_crack, 0, NUUL,
NULL), "Copy result back");
return count;
}
static int crypt_all(int *pcount, struct db_salt *salt)
{
int count = *pcount;
int index;
global_work_size = (count + local_work_size - 1) / local_work_size * local_work_size;
if (any_cracked) {
memset(cracked, 0, cracked_size);
any_cracked = 0;
}
/// Copy data to gpu
HANDLE_CLERROR(clEnqueueWriteBuffer(queue[ocl_gpu_id], mem_in, CL_FALSE, 0,
insize, inbuffer, 0, NULL, NULL), "Copy data to gpu");
/// Run kernel
HANDLE_CLERROR(clEnqueueNDRangeKernel(queue[ocl_gpu_id], crypt_kernel, 1,
NULL, &global_work_size, &local_work_size, 0, NULL, profilingEvent),
"Run kernel");
HANDLE_CLERROR(clFinish(queue[ocl_gpu_id]), "clFinish");
/// Read the result back
HANDLE_CLERROR(clEnqueueReadBuffer(queue[ocl_gpu_id], mem_out, CL_FALSE, 0,
outsize, outbuffer, 0, NULL, NULL), "Copy result back");
/// Await completion of all the above
HANDLE_CLERROR(clFinish(queue[ocl_gpu_id]), "clFinish");
#ifdef _OPENMP
#pragma omp parallel for
#endif
for (index = 0; index < count; index++)
if (!memcmp(&((unsigned char*)outbuffer[index].v)[2 * KEY_LENGTH(cur_salt->mode)], cur_salt->passverify, 2))
any_cracked = cracked[index] = 1;
return count;
}
static void crypt_all(int count) {
//memcpy(mysqlsha_plain,saved_key,PLAINTEXT_LENGTH*count);
HANDLE_CLERROR(
clEnqueueWriteBuffer(queue[gpu_id], data_info, CL_TRUE, 0,
sizeof(unsigned int) * 2, datai, 0, NULL, NULL),
"failed in clEnqueueWriteBuffer data_info");
HANDLE_CLERROR(
clEnqueueWriteBuffer(queue[gpu_id], buf_msha_keys, CL_TRUE, 0,
(PLAINTEXT_LENGTH) * max_keys_per_crypt, mysqlsha_plain, 0, NULL, NULL),
"failed in clEnqueueWriteBuffer mysqlsha_plain");
HANDLE_CLERROR(
clEnqueueNDRangeKernel(queue[gpu_id], crypt_kernel, 1, NULL,
&global_work_size, &local_work_size, 0, NULL, NULL),
"failed in clEnqueueNDRangeKernel");
HANDLE_CLERROR(clFinish(queue[gpu_id]),"failed in clFinish");
// read back partial hashes
HANDLE_CLERROR(clEnqueueReadBuffer(queue[gpu_id], buf_msha_out, CL_TRUE, 0,
sizeof(cl_uint) * max_keys_per_crypt, par_msha_hashes, 0, NULL, NULL),
"failed in reading data back");
have_full_hashes = 0;
}
static void crypt_all(int count)
{
size_t worksize = KEYS_PER_CRYPT;
size_t localworksize = local_work_size;
///Copy data to GPU memory
if (sha512_key_changed) {
HANDLE_CLERROR(clEnqueueWriteBuffer
(queue[ocl_gpu_id], mem_in, CL_FALSE, 0, insize, gkey, 0, NULL,
NULL), "Copy memin");
}
///Run kernel
HANDLE_CLERROR(clEnqueueNDRangeKernel
(queue[ocl_gpu_id], crypt_kernel, 1, NULL, &worksize, &localworksize,
0, NULL, profilingEvent), "Set ND range");
///Await completion of all the above
HANDLE_CLERROR(clFinish(queue[ocl_gpu_id]), "clFinish error");
/// Reset key to unchanged and hashes uncopy to host
sha512_key_changed = 0;
hash_copy_back = 0;
}
