/**
- * Author......: Jens Steube <jens.steube@gmail.com>
+ * Authors.....: Jens Steube <jens.steube@gmail.com>
+ * Gabriele Gristina <matrix@hashcat.net>
+ *
* License.....: MIT
*/
#endif
+/**
+ * libraries stuff
+ */
+
+#ifdef _WIN
+#define hc_dlopen LoadLibrary
+#define hc_dlclose FreeLibrary
+#define hc_dlsym GetProcAddress
+#else
+#define hc_dlopen dlopen
+#define hc_dlclose dlclose
+#define hc_dlsym dlsym
+#endif
+
/**
* system stuff
*/
#define CL_VENDOR_NV "NVIDIA Corporation"
#define CL_VENDOR_AMD "Advanced Micro Devices, Inc."
+#define CL_VENDOR_APPLE "Apple"
#define CL_VENDOR_POCL "The pocl project"
#define VENDOR_ID_AMD 4098
#define VENDOR_ID_NV 4318
+#define VENDOR_ID_APPLE 16925952
#define VENDOR_ID_GENERIC 9999
#define BLOCK_SIZE 64
#define PW_LENGTH_MAX_7400 15
/**
- * device accel / loops macro
+ * device accel macro
*/
+#ifdef OSX
+#define KERNEL_ACCEL_5000 16
+#define KERNEL_ACCEL_6100 1
+#define KERNEL_ACCEL_6211 4
+#define KERNEL_ACCEL_6231 1
+#define KERNEL_ACCEL_6241 4
+#define KERNEL_ACCEL_8200 1
+#define KERNEL_ACCEL_8700 2
+#define KERNEL_ACCEL_9500 1
+#define KERNEL_ACCEL_9600 1
+#define KERNEL_ACCEL_10500 4
+#define KERNEL_ACCEL_11300 1
+#define KERNEL_ACCEL_11600 1
+#define KERNEL_ACCEL_11700 1
+#define KERNEL_ACCEL_11800 1
+#define KERNEL_ACCEL_12200 1
+#define KERNEL_ACCEL_12400 1
+#define KERNEL_ACCEL_12500 1
+#define KERNEL_ACCEL_13000 1
+#else
+#define KERNEL_ACCEL_5000 64
+#define KERNEL_ACCEL_6100 8
+#define KERNEL_ACCEL_6211 16
+#define KERNEL_ACCEL_6231 4
+#define KERNEL_ACCEL_6241 32
+#define KERNEL_ACCEL_8200 2
+#define KERNEL_ACCEL_8700 8
+#define KERNEL_ACCEL_9500 8
+#define KERNEL_ACCEL_9600 2
+#define KERNEL_ACCEL_10500 64
+#define KERNEL_ACCEL_11300 2
+#define KERNEL_ACCEL_11600 2
+#define KERNEL_ACCEL_11700 4
+#define KERNEL_ACCEL_11800 4
+#define KERNEL_ACCEL_12200 2
+#define KERNEL_ACCEL_12400 64
+#define KERNEL_ACCEL_12500 8
+#define KERNEL_ACCEL_13000 8
+#endif // OSX
+
#define KERNEL_ACCEL_0 128
#define KERNEL_ACCEL_10 128
#define KERNEL_ACCEL_11 128
#define KERNEL_ACCEL_4700 64
#define KERNEL_ACCEL_4800 128
#define KERNEL_ACCEL_4900 64
-#define KERNEL_ACCEL_5000 64
#define KERNEL_ACCEL_5100 64
#define KERNEL_ACCEL_5200 8
#define KERNEL_ACCEL_5300 32
#define KERNEL_ACCEL_5700 64
#define KERNEL_ACCEL_5800 8
#define KERNEL_ACCEL_6000 64
-#define KERNEL_ACCEL_6100 8
-#define KERNEL_ACCEL_6211 16
#define KERNEL_ACCEL_6212 8
#define KERNEL_ACCEL_6213 8
#define KERNEL_ACCEL_6221 4
#define KERNEL_ACCEL_6222 4
#define KERNEL_ACCEL_6223 4
-#define KERNEL_ACCEL_6231 4
#define KERNEL_ACCEL_6232 4
#define KERNEL_ACCEL_6233 4
-#define KERNEL_ACCEL_6241 32
#define KERNEL_ACCEL_6242 16
#define KERNEL_ACCEL_6243 16
#define KERNEL_ACCEL_6300 8
#define KERNEL_ACCEL_7900 2
#define KERNEL_ACCEL_8000 8
#define KERNEL_ACCEL_8100 64
-#define KERNEL_ACCEL_8200 2
#define KERNEL_ACCEL_8300 64
#define KERNEL_ACCEL_8400 64
#define KERNEL_ACCEL_8500 64
#define KERNEL_ACCEL_8600 8
-#define KERNEL_ACCEL_8700 8
#define KERNEL_ACCEL_8800 8
#define KERNEL_ACCEL_8900 16
#define KERNEL_ACCEL_9000 2
#define KERNEL_ACCEL_9200 2
#define KERNEL_ACCEL_9300 2
#define KERNEL_ACCEL_9400 8
-#define KERNEL_ACCEL_9500 8
-#define KERNEL_ACCEL_9600 2
#define KERNEL_ACCEL_9700 8
#define KERNEL_ACCEL_9710 8
#define KERNEL_ACCEL_9720 8
#define KERNEL_ACCEL_10400 8
#define KERNEL_ACCEL_10410 8
#define KERNEL_ACCEL_10420 8
-#define KERNEL_ACCEL_10500 64
#define KERNEL_ACCEL_10600 64
#define KERNEL_ACCEL_10700 1
#define KERNEL_ACCEL_10800 64
#define KERNEL_ACCEL_11000 64
#define KERNEL_ACCEL_11100 64
#define KERNEL_ACCEL_11200 64
-#define KERNEL_ACCEL_11300 2
#define KERNEL_ACCEL_11400 8
#define KERNEL_ACCEL_11500 128
-#define KERNEL_ACCEL_11600 2
-#define KERNEL_ACCEL_11700 4
-#define KERNEL_ACCEL_11800 4
#define KERNEL_ACCEL_11900 2
#define KERNEL_ACCEL_12000 2
#define KERNEL_ACCEL_12100 2
-#define KERNEL_ACCEL_12200 2
#define KERNEL_ACCEL_12300 2
-#define KERNEL_ACCEL_12400 64
-#define KERNEL_ACCEL_12500 8
#define KERNEL_ACCEL_12600 32
#define KERNEL_ACCEL_12700 64
#define KERNEL_ACCEL_12800 64
#define KERNEL_ACCEL_12900 8
-#define KERNEL_ACCEL_13000 8
+/**
+ * device loops macro
+ */
+
+#ifdef OSX
+#define KERNEL_LOOPS_0 2
+#define KERNEL_LOOPS_10 2
+#define KERNEL_LOOPS_11 2
+#define KERNEL_LOOPS_12 2
+#define KERNEL_LOOPS_20 2
+#define KERNEL_LOOPS_21 2
+#define KERNEL_LOOPS_22 2
+#define KERNEL_LOOPS_23 2
+#define KERNEL_LOOPS_30 2
+#define KERNEL_LOOPS_40 2
+#define KERNEL_LOOPS_50 2
+#define KERNEL_LOOPS_60 2
+#define KERNEL_LOOPS_100 2
+#define KERNEL_LOOPS_101 2
+#define KERNEL_LOOPS_110 2
+#define KERNEL_LOOPS_111 2
+#define KERNEL_LOOPS_112 2
+#define KERNEL_LOOPS_120 2
+#define KERNEL_LOOPS_121 2
+#define KERNEL_LOOPS_122 2
+#define KERNEL_LOOPS_124 2
+#define KERNEL_LOOPS_130 2
+#define KERNEL_LOOPS_131 2
+#define KERNEL_LOOPS_132 2
+#define KERNEL_LOOPS_133 2
+#define KERNEL_LOOPS_140 2
+#define KERNEL_LOOPS_141 2
+#define KERNEL_LOOPS_150 2
+#define KERNEL_LOOPS_160 2
+#define KERNEL_LOOPS_190 2
+#define KERNEL_LOOPS_200 2
+#define KERNEL_LOOPS_300 2
+#define KERNEL_LOOPS_900 2
+#define KERNEL_LOOPS_1000 2
+#define KERNEL_LOOPS_1100 2
+#define KERNEL_LOOPS_1400 2
+#define KERNEL_LOOPS_1410 2
+#define KERNEL_LOOPS_1420 2
+#define KERNEL_LOOPS_1421 2
+#define KERNEL_LOOPS_1430 2
+#define KERNEL_LOOPS_1440 2
+#define KERNEL_LOOPS_1441 2
+#define KERNEL_LOOPS_1450 2
+#define KERNEL_LOOPS_1460 2
+#define KERNEL_LOOPS_1700 2
+#define KERNEL_LOOPS_1710 2
+#define KERNEL_LOOPS_1711 2
+#define KERNEL_LOOPS_1720 2
+#define KERNEL_LOOPS_1722 2
+#define KERNEL_LOOPS_1730 2
+#define KERNEL_LOOPS_1731 2
+#define KERNEL_LOOPS_1740 2
+#define KERNEL_LOOPS_1750 2
+#define KERNEL_LOOPS_1760 2
+#define KERNEL_LOOPS_2400 2
+#define KERNEL_LOOPS_2410 2
+#define KERNEL_LOOPS_2600 2
+#define KERNEL_LOOPS_2611 2
+#define KERNEL_LOOPS_2612 2
+#define KERNEL_LOOPS_2711 2
+#define KERNEL_LOOPS_2811 2
+#define KERNEL_LOOPS_3100 2
+#define KERNEL_LOOPS_3200 4
+#define KERNEL_LOOPS_3710 2
+#define KERNEL_LOOPS_3711 2
+#define KERNEL_LOOPS_3800 2
+#define KERNEL_LOOPS_4300 2
+#define KERNEL_LOOPS_4400 2
+#define KERNEL_LOOPS_4500 2
+#define KERNEL_LOOPS_4700 2
+#define KERNEL_LOOPS_4800 2
+#define KERNEL_LOOPS_4900 2
+#define KERNEL_LOOPS_5000 2
+#define KERNEL_LOOPS_5100 2
+#define KERNEL_LOOPS_5300 2
+#define KERNEL_LOOPS_5400 2
+#define KERNEL_LOOPS_5500 2
+#define KERNEL_LOOPS_5600 2
+#define KERNEL_LOOPS_5700 2
+#define KERNEL_LOOPS_6000 2
+#define KERNEL_LOOPS_6100 2
+#define KERNEL_LOOPS_6231 2
+#define KERNEL_LOOPS_6232 2
+#define KERNEL_LOOPS_6233 2
+#define KERNEL_LOOPS_6900 2
+#define KERNEL_LOOPS_7300 2
+#define KERNEL_LOOPS_7500 2
+#define KERNEL_LOOPS_7600 2
+#define KERNEL_LOOPS_7700 2
+#define KERNEL_LOOPS_7800 2
+#define KERNEL_LOOPS_8000 2
+#define KERNEL_LOOPS_8100 2
+#define KERNEL_LOOPS_8300 2
+#define KERNEL_LOOPS_8400 2
+#define KERNEL_LOOPS_8500 2
+#define KERNEL_LOOPS_8600 2
+#define KERNEL_LOOPS_8700 4
+#define KERNEL_LOOPS_9700 2
+#define KERNEL_LOOPS_9710 8
+#define KERNEL_LOOPS_9720 8
+#define KERNEL_LOOPS_9800 2
+#define KERNEL_LOOPS_9810 2
+#define KERNEL_LOOPS_9820 2
+#define KERNEL_LOOPS_9900 2
+#define KERNEL_LOOPS_10100 2
+#define KERNEL_LOOPS_10200 2
+#define KERNEL_LOOPS_10400 2
+#define KERNEL_LOOPS_10410 2
+#define KERNEL_LOOPS_10420 2
+#define KERNEL_LOOPS_10600 2
+#define KERNEL_LOOPS_10700 2
+#define KERNEL_LOOPS_10800 2
+#define KERNEL_LOOPS_11000 2
+#define KERNEL_LOOPS_11100 2
+#define KERNEL_LOOPS_11200 2
+#define KERNEL_LOOPS_11400 2
+#define KERNEL_LOOPS_11500 2
+#define KERNEL_LOOPS_11700 8
+#define KERNEL_LOOPS_11800 8
+#define KERNEL_LOOPS_12600 2
+#else
#define KERNEL_LOOPS_0 256
#define KERNEL_LOOPS_10 256
#define KERNEL_LOOPS_11 256
#define KERNEL_LOOPS_190 128
#define KERNEL_LOOPS_200 128
#define KERNEL_LOOPS_300 64
-#define KERNEL_LOOPS_400 256
-#define KERNEL_LOOPS_500 256
-#define KERNEL_LOOPS_501 256
#define KERNEL_LOOPS_900 256
-#define KERNEL_LOOPS_910 256
#define KERNEL_LOOPS_1000 256
#define KERNEL_LOOPS_1100 128
#define KERNEL_LOOPS_1400 64
#define KERNEL_LOOPS_1441 64
#define KERNEL_LOOPS_1450 32
#define KERNEL_LOOPS_1460 32
-#define KERNEL_LOOPS_1500 256
-#define KERNEL_LOOPS_1600 256
#define KERNEL_LOOPS_1700 32
#define KERNEL_LOOPS_1710 32
#define KERNEL_LOOPS_1711 32
#define KERNEL_LOOPS_1740 32
#define KERNEL_LOOPS_1750 16
#define KERNEL_LOOPS_1760 16
-#define KERNEL_LOOPS_1800 16
-#define KERNEL_LOOPS_2100 256
#define KERNEL_LOOPS_2400 256
#define KERNEL_LOOPS_2410 256
-#define KERNEL_LOOPS_2500 256
#define KERNEL_LOOPS_2600 128
#define KERNEL_LOOPS_2611 128
#define KERNEL_LOOPS_2612 128
#define KERNEL_LOOPS_2711 64
#define KERNEL_LOOPS_2811 64
-#define KERNEL_LOOPS_3000 256
#define KERNEL_LOOPS_3100 16
#define KERNEL_LOOPS_3200 16
#define KERNEL_LOOPS_3710 128
#define KERNEL_LOOPS_4900 128
#define KERNEL_LOOPS_5000 64
#define KERNEL_LOOPS_5100 256
-#define KERNEL_LOOPS_5200 256
#define KERNEL_LOOPS_5300 32
#define KERNEL_LOOPS_5400 32
#define KERNEL_LOOPS_5500 128
#define KERNEL_LOOPS_5600 64
#define KERNEL_LOOPS_5700 64
-#define KERNEL_LOOPS_5800 256
#define KERNEL_LOOPS_6000 64
#define KERNEL_LOOPS_6100 64
-#define KERNEL_LOOPS_6211 200
-#define KERNEL_LOOPS_6212 200
-#define KERNEL_LOOPS_6213 200
-#define KERNEL_LOOPS_6221 200
-#define KERNEL_LOOPS_6222 200
-#define KERNEL_LOOPS_6223 200
#define KERNEL_LOOPS_6231 200
#define KERNEL_LOOPS_6232 200
#define KERNEL_LOOPS_6233 200
-#define KERNEL_LOOPS_6241 200
-#define KERNEL_LOOPS_6242 200
-#define KERNEL_LOOPS_6243 200
-#define KERNEL_LOOPS_6300 256
-#define KERNEL_LOOPS_6400 256
-#define KERNEL_LOOPS_6500 256
-#define KERNEL_LOOPS_6600 200
-#define KERNEL_LOOPS_6700 256
-#define KERNEL_LOOPS_6800 200
#define KERNEL_LOOPS_6900 64
-#define KERNEL_LOOPS_7100 256
-#define KERNEL_LOOPS_7200 200
#define KERNEL_LOOPS_7300 64
-#define KERNEL_LOOPS_7400 200
#define KERNEL_LOOPS_7500 16
#define KERNEL_LOOPS_7600 128
#define KERNEL_LOOPS_7700 128
#define KERNEL_LOOPS_7800 64
-#define KERNEL_LOOPS_7900 256
#define KERNEL_LOOPS_8000 64
#define KERNEL_LOOPS_8100 128
-#define KERNEL_LOOPS_8200 200
#define KERNEL_LOOPS_8300 64
#define KERNEL_LOOPS_8400 64
#define KERNEL_LOOPS_8500 16
#define KERNEL_LOOPS_8600 16
#define KERNEL_LOOPS_8700 16
-#define KERNEL_LOOPS_8800 256
-#define KERNEL_LOOPS_8900 1
-#define KERNEL_LOOPS_9000 16
-#define KERNEL_LOOPS_9100 256
-#define KERNEL_LOOPS_9200 200
-#define KERNEL_LOOPS_9300 1
-#define KERNEL_LOOPS_9400 200
-#define KERNEL_LOOPS_9500 200
-#define KERNEL_LOOPS_9600 200
#define KERNEL_LOOPS_9700 200
#define KERNEL_LOOPS_9710 200
#define KERNEL_LOOPS_9720 200
#define KERNEL_LOOPS_9800 200
-#define KERNEL_LOOPS_9810 200
#define KERNEL_LOOPS_9820 200
+#define KERNEL_LOOPS_9810 200
#define KERNEL_LOOPS_9900 256
-#define KERNEL_LOOPS_10000 200
#define KERNEL_LOOPS_10100 512
#define KERNEL_LOOPS_10200 64
-#define KERNEL_LOOPS_10300 128
#define KERNEL_LOOPS_10400 256
#define KERNEL_LOOPS_10410 256
#define KERNEL_LOOPS_10420 256
-#define KERNEL_LOOPS_10500 64
#define KERNEL_LOOPS_10600 64
#define KERNEL_LOOPS_10700 64
#define KERNEL_LOOPS_10800 32
-#define KERNEL_LOOPS_10900 200
#define KERNEL_LOOPS_11000 256
#define KERNEL_LOOPS_11100 128
#define KERNEL_LOOPS_11200 128
-#define KERNEL_LOOPS_11300 256
#define KERNEL_LOOPS_11400 128
#define KERNEL_LOOPS_11500 256
-#define KERNEL_LOOPS_11600 512
#define KERNEL_LOOPS_11700 64
#define KERNEL_LOOPS_11800 64
+#define KERNEL_LOOPS_12600 32
+#endif // OSX
+
+#define KERNEL_LOOPS_400 256
+#define KERNEL_LOOPS_500 256
+#define KERNEL_LOOPS_501 256
+#define KERNEL_LOOPS_910 256
+#define KERNEL_LOOPS_1500 256
+#define KERNEL_LOOPS_1600 256
+#define KERNEL_LOOPS_1800 16
+#define KERNEL_LOOPS_2100 256
+#define KERNEL_LOOPS_2500 256
+#define KERNEL_LOOPS_3000 256
+#define KERNEL_LOOPS_5200 256
+#define KERNEL_LOOPS_5800 256
+#define KERNEL_LOOPS_6211 200
+#define KERNEL_LOOPS_6212 200
+#define KERNEL_LOOPS_6213 200
+#define KERNEL_LOOPS_6221 200
+#define KERNEL_LOOPS_6222 200
+#define KERNEL_LOOPS_6223 200
+#define KERNEL_LOOPS_6241 200
+#define KERNEL_LOOPS_6242 200
+#define KERNEL_LOOPS_6243 200
+#define KERNEL_LOOPS_6300 256
+#define KERNEL_LOOPS_6400 256
+#define KERNEL_LOOPS_6500 256
+#define KERNEL_LOOPS_6600 200
+#define KERNEL_LOOPS_6700 256
+#define KERNEL_LOOPS_6800 200
+#define KERNEL_LOOPS_7100 256
+#define KERNEL_LOOPS_7200 200
+#define KERNEL_LOOPS_7400 200
+#define KERNEL_LOOPS_7900 256
+#define KERNEL_LOOPS_8200 200
+#define KERNEL_LOOPS_8800 256
+#define KERNEL_LOOPS_8900 1
+#define KERNEL_LOOPS_9000 16
+#define KERNEL_LOOPS_9100 256
+#define KERNEL_LOOPS_9200 200
+#define KERNEL_LOOPS_9300 1
+#define KERNEL_LOOPS_9400 200
+#define KERNEL_LOOPS_9500 200
+#define KERNEL_LOOPS_9600 200
+#define KERNEL_LOOPS_10000 200
+#define KERNEL_LOOPS_10300 128
+#define KERNEL_LOOPS_10500 64
+#define KERNEL_LOOPS_10900 200
+#define KERNEL_LOOPS_11300 256
+#define KERNEL_LOOPS_11600 512
#define KERNEL_LOOPS_11900 200
#define KERNEL_LOOPS_12000 200
#define KERNEL_LOOPS_12100 200
#define KERNEL_LOOPS_12300 256
#define KERNEL_LOOPS_12400 256
#define KERNEL_LOOPS_12500 256
-#define KERNEL_LOOPS_12600 32
#define KERNEL_LOOPS_12700 10
#define KERNEL_LOOPS_12800 100
#define KERNEL_LOOPS_12900 64
*/
#define DGST_SIZE_0 0
-#define DGST_SIZE_4_2 (2 * sizeof (uint)) // 8
-#define DGST_SIZE_4_4 (4 * sizeof (uint)) // 16
-#define DGST_SIZE_4_5 (5 * sizeof (uint)) // 20
-#define DGST_SIZE_4_6 (6 * sizeof (uint)) // 24
-#define DGST_SIZE_4_8 (8 * sizeof (uint)) // 32
-#define DGST_SIZE_4_16 (16 * sizeof (uint)) // 64 !!!
-#define DGST_SIZE_4_32 (32 * sizeof (uint)) // 128 !!!
-#define DGST_SIZE_4_64 (64 * sizeof (uint)) // 256
+#define DGST_SIZE_4_2 (2 * sizeof (uint)) // 8
+#define DGST_SIZE_4_4 (4 * sizeof (uint)) // 16
+#define DGST_SIZE_4_5 (5 * sizeof (uint)) // 20
+#define DGST_SIZE_4_6 (6 * sizeof (uint)) // 24
+#define DGST_SIZE_4_8 (8 * sizeof (uint)) // 32
+#define DGST_SIZE_4_16 (16 * sizeof (uint)) // 64 !!!
+#define DGST_SIZE_4_32 (32 * sizeof (uint)) // 128 !!!
+#define DGST_SIZE_4_64 (64 * sizeof (uint)) // 256
#define DGST_SIZE_8_8 (8 * sizeof (u64)) // 64 !!!
#define DGST_SIZE_8_16 (16 * sizeof (u64)) // 128 !!!
#define DGST_SIZE_8_25 (25 * sizeof (u64)) // 200
/**
- * Author......: Jens Steube <jens.steube@gmail.com>
+ * Authors.....: Jens Steube <jens.steube@gmail.com>
+ * Gabriele Gristina <matrix@hashcat.net>
+ *
* License.....: MIT
*/
return "CL_UNKNOWN_ERROR";
}
-void hc_clEnqueueNDRangeKernel (cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
+#define LOAD_FUNC(ptr,name,type) \
+ ptr->name = (type) hc_dlsym (ptr->lib, #name); \
+ if (!ptr->name) { \
+ log_error ("ERROR: #name is missing from opencl shared library"); \
+ exit (-1); \
+ }
+
+void ocl_init (OCL_PTR *ocl)
+{
+ memset (ocl, 0, sizeof (hc_opencl_lib_t));
+
+ #ifdef _WIN
+ ocl->lib = hc_dlopen ("OpenCL");
+ #elif OSX
+ ocl->lib = hc_dlopen ("/System/Library/Frameworks/OpenCL.framework/OpenCL", RTLD_NOW);
+ #else
+ ocl->lib = hc_dlopen ("libOpenCL.so", RTLD_NOW);
+ #endif
+
+ if (!ocl->lib)
+ {
+ log_error ("ERROR: cannot load opencl library");
+
+ exit (-1);
+ }
+
+ LOAD_FUNC(ocl, clBuildProgram, OCL_CLBUILDPROGRAM)
+ LOAD_FUNC(ocl, clCreateBuffer, OCL_CLCREATEBUFFER)
+ LOAD_FUNC(ocl, clCreateCommandQueue, OCL_CLCREATECOMMANDQUEUE)
+ LOAD_FUNC(ocl, clCreateContext, OCL_CLCREATECONTEXT)
+ LOAD_FUNC(ocl, clCreateKernel, OCL_CLCREATEKERNEL)
+ LOAD_FUNC(ocl, clCreateProgramWithBinary, OCL_CLCREATEPROGRAMWITHBINARY)
+ LOAD_FUNC(ocl, clCreateProgramWithSource, OCL_CLCREATEPROGRAMWITHSOURCE)
+ LOAD_FUNC(ocl, clEnqueueCopyBuffer, OCL_CLENQUEUECOPYBUFFER)
+ LOAD_FUNC(ocl, clEnqueueFillBuffer, OCL_CLENQUEUEFILLBUFFER)
+ LOAD_FUNC(ocl, clEnqueueMapBuffer, OCL_CLENQUEUEMAPBUFFER)
+ LOAD_FUNC(ocl, clEnqueueNDRangeKernel, OCL_CLENQUEUENDRANGEKERNEL)
+ LOAD_FUNC(ocl, clEnqueueReadBuffer, OCL_CLENQUEUEREADBUFFER)
+ LOAD_FUNC(ocl, clEnqueueUnmapMemObject, OCL_CLENQUEUEUNMAPMEMOBJECT)
+ LOAD_FUNC(ocl, clEnqueueWriteBuffer, OCL_CLENQUEUEWRITEBUFFER)
+ LOAD_FUNC(ocl, clFinish, OCL_CLFINISH)
+ LOAD_FUNC(ocl, clFlush, OCL_CLFLUSH)
+ LOAD_FUNC(ocl, clGetDeviceIDs, OCL_CLGETDEVICEIDS)
+ LOAD_FUNC(ocl, clGetDeviceInfo, OCL_CLGETDEVICEINFO)
+ LOAD_FUNC(ocl, clGetEventInfo, OCL_CLGETEVENTINFO)
+ LOAD_FUNC(ocl, clGetKernelWorkGroupInfo, OCL_CLGETKERNELWORKGROUPINFO)
+ LOAD_FUNC(ocl, clGetPlatformIDs, OCL_CLGETPLATFORMIDS)
+ LOAD_FUNC(ocl, clGetPlatformInfo, OCL_CLGETPLATFORMINFO)
+ LOAD_FUNC(ocl, clGetProgramBuildInfo, OCL_CLGETPROGRAMBUILDINFO)
+ LOAD_FUNC(ocl, clGetProgramInfo, OCL_CLGETPROGRAMINFO)
+ LOAD_FUNC(ocl, clReleaseCommandQueue, OCL_CLRELEASECOMMANDQUEUE)
+ LOAD_FUNC(ocl, clReleaseContext, OCL_CLRELEASECONTEXT)
+ LOAD_FUNC(ocl, clReleaseKernel, OCL_CLRELEASEKERNEL)
+ LOAD_FUNC(ocl, clReleaseMemObject, OCL_CLRELEASEMEMOBJECT)
+ LOAD_FUNC(ocl, clReleaseProgram, OCL_CLRELEASEPROGRAM)
+ LOAD_FUNC(ocl, clSetKernelArg, OCL_CLSETKERNELARG)
+}
+
+void ocl_close (OCL_PTR *ocl)
{
- cl_int CL_err = clEnqueueNDRangeKernel (command_queue, kernel, work_dim, global_work_offset, global_work_size, local_work_size, num_events_in_wait_list, event_wait_list, event);
+ if (ocl)
+ {
+ if (ocl->lib)
+ hc_dlclose (ocl->lib);
+
+ free (ocl);
+ }
+}
+
+void hc_clEnqueueNDRangeKernel (OCL_PTR *ocl, cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
+{
+ cl_int CL_err = ocl->clEnqueueNDRangeKernel (command_queue, kernel, work_dim, global_work_offset, global_work_size, local_work_size, num_events_in_wait_list, event_wait_list, event);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clGetEventInfo (cl_event event, cl_event_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
+void hc_clGetEventInfo (OCL_PTR *ocl, cl_event event, cl_event_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
{
- cl_int CL_err = clGetEventInfo (event, param_name, param_value_size, param_value, param_value_size_ret);
+ cl_int CL_err = ocl->clGetEventInfo (event, param_name, param_value_size, param_value, param_value_size_ret);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clFlush (cl_command_queue command_queue)
+void hc_clFlush (OCL_PTR *ocl, cl_command_queue command_queue)
{
- cl_int CL_err = clFlush (command_queue);
+ cl_int CL_err = ocl->clFlush (command_queue);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clFinish (cl_command_queue command_queue)
+void hc_clFinish (OCL_PTR *ocl, cl_command_queue command_queue)
{
- cl_int CL_err = clFinish (command_queue);
+ cl_int CL_err = ocl->clFinish (command_queue);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clSetKernelArg (cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void *arg_value)
+void hc_clSetKernelArg (OCL_PTR *ocl, cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void *arg_value)
{
- cl_int CL_err = clSetKernelArg (kernel, arg_index, arg_size, arg_value);
+ cl_int CL_err = ocl->clSetKernelArg (kernel, arg_index, arg_size, arg_value);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clEnqueueWriteBuffer (cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write, size_t offset, size_t cb, const void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
+void hc_clEnqueueWriteBuffer (OCL_PTR *ocl, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write, size_t offset, size_t cb, const void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
{
- cl_int CL_err = clEnqueueWriteBuffer (command_queue, buffer, blocking_write, offset, cb, ptr, num_events_in_wait_list, event_wait_list, event);
+ cl_int CL_err = ocl->clEnqueueWriteBuffer (command_queue, buffer, blocking_write, offset, cb, ptr, num_events_in_wait_list, event_wait_list, event);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clEnqueueCopyBuffer (cl_command_queue command_queue, cl_mem src_buffer, cl_mem dst_buffer, size_t src_offset, size_t dst_offset, size_t cb, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
+void hc_clEnqueueCopyBuffer (OCL_PTR *ocl, cl_command_queue command_queue, cl_mem src_buffer, cl_mem dst_buffer, size_t src_offset, size_t dst_offset, size_t cb, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
{
- cl_int CL_err = clEnqueueCopyBuffer (command_queue, src_buffer, dst_buffer, src_offset, dst_offset, cb, num_events_in_wait_list, event_wait_list, event);
+ cl_int CL_err = ocl->clEnqueueCopyBuffer (command_queue, src_buffer, dst_buffer, src_offset, dst_offset, cb, num_events_in_wait_list, event_wait_list, event);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clEnqueueReadBuffer (cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read, size_t offset, size_t cb, void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
+void hc_clEnqueueReadBuffer (OCL_PTR *ocl, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read, size_t offset, size_t cb, void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
{
- cl_int CL_err = clEnqueueReadBuffer (command_queue, buffer, blocking_read, offset, cb, ptr, num_events_in_wait_list, event_wait_list, event);
+ cl_int CL_err = ocl->clEnqueueReadBuffer (command_queue, buffer, blocking_read, offset, cb, ptr, num_events_in_wait_list, event_wait_list, event);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clGetPlatformIDs (cl_uint num_entries, cl_platform_id *platforms, cl_uint *num_platforms)
+void hc_clGetPlatformIDs (OCL_PTR *ocl, cl_uint num_entries, cl_platform_id *platforms, cl_uint *num_platforms)
{
- cl_int CL_err = clGetPlatformIDs (num_entries, platforms, num_platforms);
+ cl_int CL_err = ocl->clGetPlatformIDs (num_entries, platforms, num_platforms);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clGetPlatformInfo (cl_platform_id platform, cl_platform_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
+void hc_clGetPlatformInfo (OCL_PTR *ocl, cl_platform_id platform, cl_platform_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
{
- cl_int CL_err = clGetPlatformInfo (platform, param_name, param_value_size, param_value, param_value_size_ret);
+ cl_int CL_err = ocl->clGetPlatformInfo (platform, param_name, param_value_size, param_value, param_value_size_ret);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clGetDeviceIDs (cl_platform_id platform, cl_device_type device_type, cl_uint num_entries, cl_device_id *devices, cl_uint *num_devices)
+void hc_clGetDeviceIDs (OCL_PTR *ocl, cl_platform_id platform, cl_device_type device_type, cl_uint num_entries, cl_device_id *devices, cl_uint *num_devices)
{
- cl_int CL_err = clGetDeviceIDs (platform, device_type, num_entries, devices, num_devices);
+ cl_int CL_err = ocl->clGetDeviceIDs (platform, device_type, num_entries, devices, num_devices);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clGetDeviceInfo (cl_device_id device, cl_device_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
+void hc_clGetDeviceInfo (OCL_PTR *ocl, cl_device_id device, cl_device_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
{
- cl_int CL_err = clGetDeviceInfo (device, param_name, param_value_size, param_value, param_value_size_ret);
+ cl_int CL_err = ocl->clGetDeviceInfo (device, param_name, param_value_size, param_value, param_value_size_ret);
if (CL_err != CL_SUCCESS)
{
}
}
-cl_context hc_clCreateContext (cl_context_properties *properties, cl_uint num_devices, const cl_device_id *devices, void (CL_CALLBACK *pfn_notify) (const char *, const void *, size_t, void *), void *user_data)
+cl_context hc_clCreateContext (OCL_PTR *ocl, cl_context_properties *properties, cl_uint num_devices, const cl_device_id *devices, void (CL_CALLBACK *pfn_notify) (const char *, const void *, size_t, void *), void *user_data)
{
cl_int CL_err;
- cl_context context = clCreateContext (properties, num_devices, devices, pfn_notify, user_data, &CL_err);
+ cl_context context = ocl->clCreateContext (properties, num_devices, devices, pfn_notify, user_data, &CL_err);
if (CL_err != CL_SUCCESS)
{
return (context);
}
-
-cl_command_queue hc_clCreateCommandQueue (cl_context context, cl_device_id device, cl_command_queue_properties properties)
+cl_command_queue hc_clCreateCommandQueue (OCL_PTR *ocl, cl_context context, cl_device_id device, cl_command_queue_properties properties)
{
cl_int CL_err;
- cl_command_queue command_queue = clCreateCommandQueue (context, device, properties, &CL_err);
+ cl_command_queue command_queue = ocl->clCreateCommandQueue (context, device, properties, &CL_err);
if (CL_err != CL_SUCCESS)
{
}
*/
-cl_mem hc_clCreateBuffer (cl_context context, cl_mem_flags flags, size_t size, void *host_ptr)
+cl_mem hc_clCreateBuffer (OCL_PTR *ocl, cl_context context, cl_mem_flags flags, size_t size, void *host_ptr)
{
cl_int CL_err;
- cl_mem mem = clCreateBuffer (context, flags, size, host_ptr, &CL_err);
+ cl_mem mem = ocl->clCreateBuffer (context, flags, size, host_ptr, &CL_err);
if (CL_err != CL_SUCCESS)
{
return (mem);
}
-cl_program hc_clCreateProgramWithSource (cl_context context, cl_uint count, const char **strings, const size_t *lengths)
+cl_program hc_clCreateProgramWithSource (OCL_PTR *ocl, cl_context context, cl_uint count, const char **strings, const size_t *lengths)
{
cl_int CL_err;
- cl_program program = clCreateProgramWithSource (context, count, strings, lengths, &CL_err);
+ cl_program program = ocl->clCreateProgramWithSource (context, count, strings, lengths, &CL_err);
if (CL_err != CL_SUCCESS)
{
return (program);
}
-cl_program hc_clCreateProgramWithBinary (cl_context context, cl_uint num_devices, const cl_device_id *device_list, const size_t *lengths, const unsigned char **binaries, cl_int *binary_status)
+cl_program hc_clCreateProgramWithBinary (OCL_PTR *ocl, cl_context context, cl_uint num_devices, const cl_device_id *device_list, const size_t *lengths, const unsigned char **binaries, cl_int *binary_status)
{
cl_int CL_err;
- cl_program program = clCreateProgramWithBinary (context, num_devices, device_list, lengths, binaries, binary_status, &CL_err);
+ cl_program program = ocl->clCreateProgramWithBinary (context, num_devices, device_list, lengths, binaries, binary_status, &CL_err);
if (CL_err != CL_SUCCESS)
{
return (program);
}
-void hc_clBuildProgram (cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data)
+void hc_clBuildProgram (OCL_PTR *ocl, cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data)
{
- cl_int CL_err = clBuildProgram (program, num_devices, device_list, options, pfn_notify, user_data);
+ cl_int CL_err = ocl->clBuildProgram (program, num_devices, device_list, options, pfn_notify, user_data);
if (CL_err != CL_SUCCESS)
{
}
}
-cl_kernel hc_clCreateKernel (cl_program program, const char *kernel_name)
+cl_kernel hc_clCreateKernel (OCL_PTR *ocl, cl_program program, const char *kernel_name)
{
cl_int CL_err;
- cl_kernel kernel = clCreateKernel (program, kernel_name, &CL_err);
+ cl_kernel kernel = ocl->clCreateKernel (program, kernel_name, &CL_err);
if (CL_err != CL_SUCCESS)
{
return (kernel);
}
-void hc_clReleaseMemObject (cl_mem mem)
+void hc_clReleaseMemObject (OCL_PTR *ocl, cl_mem mem)
{
- cl_int CL_err = clReleaseMemObject (mem);
+ cl_int CL_err = ocl->clReleaseMemObject (mem);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clReleaseKernel (cl_kernel kernel)
+void hc_clReleaseKernel (OCL_PTR *ocl, cl_kernel kernel)
{
- cl_int CL_err = clReleaseKernel (kernel);
+ cl_int CL_err = ocl->clReleaseKernel (kernel);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clReleaseProgram (cl_program program)
+void hc_clReleaseProgram (OCL_PTR *ocl, cl_program program)
{
- cl_int CL_err = clReleaseProgram (program);
+ cl_int CL_err = ocl->clReleaseProgram (program);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clReleaseCommandQueue (cl_command_queue command_queue)
+void hc_clReleaseCommandQueue (OCL_PTR *ocl, cl_command_queue command_queue)
{
- cl_int CL_err = clReleaseCommandQueue (command_queue);
+ cl_int CL_err = ocl->clReleaseCommandQueue (command_queue);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clReleaseContext (cl_context context)
+void hc_clReleaseContext (OCL_PTR *ocl, cl_context context)
{
- cl_int CL_err = clReleaseContext (context);
+ cl_int CL_err = ocl->clReleaseContext (context);
if (CL_err != CL_SUCCESS)
{
}
}
-void *hc_clEnqueueMapBuffer (cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read, cl_map_flags map_flags, size_t offset, size_t cb, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
+void *hc_clEnqueueMapBuffer (OCL_PTR *ocl, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read, cl_map_flags map_flags, size_t offset, size_t cb, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
{
cl_int CL_err;
- void *buf = clEnqueueMapBuffer (command_queue, buffer, blocking_read, map_flags, offset, cb, num_events_in_wait_list, event_wait_list, event, &CL_err);
+ void *buf = ocl->clEnqueueMapBuffer (command_queue, buffer, blocking_read, map_flags, offset, cb, num_events_in_wait_list, event_wait_list, event, &CL_err);
if (CL_err != CL_SUCCESS)
{
return buf;
}
-void hc_clEnqueueUnmapMemObject (cl_command_queue command_queue, cl_mem memobj, void *mapped_ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
+void hc_clEnqueueUnmapMemObject (OCL_PTR *ocl, cl_command_queue command_queue, cl_mem memobj, void *mapped_ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
{
- cl_int CL_err = clEnqueueUnmapMemObject (command_queue, memobj, mapped_ptr, num_events_in_wait_list, event_wait_list, event);
+ cl_int CL_err = ocl->clEnqueueUnmapMemObject (command_queue, memobj, mapped_ptr, num_events_in_wait_list, event_wait_list, event);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clEnqueueFillBuffer (cl_command_queue command_queue, cl_mem buffer, const void *pattern, size_t pattern_size, size_t offset, size_t size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
+void hc_clEnqueueFillBuffer (OCL_PTR *ocl, cl_command_queue command_queue, cl_mem buffer, const void *pattern, size_t pattern_size, size_t offset, size_t size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event)
{
- cl_int CL_err = clEnqueueFillBuffer (command_queue, buffer, pattern, pattern_size, offset, size, num_events_in_wait_list, event_wait_list, event);
+ cl_int CL_err = ocl->clEnqueueFillBuffer (command_queue, buffer, pattern, pattern_size, offset, size, num_events_in_wait_list, event_wait_list, event);
if (CL_err != CL_SUCCESS)
{
}
}
-void hc_clGetKernelWorkGroupInfo (cl_kernel kernel, cl_device_id device, cl_kernel_work_group_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
+void hc_clGetKernelWorkGroupInfo (OCL_PTR *ocl, cl_kernel kernel, cl_device_id device, cl_kernel_work_group_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
{
- cl_int CL_err = clGetKernelWorkGroupInfo (kernel, device, param_name, param_value_size, param_value, param_value_size_ret);
+ cl_int CL_err = ocl->clGetKernelWorkGroupInfo (kernel, device, param_name, param_value_size, param_value, param_value_size_ret);
if (CL_err != CL_SUCCESS)
{
exit (-1);
}
}
+
+void hc_clGetProgramBuildInfo (OCL_PTR *ocl, cl_program program, cl_device_id device, cl_program_build_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
+{
+ cl_int CL_err = ocl->clGetProgramBuildInfo (program, device, param_name, param_value_size, param_value, param_value_size_ret);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: %s : %d : %s\n", "clGetProgramBuildInfo()", CL_err, val2cstr_cl (CL_err));
+
+ exit (-1);
+ }
+}
+
+void hc_clGetProgramInfo (OCL_PTR *ocl, cl_program program, cl_program_info param_name, size_t param_value_size, void *param_value, size_t * param_value_size_ret)
+{
+ cl_int CL_err = ocl->clGetProgramInfo (program, param_name, param_value_size, param_value, param_value_size_ret);
+
+ if (CL_err != CL_SUCCESS)
+ {
+ log_error ("ERROR: %s : %d : %s\n", "clGetProgramInfo()", CL_err, val2cstr_cl (CL_err));
+
+ exit (-1);
+ }
+}
/**
- * Author......: Jens Steube <jens.steube@gmail.com>
+ * Authors.....: Jens Steube <jens.steube@gmail.com>
+ * Gabriele Gristina <matrix@hashcat.net>
+ *
* License.....: MIT
*/
#include <rp_kernel_on_cpu.h>
#include <getopt.h>
-const char *PROGNAME = "oclHashcat";
-const uint VERSION_BIN = 210;
-const uint RESTORE_MIN = 210;
+const char *PROGNAME = "oclHashcat";
+const uint VERSION_BIN = 210;
+const uint RESTORE_MIN = 210;
-#define INCR_RULES 10000
-#define INCR_SALTS 100000
-#define INCR_MASKS 1000
-#define INCR_POT 1000
+#define INCR_RULES 10000
+#define INCR_SALTS 100000
+#define INCR_MASKS 1000
+#define INCR_POT 1000
#define USAGE 0
#define VERSION 0
5000,
10100,
6000,
- 6100,
- 6900,
- 11700,
- 11800,
+ 6100, // broken in osx
+ 6900, // broken in osx
+ 11700, // broken in osx
+ 11800, // broken in osx
400,
8900,
11900,
12100,
23,
2500,
- 5300,
- 5400,
- 5500,
- 5600,
+ 5300, // broken in osx
+ 5400, // broken in osx
+ 5500, // broken in osx
+ 5600, // broken in osx
7300,
- 7500,
+ 7500, // broken in osx
8300,
- 11100,
+ 11100, // broken in osx
11200,
- 11400,
+ 11400, // broken in osx
121,
- 2611,
- 2711,
- 2811,
- 8400,
+ 2611, // broken in osx
+ 2711, // broken in osx
+ 2811, // broken in osx
+ 8400, // broken in osx
11,
- 2612,
+ 2612, // broken in osx
7900,
21,
11000,
124,
10000,
- 3711,
- 7600,
+ 3711, // broken in osx
+ 7600, // broken in osx
12,
131,
132,
1731,
200,
300,
- 3100,
+ 3100, // broken in osx
112,
12300,
- 8000,
+ 8000, // broken in osx
141,
1441,
1600,
- 12600,
+ 12600, // broken in osx
1421,
101,
111,
1711,
- 3000,
+ 3000, // broken in osx
1000,
1100,
2100,
12800,
- 1500,
- 12400,
+ 1500, // broken in osx
+ 12400, // broken in osx
500,
3200,
7400,
501,
5800,
8100,
- 8500,
+ 8500, // broken in osx
7200,
9900,
7700,
7800,
10300,
- 8600,
- 8700,
- 9100,
+ 8600, // broken in osx
+ 8700, // broken in osx
+ 9100, // broken in osx
133,
- 11600,
- 12500,
+ 11600, // broken in osx
+ 12500, // broken in osx
13000,
6211,
6221,
- 6231,
+ 6231, // broken in osx
6241,
- 8800,
+ 8800, // broken in osx
12900,
12200,
- 9700,
- 9710,
- 9800,
- 9810,
- 9400,
- 9500,
- 9600,
- 10400,
- 10410,
+ 9700, // broken in osx
+ 9710, // broken in osx
+ 9800, // broken in osx
+ 9810, // broken in osx
+ 9400, // broken in osx
+ 9500, // broken in osx
+ 9600, // broken in osx
+ 10400, // broken in osx
+ 10410, // broken in osx
10500,
10600,
- 10700,
+ 10700, // broken in osx
9000,
5200,
- 6800,
- 6600,
+ 6800, // broken in osx
+ 6600, // broken in osx
8200,
- 11300,
- 12700
+ 11300, // broken in osx
+ 12700 // broken in osx
};
/**
static void gidd_to_pw_t (hc_device_param_t *device_param, const u64 gidd, pw_t *pw)
{
- hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_pws_buf, CL_TRUE, gidd * sizeof (pw_t), sizeof (pw_t), pw, 0, NULL, NULL);
+ hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, gidd * sizeof (pw_t), sizeof (pw_t), pw, 0, NULL, NULL);
}
static void check_hash (hc_device_param_t *device_param, const uint salt_pos, const uint digest_pos)
plain_t plain;
- hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_plain_bufs, CL_TRUE, idx * sizeof (plain_t), sizeof (plain_t), &plain, 0, NULL, NULL);
+ hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_plain_bufs, CL_TRUE, idx * sizeof (plain_t), sizeof (plain_t), &plain, 0, NULL, NULL);
uint gidvid = plain.gidvid;
uint il_pos = plain.il_pos;
for (int i = 0, j = gidm; i < 16; i++, j++)
{
- plain_buf[i] = pw.hi1[0][j];
+ plain_buf[i] = pw.h.hi1[0][j];
}
plain_len = pw.pw_len;
for (int i = 0, j = gidm; i < 16; i++, j++)
{
- plain_buf[i] = pw.hi1[0][j];
+ plain_buf[i] = pw.h.hi1[0][j];
}
plain_len = pw.pw_len;
for (int i = 0, j = gidm; i < 16; i++, j++)
{
- plain_buf[i] = pw.hi1[0][j];
+ plain_buf[i] = pw.h.hi1[0][j];
}
plain_len = pw.pw_len;
for (int i = 0, j = gidm; i < 16; i++, j++)
{
- plain_buf[i] = pw.hi1[0][j];
+ plain_buf[i] = pw.h.hi1[0][j];
}
plain_len = pw.pw_len;
int found = 0;
- hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
+ hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
for (uint i = 0; i < KERNEL_THREADS; i++) if (device_param->result[i] == 1) found = 1;
log_info_nn ("");
- hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (uint), salt_buf->digests_cnt * sizeof (uint), &data.digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL);
+ hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (uint), salt_buf->digests_cnt * sizeof (uint), &data.digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL);
uint cpt_cracked = 0;
memset (data.digests_shown_tmp, 0, salt_buf->digests_cnt * sizeof (uint));
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (uint), salt_buf->digests_cnt * sizeof (uint), &data.digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (uint), salt_buf->digests_cnt * sizeof (uint), &data.digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL);
}
memset (device_param->result, 0, device_param->size_results);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
}
}
case KERN_RUN_3: kernel = device_param->kernel3; break;
}
- hc_clSetKernelArg (kernel, 21, sizeof (cl_uint), device_param->kernel_params[21]);
- hc_clSetKernelArg (kernel, 22, sizeof (cl_uint), device_param->kernel_params[22]);
- hc_clSetKernelArg (kernel, 23, sizeof (cl_uint), device_param->kernel_params[23]);
- hc_clSetKernelArg (kernel, 24, sizeof (cl_uint), device_param->kernel_params[24]);
- hc_clSetKernelArg (kernel, 25, sizeof (cl_uint), device_param->kernel_params[25]);
- hc_clSetKernelArg (kernel, 26, sizeof (cl_uint), device_param->kernel_params[26]);
- hc_clSetKernelArg (kernel, 27, sizeof (cl_uint), device_param->kernel_params[27]);
- hc_clSetKernelArg (kernel, 28, sizeof (cl_uint), device_param->kernel_params[28]);
- hc_clSetKernelArg (kernel, 29, sizeof (cl_uint), device_param->kernel_params[29]);
- hc_clSetKernelArg (kernel, 30, sizeof (cl_uint), device_param->kernel_params[30]);
- hc_clSetKernelArg (kernel, 31, sizeof (cl_uint), device_param->kernel_params[31]);
+ hc_clSetKernelArg (data.ocl, kernel, 21, sizeof (cl_uint), device_param->kernel_params[21]);
+ hc_clSetKernelArg (data.ocl, kernel, 22, sizeof (cl_uint), device_param->kernel_params[22]);
+ hc_clSetKernelArg (data.ocl, kernel, 23, sizeof (cl_uint), device_param->kernel_params[23]);
+ hc_clSetKernelArg (data.ocl, kernel, 24, sizeof (cl_uint), device_param->kernel_params[24]);
+ hc_clSetKernelArg (data.ocl, kernel, 25, sizeof (cl_uint), device_param->kernel_params[25]);
+ hc_clSetKernelArg (data.ocl, kernel, 26, sizeof (cl_uint), device_param->kernel_params[26]);
+ hc_clSetKernelArg (data.ocl, kernel, 27, sizeof (cl_uint), device_param->kernel_params[27]);
+ hc_clSetKernelArg (data.ocl, kernel, 28, sizeof (cl_uint), device_param->kernel_params[28]);
+ hc_clSetKernelArg (data.ocl, kernel, 29, sizeof (cl_uint), device_param->kernel_params[29]);
+ hc_clSetKernelArg (data.ocl, kernel, 30, sizeof (cl_uint), device_param->kernel_params[30]);
+ hc_clSetKernelArg (data.ocl, kernel, 31, sizeof (cl_uint), device_param->kernel_params[31]);
if ((data.opts_type & OPTS_TYPE_PT_BITSLICE) && (data.attack_mode == ATTACK_MODE_BF))
{
const size_t global_work_size[3] = { num_elements, 32, 1 };
const size_t local_work_size[3] = { kernel_threads / 32, 32, 1 };
- hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
}
else
{
const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 };
- hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
}
- hc_clFlush (device_param->command_queue);
+ hc_clFlush (data.ocl, device_param->command_queue);
- hc_clFinish (device_param->command_queue);
+ hc_clFinish (data.ocl, device_param->command_queue);
}
static void run_kernel_mp (const uint kern_run, hc_device_param_t *device_param, const uint num)
switch (kern_run)
{
- case KERN_RUN_MP: hc_clSetKernelArg (kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp[3]);
- hc_clSetKernelArg (kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp[4]);
- hc_clSetKernelArg (kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp[5]);
- hc_clSetKernelArg (kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp[6]);
- hc_clSetKernelArg (kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp[7]);
- hc_clSetKernelArg (kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp[8]);
+ case KERN_RUN_MP: hc_clSetKernelArg (data.ocl, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp[3]);
+ hc_clSetKernelArg (data.ocl, kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp[4]);
+ hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp[5]);
+ hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp[6]);
+ hc_clSetKernelArg (data.ocl, kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp[7]);
+ hc_clSetKernelArg (data.ocl, kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp[8]);
break;
- case KERN_RUN_MP_R: hc_clSetKernelArg (kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_r[3]);
- hc_clSetKernelArg (kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_r[4]);
- hc_clSetKernelArg (kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_r[5]);
- hc_clSetKernelArg (kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_r[6]);
- hc_clSetKernelArg (kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_r[7]);
- hc_clSetKernelArg (kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp_r[8]);
+ case KERN_RUN_MP_R: hc_clSetKernelArg (data.ocl, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_r[3]);
+ hc_clSetKernelArg (data.ocl, kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_r[4]);
+ hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_r[5]);
+ hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_r[6]);
+ hc_clSetKernelArg (data.ocl, kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_r[7]);
+ hc_clSetKernelArg (data.ocl, kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp_r[8]);
break;
- case KERN_RUN_MP_L: hc_clSetKernelArg (kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_l[3]);
- hc_clSetKernelArg (kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_l[4]);
- hc_clSetKernelArg (kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_l[5]);
- hc_clSetKernelArg (kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_l[6]);
- hc_clSetKernelArg (kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_l[7]);
- hc_clSetKernelArg (kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp_l[8]);
- hc_clSetKernelArg (kernel, 9, sizeof (cl_uint), device_param->kernel_params_mp_l[9]);
+ case KERN_RUN_MP_L: hc_clSetKernelArg (data.ocl, kernel, 3, sizeof (cl_ulong), device_param->kernel_params_mp_l[3]);
+ hc_clSetKernelArg (data.ocl, kernel, 4, sizeof (cl_uint), device_param->kernel_params_mp_l[4]);
+ hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_mp_l[5]);
+ hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_mp_l[6]);
+ hc_clSetKernelArg (data.ocl, kernel, 7, sizeof (cl_uint), device_param->kernel_params_mp_l[7]);
+ hc_clSetKernelArg (data.ocl, kernel, 8, sizeof (cl_uint), device_param->kernel_params_mp_l[8]);
+ hc_clSetKernelArg (data.ocl, kernel, 9, sizeof (cl_uint), device_param->kernel_params_mp_l[9]);
break;
}
const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 };
- hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
- hc_clFlush (device_param->command_queue);
+ hc_clFlush (data.ocl, device_param->command_queue);
- hc_clFinish (device_param->command_queue);
+ hc_clFinish (data.ocl, device_param->command_queue);
}
static void run_kernel_tb (hc_device_param_t *device_param, const uint num)
const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 };
- hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
- hc_clFlush (device_param->command_queue);
+ hc_clFlush (data.ocl, device_param->command_queue);
- hc_clFinish (device_param->command_queue);
+ hc_clFinish (data.ocl, device_param->command_queue);
}
static void run_kernel_tm (hc_device_param_t *device_param)
const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 };
- hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
- hc_clFlush (device_param->command_queue);
+ hc_clFlush (data.ocl, device_param->command_queue);
- hc_clFinish (device_param->command_queue);
+ hc_clFinish (data.ocl, device_param->command_queue);
}
static void run_kernel_amp (hc_device_param_t *device_param, const uint num)
cl_kernel kernel = device_param->kernel_amp;
- hc_clSetKernelArg (kernel, 5, sizeof (cl_uint), device_param->kernel_params_amp[5]);
- hc_clSetKernelArg (kernel, 6, sizeof (cl_uint), device_param->kernel_params_amp[6]);
+ hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_amp[5]);
+ hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_amp[6]);
const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 };
- hc_clEnqueueNDRangeKernel (device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
- hc_clFlush (device_param->command_queue);
+ hc_clFlush (data.ocl, device_param->command_queue);
- hc_clFinish (device_param->command_queue);
+ hc_clFinish (data.ocl, device_param->command_queue);
}
static void run_kernel_bzero (hc_device_param_t *device_param, cl_mem buf, const uint size)
const cl_uchar zero = 0;
- hc_clEnqueueFillBuffer (device_param->command_queue, buf, &zero, sizeof (cl_uchar), 0, size, 0, NULL, NULL);
+ hc_clEnqueueFillBuffer (data.ocl, device_param->command_queue, buf, &zero, sizeof (cl_uchar), 0, size, 0, NULL, NULL);
}
else
{
const int fillsz = MIN (FILLSZ, left);
- hc_clEnqueueWriteBuffer (device_param->command_queue, buf, CL_TRUE, i, fillsz, tmp, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, buf, CL_TRUE, i, fillsz, tmp, 0, NULL, NULL);
}
myfree (tmp);
{
if (data.attack_kern == ATTACK_KERN_STRAIGHT)
{
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
}
else if (data.attack_kern == ATTACK_KERN_COMBI)
{
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
}
else if (data.attack_kern == ATTACK_KERN_BF)
{
if (data.attack_mode == ATTACK_MODE_STRAIGHT)
{
- hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_rules, device_param->d_rules_c, innerloop_pos * sizeof (kernel_rule_t), 0, innerloop_left * sizeof (kernel_rule_t), 0, NULL, NULL);
+ hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_rules, device_param->d_rules_c, innerloop_pos * sizeof (kernel_rule_t), 0, innerloop_left * sizeof (kernel_rule_t), 0, NULL, NULL);
}
else if (data.attack_mode == ATTACK_MODE_COMBI)
{
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (comb_t), device_param->combs_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (comb_t), device_param->combs_buf, 0, NULL, NULL);
}
else if (data.attack_mode == ATTACK_MODE_BF)
{
- hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_bfs, device_param->d_bfs_c, 0, 0, innerloop_left * sizeof (bf_t), 0, NULL, NULL);
+ hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_bfs, device_param->d_bfs_c, 0, 0, innerloop_left * sizeof (bf_t), 0, NULL, NULL);
}
else if (data.attack_mode == ATTACK_MODE_HYBRID1)
{
- hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
+ hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
}
else if (data.attack_mode == ATTACK_MODE_HYBRID2)
{
- hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
+ hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_combs, device_param->d_combs_c, 0, 0, innerloop_left * sizeof (comb_t), 0, NULL, NULL);
}
if (data.attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
run_kernel_tm (device_param);
- hc_clEnqueueCopyBuffer (device_param->command_queue, device_param->d_tm_c, device_param->d_bfs_c, 0, 0, size_tm, 0, NULL, NULL);
+ hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_tm_c, device_param->d_bfs_c, 0, 0, size_tm, 0, NULL, NULL);
}
}
{
run_kernel (KERN_RUN_23, device_param, pws_cnt);
- hc_clEnqueueReadBuffer (device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
+ hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
// do something with data
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL);
}
run_kernel (KERN_RUN_3, device_param, pws_cnt);
static void pw_transpose_to_hi1 (const pw_t *p1, pw_t *p2)
{
- memcpy (p2->hi1, p1->hi1, 64 * sizeof (uint));
+ memcpy (p2->h.hi1, p1->h.hi1, 64 * sizeof (uint));
}
static uint pw_add_to_hc1 (hc_device_param_t *device_param, const u8 *pw_buf, const uint pw_len)
uint cache_cnt = pw_cache->cnt;
- u8 *pw_hc1 = pw_cache->pw_buf.hc1[cache_cnt];
+ u8 *pw_hc1 = pw_cache->pw_buf.h.hc1[cache_cnt];
memcpy (pw_hc1, pw_buf, pw_len);
#ifndef OSX
char *cpu_affinity = NULL;
#endif
+ OCL_PTR *ocl = NULL;
char *opencl_devices = NULL;
char *opencl_platforms = NULL;
char *opencl_device_types = NULL;
logfile_top_string (session);
logfile_top_string (truecrypt_keyfiles);
+ /**
+ * Init OpenCL library loader
+ */
+
+ if (keyspace == 0)
+ {
+ ocl = (OCL_PTR *) mymalloc (sizeof (OCL_PTR));
+
+ ocl_init(ocl);
+
+ data.ocl = ocl;
+ }
+
/**
* OpenCL platform selection
*/
kernel_accel = 16;
break;
case 6211: kernel_loops = ROUNDS_TRUECRYPT_2K;
+ #ifndef OSX
kernel_accel = 64;
+ #endif
break;
case 6212: kernel_loops = ROUNDS_TRUECRYPT_2K;
kernel_accel = 32;
kernel_accel = 8;
break;
case 6241: kernel_loops = ROUNDS_TRUECRYPT_1K;
+ #ifndef OSX
kernel_accel = 128;
+ #endif
break;
case 6242: kernel_loops = ROUNDS_TRUECRYPT_1K;
kernel_accel = 64;
kernel_accel = 8;
break;
case 7200: kernel_loops = ROUNDS_GRUB;
+ #ifndef OSX
kernel_accel = 16;
+ #endif
break;
case 7400: kernel_loops = ROUNDS_SHA256CRYPT;
kernel_accel = 8;
cl_uint platform_devices_cnt;
- hc_clGetPlatformIDs (CL_PLATFORMS_MAX, platforms, &platforms_cnt);
-
- if (platforms_cnt == 0)
+ if (keyspace == 0)
{
- log_error ("ERROR: No OpenCL compatible platform found");
+ hc_clGetPlatformIDs (data.ocl, CL_PLATFORMS_MAX, platforms, &platforms_cnt);
- return (-1);
+ if (platforms_cnt == 0)
+ {
+ log_error ("ERROR: No OpenCL compatible platform found");
+
+ return (-1);
+ }
}
/**
char platform_vendor[INFOSZ] = { 0 };
- hc_clGetPlatformInfo (platform, CL_PLATFORM_VENDOR, sizeof (platform_vendor), platform_vendor, NULL);
+ hc_clGetPlatformInfo (data.ocl, platform, CL_PLATFORM_VENDOR, sizeof (platform_vendor), platform_vendor, NULL);
#ifdef HAVE_HWMON
#if defined(HAVE_NVML) || defined(HAVE_NVAPI)
cl_platform_id platform = platforms[platform_id];
- hc_clGetDeviceIDs (platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt);
+ hc_clGetDeviceIDs (data.ocl, platform, CL_DEVICE_TYPE_ALL, DEVICES_MAX, platform_devices, &platform_devices_cnt);
for (uint platform_devices_id = 0; platform_devices_id < platform_devices_cnt; platform_devices_id++)
{
cl_device_type device_type;
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_TYPE, sizeof (device_type), &device_type, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_TYPE, sizeof (device_type), &device_type, NULL);
device_type &= ~CL_DEVICE_TYPE_DEFAULT;
cl_uint vendor_id = 0;
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_VENDOR_ID, sizeof (vendor_id), &vendor_id, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VENDOR_ID, sizeof (vendor_id), &vendor_id, NULL);
device_param->vendor_id = vendor_id;
char *device_name = (char *) mymalloc (INFOSZ);
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_NAME, INFOSZ, device_name, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NAME, INFOSZ, device_name, NULL);
device_param->device_name = device_name;
char *device_version = (char *) mymalloc (INFOSZ);
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_VERSION, INFOSZ, device_version, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VERSION, INFOSZ, device_version, NULL);
device_param->device_version = device_version;
if (opencl_vector_width == OPENCL_VECTOR_WIDTH)
{
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, sizeof (vector_width), &vector_width, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, sizeof (vector_width), &vector_width, NULL);
if ((vendor_id == VENDOR_ID_NV) && (strstr (device_name, " Ti") || strstr (device_name, " TI")))
{
cl_uint device_processors;
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof (device_processors), &device_processors, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof (device_processors), &device_processors, NULL);
device_param->device_processors = device_processors;
cl_ulong device_maxmem_alloc;
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof (device_maxmem_alloc), &device_maxmem_alloc, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof (device_maxmem_alloc), &device_maxmem_alloc, NULL);
device_param->device_maxmem_alloc = device_maxmem_alloc;
cl_ulong device_global_mem;
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof (device_global_mem), &device_global_mem, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof (device_global_mem), &device_global_mem, NULL);
device_param->device_global_mem = device_global_mem;
cl_uint device_maxclock_frequency;
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof (device_maxclock_frequency), &device_maxclock_frequency, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof (device_maxclock_frequency), &device_maxclock_frequency, NULL);
device_param->device_maxclock_frequency = device_maxclock_frequency;
char *driver_version = (char *) mymalloc (INFOSZ);
- hc_clGetDeviceInfo (device_param->device, CL_DRIVER_VERSION, INFOSZ, driver_version, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DRIVER_VERSION, INFOSZ, driver_version, NULL);
device_param->driver_version = driver_version;
#define CL_DEVICE_WAVEFRONT_WIDTH_AMD 0x4043
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_WAVEFRONT_WIDTH_AMD, sizeof (device_processor_cores), &device_processor_cores, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_WAVEFRONT_WIDTH_AMD, sizeof (device_processor_cores), &device_processor_cores, NULL);
device_param->device_processor_cores = device_processor_cores;
}
#define CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV 0x4005
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV, sizeof (kernel_exec_timeout), &kernel_exec_timeout, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV, sizeof (kernel_exec_timeout), &kernel_exec_timeout, NULL);
device_param->kernel_exec_timeout = kernel_exec_timeout;
#define CL_DEVICE_WARP_SIZE_NV 0x4003
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_WARP_SIZE_NV, sizeof (device_processor_cores), &device_processor_cores, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_WARP_SIZE_NV, sizeof (device_processor_cores), &device_processor_cores, NULL);
device_param->device_processor_cores = device_processor_cores;
#define CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV 0x4000
#define CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV 0x4001
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV, sizeof (sm_minor), &sm_minor, NULL);
- hc_clGetDeviceInfo (device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof (sm_major), &sm_major, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV, sizeof (sm_minor), &sm_minor, NULL);
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof (sm_major), &sm_major, NULL);
device_param->sm_minor = sm_minor;
device_param->sm_major = sm_major;
}
}
- if (devices_active == 0)
+ if (keyspace == 0 && devices_active == 0)
{
log_error ("ERROR: No devices found/left");
#endif // HAVE_ADK
#endif // HAVE_HWMON
+ #ifdef OSX
+ /*
+ * List of OSX kernel to fix
+ */
+ if ((hash_mode == 6100) || (hash_mode == 6900) || (hash_mode == 11700) || (hash_mode == 11800) || (hash_mode == 5300) || \
+ (hash_mode == 5400) || (hash_mode == 5500) || (hash_mode == 5600) || (hash_mode == 7500) || (hash_mode == 11100) || \
+ (hash_mode == 11400) || (hash_mode == 2611) || (hash_mode == 2711) || (hash_mode == 2811) || (hash_mode == 8400) || \
+ (hash_mode == 2612) || (hash_mode == 3711) || (hash_mode == 7600) || (hash_mode == 3100) || (hash_mode == 8000) || \
+ (hash_mode == 12600) || (hash_mode == 3000) || (hash_mode == 1500) || (hash_mode == 12400) || (hash_mode == 8500) || \
+ (hash_mode == 8600) || (hash_mode == 8700) || (hash_mode == 9100) || (hash_mode == 11600) || (hash_mode == 12500) || \
+ (hash_mode == 6231) || (hash_mode == 8800) || (hash_mode == 9700) || (hash_mode == 9710) || (hash_mode == 9800) || \
+ (hash_mode == 9810) || (hash_mode == 9400) || (hash_mode == 9500) || (hash_mode == 9600) || (hash_mode == 10400) || \
+ (hash_mode == 10410) || (hash_mode == 10700) || (hash_mode == 6800) || (hash_mode == 6600) || (hash_mode == 11300) || \
+ (hash_mode == 12700))
+ {
+ if (force == 0)
+ {
+ log_info ("");
+ log_info ("Warning: Hash mode %d is not stable in OSX.", hash_mode);
+ log_info ("You can use --force to override this but do not post error reports if you do so");
+ log_info ("");
+
+ continue;
+ }
+ }
+ #endif
+
+ #ifdef DEBUG
+ if (benchmark == 1) log_info ("Hashmode: %d", data.hash_mode);
+ #endif
+
uint kernel_blocks_all = 0;
for (uint device_id = 0; device_id < devices_cnt; device_id++)
* create context for each device
*/
- device_param->context = hc_clCreateContext (NULL, 1, &device_param->device, NULL, NULL);
+ device_param->context = hc_clCreateContext (data.ocl, NULL, 1, &device_param->device, NULL, NULL);
/**
* create command-queue
// not supported with NV
// device_param->command_queue = hc_clCreateCommandQueueWithProperties (device_param->context, device_param->device, NULL);
- device_param->command_queue = hc_clCreateCommandQueue (device_param->context, device_param->device, 0);
+ device_param->command_queue = hc_clCreateCommandQueue (data.ocl, device_param->context, device_param->device, 0);
/**
* create input buffers on device
load_kernel (source_file, 1, kernel_lengths, kernel_sources);
- device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
+ device_param->program = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
- hc_clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
size_t binary_size;
- clGetProgramInfo (device_param->program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
+ hc_clGetProgramInfo (data.ocl, device_param->program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
u8 *binary = (u8 *) mymalloc (binary_size);
- clGetProgramInfo (device_param->program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
+ hc_clGetProgramInfo (data.ocl, device_param->program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
writeProgramBin (cached_file, binary, binary_size);
load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
- device_param->program = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
+ device_param->program = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
- hc_clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
}
}
else
load_kernel (source_file, 1, kernel_lengths, kernel_sources);
- device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
+ device_param->program = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
if (force_jit_compilation == 1500)
{
snprintf (build_opts, sizeof (build_opts) - 1, "%s -DSCRYPT_N=%d -DSCRYPT_R=%d -DSCRYPT_P=%d -DSCRYPT_TMTO=%d", build_opts, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, 1 << data.salts_buf[0].scrypt_tmto);
}
- hc_clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
}
local_free (kernel_lengths);
size_t ret_val_size = 0;
- clGetProgramBuildInfo (device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
+ hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
if (ret_val_size > 2)
{
char *build_log = (char *) mymalloc (ret_val_size + 1);
- clGetProgramBuildInfo (device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
+ hc_clGetProgramBuildInfo (data.ocl, device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
puts (build_log);
load_kernel (source_file, 1, kernel_lengths, kernel_sources);
- device_param->program_mp = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
+ device_param->program_mp = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
- hc_clBuildProgram (device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
size_t binary_size;
- clGetProgramInfo (device_param->program_mp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
+ hc_clGetProgramInfo (data.ocl, device_param->program_mp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
u8 *binary = (u8 *) mymalloc (binary_size);
- clGetProgramInfo (device_param->program_mp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
+ hc_clGetProgramInfo (data.ocl, device_param->program_mp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
writeProgramBin (cached_file, binary, binary_size);
load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
- device_param->program_mp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
+ device_param->program_mp = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
- hc_clBuildProgram (device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
}
local_free (kernel_lengths);
size_t ret_val_size = 0;
- clGetProgramBuildInfo (device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
+ hc_clGetProgramBuildInfo (data.ocl, device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
if (ret_val_size > 2)
{
char *build_log = (char *) mymalloc (ret_val_size + 1);
- clGetProgramBuildInfo (device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
+ hc_clGetProgramBuildInfo (data.ocl, device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
puts (build_log);
load_kernel (source_file, 1, kernel_lengths, kernel_sources);
- device_param->program_amp = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
+ device_param->program_amp = hc_clCreateProgramWithSource (data.ocl, device_param->context, 1, (const char **) kernel_sources, NULL);
- hc_clBuildProgram (device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
size_t binary_size;
- clGetProgramInfo (device_param->program_amp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
+ hc_clGetProgramInfo (data.ocl, device_param->program_amp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
u8 *binary = (u8 *) mymalloc (binary_size);
- clGetProgramInfo (device_param->program_amp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
+ hc_clGetProgramInfo (data.ocl, device_param->program_amp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
writeProgramBin (cached_file, binary, binary_size);
load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
- device_param->program_amp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
+ device_param->program_amp = hc_clCreateProgramWithBinary (data.ocl, device_param->context, 1, &device_param->device, kernel_lengths, (const u8 **) kernel_sources, NULL);
- hc_clBuildProgram (device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
+ hc_clBuildProgram (data.ocl, device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
}
local_free (kernel_lengths);
size_t ret_val_size = 0;
- clGetProgramBuildInfo (device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
+ hc_clGetProgramBuildInfo (data.ocl, device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
if (ret_val_size > 2)
{
char *build_log = (char *) mymalloc (ret_val_size + 1);
- clGetProgramBuildInfo (device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
+ hc_clGetProgramBuildInfo (data.ocl, device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
puts (build_log);
* global buffers
*/
- device_param->d_pws_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_pws, NULL);
- device_param->d_pws_amp_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_pws, NULL);
- device_param->d_tmps = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE, size_tmps, NULL);
- device_param->d_hooks = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE, size_hooks, NULL);
- device_param->d_bitmap_s1_a = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s1_b = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s1_c = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s1_d = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s2_a = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s2_b = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s2_c = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_bitmap_s2_d = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
- device_param->d_plain_bufs = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE, size_plains, NULL);
- device_param->d_digests_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_digests, NULL);
- device_param->d_digests_shown = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE, size_shown, NULL);
- device_param->d_salt_bufs = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_salts, NULL);
- device_param->d_result = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE, size_results, NULL);
- device_param->d_scryptV_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_WRITE, size_scryptV, NULL);
-
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s1_a, CL_TRUE, 0, bitmap_size, bitmap_s1_a, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s1_b, CL_TRUE, 0, bitmap_size, bitmap_s1_b, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s1_c, CL_TRUE, 0, bitmap_size, bitmap_s1_c, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s1_d, CL_TRUE, 0, bitmap_size, bitmap_s1_d, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s2_a, CL_TRUE, 0, bitmap_size, bitmap_s2_a, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s2_b, CL_TRUE, 0, bitmap_size, bitmap_s2_b, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s2_c, CL_TRUE, 0, bitmap_size, bitmap_s2_c, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_bitmap_s2_d, CL_TRUE, 0, bitmap_size, bitmap_s2_d, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_digests_buf, CL_TRUE, 0, size_digests, data.digests_buf, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_digests_shown, CL_TRUE, 0, size_shown, data.digests_shown, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_salt_bufs, CL_TRUE, 0, size_salts, data.salts_buf, 0, NULL, NULL);
+ device_param->d_pws_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL);
+ device_param->d_pws_amp_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_pws, NULL);
+ device_param->d_tmps = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_tmps, NULL);
+ device_param->d_hooks = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_hooks, NULL);
+ device_param->d_bitmap_s1_a = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s1_b = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s1_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s1_d = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s2_a = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s2_b = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s2_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_bitmap_s2_d = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, bitmap_size, NULL);
+ device_param->d_plain_bufs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_plains, NULL);
+ device_param->d_digests_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_digests, NULL);
+ device_param->d_digests_shown = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_shown, NULL);
+ device_param->d_salt_bufs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_salts, NULL);
+ device_param->d_result = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_results, NULL);
+ device_param->d_scryptV_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE, size_scryptV, NULL);
+
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_a, CL_TRUE, 0, bitmap_size, bitmap_s1_a, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_b, CL_TRUE, 0, bitmap_size, bitmap_s1_b, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_c, CL_TRUE, 0, bitmap_size, bitmap_s1_c, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_d, CL_TRUE, 0, bitmap_size, bitmap_s1_d, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_a, CL_TRUE, 0, bitmap_size, bitmap_s2_a, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_b, CL_TRUE, 0, bitmap_size, bitmap_s2_b, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_c, CL_TRUE, 0, bitmap_size, bitmap_s2_c, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s2_d, CL_TRUE, 0, bitmap_size, bitmap_s2_d, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_buf, CL_TRUE, 0, size_digests, data.digests_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, 0, size_shown, data.digests_shown, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_salt_bufs, CL_TRUE, 0, size_salts, data.salts_buf, 0, NULL, NULL);
run_kernel_bzero (device_param, device_param->d_pws_buf, size_pws);
run_kernel_bzero (device_param, device_param->d_pws_amp_buf, size_pws);
if (attack_kern == ATTACK_KERN_STRAIGHT)
{
- device_param->d_rules = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_rules, NULL);
- device_param->d_rules_c = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_rules_c, NULL);
+ device_param->d_rules = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_rules, NULL);
+ device_param->d_rules_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_rules_c, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_rules, CL_TRUE, 0, size_rules, kernel_rules_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_rules, CL_TRUE, 0, size_rules, kernel_rules_buf, 0, NULL, NULL);
run_kernel_bzero (device_param, device_param->d_rules_c, size_rules_c);
}
else if (attack_kern == ATTACK_KERN_COMBI)
{
- device_param->d_combs = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_combs, NULL);
- device_param->d_combs_c = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_combs, NULL);
- device_param->d_root_css_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL);
- device_param->d_markov_css_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
+ device_param->d_combs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL);
+ device_param->d_combs_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL);
+ device_param->d_root_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL);
+ device_param->d_markov_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
run_kernel_bzero (device_param, device_param->d_combs, size_combs);
run_kernel_bzero (device_param, device_param->d_combs_c, size_combs);
}
else if (attack_kern == ATTACK_KERN_BF)
{
- device_param->d_bfs = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL);
- device_param->d_bfs_c = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL);
- device_param->d_tm_c = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_tm, NULL);
- device_param->d_root_css_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL);
- device_param->d_markov_css_buf = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
+ device_param->d_bfs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL);
+ device_param->d_bfs_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_bfs, NULL);
+ device_param->d_tm_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_tm, NULL);
+ device_param->d_root_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL);
+ device_param->d_markov_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
run_kernel_bzero (device_param, device_param->d_bfs, size_bfs);
run_kernel_bzero (device_param, device_param->d_bfs_c, size_bfs);
if (size_esalts)
{
- device_param->d_esalt_bufs = hc_clCreateBuffer (device_param->context, CL_MEM_READ_ONLY, size_esalts, NULL);
+ device_param->d_esalt_bufs = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_esalts, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_esalt_bufs, CL_TRUE, 0, size_esalts, data.esalts_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_esalt_bufs, CL_TRUE, 0, size_esalts, data.esalts_buf, 0, NULL, NULL);
}
/**
{
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", kern_type, 4);
- device_param->kernel1 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel1 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", kern_type, 8);
- device_param->kernel2 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel2 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_s%02d", kern_type, 16);
- device_param->kernel3 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel3 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
}
else
{
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_m%02d", kern_type, 4);
- device_param->kernel1 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel1 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_m%02d", kern_type, 8);
- device_param->kernel2 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel2 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_m%02d", kern_type, 16);
- device_param->kernel3 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel3 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
}
if (data.attack_mode == ATTACK_MODE_BF)
{
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_tb", kern_type);
- device_param->kernel_tb = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel_tb = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_tm", kern_type);
- device_param->kernel_tm = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel_tm = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
}
}
}
{
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_init", kern_type);
- device_param->kernel1 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel1 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_loop", kern_type);
- device_param->kernel2 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel2 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_comp", kern_type);
- device_param->kernel3 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel3 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
if (opts_type & OPTS_TYPE_HOOK12)
{
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook12", kern_type);
- device_param->kernel12 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel12 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
}
if (opts_type & OPTS_TYPE_HOOK23)
{
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook23", kern_type);
- device_param->kernel23 = hc_clCreateKernel (device_param->program, kernel_name);
+ device_param->kernel23 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
}
}
for (uint i = 0; i <= 20; i++)
{
- hc_clSetKernelArg (device_param->kernel1, i, sizeof (cl_mem), device_param->kernel_params[i]);
- hc_clSetKernelArg (device_param->kernel2, i, sizeof (cl_mem), device_param->kernel_params[i]);
- hc_clSetKernelArg (device_param->kernel3, i, sizeof (cl_mem), device_param->kernel_params[i]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_mem), device_param->kernel_params[i]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_mem), device_param->kernel_params[i]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel3, i, sizeof (cl_mem), device_param->kernel_params[i]);
- if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (device_param->kernel12, i, sizeof (cl_mem), device_param->kernel_params[i]);
- if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (device_param->kernel23, i, sizeof (cl_mem), device_param->kernel_params[i]);
+ if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (data.ocl, device_param->kernel12, i, sizeof (cl_mem), device_param->kernel_params[i]);
+ if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_mem), device_param->kernel_params[i]);
}
for (uint i = 21; i <= 31; i++)
{
- hc_clSetKernelArg (device_param->kernel1, i, sizeof (cl_uint), device_param->kernel_params[i]);
- hc_clSetKernelArg (device_param->kernel2, i, sizeof (cl_uint), device_param->kernel_params[i]);
- hc_clSetKernelArg (device_param->kernel3, i, sizeof (cl_uint), device_param->kernel_params[i]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_uint), device_param->kernel_params[i]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_uint), device_param->kernel_params[i]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel3, i, sizeof (cl_uint), device_param->kernel_params[i]);
- if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (device_param->kernel12, i, sizeof (cl_uint), device_param->kernel_params[i]);
- if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (device_param->kernel23, i, sizeof (cl_uint), device_param->kernel_params[i]);
+ if (opts_type & OPTS_TYPE_HOOK12) hc_clSetKernelArg (data.ocl, device_param->kernel12, i, sizeof (cl_uint), device_param->kernel_params[i]);
+ if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_uint), device_param->kernel_params[i]);
}
if (attack_mode == ATTACK_MODE_BF)
{
- device_param->kernel_mp_l = hc_clCreateKernel (device_param->program_mp, "l_markov");
- device_param->kernel_mp_r = hc_clCreateKernel (device_param->program_mp, "r_markov");
+ device_param->kernel_mp_l = hc_clCreateKernel (data.ocl, device_param->program_mp, "l_markov");
+ device_param->kernel_mp_r = hc_clCreateKernel (data.ocl, device_param->program_mp, "r_markov");
if (opts_type & OPTS_TYPE_PT_BITSLICE)
{
- hc_clSetKernelArg (device_param->kernel_tb, 0, sizeof (cl_mem), device_param->kernel_params_tb[0]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel_tb, 0, sizeof (cl_mem), device_param->kernel_params_tb[0]);
- hc_clSetKernelArg (device_param->kernel_tm, 0, sizeof (cl_mem), device_param->kernel_params_tm[0]);
- hc_clSetKernelArg (device_param->kernel_tm, 1, sizeof (cl_mem), device_param->kernel_params_tm[1]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel_tm, 0, sizeof (cl_mem), device_param->kernel_params_tm[0]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel_tm, 1, sizeof (cl_mem), device_param->kernel_params_tm[1]);
}
}
else if (attack_mode == ATTACK_MODE_HYBRID1)
{
- device_param->kernel_mp = hc_clCreateKernel (device_param->program_mp, "C_markov");
+ device_param->kernel_mp = hc_clCreateKernel (data.ocl, device_param->program_mp, "C_markov");
}
else if (attack_mode == ATTACK_MODE_HYBRID2)
{
- device_param->kernel_mp = hc_clCreateKernel (device_param->program_mp, "C_markov");
+ device_param->kernel_mp = hc_clCreateKernel (data.ocl, device_param->program_mp, "C_markov");
}
if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
}
else
{
- device_param->kernel_amp = hc_clCreateKernel (device_param->program_amp, "amp");
+ device_param->kernel_amp = hc_clCreateKernel (data.ocl, device_param->program_amp, "amp");
}
if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
{
for (uint i = 0; i < 5; i++)
{
- hc_clSetKernelArg (device_param->kernel_amp, i, sizeof (cl_mem), device_param->kernel_params_amp[i]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel_amp, i, sizeof (cl_mem), device_param->kernel_params_amp[i]);
}
for (uint i = 5; i < 7; i++)
{
- hc_clSetKernelArg (device_param->kernel_amp, i, sizeof (cl_uint), device_param->kernel_params_amp[i]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel_amp, i, sizeof (cl_uint), device_param->kernel_params_amp[i]);
}
}
device_param->kernel_params_mp_buf32[7] = 0;
}
- for (uint i = 0; i < 3; i++) hc_clSetKernelArg (device_param->kernel_mp, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp[i]);
- for (uint i = 3; i < 4; i++) hc_clSetKernelArg (device_param->kernel_mp, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp[i]);
- for (uint i = 4; i < 8; i++) hc_clSetKernelArg (device_param->kernel_mp, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp[i]);
+ for (uint i = 0; i < 3; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp[i]);
+ for (uint i = 3; i < 4; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp[i]);
+ for (uint i = 4; i < 8; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp[i]);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_root_css_buf, CL_TRUE, 0, device_param->size_root_css, root_css_buf, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, markov_css_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_root_css_buf, CL_TRUE, 0, device_param->size_root_css, root_css_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, markov_css_buf, 0, NULL, NULL);
}
}
else if (attack_mode == ATTACK_MODE_BF)
// some more resets:
- memset (device_param->pw_caches, 0, 64 * sizeof (pw_cache_t));
+ if (device_param->pw_caches) memset (device_param->pw_caches, 0, 64 * sizeof (pw_cache_t));
- memset (device_param->pws_buf, 0, device_param->size_pws);
+ if (device_param->pws_buf) memset (device_param->pws_buf, 0, device_param->size_pws);
device_param->pw_cnt = 0;
device_param->pws_cnt = 0;
device_param->kernel_params_mp_r_buf32[6] = 0;
device_param->kernel_params_mp_r_buf32[7] = 0;
- for (uint i = 0; i < 3; i++) hc_clSetKernelArg (device_param->kernel_mp_l, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp_l[i]);
- for (uint i = 3; i < 4; i++) hc_clSetKernelArg (device_param->kernel_mp_l, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_l[i]);
- for (uint i = 4; i < 9; i++) hc_clSetKernelArg (device_param->kernel_mp_l, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp_l[i]);
+ for (uint i = 0; i < 3; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp_l[i]);
+ for (uint i = 3; i < 4; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_l[i]);
+ for (uint i = 4; i < 9; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_l, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp_l[i]);
- for (uint i = 0; i < 3; i++) hc_clSetKernelArg (device_param->kernel_mp_r, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp_r[i]);
- for (uint i = 3; i < 4; i++) hc_clSetKernelArg (device_param->kernel_mp_r, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_r[i]);
- for (uint i = 4; i < 8; i++) hc_clSetKernelArg (device_param->kernel_mp_r, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp_r[i]);
+ for (uint i = 0; i < 3; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_mem), (void *) device_param->kernel_params_mp_r[i]);
+ for (uint i = 3; i < 4; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_ulong), (void *) device_param->kernel_params_mp_r[i]);
+ for (uint i = 4; i < 8; i++) hc_clSetKernelArg (data.ocl, device_param->kernel_mp_r, i, sizeof (cl_uint), (void *) device_param->kernel_params_mp_r[i]);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_root_css_buf, CL_TRUE, 0, device_param->size_root_css, root_css_buf, 0, NULL, NULL);
- hc_clEnqueueWriteBuffer (device_param->command_queue, device_param->d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, markov_css_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_root_css_buf, CL_TRUE, 0, device_param->size_root_css, root_css_buf, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_markov_css_buf, CL_TRUE, 0, device_param->size_markov_css, markov_css_buf, 0, NULL, NULL);
}
}
local_free (device_param->driver_version);
if (device_param->pws_buf) myfree (device_param->pws_buf);
- if (device_param->d_pws_buf) hc_clReleaseMemObject (device_param->d_pws_buf);
- if (device_param->d_pws_amp_buf) hc_clReleaseMemObject (device_param->d_pws_amp_buf);
- if (device_param->d_rules) hc_clReleaseMemObject (device_param->d_rules);
- if (device_param->d_rules_c) hc_clReleaseMemObject (device_param->d_rules_c);
- if (device_param->d_combs) hc_clReleaseMemObject (device_param->d_combs);
- if (device_param->d_combs_c) hc_clReleaseMemObject (device_param->d_combs_c);
- if (device_param->d_bfs) hc_clReleaseMemObject (device_param->d_bfs);
- if (device_param->d_bfs_c) hc_clReleaseMemObject (device_param->d_bfs_c);
- if (device_param->d_bitmap_s1_a) hc_clReleaseMemObject (device_param->d_bitmap_s1_a);
- if (device_param->d_bitmap_s1_b) hc_clReleaseMemObject (device_param->d_bitmap_s1_b);
- if (device_param->d_bitmap_s1_c) hc_clReleaseMemObject (device_param->d_bitmap_s1_c);
- if (device_param->d_bitmap_s1_d) hc_clReleaseMemObject (device_param->d_bitmap_s1_d);
- if (device_param->d_bitmap_s2_a) hc_clReleaseMemObject (device_param->d_bitmap_s2_a);
- if (device_param->d_bitmap_s2_b) hc_clReleaseMemObject (device_param->d_bitmap_s2_b);
- if (device_param->d_bitmap_s2_c) hc_clReleaseMemObject (device_param->d_bitmap_s2_c);
- if (device_param->d_bitmap_s2_d) hc_clReleaseMemObject (device_param->d_bitmap_s2_d);
- if (device_param->d_plain_bufs) hc_clReleaseMemObject (device_param->d_plain_bufs);
- if (device_param->d_digests_buf) hc_clReleaseMemObject (device_param->d_digests_buf);
- if (device_param->d_digests_shown) hc_clReleaseMemObject (device_param->d_digests_shown);
- if (device_param->d_salt_bufs) hc_clReleaseMemObject (device_param->d_salt_bufs);
- if (device_param->d_esalt_bufs) hc_clReleaseMemObject (device_param->d_esalt_bufs);
- if (device_param->d_tmps) hc_clReleaseMemObject (device_param->d_tmps);
- if (device_param->d_hooks) hc_clReleaseMemObject (device_param->d_hooks);
- if (device_param->d_result) hc_clReleaseMemObject (device_param->d_result);
- if (device_param->d_scryptV_buf) hc_clReleaseMemObject (device_param->d_scryptV_buf);
- if (device_param->d_root_css_buf) hc_clReleaseMemObject (device_param->d_root_css_buf);
- if (device_param->d_markov_css_buf) hc_clReleaseMemObject (device_param->d_markov_css_buf);
- if (device_param->d_tm_c) hc_clReleaseMemObject (device_param->d_tm_c);
-
- if (device_param->kernel1) hc_clReleaseKernel (device_param->kernel1);
- if (device_param->kernel12) hc_clReleaseKernel (device_param->kernel12);
- if (device_param->kernel2) hc_clReleaseKernel (device_param->kernel2);
- if (device_param->kernel23) hc_clReleaseKernel (device_param->kernel23);
- if (device_param->kernel3) hc_clReleaseKernel (device_param->kernel3);
- if (device_param->kernel_mp) hc_clReleaseKernel (device_param->kernel_mp);
- if (device_param->kernel_mp_l) hc_clReleaseKernel (device_param->kernel_mp_l);
- if (device_param->kernel_mp_r) hc_clReleaseKernel (device_param->kernel_mp_r);
- if (device_param->kernel_tb) hc_clReleaseKernel (device_param->kernel_tb);
- if (device_param->kernel_tm) hc_clReleaseKernel (device_param->kernel_tm);
- if (device_param->kernel_amp) hc_clReleaseKernel (device_param->kernel_amp);
-
- if (device_param->program) hc_clReleaseProgram (device_param->program);
- if (device_param->program_mp) hc_clReleaseProgram (device_param->program_mp);
- if (device_param->program_amp) hc_clReleaseProgram (device_param->program_amp);
-
- if (device_param->command_queue) hc_clReleaseCommandQueue (device_param->command_queue);
- if (device_param->context) hc_clReleaseContext (device_param->context);
+ if (device_param->d_pws_buf) hc_clReleaseMemObject (data.ocl, device_param->d_pws_buf);
+ if (device_param->d_pws_amp_buf) hc_clReleaseMemObject (data.ocl, device_param->d_pws_amp_buf);
+ if (device_param->d_rules) hc_clReleaseMemObject (data.ocl, device_param->d_rules);
+ if (device_param->d_rules_c) hc_clReleaseMemObject (data.ocl, device_param->d_rules_c);
+ if (device_param->d_combs) hc_clReleaseMemObject (data.ocl, device_param->d_combs);
+ if (device_param->d_combs_c) hc_clReleaseMemObject (data.ocl, device_param->d_combs_c);
+ if (device_param->d_bfs) hc_clReleaseMemObject (data.ocl, device_param->d_bfs);
+ if (device_param->d_bfs_c) hc_clReleaseMemObject (data.ocl, device_param->d_bfs_c);
+ if (device_param->d_bitmap_s1_a) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_a);
+ if (device_param->d_bitmap_s1_b) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_b);
+ if (device_param->d_bitmap_s1_c) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_c);
+ if (device_param->d_bitmap_s1_d) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s1_d);
+ if (device_param->d_bitmap_s2_a) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_a);
+ if (device_param->d_bitmap_s2_b) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_b);
+ if (device_param->d_bitmap_s2_c) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_c);
+ if (device_param->d_bitmap_s2_d) hc_clReleaseMemObject (data.ocl, device_param->d_bitmap_s2_d);
+ if (device_param->d_plain_bufs) hc_clReleaseMemObject (data.ocl, device_param->d_plain_bufs);
+ if (device_param->d_digests_buf) hc_clReleaseMemObject (data.ocl, device_param->d_digests_buf);
+ if (device_param->d_digests_shown) hc_clReleaseMemObject (data.ocl, device_param->d_digests_shown);
+ if (device_param->d_salt_bufs) hc_clReleaseMemObject (data.ocl, device_param->d_salt_bufs);
+ if (device_param->d_esalt_bufs) hc_clReleaseMemObject (data.ocl, device_param->d_esalt_bufs);
+ if (device_param->d_tmps) hc_clReleaseMemObject (data.ocl, device_param->d_tmps);
+ if (device_param->d_hooks) hc_clReleaseMemObject (data.ocl, device_param->d_hooks);
+ if (device_param->d_result) hc_clReleaseMemObject (data.ocl, device_param->d_result);
+ if (device_param->d_scryptV_buf) hc_clReleaseMemObject (data.ocl, device_param->d_scryptV_buf);
+ if (device_param->d_root_css_buf) hc_clReleaseMemObject (data.ocl, device_param->d_root_css_buf);
+ if (device_param->d_markov_css_buf) hc_clReleaseMemObject (data.ocl, device_param->d_markov_css_buf);
+ if (device_param->d_tm_c) hc_clReleaseMemObject (data.ocl, device_param->d_tm_c);
+
+ if (device_param->kernel1) hc_clReleaseKernel (data.ocl, device_param->kernel1);
+ if (device_param->kernel12) hc_clReleaseKernel (data.ocl, device_param->kernel12);
+ if (device_param->kernel2) hc_clReleaseKernel (data.ocl, device_param->kernel2);
+ if (device_param->kernel23) hc_clReleaseKernel (data.ocl, device_param->kernel23);
+ if (device_param->kernel3) hc_clReleaseKernel (data.ocl, device_param->kernel3);
+ if (device_param->kernel_mp) hc_clReleaseKernel (data.ocl, device_param->kernel_mp);
+ if (device_param->kernel_mp_l) hc_clReleaseKernel (data.ocl, device_param->kernel_mp_l);
+ if (device_param->kernel_mp_r) hc_clReleaseKernel (data.ocl, device_param->kernel_mp_r);
+ if (device_param->kernel_tb) hc_clReleaseKernel (data.ocl, device_param->kernel_tb);
+ if (device_param->kernel_tm) hc_clReleaseKernel (data.ocl, device_param->kernel_tm);
+ if (device_param->kernel_amp) hc_clReleaseKernel (data.ocl, device_param->kernel_amp);
+
+ if (device_param->program) hc_clReleaseProgram (data.ocl, device_param->program);
+ if (device_param->program_mp) hc_clReleaseProgram (data.ocl, device_param->program_mp);
+ if (device_param->program_amp) hc_clReleaseProgram (data.ocl, device_param->program_amp);
+
+ if (device_param->command_queue) hc_clReleaseCommandQueue (data.ocl, device_param->command_queue);
+ if (device_param->context) hc_clReleaseContext (data.ocl, device_param->context);
}
// reset default fan speed
if (quiet == 0) log_info_nn ("Started: %s", ctime (&proc_start));
if (quiet == 0) log_info_nn ("Stopped: %s", ctime (&proc_stop));
+ if (data.ocl) ocl_close (data.ocl);
+
if (data.devices_status == STATUS_ABORTED) return 2;
if (data.devices_status == STATUS_QUIT) return 2;
if (data.devices_status == STATUS_STOP_AT_CHECKPOINT) return 2;
projects / hashcat.git / commitdiff