+++ /dev/null
-/**
- * Author......: Jens Steube <jens.steube@gmail.com>
- * License.....: MIT
- */
-
-#include <ext_cuda.h>
-
-void hc_cuDeviceGetCount (int *count)
-{
- CUresult CU_err = cuDeviceGetCount (count);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuDeviceGetCount()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuDeviceGet (CUdevice *device, int ordinal)
-{
- CUresult CU_err = cuDeviceGet (device, ordinal);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuDeviceGet()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuDeviceGetName (char *name, int len, CUdevice dev)
-{
- CUresult CU_err = cuDeviceGetName (name, len, dev);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuDeviceGetName()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuDeviceTotalMem (size_t *bytes, CUdevice dev)
-{
- CUresult CU_err = cuDeviceTotalMem (bytes, dev);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuDeviceTotalMem()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuDeviceGetAttribute (int *pi, CUdevice_attribute attrib, CUdevice dev)
-{
- CUresult CU_err = cuDeviceGetAttribute (pi, attrib, dev);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuDeviceGetAttribute()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuCtxCreate (CUcontext *pctx, unsigned int flags, CUdevice dev)
-{
- CUresult CU_err = cuCtxCreate (pctx, flags, dev);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuCtxCreate()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuMemAllocHost (void **pp, size_t bytesize)
-{
- CUresult CU_err = cuMemAllocHost (pp, bytesize);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuMemAllocHost()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuMemAlloc (CUdeviceptr *dptr, size_t bytesize)
-{
- CUresult CU_err = cuMemAlloc (dptr, bytesize);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuMemAlloc()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuMemsetD8 (CUdeviceptr dstDevice, unsigned char uc, unsigned int N)
-{
- CUresult CU_err = cuMemsetD8 (dstDevice, uc, N);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuMemsetD8()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuMemsetD32 (CUdeviceptr dstDevice, unsigned int ui, unsigned int N)
-{
- CUresult CU_err = cuMemsetD32 (dstDevice, ui, N);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuMemsetD32()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuMemcpyDtoD (CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount)
-{
- CUresult CU_err = cuMemcpyDtoD (dstDevice, srcDevice, ByteCount);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuMemcpyDtoD()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuMemcpyDtoDAsync (CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream)
-{
- CUresult CU_err = cuMemcpyDtoDAsync (dstDevice, srcDevice, ByteCount, hStream);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuMemcpyDtoDAsync()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuMemcpyHtoD (CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount)
-{
- CUresult CU_err = cuMemcpyHtoD (dstDevice, srcHost, ByteCount);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuMemcpyHtoD()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuMemcpyHtoDAsync (CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount, CUstream hStream)
-{
- CUresult CU_err = cuMemcpyHtoDAsync (dstDevice, srcHost, ByteCount, hStream);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuMemcpyHtoDAsync()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuMemcpyDtoH (void *dstHost, CUdeviceptr srcDevice, size_t ByteCount)
-{
- CUresult CU_err = cuMemcpyDtoH (dstHost, srcDevice, ByteCount);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuMemcpyDtoH()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuMemcpyDtoHAsync (void *dstHost, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream)
-{
- CUresult CU_err = cuMemcpyDtoHAsync (dstHost, srcDevice, ByteCount, hStream);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuMemcpyDtoHAsync()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuEventCreate (CUevent *phEvent, unsigned int Flags)
-{
- CUresult CU_err = cuEventCreate (phEvent, Flags);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuEventCreate()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuStreamCreate (CUstream *phStream, unsigned int Flags)
-{
- CUresult CU_err = cuStreamCreate (phStream, Flags);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuStreamCreate()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuDeviceComputeCapability (int *major, int *minor, CUdevice dev)
-{
- CUresult CU_err = cuDeviceComputeCapability (major, minor, dev);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuDeviceComputeCapability()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuModuleLoad (CUmodule *module, const char *fname)
-{
- CUresult CU_err = cuModuleLoad (module, fname);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuModuleLoad()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuModuleLoadData (CUmodule *module, const void *image)
-{
- CUresult CU_err = cuModuleLoadData (module, image);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuModuleLoadData()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuModuleGetFunction (CUfunction *hfunc, CUmodule hmod, const char *name)
-{
- CUresult CU_err = cuModuleGetFunction (hfunc, hmod, name);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuModuleGetFunction()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuFuncSetBlockShape (CUfunction hfunc, int x, int y, int z)
-{
- CUresult CU_err = cuFuncSetBlockShape (hfunc, x, y, z);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuFuncSetBlockShape()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuParamSetSize (CUfunction hfunc, unsigned int numbytes)
-{
- CUresult CU_err = cuParamSetSize (hfunc, numbytes);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuParamSetSize()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuParamSetv (CUfunction hfunc, int offset, void *ptr, unsigned int numbytes)
-{
- CUresult CU_err = cuParamSetv (hfunc, offset, ptr, numbytes);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuParamSetv()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuParamSeti (CUfunction hfunc, int offset, unsigned int value)
-{
- CUresult CU_err = cuParamSeti (hfunc, offset, value);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuParamSeti()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuModuleGetGlobal (CUdeviceptr *dptr, size_t *bytes, CUmodule hmod, const char *name)
-{
- CUresult CU_err = cuModuleGetGlobal (dptr, bytes, hmod, name);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuModuleGetGlobal()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuCtxPopCurrent (CUcontext *pctx)
-{
- CUresult CU_err = cuCtxPopCurrent (pctx);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuCtxPopCurrent()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuCtxPushCurrent (CUcontext ctx)
-{
- CUresult CU_err = cuCtxPushCurrent (ctx);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuCtxPushCurrent()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuLaunchKernel (CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void **kernelParams, void **extra)
-{
- CUresult CU_err = cuLaunchKernel (f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuLaunchKernel()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuLaunchGrid (CUfunction f, int grid_width, int grid_height)
-{
- CUresult CU_err = cuLaunchGrid (f, grid_width, grid_height);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuLaunchGrid()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuLaunchGridAsync (CUfunction f, int grid_width, int grid_height, CUstream hStream)
-{
- CUresult CU_err = cuLaunchGridAsync (f, grid_width, grid_height, hStream);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuLaunchGridAsync()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuEventSynchronize (CUevent hEvent)
-{
- CUresult CU_err = cuEventSynchronize (hEvent);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuEventSynchronize()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuStreamSynchronize (CUstream hStream)
-{
- CUresult CU_err = cuStreamSynchronize (hStream);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuStreamSynchronize()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuMemFreeHost (void *p)
-{
- CUresult CU_err = cuMemFreeHost (p);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuMemFreeHost()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuMemFree (CUdeviceptr dptr)
-{
- CUresult CU_err = cuMemFree (dptr);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuMemFree()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuEventDestroy (CUevent hEvent)
-{
- CUresult CU_err = cuEventDestroy (hEvent);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuEventDestroy()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuStreamDestroy (CUstream hStream)
-{
- CUresult CU_err = cuStreamDestroy (hStream);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuStreamDestroy()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuModuleUnload (CUmodule hmod)
-{
- CUresult CU_err = cuModuleUnload (hmod);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuModuleUnload()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuCtxDestroy (CUcontext ctx)
-{
- CUresult CU_err = cuCtxDestroy (ctx);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuCtxDestroy()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuCtxAttach (CUcontext *pctx, unsigned int flags)
-{
- CUresult CU_err = cuCtxAttach (pctx, flags);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuCtxAttach()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuCtxDetach (CUcontext ctx)
-{
- CUresult CU_err = cuCtxDetach (ctx);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuCtxDetach()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuCtxSynchronize (void)
-{
- CUresult CU_err = cuCtxSynchronize ();
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuCtxSynchronize()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuCtxSetCacheConfig (CUfunc_cache config)
-{
- CUresult CU_err = cuCtxSetCacheConfig (config);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuCtxSetCacheConfig()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuDriverGetVersion (int *driverVersion)
-{
- CUresult CU_err = cuDriverGetVersion (driverVersion);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuDriverGetVersion()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuModuleLoadDataEx (CUmodule *module, const void *image, unsigned int numOptions, CUjit_option *options, void **optionValues)
-{
- CUresult CU_err = cuModuleLoadDataEx (module, image, numOptions, options, optionValues);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuModuleLoadDataEx()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuLinkAddFile (CUlinkState state, CUjitInputType type, const char *path, unsigned int numOptions, CUjit_option *options, void **optionValues)
-{
- CUresult CU_err = cuLinkAddFile (state, type, path, numOptions, options, optionValues);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuLinkAddFile()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuLinkComplete (CUlinkState state, void **cubinOut, size_t *sizeOut)
-{
- CUresult CU_err = cuLinkComplete (state, cubinOut, sizeOut);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuLinkComplete()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuLinkCreate (unsigned int numOptions, CUjit_option *options, void **optionValues, CUlinkState *stateOut)
-{
- CUresult CU_err = cuLinkCreate (numOptions, options, optionValues, stateOut);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuLinkCreate()", CU_err);
-
- exit (-1);
- }
-}
-
-void hc_cuLinkDestroy (CUlinkState state)
-{
- CUresult CU_err = cuLinkDestroy (state);
-
- if (CU_err != CUDA_SUCCESS)
- {
- log_error ("ERROR: %s %d\n", "cuLinkDestroy()", CU_err);
-
- exit (-1);
- }
-}
#define INCR_MASKS 1000
#define INCR_POT 1000
-// comment-out for kernel source mode
-
-//#define BINARY_KERNEL
-
#define USAGE 0
#define VERSION 0
#define QUIET 0
* oclHashcat -only- functions
*/
-static void generate_source_kernel_filename (const uint attack_exec, const uint attack_kern, const uint kern_type, char *install_dir, char *kernel_file)
+static void generate_source_kernel_filename (const uint attack_exec, const uint attack_kern, const uint kern_type, char *install_dir, char *source_file)
+{
+ if (attack_exec == ATTACK_EXEC_ON_GPU)
+ {
+ if (attack_kern == ATTACK_KERN_STRAIGHT)
+ snprintf (source_file, 255, "%s/OpenCL/m%05d_a0.cl", install_dir, (int) kern_type);
+ else if (attack_kern == ATTACK_KERN_COMBI)
+ snprintf (source_file, 255, "%s/OpenCL/m%05d_a1.cl", install_dir, (int) kern_type);
+ else if (attack_kern == ATTACK_KERN_BF)
+ snprintf (source_file, 255, "%s/OpenCL/m%05d_a3.cl", install_dir, (int) kern_type);
+ }
+ else
+ snprintf (source_file, 255, "%s/OpenCL/m%05d.cl", install_dir, (int) kern_type);
+}
+
+static void generate_cached_kernel_filename (const uint attack_exec, const uint attack_kern, const uint kern_type, char *install_dir, char *device_name, char *device_version, char *driver_version, int vendor_id, char *cached_file)
{
if (attack_exec == ATTACK_EXEC_ON_GPU)
{
if (attack_kern == ATTACK_KERN_STRAIGHT)
- snprintf (kernel_file, 255, "%s/OpenCL/m%05d_a0.cl", install_dir, (int) kern_type);
+ snprintf (cached_file, 255, "%s/kernels/%d/m%05d_a0.%s_%s_%s_%d.kernel", install_dir, vendor_id, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
else if (attack_kern == ATTACK_KERN_COMBI)
- snprintf (kernel_file, 255, "%s/OpenCL/m%05d_a1.cl", install_dir, (int) kern_type);
+ snprintf (cached_file, 255, "%s/kernels/%d/m%05d_a1.%s_%s_%s_%d.kernel", install_dir, vendor_id, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
else if (attack_kern == ATTACK_KERN_BF)
- snprintf (kernel_file, 255, "%s/OpenCL/m%05d_a3.cl", install_dir, (int) kern_type);
+ snprintf (cached_file, 255, "%s/kernels/%d/m%05d_a3.%s_%s_%s_%d.kernel", install_dir, vendor_id, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
+ }
+ else
+ {
+ snprintf (cached_file, 255, "%s/kernels/%d/m%05d.%s_%s_%s_%d.kernel", install_dir, vendor_id, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
+ }
+}
+
+static void generate_source_kernel_mp_filename (const uint opti_type, const uint opts_type, char *install_dir, char *source_file)
+{
+ if ((opti_type & OPTI_TYPE_BRUTE_FORCE) && (opts_type & OPTS_TYPE_PT_GENERATE_BE))
+ {
+ snprintf (source_file, 255, "%s/OpenCL/markov_be.cl", install_dir);
+ }
+ else
+ {
+ snprintf (source_file, 255, "%s/OpenCL/markov_le.cl", install_dir);
+ }
+}
+
+static void generate_cached_kernel_mp_filename (const uint opti_type, const uint opts_type, char *install_dir, char *device_name, char *device_version, char *driver_version, int vendor_id, char *cached_file)
+{
+ if ((opti_type & OPTI_TYPE_BRUTE_FORCE) && (opts_type & OPTS_TYPE_PT_GENERATE_BE))
+ {
+ snprintf (cached_file, 255, "%s/kernels/%d/markov_be.%s_%s_%s_%d.kernel", install_dir, vendor_id, device_name, device_version, driver_version, COMPTIME);
}
else
- snprintf (kernel_file, 255, "%s/OpenCL/m%05d.cl", install_dir, (int) kern_type);
+ {
+ snprintf (cached_file, 255, "%s/kernels/%d/markov_le.%s_%s_%s_%d.kernel", install_dir, vendor_id, device_name, device_version, driver_version, COMPTIME);
+ }
+}
+
+static void generate_source_kernel_amp_filename (const uint attack_kern, char *install_dir, char *source_file)
+{
+ snprintf (source_file, 255, "%s/OpenCL/amp_a%d.cl", install_dir, attack_kern);
+}
+
+static void generate_cached_kernel_amp_filename (const uint attack_kern, char *install_dir, char *device_name, char *device_version, char *driver_version, int vendor_id, char *cached_file)
+{
+ snprintf (cached_file, 255, "%s/kernels/%d/amp_a%d.%s_%s_%s_%d.kernel", install_dir, vendor_id, attack_kern, device_name, device_version, driver_version, COMPTIME);
}
static uint convert_from_hex (char *line_buf, const uint line_len)
}
/**
- * kernel find
+ * default building options
*/
char build_opts[1024];
sprintf (build_opts, "-I. -IOpenCL/ -DVENDOR_ID=%d -DCUDA_ARCH=%d", vendor_id, (device_param->sm_major * 100) + device_param->sm_minor);
- struct stat st;
+ /**
+ * main kernel
+ */
- char kernel_file[256];
+ {
+ /**
+ * kernel source filename
+ */
- memset (kernel_file, 0, sizeof (kernel_file));
+ char source_file[256];
- size_t *kernel_lengths = (size_t *) mymalloc (sizeof (size_t));
+ memset (source_file, 0, sizeof (source_file));
- const unsigned char **kernel_sources = (const unsigned char **) mymalloc (sizeof (unsigned char *));
+ generate_source_kernel_filename (attack_exec, attack_kern, kern_type, install_dir, source_file);
- #ifdef BINARY_KERNEL
- if (force_jit_compilation == 0)
- {
- if (attack_exec == ATTACK_EXEC_ON_GPU)
- {
- if (attack_kern == ATTACK_KERN_STRAIGHT)
- snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d_a0.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
- else if (attack_kern == ATTACK_KERN_COMBI)
- snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d_a1.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
- else if (attack_kern == ATTACK_KERN_BF)
- snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d_a3.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
- }
- else
+ struct stat sst;
+
+ if (stat (source_file, &sst) == -1)
{
- snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
+ log_error ("ERROR: %s: %s", source_file, strerror (errno));
- if ((hash_mode == 8900) || (hash_mode == 9300))
- {
- snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d_%d_%d_%d_%d.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, data.salts_buf[0].scrypt_tmto, device_name, device_version, driver_version, COMPTIME);
- }
+ return -1;
}
- if (stat (kernel_file, &st) == -1)
+ /**
+ * kernel cached filename
+ */
+
+ char cached_file[256];
+
+ memset (cached_file, 0, sizeof (cached_file));
+
+ generate_cached_kernel_filename (attack_exec, attack_kern, kern_type, install_dir, device_name, device_version, driver_version, vendor_id, cached_file);
+
+ int cached = 1;
+
+ struct stat cst;
+
+ if (stat (cached_file, &cst) == -1)
{
- if (quiet == 0) log_info ("Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, kernel_file);
+ cached = 0;
+ }
- char module_file[256];
+ /**
+ * kernel compile or load
+ */
- memset (module_file, 0, sizeof (module_file));
+ size_t *kernel_lengths = (size_t *) mymalloc (sizeof (size_t));
- if (attack_exec == ATTACK_EXEC_ON_GPU)
- {
- if (attack_kern == ATTACK_KERN_STRAIGHT)
- snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d_a0.llvmir", install_dir, (int) kern_type);
- else if (attack_kern == ATTACK_KERN_COMBI)
- snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d_a1.llvmir", install_dir, (int) kern_type);
- else if (attack_kern == ATTACK_KERN_BF)
- snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d_a3.llvmir", install_dir, (int) kern_type);
- }
- else
+ const unsigned char **kernel_sources = (const unsigned char **) mymalloc (sizeof (unsigned char *));
+
+ if (force_jit_compilation == 0)
+ {
+ if (cached == 0)
{
- snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d.llvmir", install_dir, (int) kern_type);
+ if (quiet == 0) log_info ("Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, cached_file);
- if ((hash_mode == 8900) || (hash_mode == 9300))
- {
- snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d_%d_%d_%d_%d.llvmir", install_dir, (int) kern_type, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, data.salts_buf[0].scrypt_tmto);
- }
- }
+ load_kernel (source_file, 1, kernel_lengths, kernel_sources);
- load_kernel (module_file, 1, kernel_lengths, kernel_sources);
+ device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
- cl_program program = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const unsigned char **) kernel_sources, NULL);
+ hc_clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
- local_free (kernel_sources[0]);
+ size_t binary_size;
- hc_clBuildProgram (program, 1, &device_param->device, build_opts, NULL, NULL);
+ clGetProgramInfo (device_param->program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
- size_t binary_size;
+ unsigned char *binary = (unsigned char *) mymalloc (binary_size);
- clGetProgramInfo (program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
+ clGetProgramInfo (device_param->program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
- unsigned char *binary = (unsigned char *) mymalloc (binary_size);
+ writeProgramBin (cached_file, binary, binary_size);
- clGetProgramInfo (program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
+ local_free (binary);
+ }
+ else
+ {
+ if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size);
- writeProgramBin (kernel_file, binary, binary_size);
+ load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
- local_free (binary);
+ device_param->program = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const unsigned char **) kernel_sources, NULL);
- stat (kernel_file, &st); // to reload filesize
+ hc_clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
+ }
}
- }
- else
- {
- generate_source_kernel_filename (attack_exec, attack_kern, kern_type, install_dir, kernel_file);
-
- if (stat (kernel_file, &st) == -1)
+ else
{
- log_error ("ERROR: %s: %s", kernel_file, strerror (errno));
+ if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, source_file, sst.st_size);
- return -1;
- }
- }
+ load_kernel (source_file, 1, kernel_lengths, kernel_sources);
- #else
+ device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
- generate_source_kernel_filename (attack_exec, attack_kern, kern_type, install_dir, kernel_file);
+ if (force_jit_compilation == 1500)
+ {
+ sprintf (build_opts, "%s -DDESCRYPT_SALT=%d", build_opts, data.salts_buf[0].salt_buf[0]);
+ }
+ else if (force_jit_compilation == 8900)
+ {
+ sprintf (build_opts, "%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, data.salts_buf[0].scrypt_tmto);
+ }
- if (stat (kernel_file, &st) == -1)
- {
- log_error ("ERROR: %s: %s", kernel_file, strerror (errno));
+ hc_clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
+ }
- return -1;
- }
+ local_free (kernel_lengths);
+ local_free (kernel_sources[0]);
+ local_free (kernel_sources);
- #endif
+ // this is mostly for debug
- load_kernel (kernel_file, 1, kernel_lengths, kernel_sources);
+ size_t ret_val_size = 0;
- if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, kernel_file, st.st_size);
+ clGetProgramBuildInfo (device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
- #ifdef BINARY_KERNEL
- if (force_jit_compilation == 0)
- {
- device_param->program = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const unsigned char **) kernel_sources, NULL);
- }
- else
- {
- device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
- }
- #else
- device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
- #endif
+ if (ret_val_size > 2)
+ {
+ char *build_log = (char *) mymalloc (ret_val_size + 1);
- local_free (kernel_lengths);
+ memset (build_log, 0, ret_val_size + 1);
- local_free (kernel_sources[0]);
+ clGetProgramBuildInfo (device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
- local_free (kernel_sources)
+ puts (build_log);
+
+ myfree (build_log);
+ }
+ }
/**
- * kernel mp find
+ * word generator kernel
*/
if (attack_mode != ATTACK_MODE_STRAIGHT)
{
- char kernel_mp_file[256];
+ /**
+ * kernel mp source filename
+ */
- memset (kernel_mp_file, 0, sizeof (kernel_mp_file));
+ char source_file[256];
- size_t *kernel_mp_lengths = (size_t *) mymalloc (sizeof (size_t));
+ memset (source_file, 0, sizeof (source_file));
- const unsigned char **kernel_mp_sources = (const unsigned char **) mymalloc (sizeof (unsigned char *));
+ generate_source_kernel_mp_filename (opti_type, opts_type, install_dir, source_file);
- #ifdef BINARY_KERNEL
- if ((opti_type & OPTI_TYPE_BRUTE_FORCE) && (opts_type & OPTS_TYPE_PT_GENERATE_BE))
- {
- snprintf (kernel_mp_file, sizeof (kernel_mp_file) - 1, "%s/kernels/4098/markov_be.%s_%s_%s_%d.kernel", install_dir, device_name, device_version, driver_version, COMPTIME);
- }
- else
+ struct stat sst;
+
+ if (stat (source_file, &sst) == -1)
{
- snprintf (kernel_mp_file, sizeof (kernel_mp_file) - 1, "%s/kernels/4098/markov_le.%s_%s_%s_%d.kernel", install_dir, device_name, device_version, driver_version, COMPTIME);
+ log_error ("ERROR: %s: %s", source_file, strerror (errno));
+
+ return -1;
}
- if (stat (kernel_mp_file, &st) == -1)
- {
- if (quiet == 0) log_info ("Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, kernel_mp_file);
+ /**
+ * kernel mp cached filename
+ */
- char module_mp_file[256];
+ char cached_file[256];
- memset (module_mp_file, 0, sizeof (module_mp_file));
+ memset (cached_file, 0, sizeof (cached_file));
- if ((opti_type & OPTI_TYPE_BRUTE_FORCE) && (opts_type & OPTS_TYPE_PT_GENERATE_BE))
- {
- snprintf (module_mp_file, sizeof (module_mp_file) - 1, "%s/kernels/4098/markov_be.llvmir", install_dir);
- }
- else
- {
- snprintf (module_mp_file, sizeof (module_mp_file) - 1, "%s/kernels/4098/markov_le.llvmir", install_dir);
- }
+ generate_cached_kernel_mp_filename (opti_type, opts_type, install_dir, device_name, device_version, driver_version, vendor_id, cached_file);
+
+ int cached = 1;
+
+ struct stat cst;
+
+ if (stat (cached_file, &cst) == -1)
+ {
+ cached = 0;
+ }
- load_kernel (module_mp_file, 1, kernel_mp_lengths, kernel_mp_sources);
+ /**
+ * kernel compile or load
+ */
- cl_program program_mp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_mp_lengths, (const unsigned char **) kernel_mp_sources, NULL);
+ size_t *kernel_lengths = (size_t *) mymalloc (sizeof (size_t));
- local_free (kernel_mp_sources[0]);
+ const unsigned char **kernel_sources = (const unsigned char **) mymalloc (sizeof (unsigned char *));
- hc_clBuildProgram (program_mp, 1, &device_param->device, build_opts, NULL, NULL);
+ if (cached == 0)
+ {
+ if (quiet == 0) log_info ("Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, cached_file);
+
+ load_kernel (source_file, 1, kernel_lengths, kernel_sources);
- size_t binary_mp_size;
+ device_param->program_mp = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
- clGetProgramInfo (program_mp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_mp_size, NULL);
+ hc_clBuildProgram (device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
- unsigned char *binary_mp = (unsigned char *) mymalloc (binary_mp_size);
+ size_t binary_size;
- clGetProgramInfo (program_mp, CL_PROGRAM_BINARIES, sizeof (binary_mp), &binary_mp, NULL);
+ clGetProgramInfo (device_param->program_mp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
- writeProgramBin (kernel_mp_file, binary_mp, binary_mp_size);
+ unsigned char *binary = (unsigned char *) mymalloc (binary_size);
- local_free (binary_mp);
+ clGetProgramInfo (device_param->program_mp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
- stat (kernel_mp_file, &st); // to reload filesize
- }
+ writeProgramBin (cached_file, binary, binary_size);
- #else
- if ((opti_type & OPTI_TYPE_BRUTE_FORCE) && (opts_type & OPTS_TYPE_PT_GENERATE_BE))
- {
- snprintf (kernel_mp_file, sizeof (kernel_mp_file) - 1, "%s/OpenCL/markov_be.cl", install_dir);
+ local_free (binary);
}
else
{
- snprintf (kernel_mp_file, sizeof (kernel_mp_file) - 1, "%s/OpenCL/markov_le.cl", install_dir);
- }
+ if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size);
- if (stat (kernel_mp_file, &st) == -1)
- {
- log_error ("ERROR: %s: %s", kernel_mp_file, strerror (errno));
+ load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
- return -1;
+ device_param->program_mp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const unsigned char **) kernel_sources, NULL);
+
+ hc_clBuildProgram (device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
}
- #endif
+ local_free (kernel_lengths);
+ local_free (kernel_sources[0]);
+ local_free (kernel_sources);
- load_kernel (kernel_mp_file, 1, kernel_mp_lengths, kernel_mp_sources);
+ // this is mostly for debug
- if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, kernel_mp_file, st.st_size);
+ size_t ret_val_size = 0;
- #ifdef BINARY_KERNEL
- device_param->program_mp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_mp_lengths, (const unsigned char **) kernel_mp_sources, NULL);
- #else
- device_param->program_mp = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_mp_sources, NULL);
- #endif
+ clGetProgramBuildInfo (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);
- local_free (kernel_mp_lengths);
+ memset (build_log, 0, ret_val_size + 1);
- local_free (kernel_mp_sources[0]);
+ clGetProgramBuildInfo (device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
- local_free (kernel_mp_sources);
+ puts (build_log);
+
+ myfree (build_log);
+ }
}
/**
- * kernel amp find
+ * amplifier kernel
*/
if (attack_exec == ATTACK_EXEC_ON_GPU)
{
- // nothing to do
+
}
else
{
- char kernel_amp_file[256];
+ /**
+ * kernel amp source filename
+ */
- memset (kernel_amp_file, 0, sizeof (kernel_amp_file));
+ char source_file[256];
- size_t *kernel_amp_lengths = (size_t *) mymalloc (sizeof (size_t));
+ memset (source_file, 0, sizeof (source_file));
- const unsigned char **kernel_amp_sources = (const unsigned char **) mymalloc (sizeof (unsigned char *));
+ generate_source_kernel_amp_filename (attack_kern, install_dir, source_file);
- #ifdef BINARY_KERNEL
- snprintf (kernel_amp_file, sizeof (kernel_amp_file) - 1, "%s/kernels/4098/amp_a%d.%s_%s_%s_%d.kernel", install_dir, attack_kern, device_name, device_version, driver_version, COMPTIME);
+ struct stat sst;
- if (stat (kernel_amp_file, &st) == -1)
+ if (stat (source_file, &sst) == -1)
{
- if (quiet == 0) log_info ("Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, kernel_amp_file);
-
- char module_amp_file[256];
+ log_error ("ERROR: %s: %s", source_file, strerror (errno));
- memset (module_amp_file, 0, sizeof (module_amp_file));
-
- snprintf (module_amp_file, sizeof (module_amp_file) - 1, "%s/kernels/4098/amp_a%d.llvmir", install_dir, attack_kern);
-
- load_kernel (module_amp_file, 1, kernel_amp_lengths, kernel_amp_sources);
+ return -1;
+ }
- cl_program program_amp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_amp_lengths, (const unsigned char **) kernel_amp_sources, NULL);
+ /**
+ * kernel amp cached filename
+ */
- local_free (kernel_amp_sources[0]);
+ char cached_file[256];
- hc_clBuildProgram (program_amp, 1, &device_param->device, build_opts, NULL, NULL);
+ memset (cached_file, 0, sizeof (cached_file));
- size_t binary_amp_size;
+ generate_cached_kernel_amp_filename (attack_kern, install_dir, device_name, device_version, driver_version, vendor_id, cached_file);
- clGetProgramInfo (program_amp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_amp_size, NULL);
+ int cached = 1;
- unsigned char *binary_amp = (unsigned char *) mymalloc (binary_amp_size);
+ struct stat cst;
- clGetProgramInfo (program_amp, CL_PROGRAM_BINARIES, sizeof (binary_amp), &binary_amp, NULL);
+ if (stat (cached_file, &cst) == -1)
+ {
+ cached = 0;
+ }
- writeProgramBin (kernel_amp_file, binary_amp, binary_amp_size);
+ /**
+ * kernel compile or load
+ */
- local_free (binary_amp);
+ size_t *kernel_lengths = (size_t *) mymalloc (sizeof (size_t));
- stat (kernel_amp_file, &st); // to reload filesize
- }
- #else
- snprintf (kernel_amp_file, sizeof (kernel_amp_file) - 1, "%s/OpenCL/amp_a%d.cl", install_dir, attack_kern);
+ const unsigned char **kernel_sources = (const unsigned char **) mymalloc (sizeof (unsigned char *));
- if (stat (kernel_amp_file, &st) == -1)
+ if (cached == 0)
{
- log_error ("ERROR: %s: %s", kernel_amp_file, strerror (errno));
-
- return -1;
- }
- #endif
+ if (quiet == 0) log_info ("Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, cached_file);
- load_kernel (kernel_amp_file, 1, kernel_amp_lengths, kernel_amp_sources);
+ load_kernel (source_file, 1, kernel_lengths, kernel_sources);
- if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, kernel_amp_file, st.st_size);
+ device_param->program_amp = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
- #ifdef BINARY_KERNEL
- device_param->program_amp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_amp_lengths, (const unsigned char **) kernel_amp_sources, NULL);
- #else
- device_param->program_amp = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_amp_sources, NULL);
- #endif
+ hc_clBuildProgram (device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
- local_free (kernel_amp_lengths);
+ size_t binary_size;
- local_free (kernel_amp_sources[0]);
+ clGetProgramInfo (device_param->program_amp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
- local_free (kernel_amp_sources);
- }
+ unsigned char *binary = (unsigned char *) mymalloc (binary_size);
- /**
- * kernel compile
- */
+ clGetProgramInfo (device_param->program_amp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
- //#ifdef BINARY_KERNEL
+ writeProgramBin (cached_file, binary, binary_size);
- if (force_jit_compilation == 0)
- {
- // nothing to do
- }
- else if (force_jit_compilation == 1500)
- {
- sprintf (build_opts, "%s -DDESCRYPT_SALT=%d", build_opts, data.salts_buf[0].salt_buf[0]);
- }
- else if (force_jit_compilation == 8900)
- {
- sprintf (build_opts, "%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, data.salts_buf[0].scrypt_tmto);
- }
+ local_free (binary);
+ }
+ else
+ {
+ if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size);
- //#endif
+ load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
- clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
+ device_param->program_amp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const unsigned char **) kernel_sources, NULL);
- size_t ret_val_size = 0;
+ hc_clBuildProgram (device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
+ }
- clGetProgramBuildInfo (device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
+ local_free (kernel_lengths);
+ local_free (kernel_sources[0]);
+ local_free (kernel_sources);
- if (ret_val_size > 2)
- {
- char *build_log = (char *) malloc (ret_val_size + 1);
+ // this is mostly for debug
- memset (build_log, 0, ret_val_size + 1);
+ size_t ret_val_size = 0;
- clGetProgramBuildInfo (device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
+ clGetProgramBuildInfo (device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
- puts (build_log);
+ if (ret_val_size > 2)
+ {
+ char *build_log = (char *) mymalloc (ret_val_size + 1);
- free (build_log);
- }
+ memset (build_log, 0, ret_val_size + 1);
- if (attack_mode != ATTACK_MODE_STRAIGHT)
- {
- hc_clBuildProgram (device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
- }
+ clGetProgramBuildInfo (device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
- if (attack_exec == ATTACK_EXEC_ON_GPU)
- {
- // nothing to do
- }
- else
- {
- hc_clBuildProgram (device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
- }
+ puts (build_log);
- /**
- * amp is not independent
- */
-
- if (attack_exec == ATTACK_EXEC_ON_GPU)
- {
- // nothing to do
- }
- else
- {
- device_param->kernel_amp = hc_clCreateKernel (device_param->program_amp, "amp");
+ myfree (build_log);
+ }
}
/**
device_param->kernel_mp = hc_clCreateKernel (device_param->program_mp, "C_markov");
}
+ if (attack_exec == ATTACK_EXEC_ON_GPU)
+ {
+ // nothing to do
+ }
+ else
+ {
+ device_param->kernel_amp = hc_clCreateKernel (device_param->program_amp, "amp");
+ }
+
if (attack_exec == ATTACK_EXEC_ON_GPU)
{
// nothing to do
projects / hashcat.git / commitdiff