arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-22 08:08:10 +00:00
Code Issues Releases Wiki Activity

Fix caching system for use with AMD and NV, drop BINARY_KERNEL define

Browse Source
This commit is contained in:
jsteube 2015-12-21 12:01:38 +01:00
parent c8f7b7b5d3
commit 378258d789
4 changed files with 359 additions and 999 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

64
include/ext_cuda.h
View File

@ -1,64 +0,0 @@
/**
* Author......: Jens Steube <jens.steube@gmail.com>
* License.....: MIT
*/
#ifndef EXT_CUDA_H
#define EXT_CUDA_H
#include <common.h>
#include <cuda.h>
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

2
src/ext_OpenCL.c
View File

@ -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);
}

606
src/ext_cuda.c
View File

@ -1,606 +0,0 @@
/**
* 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);
}
}

686
src/oclHashcat.c
View File

@ -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 (kernel_file, 255, "%s/OpenCL/m%05d_a0.cl", install_dir, (int) kern_type);
snprintf (source_file, 255, "%s/OpenCL/m%05d_a0.cl", install_dir, (int) kern_type);
else if (attack_kern == ATTACK_KERN_COMBI)
snprintf (kernel_file, 255, "%s/OpenCL/m%05d_a1.cl", install_dir, (int) kern_type);
snprintf (source_file, 255, "%s/OpenCL/m%05d_a1.cl", install_dir, (int) kern_type);
else if (attack_kern == ATTACK_KERN_BF)
snprintf (kernel_file, 255, "%s/OpenCL/m%05d_a3.cl", install_dir, (int) kern_type);
snprintf (source_file, 255, "%s/OpenCL/m%05d_a3.cl", install_dir, (int) kern_type);
}
else
snprintf (kernel_file, 255, "%s/OpenCL/m%05d.cl", install_dir, (int) kern_type);
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 (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 (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 (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 (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];
memset (kernel_file, 0, sizeof (kernel_file));
size_t *kernel_lengths = (size_t *) mymalloc (sizeof (size_t));
const unsigned char **kernel_sources = (const unsigned char **) mymalloc (sizeof (unsigned char *));
#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
{
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);
/**
* kernel source filename
*/
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);
}
char source_file[256];
memset (source_file, 0, sizeof (source_file));
generate_source_kernel_filename (attack_exec, attack_kern, kern_type, install_dir, source_file);
struct stat sst;
if (stat (source_file, &sst) == -1)
{
log_error ("ERROR: %s: %s", source_file, strerror (errno));
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)
const unsigned char **kernel_sources = (const unsigned char **) mymalloc (sizeof (unsigned char *));
if (force_jit_compilation == 0)
{
if (cached == 0)
{
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);
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);
device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
hc_clBuildProgram (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);
unsigned char *binary = (unsigned char *) mymalloc (binary_size);
clGetProgramInfo (device_param->program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
writeProgramBin (cached_file, binary, binary_size);
local_free (binary);
}
else
{
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 (%ld bytes)", device_id + 1, cached_file, cst.st_size);
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 (cached_file, 1, kernel_lengths, kernel_sources);
device_param->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);
}
}
else
{
if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, source_file, sst.st_size);
load_kernel (source_file, 1, kernel_lengths, kernel_sources);
device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
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);
}
load_kernel (module_file, 1, kernel_lengths, kernel_sources);
hc_clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
}
cl_program program = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const unsigned char **) kernel_sources, NULL);
local_free (kernel_lengths);
local_free (kernel_sources[0]);
local_free (kernel_sources);
local_free (kernel_sources[0]);
// this is mostly for debug
hc_clBuildProgram (program, 1, &device_param->device, build_opts, NULL, NULL);
size_t ret_val_size = 0;
clGetProgramBuildInfo (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);
memset (build_log, 0, ret_val_size + 1);
clGetProgramBuildInfo (device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
puts (build_log);
myfree (build_log);
}
}
/**
* word generator kernel
*/
if (attack_mode != ATTACK_MODE_STRAIGHT)
{
/**
* kernel mp source filename
*/
char source_file[256];
memset (source_file, 0, sizeof (source_file));
generate_source_kernel_mp_filename (opti_type, opts_type, install_dir, source_file);
struct stat sst;
if (stat (source_file, &sst) == -1)
{
log_error ("ERROR: %s: %s", source_file, strerror (errno));
return -1;
}
/**
* kernel mp cached filename
*/
char cached_file[256];
memset (cached_file, 0, sizeof (cached_file));
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;
}
/**
* kernel compile or load
*/
size_t *kernel_lengths = (size_t *) mymalloc (sizeof (size_t));
const unsigned char **kernel_sources = (const unsigned char **) mymalloc (sizeof (unsigned char *));
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);
device_param->program_mp = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
hc_clBuildProgram (device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
size_t binary_size;
clGetProgramInfo (program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
clGetProgramInfo (device_param->program_mp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
unsigned char *binary = (unsigned char *) mymalloc (binary_size);
clGetProgramInfo (program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
clGetProgramInfo (device_param->program_mp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
writeProgramBin (kernel_file, binary, binary_size);
writeProgramBin (cached_file, binary, binary_size);
local_free (binary);
stat (kernel_file, &st); // to reload filesize
}
}
else
{
generate_source_kernel_filename (attack_exec, attack_kern, kern_type, install_dir, kernel_file);
if (stat (kernel_file, &st) == -1)
{
log_error ("ERROR: %s: %s", kernel_file, strerror (errno));
return -1;
}
}
#else
generate_source_kernel_filename (attack_exec, attack_kern, kern_type, install_dir, kernel_file);
if (stat (kernel_file, &st) == -1)
{
log_error ("ERROR: %s: %s", kernel_file, strerror (errno));
return -1;
}
#endif
load_kernel (kernel_file, 1, kernel_lengths, kernel_sources);
if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, kernel_file, st.st_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
local_free (kernel_lengths);
local_free (kernel_sources[0]);
local_free (kernel_sources)
/**
* kernel mp find
*/
if (attack_mode != ATTACK_MODE_STRAIGHT)
{
char kernel_mp_file[256];
memset (kernel_mp_file, 0, sizeof (kernel_mp_file));
size_t *kernel_mp_lengths = (size_t *) mymalloc (sizeof (size_t));
const unsigned char **kernel_mp_sources = (const unsigned char **) mymalloc (sizeof (unsigned char *));
#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
{
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);
if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_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 unsigned char **) kernel_sources, NULL);
hc_clBuildProgram (device_param->program_mp, 1, &device_param->device, build_opts, NULL, NULL);
}
if (stat (kernel_mp_file, &st) == -1)
local_free (kernel_lengths);
local_free (kernel_sources[0]);
local_free (kernel_sources);
// this is mostly for debug
size_t ret_val_size = 0;
clGetProgramBuildInfo (device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
if (ret_val_size > 2)
{
if (quiet == 0) log_info ("Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, kernel_mp_file);
char *build_log = (char *) mymalloc (ret_val_size + 1);
char module_mp_file[256];
memset (build_log, 0, ret_val_size + 1);
memset (module_mp_file, 0, sizeof (module_mp_file));
clGetProgramBuildInfo (device_param->program_mp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
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);
}
puts (build_log);
load_kernel (module_mp_file, 1, kernel_mp_lengths, kernel_mp_sources);
myfree (build_log);
}
}
cl_program program_mp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_mp_lengths, (const unsigned char **) kernel_mp_sources, NULL);
/**
* amplifier kernel
*/
local_free (kernel_mp_sources[0]);
if (attack_exec == ATTACK_EXEC_ON_GPU)
{
hc_clBuildProgram (program_mp, 1, &device_param->device, build_opts, NULL, NULL);
}
else
{
/**
* kernel amp source filename
*/
size_t binary_mp_size;
char source_file[256];
clGetProgramInfo (program_mp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_mp_size, NULL);
memset (source_file, 0, sizeof (source_file));
unsigned char *binary_mp = (unsigned char *) mymalloc (binary_mp_size);
generate_source_kernel_amp_filename (attack_kern, install_dir, source_file);
clGetProgramInfo (program_mp, CL_PROGRAM_BINARIES, sizeof (binary_mp), &binary_mp, NULL);
struct stat sst;
writeProgramBin (kernel_mp_file, binary_mp, binary_mp_size);
if (stat (source_file, &sst) == -1)
{
log_error ("ERROR: %s: %s", source_file, strerror (errno));
local_free (binary_mp);
stat (kernel_mp_file, &st); // to reload filesize
return -1;
}
#else
if ((opti_type & OPTI_TYPE_BRUTE_FORCE) && (opts_type & OPTS_TYPE_PT_GENERATE_BE))
/**
* kernel amp cached filename
*/
char cached_file[256];
memset (cached_file, 0, sizeof (cached_file));
generate_cached_kernel_amp_filename (attack_kern, install_dir, device_name, device_version, driver_version, vendor_id, cached_file);
int cached = 1;
struct stat cst;
if (stat (cached_file, &cst) == -1)
{
snprintf (kernel_mp_file, sizeof (kernel_mp_file) - 1, "%s/OpenCL/markov_be.cl", install_dir);
cached = 0;
}
/**
* kernel compile or load
*/
size_t *kernel_lengths = (size_t *) mymalloc (sizeof (size_t));
const unsigned char **kernel_sources = (const unsigned char **) mymalloc (sizeof (unsigned char *));
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);
device_param->program_amp = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
hc_clBuildProgram (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);
unsigned char *binary = (unsigned char *) mymalloc (binary_size);
clGetProgramInfo (device_param->program_amp, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
writeProgramBin (cached_file, binary, binary_size);
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);
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 unsigned char **) kernel_sources, NULL);
hc_clBuildProgram (device_param->program_amp, 1, &device_param->device, build_opts, NULL, NULL);
}
if (stat (kernel_mp_file, &st) == -1)
local_free (kernel_lengths);
local_free (kernel_sources[0]);
local_free (kernel_sources);
// this is mostly for debug
size_t ret_val_size = 0;
clGetProgramBuildInfo (device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
if (ret_val_size > 2)
{
log_error ("ERROR: %s: %s", kernel_mp_file, strerror (errno));
char *build_log = (char *) mymalloc (ret_val_size + 1);
return -1;
memset (build_log, 0, ret_val_size + 1);
clGetProgramBuildInfo (device_param->program_amp, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
puts (build_log);
myfree (build_log);
}
#endif
load_kernel (kernel_mp_file, 1, kernel_mp_lengths, kernel_mp_sources);
if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, kernel_mp_file, st.st_size);
#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
local_free (kernel_mp_lengths);
local_free (kernel_mp_sources[0]);
local_free (kernel_mp_sources);
}
/**
* kernel amp find
*/
if (attack_exec == ATTACK_EXEC_ON_GPU)
{
// nothing to do
}
else
{
char kernel_amp_file[256];
memset (kernel_amp_file, 0, sizeof (kernel_amp_file));
size_t *kernel_amp_lengths = (size_t *) mymalloc (sizeof (size_t));
const unsigned char **kernel_amp_sources = (const unsigned char **) mymalloc (sizeof (unsigned char *));
#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);
if (stat (kernel_amp_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_amp_file);
char module_amp_file[256];
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);
cl_program program_amp = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_amp_lengths, (const unsigned char **) kernel_amp_sources, NULL);
local_free (kernel_amp_sources[0]);
hc_clBuildProgram (program_amp, 1, &device_param->device, build_opts, NULL, NULL);
size_t binary_amp_size;
clGetProgramInfo (program_amp, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_amp_size, NULL);
unsigned char *binary_amp = (unsigned char *) mymalloc (binary_amp_size);
clGetProgramInfo (program_amp, CL_PROGRAM_BINARIES, sizeof (binary_amp), &binary_amp, NULL);
writeProgramBin (kernel_amp_file, binary_amp, binary_amp_size);
local_free (binary_amp);
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);
if (stat (kernel_amp_file, &st) == -1)
{
log_error ("ERROR: %s: %s", kernel_amp_file, strerror (errno));
return -1;
}
#endif
load_kernel (kernel_amp_file, 1, kernel_amp_lengths, kernel_amp_sources);
if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, kernel_amp_file, st.st_size);
#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
local_free (kernel_amp_lengths);
local_free (kernel_amp_sources[0]);
local_free (kernel_amp_sources);
}
/**
* kernel compile
*/
//#ifdef BINARY_KERNEL
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);
}
//#endif
clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
size_t ret_val_size = 0;
clGetProgramBuildInfo (device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
if (ret_val_size > 2)
{
char *build_log = (char *) malloc (ret_val_size + 1);
memset (build_log, 0, ret_val_size + 1);
clGetProgramBuildInfo (device_param->program, device_param->device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
puts (build_log);
free (build_log);
}
if (attack_mode != ATTACK_MODE_STRAIGHT)
{
hc_clBuildProgram (device_param->program_mp, 1, &device_param->device, build_opts, NULL, 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);
}
/**
* 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");
}
/**
@ -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