Openwall GNU/*/Linux - a small security-enhanced Linux distro for servers
[<prev] [next>] [thread-next>] [day] [month] [year] [list]
Date: Tue, 30 Oct 2012 22:09:47 +0100
From: magnum <john.magnum@...hmail.com>
To: john-dev@...ts.openwall.com
Subject: mscash2-opencl problems with GT650M

Sayantan,

On some weaker cards (GT650M for one) I get this from mscash2-opencl:

OpenCL error (UNKNOWN ERROR :() in file (common_opencl_pbkdf2.c) at line 
(313) - (SYNC FAILED)

That's not really a good error message. After some looks at the code I 
came up with this patch (not applied to git yet):

diff --git a/src/common_opencl_pbkdf2.c b/src/common_opencl_pbkdf2.c
index aa43382..07c749c 100644
--- a/src/common_opencl_pbkdf2.c
+++ b/src/common_opencl_pbkdf2.c
@@ -291,7 +291,8 @@ static gpu_mem_buffer exec_pbkdf2(cl_uint 
*pass_api,cl_uint *salt_api,cl_uint sa

 
HANDLE_CLERROR(clSetKernelArg(krnl[platform_no][dev_no],3,sizeof(cl_uint),&num),"Set 
Kernel Arg FAILED arg3");

- 
err=clEnqueueNDRangeKernel(cmdq[platform_no][dev_no],krnl[platform_no][dev_no],1,NULL,&N,&M,0,NULL,&evnt);
+ 
clEnqueueNDRangeKernel(cmdq[platform_no][dev_no],krnl[platform_no][dev_no],1,NULL,&N,&M,0,NULL,&evnt);
+       err=clFinish(cmdq[platform_no][dev_no]);

         if(err){
                 if(PROFILE){


This properly sets 'err' with any failure from the kernel and lets the 
subsequent if clause halve the LWS. However, the problem moves. I now 
get this:

Optimal Work Group Size:8
Kernel Execution Speed (Higher is better):-1.000000
Benchmarking: M$ Cache Hash 2 (DCC2) PBKDF2-HMAC-SHA-1 [OpenCL]... 
OpenCL error (CL_INVALID_COMMAND_QUEUE) in file (common_opencl_pbkdf2.c) 
at line (286) - (Copy data to gpu)


I could look further into this but it may be quicker for you. I believe 
my patch above should be applied but some more seem to be needed to 
adopt for available resources. Do you have any immediate ideas or should 
I go on trying to fix this?

magnum

Powered by blists - more mailing lists

Your e-mail address:

Powered by Openwall GNU/*/Linux - Powered by OpenVZ