自学教程：C++ HANDLE_CLERROR函数代码示例

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 : ");		for (j = 0; j < 64; j++) {			fprintf(stderr, "%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, &currentsalt, 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, &currentsalt, 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, &currentsalt, 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, &currentsalt, 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, &currentsalt, 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, &currentsalt, 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, &currentsalt, 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;}