From 78c7ee2af67c885d5464b61461435d09541c7872 Mon Sep 17 00:00:00 2001 From: Gabriele Gristina Date: Sun, 2 Jan 2022 19:12:41 +0100 Subject: [PATCH] HIP Backend: moved functions to ext_hip.c/ext_hiprtc.c and includes to ext_hip.h/ext_hiprtc.h --- docs/changes.txt | 1 + include/backend.h | 51 -- include/ext_hip.h | 41 ++ include/ext_hiprtc.h | 11 + src/backend.c | 1261 ------------------------------------------ src/ext_hip.c | 1080 ++++++++++++++++++++++++++++++++++++ src/ext_hiprtc.c | 185 +++++++ 7 files changed, 1318 insertions(+), 1312 deletions(-) diff --git a/docs/changes.txt b/docs/changes.txt index 877e6bbf4..94f7ef326 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -38,6 +38,7 @@ - OpenCL Runtime: Set default device-type to GPU with Apple Silicon compute devices - Unit tests: Updated test.sh to set default device-type to CPU with Apple Intel and added -f (--force) option - OpenCL Backend: moved functions to ext_OpenCL.c and includes to ext_OpenCL.h +- HIP Backend: moved functions to ext_hip.c/ext_hiprtc.c and includes to ext_hip.h/ext_hiprtc.h * changes v6.2.4 -> v6.2.5 diff --git a/include/backend.h b/include/backend.h index 72e7ca2ee..0f1e23133 100644 --- a/include/backend.h +++ b/include/backend.h @@ -26,15 +26,9 @@ static const char CL_VENDOR_POCL[] = "The pocl project"; int cuda_init (hashcat_ctx_t *hashcat_ctx); void cuda_close (hashcat_ctx_t *hashcat_ctx); -int hip_init (hashcat_ctx_t *hashcat_ctx); -void hip_close (hashcat_ctx_t *hashcat_ctx); - int nvrtc_init (hashcat_ctx_t *hashcat_ctx); void nvrtc_close (hashcat_ctx_t *hashcat_ctx); -int hiprtc_init (hashcat_ctx_t *hashcat_ctx); -void hiprtc_close (hashcat_ctx_t *hashcat_ctx); - int hc_nvrtcCreateProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram *prog, const char *src, const char *name, int numHeaders, const char * const *headers, const char * const *includeNames); int hc_nvrtcDestroyProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram *prog); int hc_nvrtcCompileProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, int numOptions, const char * const *options); @@ -85,51 +79,6 @@ int hc_cuLinkAddData (hashcat_ctx_t *hashcat_ctx, CUlinkState state, int hc_cuLinkDestroy (hashcat_ctx_t *hashcat_ctx, CUlinkState state); int hc_cuLinkComplete (hashcat_ctx_t *hashcat_ctx, CUlinkState state, void **cubinOut, size_t *sizeOut); -int hc_hipCreateProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram *prog, const char *src, const char *name, int numHeaders, const char * const *headers, const char * const *includeNames); -int hc_hipDestroyProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram *prog); -int hc_hipCompileProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, int numOptions, const char * const *options); -int hc_hipGetProgramLogSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *logSizeRet); -int hc_hipGetProgramLog (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *log); -int hc_hipGetCodeSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *codeSizeRet); -int hc_hipGetCode (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *code); - -int hc_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, hipCtx_t *pctx, unsigned int flags, hipDevice_t dev); -int hc_hipCtxDestroy (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx); -int hc_hipCtxPopCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t *pctx); -int hc_hipCtxPushCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx); -int hc_hipCtxSetCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx); -int hc_hipCtxSynchronize (hashcat_ctx_t *hashcat_ctx); -int hc_hipDeviceGet (hashcat_ctx_t *hashcat_ctx, hipDevice_t *device, int ordinal); -int hc_hipDeviceGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, hipDeviceAttribute_t attrib, hipDevice_t dev); -int hc_hipDeviceGetCount (hashcat_ctx_t *hashcat_ctx, int *count); -int hc_hipDeviceGetName (hashcat_ctx_t *hashcat_ctx, char *name, int len, hipDevice_t dev); -int hc_hipDeviceTotalMem (hashcat_ctx_t *hashcat_ctx, size_t *bytes, hipDevice_t dev); -int hc_hipDriverGetVersion (hashcat_ctx_t *hashcat_ctx, int *driverVersion); -int hc_hipEventCreate (hashcat_ctx_t *hashcat_ctx, hipEvent_t *phEvent, unsigned int Flags); -int hc_hipEventDestroy (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent); -int hc_hipEventElapsedTime (hashcat_ctx_t *hashcat_ctx, float *pMilliseconds, hipEvent_t hStart, hipEvent_t hEnd); -int hc_hipEventQuery (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent); -int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent, hipStream_t hStream); -int hc_hipEventSynchronize (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent); -int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, hipFunction_attribute attrib, hipFunction_t hfunc); -int hc_hipInit (hashcat_ctx_t *hashcat_ctx, unsigned int Flags); -int hc_hipLaunchKernel (hashcat_ctx_t *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); -int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize); -int hc_hipMemFree (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dptr); -int hc_hipMemcpyDtoDAsync (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream); -int hc_hipMemcpyDtoHAsync (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream); -int hc_hipMemcpyHtoDAsync (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount, hipStream_t hStream); -int hc_hipMemsetD32Async (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned int ui, size_t N, hipStream_t hStream); -int hc_hipMemsetD8Async (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned char uc, size_t N, hipStream_t hStream); -int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, hipFunction_t *hfunc, hipModule_t hmod, const char *name); -int hc_hipModuleGetGlobal (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name); -int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, hipModule_t *module, const void *image, unsigned int numOptions, hipJitOption *options, void **optionValues); -int hc_hipModuleUnload (hashcat_ctx_t *hashcat_ctx, hipModule_t hmod); -int hc_hipRuntimeGetVersion (hashcat_ctx_t *hashcat_ctx, int *runtimeVersion); -int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, hipStream_t *phStream, unsigned int Flags); -int hc_hipStreamDestroy (hashcat_ctx_t *hashcat_ctx, hipStream_t hStream); -int hc_hipStreamSynchronize (hashcat_ctx_t *hashcat_ctx, hipStream_t hStream); - int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u64 gidd, pw_t *pw); int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 highest_pw_len, const u64 pws_pos, const u64 pws_cnt, const u32 fast_iteration, const u32 salt_pos); diff --git a/include/ext_hip.h b/include/ext_hip.h index 7d9095588..a58d2ef28 100644 --- a/include/ext_hip.h +++ b/include/ext_hip.h @@ -389,4 +389,45 @@ typedef struct hc_hip_lib typedef hc_hip_lib_t HIP_PTR; +int hip_init (void *hashcat_ctx); +void hip_close (void *hashcat_ctx); + +int hc_hipCtxCreate (void *hashcat_ctx, hipCtx_t *pctx, unsigned int flags, hipDevice_t dev); +int hc_hipCtxDestroy (void *hashcat_ctx, hipCtx_t ctx); +int hc_hipCtxPopCurrent (void *hashcat_ctx, hipCtx_t *pctx); +int hc_hipCtxPushCurrent (void *hashcat_ctx, hipCtx_t ctx); +int hc_hipCtxSetCurrent (void *hashcat_ctx, hipCtx_t ctx); +int hc_hipCtxSynchronize (void *hashcat_ctx); +int hc_hipDeviceGet (void *hashcat_ctx, hipDevice_t *device, int ordinal); +int hc_hipDeviceGetAttribute (void *hashcat_ctx, int *pi, hipDeviceAttribute_t attrib, hipDevice_t dev); +int hc_hipDeviceGetCount (void *hashcat_ctx, int *count); +int hc_hipDeviceGetName (void *hashcat_ctx, char *name, int len, hipDevice_t dev); +int hc_hipDeviceTotalMem (void *hashcat_ctx, size_t *bytes, hipDevice_t dev); +int hc_hipDriverGetVersion (void *hashcat_ctx, int *driverVersion); +int hc_hipEventCreate (void *hashcat_ctx, hipEvent_t *phEvent, unsigned int Flags); +int hc_hipEventDestroy (void *hashcat_ctx, hipEvent_t hEvent); +int hc_hipEventElapsedTime (void *hashcat_ctx, float *pMilliseconds, hipEvent_t hStart, hipEvent_t hEnd); +int hc_hipEventQuery (void *hashcat_ctx, hipEvent_t hEvent); +int hc_hipEventRecord (void *hashcat_ctx, hipEvent_t hEvent, hipStream_t hStream); +int hc_hipEventSynchronize (void *hashcat_ctx, hipEvent_t hEvent); +int hc_hipFuncGetAttribute (void *hashcat_ctx, int *pi, hipFunction_attribute attrib, hipFunction_t hfunc); +int hc_hipInit (void *hashcat_ctx, unsigned int Flags); +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); +int hc_hipMemAlloc (void *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize); +int hc_hipMemFree (void *hashcat_ctx, hipDeviceptr_t dptr); +int hc_hipMemGetInfo (void *hashcat_ctx, size_t *free, size_t *total); +int hc_hipMemcpyDtoDAsync (void *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream); +int hc_hipMemcpyDtoHAsync (void *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream); +int hc_hipMemcpyHtoDAsync (void *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount, hipStream_t hStream); +int hc_hipMemsetD32Async (void *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned int ui, size_t N, hipStream_t hStream); +int hc_hipMemsetD8Async (void *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned char uc, size_t N, hipStream_t hStream); +int hc_hipModuleGetFunction (void *hashcat_ctx, hipFunction_t *hfunc, hipModule_t hmod, const char *name); +int hc_hipModuleGetGlobal (void *hashcat_ctx, hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name); +int hc_hipModuleLoadDataEx (void *hashcat_ctx, hipModule_t *module, const void *image, unsigned int numOptions, hipJitOption *options, void **optionValues); +int hc_hipModuleUnload (void *hashcat_ctx, hipModule_t hmod); +int hc_hipRuntimeGetVersion (void *hashcat_ctx, int *runtimeVersion); +int hc_hipStreamCreate (void *hashcat_ctx, hipStream_t *phStream, unsigned int Flags); +int hc_hipStreamDestroy (void *hashcat_ctx, hipStream_t hStream); +int hc_hipStreamSynchronize (void *hashcat_ctx, hipStream_t hStream); + #endif // _EXT_HIP_H diff --git a/include/ext_hiprtc.h b/include/ext_hiprtc.h index 347239c38..8f0b423d0 100644 --- a/include/ext_hiprtc.h +++ b/include/ext_hiprtc.h @@ -67,4 +67,15 @@ typedef hc_hiprtc_lib_t HIPRTC_PTR; int hiprtc_make_options_array_from_string (char *string, char **options); +int hiprtc_init (void *hashcat_ctx); +void hiprtc_close (void *hashcat_ctx); + +int hc_hiprtcCreateProgram (void *hashcat_ctx, hiprtcProgram *prog, const char *src, const char *name, int numHeaders, const char * const *headers, const char * const *includeNames); +int hc_hiprtcDestroyProgram (void *hashcat_ctx, hiprtcProgram *prog); +int hc_hiprtcCompileProgram (void *hashcat_ctx, hiprtcProgram prog, int numOptions, const char * const *options); +int hc_hiprtcGetProgramLogSize (void *hashcat_ctx, hiprtcProgram prog, size_t *logSizeRet); +int hc_hiprtcGetProgramLog (void *hashcat_ctx, hiprtcProgram prog, char *log); +int hc_hiprtcGetCodeSize (void *hashcat_ctx, hiprtcProgram prog, size_t *codeSizeRet); +int hc_hiprtcGetCode (void *hashcat_ctx, hiprtcProgram prog, char *code); + #endif // _EXT_HIPRTC_H diff --git a/src/backend.c b/src/backend.c index c90b51315..cbc81f84a 100644 --- a/src/backend.c +++ b/src/backend.c @@ -990,189 +990,6 @@ int hc_nvrtcVersion (hashcat_ctx_t *hashcat_ctx, int *major, int *minor) return 0; } -// HIPRTC - -int hiprtc_init (hashcat_ctx_t *hashcat_ctx) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; - - memset (hiprtc, 0, sizeof (HIPRTC_PTR)); - - #if defined (_WIN) - hiprtc->lib = hc_dlopen ("amdhip64.dll"); - #elif defined (__APPLE__) - hiprtc->lib = hc_dlopen ("fixme.dylib"); - #elif defined (__CYGWIN__) - hiprtc->lib = hc_dlopen ("amdhip64.dll"); - #else - hiprtc->lib = hc_dlopen ("libamdhip64.so"); - - if (hiprtc->lib == NULL) hiprtc->lib = hc_dlopen ("libamdhip64.so.4"); - #endif - - if (hiprtc->lib == NULL) return -1; - - HC_LOAD_FUNC (hiprtc, hiprtcAddNameExpression, HIPRTC_HIPRTCADDNAMEEXPRESSION, HIPRTC, 1); - HC_LOAD_FUNC (hiprtc, hiprtcCompileProgram, HIPRTC_HIPRTCCOMPILEPROGRAM, HIPRTC, 1); - HC_LOAD_FUNC (hiprtc, hiprtcCreateProgram, HIPRTC_HIPRTCCREATEPROGRAM, HIPRTC, 1); - HC_LOAD_FUNC (hiprtc, hiprtcDestroyProgram, HIPRTC_HIPRTCDESTROYPROGRAM, HIPRTC, 1); - HC_LOAD_FUNC (hiprtc, hiprtcGetLoweredName, HIPRTC_HIPRTCGETLOWEREDNAME, HIPRTC, 1); - HC_LOAD_FUNC (hiprtc, hiprtcGetCode, HIPRTC_HIPRTCGETCODE, HIPRTC, 1); - HC_LOAD_FUNC (hiprtc, hiprtcGetCodeSize, HIPRTC_HIPRTCGETCODESIZE, HIPRTC, 1); - HC_LOAD_FUNC (hiprtc, hiprtcGetProgramLog, HIPRTC_HIPRTCGETPROGRAMLOG, HIPRTC, 1); - HC_LOAD_FUNC (hiprtc, hiprtcGetProgramLogSize, HIPRTC_HIPRTCGETPROGRAMLOGSIZE, HIPRTC, 1); - HC_LOAD_FUNC (hiprtc, hiprtcGetErrorString, HIPRTC_HIPRTCGETERRORSTRING, HIPRTC, 1); - - return 0; -} - -void hiprtc_close (hashcat_ctx_t *hashcat_ctx) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; - - if (hiprtc) - { - if (hiprtc->lib) - { - hc_dlclose (hiprtc->lib); - } - - hcfree (backend_ctx->hiprtc); - - backend_ctx->hiprtc = NULL; - } -} - -int hc_hiprtcCreateProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram *prog, const char *src, const char *name, int numHeaders, const char * const *headers, const char * const *includeNames) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; - - const hiprtcResult HIPRTC_err = hiprtc->hiprtcCreateProgram (prog, src, name, numHeaders, headers, includeNames); - - if (HIPRTC_err != HIPRTC_SUCCESS) - { - event_log_error (hashcat_ctx, "hiprtcCreateProgram(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); - - return -1; - } - - return 0; -} - -int hc_hiprtcDestroyProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram *prog) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; - - const hiprtcResult HIPRTC_err = hiprtc->hiprtcDestroyProgram (prog); - - if (HIPRTC_err != HIPRTC_SUCCESS) - { - event_log_error (hashcat_ctx, "hiprtcDestroyProgram(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); - - return -1; - } - - return 0; -} - -int hc_hiprtcCompileProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, int numOptions, const char * const *options) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; - - const hiprtcResult HIPRTC_err = hiprtc->hiprtcCompileProgram (prog, numOptions, options); - - if (HIPRTC_err != HIPRTC_SUCCESS) - { - event_log_error (hashcat_ctx, "hiprtcCompileProgram(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); - - return -1; - } - - return 0; -} - -int hc_hiprtcGetProgramLogSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *logSizeRet) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; - - const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetProgramLogSize (prog, logSizeRet); - - if (HIPRTC_err != HIPRTC_SUCCESS) - { - event_log_error (hashcat_ctx, "hiprtcGetProgramLogSize(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); - - return -1; - } - - return 0; -} - -int hc_hiprtcGetProgramLog (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *log) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; - - const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetProgramLog (prog, log); - - if (HIPRTC_err != HIPRTC_SUCCESS) - { - event_log_error (hashcat_ctx, "hiprtcGetProgramLog(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); - - return -1; - } - - return 0; -} - -int hc_hiprtcGetCodeSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *codeSizeRet) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; - - const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetCodeSize (prog, codeSizeRet); - - if (HIPRTC_err != HIPRTC_SUCCESS) - { - event_log_error (hashcat_ctx, "hiprtcGetCodeSize(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); - - return -1; - } - - return 0; -} - -int hc_hiprtcGetCode (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *code) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; - - const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetCode (prog, code); - - if (HIPRTC_err != HIPRTC_SUCCESS) - { - event_log_error (hashcat_ctx, "hiprtcGetCode(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); - - return -1; - } - - return 0; -} - // CUDA int cuda_init (hashcat_ctx_t *hashcat_ctx) @@ -2434,1084 +2251,6 @@ int hc_cuLinkComplete (hashcat_ctx_t *hashcat_ctx, CUlinkState state, void **cub return 0; } -// HIP - -int hip_init (hashcat_ctx_t *hashcat_ctx) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipCtx_t *pctx, unsigned int flags, hipDevice_t dev) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipCtx_t *pctx) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipDevice_t* device, int ordinal) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, int *pi, hipDeviceAttribute_t attrib, hipDevice_t dev) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, int *count) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, char *name, int len, hipDevice_t dev) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, size_t *bytes, hipDevice_t dev) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, int *driverVersion) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipEvent_t *phEvent, unsigned int Flags) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, float *pMilliseconds, hipEvent_t hStart, hipEvent_t hEnd) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent, hipStream_t hStream) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, int *pi, hipFunction_attribute attrib, hipFunction_t hfunc) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *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->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 (hashcat_ctx_t *hashcat_ctx, unsigned int Flags) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dptr) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, size_t *free, size_t *total) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount, hipStream_t hStream) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned int ui, size_t N, hipStream_t hStream) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned char uc, size_t N, hipStream_t hStream) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipFunction_t *hfunc, hipModule_t hmod, const char *name) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipModule_t *module, const void *image, unsigned int numOptions, hipJitOption *options, void **optionValues) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipModule_t hmod) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, int *runtimeVersion) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipStream_t *phStream, unsigned int Flags) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipStream_t hStream) -{ - backend_ctx_t *backend_ctx = 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 (hashcat_ctx_t *hashcat_ctx, hipStream_t hStream) -{ - backend_ctx_t *backend_ctx = 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; -} - // Backend int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u64 gidd, pw_t *pw) diff --git a/src/ext_hip.c b/src/ext_hip.c index 72fb2fbfe..95bec69e1 100644 --- a/src/ext_hip.c +++ b/src/ext_hip.c @@ -5,4 +5,1084 @@ #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; +} diff --git a/src/ext_hiprtc.c b/src/ext_hiprtc.c index 1ec099ae7..2d2f86566 100644 --- a/src/ext_hiprtc.c +++ b/src/ext_hiprtc.c @@ -5,8 +5,12 @@ #include "common.h" #include "types.h" +#include "memory.h" +#include "event.h" #include "ext_hiprtc.h" +#include "dynloader.h" + int hiprtc_make_options_array_from_string (char *string, char **options) { char *saveptr = NULL; @@ -25,3 +29,184 @@ int hiprtc_make_options_array_from_string (char *string, char **options) return cnt; } + +int hiprtc_init (void *hashcat_ctx) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; + + memset (hiprtc, 0, sizeof (HIPRTC_PTR)); + + #if defined (_WIN) + hiprtc->lib = hc_dlopen ("amdhip64.dll"); + #elif defined (__APPLE__) + hiprtc->lib = hc_dlopen ("fixme.dylib"); + #elif defined (__CYGWIN__) + hiprtc->lib = hc_dlopen ("amdhip64.dll"); + #else + hiprtc->lib = hc_dlopen ("libamdhip64.so"); + + if (hiprtc->lib == NULL) hiprtc->lib = hc_dlopen ("libamdhip64.so.4"); + #endif + + if (hiprtc->lib == NULL) return -1; + + HC_LOAD_FUNC (hiprtc, hiprtcAddNameExpression, HIPRTC_HIPRTCADDNAMEEXPRESSION, HIPRTC, 1); + HC_LOAD_FUNC (hiprtc, hiprtcCompileProgram, HIPRTC_HIPRTCCOMPILEPROGRAM, HIPRTC, 1); + HC_LOAD_FUNC (hiprtc, hiprtcCreateProgram, HIPRTC_HIPRTCCREATEPROGRAM, HIPRTC, 1); + HC_LOAD_FUNC (hiprtc, hiprtcDestroyProgram, HIPRTC_HIPRTCDESTROYPROGRAM, HIPRTC, 1); + HC_LOAD_FUNC (hiprtc, hiprtcGetLoweredName, HIPRTC_HIPRTCGETLOWEREDNAME, HIPRTC, 1); + HC_LOAD_FUNC (hiprtc, hiprtcGetCode, HIPRTC_HIPRTCGETCODE, HIPRTC, 1); + HC_LOAD_FUNC (hiprtc, hiprtcGetCodeSize, HIPRTC_HIPRTCGETCODESIZE, HIPRTC, 1); + HC_LOAD_FUNC (hiprtc, hiprtcGetProgramLog, HIPRTC_HIPRTCGETPROGRAMLOG, HIPRTC, 1); + HC_LOAD_FUNC (hiprtc, hiprtcGetProgramLogSize, HIPRTC_HIPRTCGETPROGRAMLOGSIZE, HIPRTC, 1); + HC_LOAD_FUNC (hiprtc, hiprtcGetErrorString, HIPRTC_HIPRTCGETERRORSTRING, HIPRTC, 1); + + return 0; +} + +void hiprtc_close (void *hashcat_ctx) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; + + if (hiprtc) + { + if (hiprtc->lib) + { + hc_dlclose (hiprtc->lib); + } + + hcfree (backend_ctx->hiprtc); + + backend_ctx->hiprtc = NULL; + } +} + +int hc_hiprtcCreateProgram (void *hashcat_ctx, hiprtcProgram *prog, const char *src, const char *name, int numHeaders, const char * const *headers, const char * const *includeNames) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; + + const hiprtcResult HIPRTC_err = hiprtc->hiprtcCreateProgram (prog, src, name, numHeaders, headers, includeNames); + + if (HIPRTC_err != HIPRTC_SUCCESS) + { + event_log_error (hashcat_ctx, "hiprtcCreateProgram(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); + + return -1; + } + + return 0; +} + +int hc_hiprtcDestroyProgram (void *hashcat_ctx, hiprtcProgram *prog) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; + + const hiprtcResult HIPRTC_err = hiprtc->hiprtcDestroyProgram (prog); + + if (HIPRTC_err != HIPRTC_SUCCESS) + { + event_log_error (hashcat_ctx, "hiprtcDestroyProgram(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); + + return -1; + } + + return 0; +} + +int hc_hiprtcCompileProgram (void *hashcat_ctx, hiprtcProgram prog, int numOptions, const char * const *options) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; + + const hiprtcResult HIPRTC_err = hiprtc->hiprtcCompileProgram (prog, numOptions, options); + + if (HIPRTC_err != HIPRTC_SUCCESS) + { + event_log_error (hashcat_ctx, "hiprtcCompileProgram(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); + + return -1; + } + + return 0; +} + +int hc_hiprtcGetProgramLogSize (void *hashcat_ctx, hiprtcProgram prog, size_t *logSizeRet) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; + + const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetProgramLogSize (prog, logSizeRet); + + if (HIPRTC_err != HIPRTC_SUCCESS) + { + event_log_error (hashcat_ctx, "hiprtcGetProgramLogSize(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); + + return -1; + } + + return 0; +} + +int hc_hiprtcGetProgramLog (void *hashcat_ctx, hiprtcProgram prog, char *log) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; + + const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetProgramLog (prog, log); + + if (HIPRTC_err != HIPRTC_SUCCESS) + { + event_log_error (hashcat_ctx, "hiprtcGetProgramLog(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); + + return -1; + } + + return 0; +} + +int hc_hiprtcGetCodeSize (void *hashcat_ctx, hiprtcProgram prog, size_t *codeSizeRet) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; + + const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetCodeSize (prog, codeSizeRet); + + if (HIPRTC_err != HIPRTC_SUCCESS) + { + event_log_error (hashcat_ctx, "hiprtcGetCodeSize(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); + + return -1; + } + + return 0; +} + +int hc_hiprtcGetCode (void *hashcat_ctx, hiprtcProgram prog, char *code) +{ + backend_ctx_t *backend_ctx = ((hashcat_ctx_t *) hashcat_ctx)->backend_ctx; + + HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; + + const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetCode (prog, code); + + if (HIPRTC_err != HIPRTC_SUCCESS) + { + event_log_error (hashcat_ctx, "hiprtcGetCode(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); + + return -1; + } + + return 0; +}