From: jsteube Date: Mon, 21 Dec 2015 11:01:38 +0000 (+0100) Subject: Fix caching system for use with AMD and NV, drop BINARY_KERNEL define X-Git-Tag: v3.00-beta~584^2~52 X-Git-Url: https://www.flypig.org.uk/git/?a=commitdiff_plain;h=378258d789929a3444d36582a418d9cbaaef0b66;p=hashcat.git Fix caching system for use with AMD and NV, drop BINARY_KERNEL define --- diff --git a/include/ext_cuda.h b/include/ext_cuda.h deleted file mode 100644 index 18692de..0000000 --- a/include/ext_cuda.h +++ /dev/null @@ -1,64 +0,0 @@ -/** - * Author......: Jens Steube - * License.....: MIT - */ - -#ifndef EXT_CUDA_H -#define EXT_CUDA_H - -#include - -#include - -void hc_cuDeviceGetCount (int *count); -void hc_cuDeviceGet (CUdevice *device, int ordinal); -void hc_cuDeviceGetName (char *name, int len, CUdevice dev); -void hc_cuDeviceTotalMem (size_t *bytes, CUdevice dev); -void hc_cuDeviceGetAttribute (int *pi, CUdevice_attribute attrib, CUdevice dev); -void hc_cuCtxCreate (CUcontext *pctx, unsigned int flags, CUdevice dev); -void hc_cuMemAllocHost (void **pp, size_t bytesize); -void hc_cuMemAlloc (CUdeviceptr *dptr, size_t bytesize); -void hc_cuMemsetD8 (CUdeviceptr dstDevice, unsigned char uc, unsigned int N); -void hc_cuMemsetD32 (CUdeviceptr dstDevice, unsigned int ui, unsigned int N); -void hc_cuMemcpyDtoD (CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount); -void hc_cuMemcpyDtoDAsync (CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream); -void hc_cuMemcpyHtoD (CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount); -void hc_cuMemcpyHtoDAsync (CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount, CUstream hStream); -void hc_cuMemcpyDtoH (void *dstHost, CUdeviceptr srcDevice, size_t ByteCount); -void hc_cuMemcpyDtoHAsync (void *dstHost, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream); -void hc_cuEventCreate (CUevent *phEvent, unsigned int Flags); -void hc_cuStreamCreate (CUstream *phStream, unsigned int Flags); -void hc_cuDeviceComputeCapability (int *major, int *minor, CUdevice dev); -void hc_cuModuleLoad (CUmodule *module, const char *fname); -void hc_cuModuleLoadData (CUmodule *module, const void *image); -void hc_cuModuleGetFunction (CUfunction *hfunc, CUmodule hmod, const char *name); -void hc_cuFuncSetBlockShape (CUfunction hfunc, int x, int y, int z); -void hc_cuParamSetSize (CUfunction hfunc, unsigned int numbytes); -void hc_cuParamSetv (CUfunction hfunc, int offset, void *ptr, unsigned int numbytes); -void hc_cuParamSeti (CUfunction hfunc, int offset, unsigned int value); -void hc_cuModuleGetGlobal (CUdeviceptr *dptr, size_t *bytes, CUmodule hmod, const char *name); -void hc_cuCtxPopCurrent (CUcontext *pctx); -void hc_cuCtxPushCurrent (CUcontext ctx); -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); -void hc_cuLaunchGrid (CUfunction f, int grid_width, int grid_height); -void hc_cuLaunchGridAsync (CUfunction f, int grid_width, int grid_height, CUstream hStream); -void hc_cuEventSynchronize (CUevent hEvent); -void hc_cuStreamSynchronize (CUstream hStream); -void hc_cuMemFreeHost (void *p); -void hc_cuMemFree (CUdeviceptr dptr); -void hc_cuEventDestroy (CUevent hEvent); -void hc_cuStreamDestroy (CUstream hStream); -void hc_cuModuleUnload (CUmodule hmod); -void hc_cuCtxDestroy (CUcontext ctx); -void hc_cuCtxAttach (CUcontext *pctx, unsigned int flags); -void hc_cuCtxDetach (CUcontext ctx); -void hc_cuCtxSynchronize (void); -void hc_cuCtxSetCacheConfig (CUfunc_cache config); -void hc_cuDriverGetVersion (int *driverVersion); -void hc_cuModuleLoadDataEx (CUmodule *module, const void *image, unsigned int numOptions, CUjit_option *options, void **optionValues); -void hc_cuLinkAddFile (CUlinkState state, CUjitInputType type, const char *path, unsigned int numOptions, CUjit_option *options, void **optionValues); -void hc_cuLinkComplete (CUlinkState state, void **cubinOut, size_t *sizeOut); -void hc_cuLinkCreate (unsigned int numOptions, CUjit_option *options, void **optionValues, CUlinkState *stateOut); -void hc_cuLinkDestroy (CUlinkState state); - -#endif diff --git a/src/ext_OpenCL.c b/src/ext_OpenCL.c index 17f287e..a632545 100644 --- a/src/ext_OpenCL.c +++ b/src/ext_OpenCL.c @@ -268,7 +268,7 @@ cl_kernel hc_clCreateKernel (cl_program program, const char *kernel_name) if (CL_err != CL_SUCCESS) { - log_error ("ERROR: %s %d\n", "clCreateKernel()", CL_err); + log_error ("ERROR: %s %d - %s\n", "clCreateKernel()", CL_err, kernel_name); exit (-1); } diff --git a/src/ext_cuda.c b/src/ext_cuda.c deleted file mode 100644 index 74b4b1d..0000000 --- a/src/ext_cuda.c +++ /dev/null @@ -1,606 +0,0 @@ -/** - * Author......: Jens Steube - * License.....: MIT - */ - -#include - -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); - } -} diff --git a/src/oclHashcat.c b/src/oclHashcat.c index b43d1d3..98398cd 100644 --- a/src/oclHashcat.c +++ b/src/oclHashcat.c @@ -18,10 +18,6 @@ const uint RESTORE_MIN = 201; #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 @@ -1597,19 +1593,70 @@ static void status_benchmark () * 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) @@ -12958,7 +13005,7 @@ int main (int argc, char **argv) } /** - * kernel find + * default building options */ char build_opts[1024]; @@ -12967,382 +13014,356 @@ int main (int argc, char **argv) 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); + } } /** @@ -13715,6 +13736,15 @@ int main (int argc, char **argv) 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