Update

Author: knabberknusperhaus

cpyrit/cpyrit.c

@@ -22,6 +22,12 @@
 
 int numThreads = 4;
 
+#ifdef HAVE_CUDA
+    int cudadev;
+    int cudadevcount;
+    struct cudaDeviceProp cuda_devprop;
+#endif
+
 #ifdef HAVE_PADLOCK
 
 pthread_mutex_t padlock_sigmutex;
@@ -339,7 +345,16 @@ cpyrit_pmklist(PyObject *self, PyObject *args)
 }
 
 #ifdef HAVE_CUDA
+    //Will be linked as object file
     PyObject *cpyrit_cuda(PyObject *self, PyObject *args);
+    
+    PyObject *cpyrit_devprops(PyObject *self, PyObject *args)
+    {
+        if (!PyArg_ParseTuple(args, "")) return NULL;
+    
+        return Py_BuildValue("iisiiii", cudadevcount, cudadev, &cuda_devprop.name, cuda_devprop.totalGlobalMem,
+                                        cuda_devprop.regsPerBlock, cuda_devprop.clockRate, cuda_devprop.multiProcessorCount);
+    }
 #endif
 
 static PyMethodDef SpamMethods[] = {
@@ -354,6 +369,7 @@ static PyMethodDef SpamMethods[] = {
 
     #ifdef HAVE_CUDA
         {"calc_cuda", cpyrit_cuda, METH_VARARGS, "Calculate PMKs from ESSID and list of strings"},
+        {"cudaprops", cpyrit_devprops, METH_VARARGS, "Returns a tuple with some properties about main device"},
     #endif
     {NULL, NULL, 0, NULL}
 };
@@ -362,25 +378,41 @@ PyMODINIT_FUNC
 init_cpyrit(void)
 {
     (void) Py_InitModule("_cpyrit", SpamMethods);
-}
-
-int
-main(int argc, char *argv[])
-{
-    /* Initialize the Python interpreter.  Required. */
-    Py_Initialize();
 
-    init_cpyrit();
 
     #ifdef HAVE_CUDA
+    
+        // This is a somewhat awkward way to initialize the GPU as a CUDA-context
+        // is thread specific. There is no guarantee that the cuda functions above will be called
+        // by the same thread which called this code. However we will not force that as at
+        // some point in the future we need multi-gpu functionality anyway.
         char* buffer;
-        cudaMallocHost( (void**) &buffer, 4 );
-        cudaFreeHost( buffer );
+        cudaMallocHost((void**) &buffer, 4);
+        cudaFreeHost(buffer);
+        
+        cudaGetDeviceCount(&cudadevcount);
+        cudaGetDevice(&cudadev);
+        cudaGetDeviceProperties(&cuda_devprop, cudadev);
+        int ret = cudaGetLastError();
+        if (ret != cudaSuccess)
+        {
+            PyErr_SetString(PyExc_SystemError, cudaGetErrorString(ret));
+            return;
+        }
+
     #endif
 
     #ifdef HAVE_PADLOCK
         pthread_mutex_init(&padlock_sigmutex, NULL);
     #endif
+}
+
+int
+main(int argc, char *argv[])
+{
+    Py_Initialize();
+
+    init_cpyrit();
 
     return -1;
 }

cpyrit/cpyrit.py

@@ -24,32 +24,53 @@
 
 class CUDACore(object):
     name = "Nvidia CUDA"
-    description = "Yeah"
+    ctype = "GPU"
     def __init__(self):
         assert 'calc_cuda' in dir(_cpyrit)
-        self.gpu_perf = (0, 0)
-        self.cpu_perf = (0, 0)
+        self.buffersize = 1024
+        props = _cpyrit.cudaprops()
+        self.devicename = props[2]
+        self.devicemem = int(props[3] / 1024.0 / 1024.0)
+        self.deviceclock = int(props[5] / 1024.0)
 
     def solve(self, essid, password):
         assert isinstance(essid, str)
         if isinstance(password, str):
             return _cpyrit.calc_pmk(essid, password)
         assert isinstance(password, list)
 
-        return _cpyrit.calc_cuda(essid, password)
+        # The kernel allows a max. execution time of 5 seconds per call (usually) so we have to
+        # limit the input-buffer to some degree. However the size of the input-buffer is crucial
+        # to overall performance. Therefor buffersize is to be calibrated somewhere near
+        # it's maximum value allowed. Target is 3.0 seconds execution time.
+        res = []
+        i = 0
+        while i < len(password):
+            t = time.time()
+            res.extend(_cpyrit.calc_cuda(essid, password[i:i+self.buffersize]))
+            i += self.buffersize
+            if len(password[i:i+self.buffersize]) > self.buffersize:
+                self.buffersize = int(max(1024, min(10240, (2 * self.buffersize + (3.0 / (time.time() - t) * self.buffersize)) / 3)))
+        return res
+
 
 class CPUCore(object):
     name = "Standard CPU"
-    description = "Yeah!"
+    ctype = "CPU"
     def __init__(self):
         assert 'calc_pmklist' in dir(_cpyrit)
 
     def solve(self, essid, password):
         assert isinstance(essid, str)
         if isinstance(password, list):
-            return _cpyrit.calc_pmklist(essid, password)
+            # slicing gives better interactivity to signals as _cpyrit's functions won't listen to them
+            res = []
+            for pwslice in xrange(0, len(password), 1000):
+                res.extend(_cpyrit.calc_pmklist(essid, password[pwslice:pwslice+1000]))
+            return res
         elif isinstance(password, str):
             return _cpyrit.calc_pmk(essid, password)
+
         else:
             raise TypeError, "Password parameter must be string or list"
 

pyrit_cli.py

@@ -38,7 +38,7 @@ def __init__(self):
                         "core_name": None,
                         "essid": None,
                         "file": None,
-                        "ncpus": 1}
+                        "ncpus": None}
         self.pyrit_obj = None
 
     def progressbar(self, idx, max_idx):
