mirror of https://github.com/hashcat/hashcat.git
You can not select more than 25 topics
Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
1089 lines
28 KiB
1089 lines
28 KiB
/**
|
|
* Author......: See docs/credits.txt
|
|
* License.....: MIT
|
|
*/
|
|
|
|
#include "common.h"
|
|
#include "types.h"
|
|
#include "memory.h"
|
|
#include "event.h"
|
|
#include "ext_hip.h"
|
|
|
|
#include "dynloader.h"
|
|
|
|
int hip_init (void *hashcat_ctx)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
memset (hip, 0, sizeof (HIP_PTR));
|
|
|
|
#if defined (_WIN)
|
|
hip->lib = hc_dlopen ("amdhip64.dll");
|
|
#elif defined (__APPLE__)
|
|
hip->lib = hc_dlopen ("fixme.dylib");
|
|
#elif defined (__CYGWIN__)
|
|
hip->lib = hc_dlopen ("amdhip64.dll");
|
|
#else
|
|
hip->lib = hc_dlopen ("libamdhip64.so");
|
|
#endif
|
|
|
|
if (hip->lib == NULL) return -1;
|
|
|
|
// finding the right symbol is a PITA,
|
|
#define HC_LOAD_FUNC_HIP(ptr,name,hipname,type,libname,noerr) \
|
|
do { \
|
|
ptr->name = (type) hc_dlsym ((ptr)->lib, #hipname); \
|
|
if ((noerr) != -1) { \
|
|
if (!(ptr)->name) { \
|
|
if ((noerr) == 1) { \
|
|
event_log_error (hashcat_ctx, "%s is missing from %s shared library.", #name, #libname); \
|
|
return -1; \
|
|
} \
|
|
if ((noerr) != 1) { \
|
|
event_log_warning (hashcat_ctx, "%s is missing from %s shared library.", #name, #libname); \
|
|
return 0; \
|
|
} \
|
|
} \
|
|
} \
|
|
} while (0)
|
|
|
|
// finding the right symbol is a PITA, because of the _v2 suffix
|
|
// a good reference is cuda.h itself
|
|
// this needs to be verified for each new cuda release
|
|
|
|
HC_LOAD_FUNC_HIP (hip, hipCtxCreate, hipCtxCreate, HIP_HIPCTXCREATE, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipCtxDestroy, hipCtxDestroy, HIP_HIPCTXDESTROY, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipCtxPopCurrent, hipCtxPopCurrent, HIP_HIPCTXPOPCURRENT, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipCtxPushCurrent, hipCtxPushCurrent, HIP_HIPCTXPUSHCURRENT, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipCtxSetCurrent, hipCtxSetCurrent, HIP_HIPCTXSETCURRENT, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipCtxSynchronize, hipCtxSynchronize, HIP_HIPCTXSYNCHRONIZE, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipDeviceGet, hipDeviceGet, HIP_HIPDEVICEGET, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipDeviceGetAttribute, hipDeviceGetAttribute, HIP_HIPDEVICEGETATTRIBUTE, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipDeviceGetCount, hipGetDeviceCount, HIP_HIPDEVICEGETCOUNT, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipDeviceGetName, hipDeviceGetName, HIP_HIPDEVICEGETNAME, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipDeviceTotalMem, hipDeviceTotalMem, HIP_HIPDEVICETOTALMEM, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipDriverGetVersion, hipDriverGetVersion, HIP_HIPDRIVERGETVERSION, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipEventCreate, hipEventCreateWithFlags, HIP_HIPEVENTCREATE, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipEventDestroy, hipEventDestroy, HIP_HIPEVENTDESTROY, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipEventElapsedTime, hipEventElapsedTime, HIP_HIPEVENTELAPSEDTIME, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipEventRecord, hipEventRecord, HIP_HIPEVENTRECORD, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipEventSynchronize, hipEventSynchronize, HIP_HIPEVENTSYNCHRONIZE, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipFuncGetAttribute, hipFuncGetAttribute, HIP_HIPFUNCGETATTRIBUTE, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipGetErrorName, hipGetErrorName, HIP_HIPGETERRORNAME, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipGetErrorString, hipGetErrorString, HIP_HIPGETERRORSTRING, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipInit, hipInit, HIP_HIPINIT, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipLaunchKernel, hipModuleLaunchKernel, HIP_HIPLAUNCHKERNEL, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipMemAlloc, hipMalloc, HIP_HIPMEMALLOC, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipMemFree, hipFree, HIP_HIPMEMFREE, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipMemGetInfo, hipMemGetInfo, HIP_HIPMEMGETINFO, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoDAsync, hipMemcpyDtoDAsync, HIP_HIPMEMCPYDTODASYNC, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoHAsync, hipMemcpyDtoHAsync, HIP_HIPMEMCPYDTOHASYNC, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipMemcpyHtoDAsync, hipMemcpyHtoDAsync, HIP_HIPMEMCPYHTODASYNC, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipMemsetD32Async, hipMemsetD32Async, HIP_HIPMEMSETD32ASYNC, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipMemsetD8Async, hipMemsetD8Async, HIP_HIPMEMSETD8ASYNC, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipMemcpyHtoDAsync, hipMemcpyHtoDAsync, HIP_HIPMEMCPYHTODASYNC, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipModuleGetFunction, hipModuleGetFunction, HIP_HIPMODULEGETFUNCTION, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipModuleGetGlobal, hipModuleGetGlobal, HIP_HIPMODULEGETGLOBAL, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipModuleLoadDataEx, hipModuleLoadDataEx, HIP_HIPMODULELOADDATAEX, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipModuleUnload, hipModuleUnload, HIP_HIPMODULEUNLOAD, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipRuntimeGetVersion, hipRuntimeGetVersion, HIP_HIPRUNTIMEGETVERSION, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipStreamCreate, hipStreamCreate, HIP_HIPSTREAMCREATE, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipStreamDestroy, hipStreamDestroy, HIP_HIPSTREAMDESTROY, HIP, 1);
|
|
HC_LOAD_FUNC_HIP (hip, hipStreamSynchronize, hipStreamSynchronize, HIP_HIPSTREAMSYNCHRONIZE, HIP, 1);
|
|
|
|
return 0;
|
|
}
|
|
|
|
void hip_close (void *hashcat_ctx)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
if (hip)
|
|
{
|
|
if (hip->lib)
|
|
{
|
|
hc_dlclose (hip->lib);
|
|
}
|
|
|
|
hcfree (backend_ctx->hip);
|
|
|
|
backend_ctx->hip = NULL;
|
|
}
|
|
}
|
|
|
|
int hc_hipCtxCreate (void *hashcat_ctx, hipCtx_t *pctx, unsigned int flags, hipDevice_t dev)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipCtxCreate (pctx, flags, dev);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipCtxCreate(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipCtxCreate(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipCtxDestroy (void *hashcat_ctx, hipCtx_t ctx)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipCtxDestroy (ctx);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipCtxDestroy(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipCtxDestroy(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipCtxPopCurrent (void *hashcat_ctx, hipCtx_t *pctx)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipCtxPopCurrent (pctx);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipCtxPopCurrent(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipCtxPopCurrent(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipCtxPushCurrent (void *hashcat_ctx, hipCtx_t ctx)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipCtxPushCurrent (ctx);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipCtxPushCurrent(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipCtxPushCurrent(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipCtxSetCurrent (void *hashcat_ctx, hipCtx_t ctx)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipCtxSetCurrent (ctx);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipCtxSetCurrent(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipCtxSetCurrent(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipCtxSynchronize (void *hashcat_ctx)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipCtxSynchronize ();
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipCtxSynchronize(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipCtxSynchronize(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipDeviceGet (void *hashcat_ctx, hipDevice_t* device, int ordinal)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipDeviceGet (device, ordinal);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipDeviceGet(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipDeviceGet(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipDeviceGetAttribute (void *hashcat_ctx, int *pi, hipDeviceAttribute_t attrib, hipDevice_t dev)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipDeviceGetAttribute (pi, attrib, dev);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipDeviceGetAttribute(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipDeviceGetAttribute(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipDeviceGetCount (void *hashcat_ctx, int *count)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipDeviceGetCount (count);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipDeviceGetCount(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipDeviceGetCount(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipDeviceGetName (void *hashcat_ctx, char *name, int len, hipDevice_t dev)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipDeviceGetName (name, len, dev);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipDeviceGetName(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipDeviceGetName(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipDeviceTotalMem (void *hashcat_ctx, size_t *bytes, hipDevice_t dev)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipDeviceTotalMem (bytes, dev);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipDeviceTotalMem(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipDeviceTotalMem(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipDriverGetVersion (void *hashcat_ctx, int *driverVersion)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipDriverGetVersion (driverVersion);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipDriverGetVersion(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipDriverGetVersion(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipEventCreate (void *hashcat_ctx, hipEvent_t *phEvent, unsigned int Flags)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipEventCreate (phEvent, Flags);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipEventCreate(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipEventCreate(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipEventDestroy (void *hashcat_ctx, hipEvent_t hEvent)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipEventDestroy (hEvent);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipEventDestroy(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipEventDestroy(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipEventElapsedTime (void *hashcat_ctx, float *pMilliseconds, hipEvent_t hStart, hipEvent_t hEnd)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipEventElapsedTime (pMilliseconds, hStart, hEnd);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipEventElapsedTime(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipEventElapsedTime(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipEventRecord (void *hashcat_ctx, hipEvent_t hEvent, hipStream_t hStream)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipEventRecord (hEvent, hStream);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipEventRecord(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipEventRecord(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipEventSynchronize (void *hashcat_ctx, hipEvent_t hEvent)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipEventSynchronize (hEvent);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipEventSynchronize(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipEventSynchronize(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipFuncGetAttribute (void *hashcat_ctx, int *pi, hipFunction_attribute attrib, hipFunction_t hfunc)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipFuncGetAttribute (pi, attrib, hfunc);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipFuncGetAttribute(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipFuncGetAttribute(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipLaunchKernel (void *hashcat_ctx, hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipLaunchKernel (f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipLaunchKernel(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipLaunchKernel(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipInit (void *hashcat_ctx, unsigned int Flags)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipInit (Flags);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipInit(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipInit(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipMemAlloc (void *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipMemAlloc (dptr, bytesize);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemAlloc(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemAlloc(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipMemFree (void *hashcat_ctx, hipDeviceptr_t dptr)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipMemFree (dptr);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemFree(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemFree(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipMemGetInfo (void *hashcat_ctx, size_t *free, size_t *total)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipMemGetInfo (free, total);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemGetInfo(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemGetInfo(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipMemcpyDtoHAsync (void *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipMemcpyDtoHAsync (dstHost, srcDevice, ByteCount, hStream);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemcpyDtoHAsync(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemcpyDtoHAsync(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipMemcpyDtoDAsync (void *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipMemcpyDtoDAsync (dstDevice, srcDevice, ByteCount, hStream);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemcpyDtoDAsync(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemcpyDtoDAsync(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipMemcpyHtoDAsync (void *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount, hipStream_t hStream)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipMemcpyHtoDAsync (dstDevice, srcHost, ByteCount, hStream);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemcpyHtoDAsync(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemcpyHtoDAsync(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipMemsetD32Async (void *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned int ui, size_t N, hipStream_t hStream)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipMemsetD32Async (dstDevice, ui, N, hStream);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemsetD32Async(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemsetD32Async(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipMemsetD8Async (void *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned char uc, size_t N, hipStream_t hStream)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipMemsetD8Async (dstDevice, uc, N, hStream);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemsetD8Async(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipMemsetD8Async(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipModuleGetFunction (void *hashcat_ctx, hipFunction_t *hfunc, hipModule_t hmod, const char *name)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipModuleGetFunction (hfunc, hmod, name);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipModuleGetFunction(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipModuleGetFunction(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipModuleGetGlobal (void *hashcat_ctx, hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipModuleGetGlobal (dptr, bytes, hmod, name);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipModuleGetGlobal(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipModuleGetGlobal(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipModuleLoadDataEx (void *hashcat_ctx, hipModule_t *module, const void *image, unsigned int numOptions, hipJitOption *options, void **optionValues)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipModuleLoadDataEx (module, image, numOptions, options, optionValues);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipModuleLoadDataEx(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipModuleLoadDataEx(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipModuleUnload (void *hashcat_ctx, hipModule_t hmod)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipModuleUnload (hmod);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipModuleUnload(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipModuleUnload(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipRuntimeGetVersion (void *hashcat_ctx, int *runtimeVersion)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipRuntimeGetVersion (runtimeVersion);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipRuntimeGetVersion(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipRuntimeGetVersion(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipStreamCreate (void *hashcat_ctx, hipStream_t *phStream, unsigned int Flags)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipStreamCreate (phStream, Flags);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipStreamCreate(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipStreamCreate(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipStreamDestroy (void *hashcat_ctx, hipStream_t hStream)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipStreamDestroy (hStream);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipStreamDestroy(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipStreamDestroy(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
int hc_hipStreamSynchronize (void *hashcat_ctx, hipStream_t hStream)
|
|
{
|
|
backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx;
|
|
|
|
HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip;
|
|
|
|
const hipError_t HIP_err = hip->hipStreamSynchronize (hStream);
|
|
|
|
if (HIP_err != hipSuccess)
|
|
{
|
|
const char *pStr = NULL;
|
|
|
|
if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess)
|
|
{
|
|
event_log_error (hashcat_ctx, "hipStreamSynchronize(): %s", pStr);
|
|
}
|
|
else
|
|
{
|
|
event_log_error (hashcat_ctx, "hipStreamSynchronize(): %d", HIP_err);
|
|
}
|
|
|
|
return -1;
|
|
}
|
|
|
|
return 0;
|
|
}
|