@@ -208,10 +208,15 @@ def batchprocess(self):
         cp = cpyrit.CPyrit(ncpus = self.options["ncpus"])
         if self.options["core_name"] is not None:
             core = cp.getCore(self.options["core_name"])
-            print "Selected core '%s' (%i CPUs)" % (core.name, cp.ncpus)
+            print "Selected core '%s'" % core.name,
         else:
             core = cp.getCore()
-            print "Using default core '%s' (%i CPUs)" % (core.name, cp.ncpus)
+            print "Using default core '%s'" % core.name,
+        if core.ctype == 'GPU':
+            print "(Device '%s')" % core.devicename
+        else:
+            print "(%i CPUs)" % cp.ncpus
+
         comptime = 0
         rescount = 0    
         essids = self.pyrit_obj.list_essids()
@@ -240,8 +245,6 @@ def batchprocess(self):
                     print "%i PMKs to do." % len(passwords)
 
                     if len(passwords) > 0:
-                        #We slice the workunit to smaller parts since calc_pmklist won't return on KeyboardInterrupt
-                        #the overhead of slicing is minimal
                         for pwslice in xrange(0,len(passwords), 15000):
                             pwset = passwords[pwslice:pwslice+15000]
                             t = time.time()
@@ -276,19 +279,19 @@ def benchmark(self):
         print "%i PMKs in %.2f seconds: %.2f PMKs/s" % (len(pws), t, len(pws) / t)
         md = md5.new()
         map(md.update, [x[1] for x in res])
-        print "Result hash: %s" % md.hexdigest(), {True: "OK", False: "FAILED"}[md.hexdigest() == "ef747d123821851a9bd1d1e94ba048ac"]
+        print "Result hash:", {True: "OK", False: "FAILED"}[md.hexdigest() == "ef747d123821851a9bd1d1e94ba048ac"]
         print ""
 
         if 'Nvidia CUDA' in [x[0] for x in c.listCores()]:
             core = c.getCore('Nvidia CUDA')
-            print "Testing GPU core '%s'..." % core.name
+            print "Testing GPU core '%s' (Device '%s')" % (core.name, core.devicename)
             t = time.time()
             res = sorted(core.solve('foo', pws))
             t = time.time() - t
             print "%i PMKs in %.2f seconds: %.2f PMKs/s" % (len(pws), t, len(pws) / t)
             md = md5.new()
             map(md.update, [x[1] for x in res])
-            print "Result hash: %s" % md.hexdigest(), {True: "OK", False: "FAILED"}[md.hexdigest() == "ef747d123821851a9bd1d1e94ba048ac"]
+            print "Result hash:", {True: "OK", False: "FAILED"}[md.hexdigest() == "ef747d123821851a9bd1d1e94ba048ac"]
             print ""
 
 if __name__ == "__main__":