From 1b84a9e53bf3be185c2ba49a98e12c23f0e162f2 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Sun, 11 Jul 2021 12:38:59 +0200 Subject: [PATCH] Add missing backports from code base v6.2.2 Fix context to thread management Fix missing code in selftest.c, autotune.c, hashes.c, dispatch.c and backend.c Use IS_HIP depending code makes it easier for future optimization related to inline assembly calls - instead of using IS_CUDA || IS_HIP See TODO markers for more optimizations / next steps --- OpenCL/inc_common.cl | 4 - OpenCL/inc_common.h | 1 - OpenCL/inc_platform.cl | 109 +- OpenCL/inc_platform.h | 25 +- OpenCL/inc_types.h | 6 +- OpenCL/inc_vendor.h | 38 +- OpenCL/shared.cl | 5 + include/backend.h | 133 +- include/types.h | 308 ++-- src/Makefile | 2 +- src/autotune.c | 48 +- src/backend.c | 3892 +++++++++++++++++++++++++++++++++++++++- src/dispatch.c | 20 + src/hashes.c | 34 + src/selftest.c | 137 +- 15 files changed, 4485 insertions(+), 277 deletions(-) diff --git a/OpenCL/inc_common.cl b/OpenCL/inc_common.cl index 3aed1ceff..26df19a2b 100644 --- a/OpenCL/inc_common.cl +++ b/OpenCL/inc_common.cl @@ -3,10 +3,6 @@ * License.....: MIT */ -#ifdef IS_HIP -#include -#endif - #include "inc_vendor.h" #include "inc_types.h" #include "inc_platform.h" diff --git a/OpenCL/inc_common.h b/OpenCL/inc_common.h index c854bb1ca..ebd0107c1 100644 --- a/OpenCL/inc_common.h +++ b/OpenCL/inc_common.h @@ -105,7 +105,6 @@ MAYBE_UNUSED const u64 pws_pos, \ MAYBE_UNUSED const u64 gid_max #endif - /* * Shortcut macros for usage in the actual kernels * diff --git a/OpenCL/inc_platform.cl b/OpenCL/inc_platform.cl index 5c30cb6ed..40002c3eb 100644 --- a/OpenCL/inc_platform.cl +++ b/OpenCL/inc_platform.cl @@ -2,9 +2,6 @@ * Author......: See docs/credits.txt * License.....: MIT */ -#ifdef IS_HIP -#include -#endif #include "inc_vendor.h" #include "inc_types.h" @@ -63,7 +60,111 @@ DECLSPEC u64 rotr64_S (const u64 a, const int n) #endif -#if defined IS_CUDA || defined IS_HIP +#if defined IS_CUDA + +#if ATTACK_EXEC == 11 + +CONSTANT_VK u32 generic_constant[8192]; // 32k + +#if ATTACK_KERN == 0 +#define bfs_buf g_bfs_buf +#define rules_buf ((const kernel_rule_t *) generic_constant) +#define words_buf_s g_words_buf_s +#define words_buf_r g_words_buf_r +#elif ATTACK_KERN == 1 +#define bfs_buf g_bfs_buf +#define rules_buf g_rules_buf +#define words_buf_s g_words_buf_s +#define words_buf_r g_words_buf_r +#elif ATTACK_KERN == 3 +#define rules_buf g_rules_buf +#define bfs_buf ((const bf_t *) generic_constant) +#define words_buf_s ((const bs_word_t *) generic_constant) +#define words_buf_r ((const u32x *) generic_constant) +#endif + +#endif + +DECLSPEC u32 hc_atomic_dec (GLOBAL_AS u32 *p) +{ + volatile const u32 val = 1; + + return atomicSub (p, val); +} + +DECLSPEC u32 hc_atomic_inc (GLOBAL_AS u32 *p) +{ + volatile const u32 val = 1; + + return atomicAdd (p, val); +} + +DECLSPEC u32 hc_atomic_or (GLOBAL_AS u32 *p, volatile const u32 val) +{ + return atomicOr (p, val); +} + +DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused))) +{ + return (blockIdx.x * blockDim.x) + threadIdx.x; +} + +DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused))) +{ + return threadIdx.x; +} + +DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused))) +{ + // verify + return blockDim.x; +} + +DECLSPEC u32x rotl32 (const u32x a, const int n) +{ + return ((a << n) | ((a >> (32 - n)))); +} + +DECLSPEC u32x rotr32 (const u32x a, const int n) +{ + return ((a >> n) | ((a << (32 - n)))); +} + +DECLSPEC u32 rotl32_S (const u32 a, const int n) +{ + return ((a << n) | ((a >> (32 - n)))); +} + +DECLSPEC u32 rotr32_S (const u32 a, const int n) +{ + return ((a >> n) | ((a << (32 - n)))); +} + +DECLSPEC u64x rotl64 (const u64x a, const int n) +{ + return ((a << n) | ((a >> (64 - n)))); +} + +DECLSPEC u64x rotr64 (const u64x a, const int n) +{ + return ((a >> n) | ((a << (64 - n)))); +} + +DECLSPEC u64 rotl64_S (const u64 a, const int n) +{ + return ((a << n) | ((a >> (64 - n)))); +} + +DECLSPEC u64 rotr64_S (const u64 a, const int n) +{ + return ((a >> n) | ((a << (64 - n)))); +} + +#define FIXED_THREAD_COUNT(n) __launch_bounds__((n), 0) +#define SYNC_THREADS() __syncthreads () +#endif + +#if defined IS_HIP #if ATTACK_EXEC == 11 diff --git a/OpenCL/inc_platform.h b/OpenCL/inc_platform.h index 50aaeb7d0..c65891a74 100644 --- a/OpenCL/inc_platform.h +++ b/OpenCL/inc_platform.h @@ -21,7 +21,7 @@ DECLSPEC u64 rotl64_S (const u64 a, const int n); DECLSPEC u64 rotr64_S (const u64 a, const int n); #endif -#if defined IS_CUDA || defined IS_HIP +#ifdef IS_CUDA DECLSPEC u32 hc_atomic_dec (volatile GLOBAL_AS u32 *p); DECLSPEC u32 hc_atomic_inc (volatile GLOBAL_AS u32 *p); DECLSPEC u32 hc_atomic_or (volatile GLOBAL_AS u32 *p, volatile const u32 val); @@ -39,10 +39,29 @@ DECLSPEC u64x rotr64 (const u64x a, const int n); DECLSPEC u64 rotl64_S (const u64 a, const int n); DECLSPEC u64 rotr64_S (const u64 a, const int n); -#ifdef IS_HIP -#define rotate(a,n) (((a) << (n)) | ((a) >> (32 - (n)))) +//#define rotate(a,n) (((a) << (n)) | ((a) >> (32 - (n)))) +#define bitselect(a,b,c) ((a) ^ ((c) & ((b) ^ (a)))) #endif +#ifdef IS_HIP +DECLSPEC u32 hc_atomic_dec (volatile GLOBAL_AS u32 *p); +DECLSPEC u32 hc_atomic_inc (volatile GLOBAL_AS u32 *p); +DECLSPEC u32 hc_atomic_or (volatile GLOBAL_AS u32 *p, volatile const u32 val); + +DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused))); +DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused))); +DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused))); + +DECLSPEC u32x rotl32 (const u32x a, const int n); +DECLSPEC u32x rotr32 (const u32x a, const int n); +DECLSPEC u32 rotl32_S (const u32 a, const int n); +DECLSPEC u32 rotr32_S (const u32 a, const int n); +DECLSPEC u64x rotl64 (const u64x a, const int n); +DECLSPEC u64x rotr64 (const u64x a, const int n); +DECLSPEC u64 rotl64_S (const u64 a, const int n); +DECLSPEC u64 rotr64_S (const u64 a, const int n); + +//#define rotate(a,n) (((a) << (n)) | ((a) >> (32 - (n)))) #define bitselect(a,b,c) ((a) ^ ((c) & ((b) ^ (a)))) #endif diff --git a/OpenCL/inc_types.h b/OpenCL/inc_types.h index 86353c087..1041a0d7f 100644 --- a/OpenCL/inc_types.h +++ b/OpenCL/inc_types.h @@ -16,12 +16,12 @@ #define DIGESTS_OFFSET digests_offset_host #endif -#if defined IS_CUDA || defined IS_HIP +#ifdef IS_CUDA //https://docs.nvidia.com/cuda/nvrtc/index.html#integer-size typedef unsigned char uchar; typedef unsigned short ushort; typedef unsigned int uint; -typedef unsigned long long xulong; +typedef unsigned long long ulong; #endif #ifdef KERNEL_STATIC @@ -68,7 +68,7 @@ typedef u64 u64x; #define make_u64x (u64) #else -#if defined IS_CUDA || defined IS_HIP +#if defined IS_CUDA #if VECT_SIZE == 2 diff --git a/OpenCL/inc_vendor.h b/OpenCL/inc_vendor.h index 0ad5de23b..dc6a41d4a 100644 --- a/OpenCL/inc_vendor.h +++ b/OpenCL/inc_vendor.h @@ -16,6 +16,10 @@ #define IS_OPENCL #endif +#ifdef IS_HIP +#include +#endif + #if defined IS_NATIVE #define CONSTANT_VK #define CONSTANT_AS @@ -23,7 +27,14 @@ #define LOCAL_VK #define LOCAL_AS #define KERNEL_FQ -#elif (defined IS_CUDA) || (defined IS_HIP) +#elif defined IS_CUDA +#define CONSTANT_VK __constant__ +#define CONSTANT_AS +#define GLOBAL_AS +#define LOCAL_VK __shared__ +#define LOCAL_AS +#define KERNEL_FQ extern "C" __global__ +#elif defined IS_HIP #define CONSTANT_VK __constant__ #define CONSTANT_AS #define GLOBAL_AS @@ -78,12 +89,14 @@ #define IS_MESA #define IS_GENERIC #elif VENDOR_ID == (1 << 5) -//#define IS_NV //TODO: FIX ME HIP -#define IS_POCL -#define IS_GENERIC +#define IS_NV #elif VENDOR_ID == (1 << 6) #define IS_POCL #define IS_GENERIC +#elif VENDOR_ID == (1 << 8) +#define IS_AMD_USE_HIP +// TODO HIP optimization potential +#define IS_GENERIC #else #define IS_GENERIC #endif @@ -116,14 +129,12 @@ */ #if defined IS_AMD && defined IS_GPU -#define DECLSPEC inline static __device__ -#else -#ifdef IS_HIP -#define DECLSPEC __device__ +#define DECLSPEC inline static +#elif defined IS_HIP +#define DECLSPEC __device__ #else #define DECLSPEC #endif -#endif /** * AMD specific @@ -141,11 +152,18 @@ // Whitelist some OpenCL specific functions // This could create more stable kernels on systems with bad OpenCL drivers -#if defined IS_CUDA || defined IS_HIP +#ifdef IS_CUDA #define USE_BITSELECT #define USE_ROTATE #endif +#ifdef IS_HIP +//TODO HIP +//#define USE_BITSELECT +//#define USE_ROTATE +//#define USE_SWIZZLE +#endif + #ifdef IS_ROCM #define USE_BITSELECT #define USE_ROTATE diff --git a/OpenCL/shared.cl b/OpenCL/shared.cl index a6ae38988..3cc96e79f 100644 --- a/OpenCL/shared.cl +++ b/OpenCL/shared.cl @@ -126,6 +126,11 @@ KERNEL_FQ void gpu_memset (GLOBAL_AS uint4 *buf, const u32 value, const u64 gid_ r.y = value; r.z = value; r.w = value; + #elif defined IS_HIP + r.x = value; + r.y = value; + r.z = value; + r.w = value; #endif buf[gid] = r; diff --git a/include/backend.h b/include/backend.h index c73c512f6..a024aa37d 100644 --- a/include/backend.h +++ b/include/backend.h @@ -22,21 +22,21 @@ static const char CL_VENDOR_MESA[] = "Mesa"; static const char CL_VENDOR_NV[] = "NVIDIA Corporation"; 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 cuda_init (hashcat_ctx_t *hashcat_ctx); +void cuda_close (hashcat_ctx_t *hashcat_ctx); -int nvrtc_init (hashcat_ctx_t *hashcat_ctx); -void nvrtc_close (hashcat_ctx_t *hashcat_ctx); +int hip_init (hashcat_ctx_t *hashcat_ctx); +void hip_close (hashcat_ctx_t *hashcat_ctx); -int hip_init (hashcat_ctx_t *hashcat_ctx); -void hip_close (hashcat_ctx_t *hashcat_ctx); +int ocl_init (hashcat_ctx_t *hashcat_ctx); +void ocl_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 ocl_init (hashcat_ctx_t *hashcat_ctx); -void ocl_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,55 +85,53 @@ 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_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); +int hc_nvrtcGetProgramLogSize (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *logSizeRet); +int hc_nvrtcGetProgramLog (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *log); +int hc_nvrtcGetPTXSize (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *ptxSizeRet); +int hc_nvrtcGetPTX (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *ptx); +int hc_nvrtcVersion (hashcat_ctx_t *hashcat_ctx, int *major, int *minor); -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); -int hc_hiprtcDestroyProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram *prog); -int hc_hiprtcCompileProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, int numOptions, const char * const *options); -int hc_hiprtcGetProgramLogSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *logSizeRet); -int hc_hiprtcGetProgramLog (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *log); -int hc_hiprtcGetCodeSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *ptxSizeRet); -int hc_hiprtcGetCode (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *ptx); -int hc_hiprtcVersion (hashcat_ctx_t *hashcat_ctx, int *major, int *minor); - -int hc_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx, unsigned int flags, HIPdevice dev); -int hc_hipCtxDestroy (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx); -int hc_hipCtxSetCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx); -int hc_hipCtxSetCacheConfig (hashcat_ctx_t *hashcat_ctx, HIPfunc_cache config); -int hc_hipCtxSynchronize (hashcat_ctx_t *hashcat_ctx); -int hc_hipDeviceGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPdevice_attribute attrib, HIPdevice dev); -int hc_hipDeviceGetCount (hashcat_ctx_t *hashcat_ctx, int *count); -int hc_hipDeviceGet (hashcat_ctx_t *hashcat_ctx, HIPdevice *device, int ordinal); -int hc_hipDeviceGetName (hashcat_ctx_t *hashcat_ctx, char *name, int len, HIPdevice dev); -int hc_hipDeviceTotalMem (hashcat_ctx_t *hashcat_ctx, size_t *bytes, HIPdevice dev); -int hc_hipDriverGetVersion (hashcat_ctx_t *hashcat_ctx, int *driverVersion); -int hc_hipEventCreate (hashcat_ctx_t *hashcat_ctx, HIPevent *phEvent, unsigned int Flags); -int hc_hipEventDestroy (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); -int hc_hipEventElapsedTime (hashcat_ctx_t *hashcat_ctx, float *pMilliseconds, HIPevent hStart, HIPevent hEnd); -int hc_hipEventQuery (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); -int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent, HIPstream hStream); -int hc_hipEventSynchronize (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); -int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPfunction_attribute attrib, HIPfunction hfunc); -int hc_hipFuncSetAttribute (hashcat_ctx_t *hashcat_ctx, HIPfunction hfunc, HIPfunction_attribute attrib, int value); -int hc_hipInit (hashcat_ctx_t *hashcat_ctx, unsigned int Flags); -int hc_hipLaunchKernel (hashcat_ctx_t *hashcat_ctx, HIPfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, HIPstream hStream, void **kernelParams, void **extra); -int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr *dptr, size_t bytesize); -int hc_hipMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, HIPdeviceptr srcDevice, size_t ByteCount); -int hc_hipMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, HIPdeviceptr srcDevice, size_t ByteCount); -int hc_hipMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, const void *srcHost, size_t ByteCount); -int hc_hipMemFree (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dptr); -int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, HIPfunction *hfunc, HIPmodule hmod, const char *name); -int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, HIPmodule *module, const void *image, unsigned int numOptions, HIPjit_option *options, void **optionValues); -int hc_hipModuleUnload (hashcat_ctx_t *hashcat_ctx, HIPmodule hmod); -int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, HIPstream *phStream, unsigned int Flags); -int hc_hipStreamDestroy (hashcat_ctx_t *hashcat_ctx, HIPstream hStream); -int hc_hipStreamSynchronize (hashcat_ctx_t *hashcat_ctx, HIPstream hStream); -int hc_hipCtxPushCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx); -int hc_hipCtxPopCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx); -int hc_hipLinkCreate (hashcat_ctx_t *hashcat_ctx, unsigned int numOptions, HIPjit_option *options, void **optionValues, HIPlinkState *stateOut); -int hc_hipLinkAddData (hashcat_ctx_t *hashcat_ctx, HIPlinkState state, HIPjitInputType type, void *data, size_t size, const char *name, unsigned int numOptions, HIPjit_option *options, void **optionValues); -int hc_hipLinkDestroy (hashcat_ctx_t *hashcat_ctx, HIPlinkState state); -int hc_hipLinkComplete (hashcat_ctx_t *hashcat_ctx, HIPlinkState state, void **hipbinOut, size_t *sizeOut); - +int hc_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx, unsigned int flags, HIPdevice dev); +int hc_hipCtxDestroy (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx); +int hc_hipCtxSetCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx); +int hc_hipCtxSetCacheConfig (hashcat_ctx_t *hashcat_ctx, HIPfunc_cache config); +int hc_hipCtxSynchronize (hashcat_ctx_t *hashcat_ctx); +int hc_hipDeviceGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPdevice_attribute attrib, HIPdevice dev); +int hc_hipDeviceGetCount (hashcat_ctx_t *hashcat_ctx, int *count); +int hc_hipDeviceGet (hashcat_ctx_t *hashcat_ctx, HIPdevice *device, int ordinal); +int hc_hipDeviceGetName (hashcat_ctx_t *hashcat_ctx, char *name, int len, HIPdevice dev); +int hc_hipDeviceTotalMem (hashcat_ctx_t *hashcat_ctx, size_t *bytes, HIPdevice dev); +int hc_hipDriverGetVersion (hashcat_ctx_t *hashcat_ctx, int *driverVersion); +int hc_hipEventCreate (hashcat_ctx_t *hashcat_ctx, HIPevent *phEvent, unsigned int Flags); +int hc_hipEventDestroy (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); +int hc_hipEventElapsedTime (hashcat_ctx_t *hashcat_ctx, float *pMilliseconds, HIPevent hStart, HIPevent hEnd); +int hc_hipEventQuery (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); +int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent, HIPstream hStream); +int hc_hipEventSynchronize (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); +int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPfunction_attribute attrib, HIPfunction hfunc); +int hc_hipFuncSetAttribute (hashcat_ctx_t *hashcat_ctx, HIPfunction hfunc, HIPfunction_attribute attrib, int value); +int hc_hipInit (hashcat_ctx_t *hashcat_ctx, unsigned int Flags); +int hc_hipLaunchKernel (hashcat_ctx_t *hashcat_ctx, HIPfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, HIPstream hStream, void **kernelParams, void **extra); +int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr *dptr, size_t bytesize); +int hc_hipMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, HIPdeviceptr srcDevice, size_t ByteCount); +int hc_hipMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, HIPdeviceptr srcDevice, size_t ByteCount); +int hc_hipMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, const void *srcHost, size_t ByteCount); +int hc_hipMemFree (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dptr); +int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, HIPfunction *hfunc, HIPmodule hmod, const char *name); +int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, HIPmodule *module, const void *image, unsigned int numOptions, HIPjit_option *options, void **optionValues); +int hc_hipModuleUnload (hashcat_ctx_t *hashcat_ctx, HIPmodule hmod); +int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, HIPstream *phStream, unsigned int Flags); +int hc_hipStreamDestroy (hashcat_ctx_t *hashcat_ctx, HIPstream hStream); +int hc_hipStreamSynchronize (hashcat_ctx_t *hashcat_ctx, HIPstream hStream); +int hc_hipCtxPushCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx); +int hc_hipCtxPopCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx); +int hc_hipLinkCreate (hashcat_ctx_t *hashcat_ctx, unsigned int numOptions, HIPjit_option *options, void **optionValues, HIPlinkState *stateOut); +int hc_hipLinkAddData (hashcat_ctx_t *hashcat_ctx, HIPlinkState state, HIPjitInputType type, void *data, size_t size, const char *name, unsigned int numOptions, HIPjit_option *options, void **optionValues); +int hc_hipLinkDestroy (hashcat_ctx_t *hashcat_ctx, HIPlinkState state); +int hc_hipLinkComplete (hashcat_ctx_t *hashcat_ctx, HIPlinkState state, void **cubinOut, size_t *sizeOut); int hc_clBuildProgram (hashcat_ctx_t *hashcat_ctx, cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data); int hc_clCompileProgram (hashcat_ctx_t *hashcat_ctx, cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, cl_uint num_input_headers, const cl_program *input_headers, const char **header_include_names, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data); @@ -177,19 +175,20 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, void rebuild_pws_compressed_append (hc_device_param_t *device_param, const u64 pws_cnt, const u8 chr); -int run_cuda_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 num); -int run_cuda_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 num); -int run_cuda_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u32 value, const u64 size); -int run_cuda_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 size); +int run_cuda_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 num); +int run_cuda_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 num); +int run_cuda_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u32 value, const u64 size); +int run_cuda_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 size); -int run_hip_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 num); -int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u32 value, const u64 size); -int run_hip_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 size); +int run_hip_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 num); +int run_hip_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 num); +int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u32 value, const u64 size); +int run_hip_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 size); -int run_opencl_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 num); +int run_opencl_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 num); int run_opencl_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 num); -int run_opencl_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u32 value, const u64 size); -int run_opencl_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 size); +int run_opencl_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u32 value, const u64 size); +int run_opencl_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 size); int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kern_run, const u64 pws_pos, const u64 num, const u32 event_update, const u32 iteration); int run_kernel_mp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kern_run, const u64 num); diff --git a/include/types.h b/include/types.h index 5d9d611c7..037e23d34 100644 --- a/include/types.h +++ b/include/types.h @@ -184,6 +184,7 @@ typedef enum vendor_id VENDOR_ID_NV = (1U << 5), VENDOR_ID_POCL = (1U << 6), VENDOR_ID_AMD_USE_INTEL = (1U << 7), + VENDOR_ID_AMD_USE_HIP = (1U << 8), VENDOR_ID_GENERIC = (1U << 31) } vendor_id_t; @@ -696,114 +697,114 @@ typedef enum user_options_map IDX_ATTACK_MODE = 'a', IDX_BACKEND_DEVICES = 'd', IDX_BACKEND_IGNORE_CUDA = 0xff01, - IDX_BACKEND_IGNORE_HIP = 0xff4d, - IDX_BACKEND_IGNORE_OPENCL = 0xff02, + IDX_BACKEND_IGNORE_HIP = 0xff02, + IDX_BACKEND_IGNORE_OPENCL = 0xff03, IDX_BACKEND_INFO = 'I', - IDX_BACKEND_VECTOR_WIDTH = 0xff03, - IDX_BENCHMARK_ALL = 0xff04, + IDX_BACKEND_VECTOR_WIDTH = 0xff04, + IDX_BENCHMARK_ALL = 0xff05, IDX_BENCHMARK = 'b', - IDX_BITMAP_MAX = 0xff05, - IDX_BITMAP_MIN = 0xff06, + IDX_BITMAP_MAX = 0xff06, + IDX_BITMAP_MIN = 0xff07, #ifdef WITH_BRAIN IDX_BRAIN_CLIENT = 'z', - IDX_BRAIN_CLIENT_FEATURES = 0xff07, - IDX_BRAIN_HOST = 0xff08, - IDX_BRAIN_PASSWORD = 0xff09, - IDX_BRAIN_PORT = 0xff0a, - IDX_BRAIN_SERVER = 0xff0b, - IDX_BRAIN_SERVER_TIMER = 0xff0c, - IDX_BRAIN_SESSION = 0xff0d, - IDX_BRAIN_SESSION_WHITELIST = 0xff0e, + IDX_BRAIN_CLIENT_FEATURES = 0xff08, + IDX_BRAIN_HOST = 0xff09, + IDX_BRAIN_PASSWORD = 0xff0a, + IDX_BRAIN_PORT = 0xff0b, + IDX_BRAIN_SERVER = 0xff0c, + IDX_BRAIN_SERVER_TIMER = 0xff0d, + IDX_BRAIN_SESSION = 0xff0e, + IDX_BRAIN_SESSION_WHITELIST = 0xff0f, #endif - IDX_CPU_AFFINITY = 0xff0f, + IDX_CPU_AFFINITY = 0xff10, IDX_CUSTOM_CHARSET_1 = '1', IDX_CUSTOM_CHARSET_2 = '2', IDX_CUSTOM_CHARSET_3 = '3', IDX_CUSTOM_CHARSET_4 = '4', - IDX_DEBUG_FILE = 0xff10, - IDX_DEBUG_MODE = 0xff11, - IDX_ENCODING_FROM = 0xff12, - IDX_ENCODING_TO = 0xff13, - IDX_HASH_INFO = 0xff14, - IDX_FORCE = 0xff15, - IDX_HWMON_DISABLE = 0xff16, - IDX_HWMON_TEMP_ABORT = 0xff17, + IDX_DEBUG_FILE = 0xff11, + IDX_DEBUG_MODE = 0xff12, + IDX_ENCODING_FROM = 0xff13, + IDX_ENCODING_TO = 0xff14, + IDX_HASH_INFO = 0xff15, + IDX_FORCE = 0xff16, + IDX_HWMON_DISABLE = 0xff17, + IDX_HWMON_TEMP_ABORT = 0xff18, IDX_HASH_MODE = 'm', - IDX_HCCAPX_MESSAGE_PAIR = 0xff18, + IDX_HCCAPX_MESSAGE_PAIR = 0xff19, IDX_HELP = 'h', - IDX_HEX_CHARSET = 0xff19, - IDX_HEX_SALT = 0xff1a, - IDX_HEX_WORDLIST = 0xff1b, - IDX_HOOK_THREADS = 0xff1c, - IDX_IDENTIFY = 0xff1d, + IDX_HEX_CHARSET = 0xff1a, + IDX_HEX_SALT = 0xff1b, + IDX_HEX_WORDLIST = 0xff1c, + IDX_HOOK_THREADS = 0xff1d, + IDX_IDENTIFY = 0xff1e, IDX_INCREMENT = 'i', - IDX_INCREMENT_MAX = 0xff1e, - IDX_INCREMENT_MIN = 0xff1f, - IDX_INDUCTION_DIR = 0xff20, - IDX_KEEP_GUESSING = 0xff21, + IDX_INCREMENT_MAX = 0xff1f, + IDX_INCREMENT_MIN = 0xff20, + IDX_INDUCTION_DIR = 0xff21, + IDX_KEEP_GUESSING = 0xff22, IDX_KERNEL_ACCEL = 'n', IDX_KERNEL_LOOPS = 'u', IDX_KERNEL_THREADS = 'T', - IDX_KEYBOARD_LAYOUT_MAPPING = 0xff22, - IDX_KEYSPACE = 0xff23, - IDX_LEFT = 0xff24, + IDX_KEYBOARD_LAYOUT_MAPPING = 0xff23, + IDX_KEYSPACE = 0xff24, + IDX_LEFT = 0xff25, IDX_LIMIT = 'l', - IDX_LOGFILE_DISABLE = 0xff25, - IDX_LOOPBACK = 0xff26, - IDX_MACHINE_READABLE = 0xff27, - IDX_MARKOV_CLASSIC = 0xff28, - IDX_MARKOV_DISABLE = 0xff29, - IDX_MARKOV_HCSTAT2 = 0xff2a, - IDX_MARKOV_INVERSE = 0xff2b, + IDX_LOGFILE_DISABLE = 0xff26, + IDX_LOOPBACK = 0xff27, + IDX_MACHINE_READABLE = 0xff28, + IDX_MARKOV_CLASSIC = 0xff29, + IDX_MARKOV_DISABLE = 0xff2a, + IDX_MARKOV_HCSTAT2 = 0xff2b, + IDX_MARKOV_INVERSE = 0xff2c, IDX_MARKOV_THRESHOLD = 't', - IDX_NONCE_ERROR_CORRECTIONS = 0xff2c, + IDX_NONCE_ERROR_CORRECTIONS = 0xff2d, IDX_OPENCL_DEVICE_TYPES = 'D', IDX_OPTIMIZED_KERNEL_ENABLE = 'O', - IDX_OUTFILE_AUTOHEX_DISABLE = 0xff2d, - IDX_OUTFILE_CHECK_DIR = 0xff2e, - IDX_OUTFILE_CHECK_TIMER = 0xff2f, - IDX_OUTFILE_FORMAT = 0xff30, + IDX_OUTFILE_AUTOHEX_DISABLE = 0xff2e, + IDX_OUTFILE_CHECK_DIR = 0xff2f, + IDX_OUTFILE_CHECK_TIMER = 0xff30, + IDX_OUTFILE_FORMAT = 0xff31, IDX_OUTFILE = 'o', - IDX_POTFILE_DISABLE = 0xff31, - IDX_POTFILE_PATH = 0xff32, - IDX_PROGRESS_ONLY = 0xff33, - IDX_QUIET = 0xff34, - IDX_REMOVE = 0xff35, - IDX_REMOVE_TIMER = 0xff36, - IDX_RESTORE = 0xff37, - IDX_RESTORE_DISABLE = 0xff38, - IDX_RESTORE_FILE_PATH = 0xff39, + IDX_POTFILE_DISABLE = 0xff32, + IDX_POTFILE_PATH = 0xff33, + IDX_PROGRESS_ONLY = 0xff34, + IDX_QUIET = 0xff35, + IDX_REMOVE = 0xff36, + IDX_REMOVE_TIMER = 0xff37, + IDX_RESTORE = 0xff38, + IDX_RESTORE_DISABLE = 0xff39, + IDX_RESTORE_FILE_PATH = 0xff3a, IDX_RP_FILE = 'r', - IDX_RP_GEN_FUNC_MAX = 0xff3a, - IDX_RP_GEN_FUNC_MIN = 0xff3b, + IDX_RP_GEN_FUNC_MAX = 0xff3b, + IDX_RP_GEN_FUNC_MIN = 0xff3c, IDX_RP_GEN = 'g', - IDX_RP_GEN_SEED = 0xff3c, + IDX_RP_GEN_SEED = 0xff3d, IDX_RULE_BUF_L = 'j', IDX_RULE_BUF_R = 'k', - IDX_RUNTIME = 0xff3d, - IDX_SCRYPT_TMTO = 0xff3e, + IDX_RUNTIME = 0xff3e, + IDX_SCRYPT_TMTO = 0xff3f, IDX_SEGMENT_SIZE = 'c', - IDX_SELF_TEST_DISABLE = 0xff3f, + IDX_SELF_TEST_DISABLE = 0xff40, IDX_SEPARATOR = 'p', - IDX_SESSION = 0xff40, - IDX_SHOW = 0xff41, + IDX_SESSION = 0xff41, + IDX_SHOW = 0xff42, IDX_SKIP = 's', IDX_SLOW_CANDIDATES = 'S', - IDX_SPEED_ONLY = 0xff42, - IDX_SPIN_DAMP = 0xff43, - IDX_STATUS = 0xff44, - IDX_STATUS_JSON = 0xff45, - IDX_STATUS_TIMER = 0xff46, - IDX_STDOUT_FLAG = 0xff47, - IDX_STDIN_TIMEOUT_ABORT = 0xff48, - IDX_TRUECRYPT_KEYFILES = 0xff49, - IDX_USERNAME = 0xff4a, - IDX_VERACRYPT_KEYFILES = 0xff4b, - IDX_VERACRYPT_PIM_START = 0xff4c, - IDX_VERACRYPT_PIM_STOP = 0xff4d, + IDX_SPEED_ONLY = 0xff43, + IDX_SPIN_DAMP = 0xff44, + IDX_STATUS = 0xff45, + IDX_STATUS_JSON = 0xff46, + IDX_STATUS_TIMER = 0xff47, + IDX_STDOUT_FLAG = 0xff48, + IDX_STDIN_TIMEOUT_ABORT = 0xff49, + IDX_TRUECRYPT_KEYFILES = 0xff4a, + IDX_USERNAME = 0xff4b, + IDX_VERACRYPT_KEYFILES = 0xff4c, + IDX_VERACRYPT_PIM_START = 0xff4d, + IDX_VERACRYPT_PIM_STOP = 0xff4e, IDX_VERSION_LOWER = 'v', IDX_VERSION = 'V', - IDX_WORDLIST_AUTOHEX_DISABLE = 0xff4e, + IDX_WORDLIST_AUTOHEX_DISABLE = 0xff4f, IDX_WORKLOAD_PROFILE = 'w', } user_options_map_t; @@ -1485,82 +1486,83 @@ typedef struct hc_device_param // API: hip - bool is_hip; + bool is_hip; - int hip_warp_size; + int hip_warp_size; - HIPdevice hip_device; - HIPcontext hip_context; - HIPstream hip_stream; + HIPdevice hip_device; + HIPcontext hip_context; + HIPstream hip_stream; - HIPevent hip_event1; - HIPevent hip_event2; + HIPevent hip_event1; + HIPevent hip_event2; - HIPmodule hip_module; - HIPmodule hip_module_shared; - HIPmodule hip_module_mp; - HIPmodule hip_module_amp; + HIPmodule hip_module; + HIPmodule hip_module_shared; + HIPmodule hip_module_mp; + HIPmodule hip_module_amp; - HIPfunction hip_function1; - HIPfunction hip_function12; - HIPfunction hip_function2; - HIPfunction hip_function2e; - HIPfunction hip_function23; - HIPfunction hip_function3; - HIPfunction hip_function4; - HIPfunction hip_function_init2; - HIPfunction hip_function_loop2; - HIPfunction hip_function_mp; - HIPfunction hip_function_mp_l; - HIPfunction hip_function_mp_r; - HIPfunction hip_function_amp; - HIPfunction hip_function_tm; - HIPfunction hip_function_memset; - HIPfunction hip_function_atinit; - HIPfunction hip_function_decompress; - HIPfunction hip_function_aux1; - HIPfunction hip_function_aux2; - HIPfunction hip_function_aux3; - HIPfunction hip_function_aux4; + HIPfunction hip_function1; + HIPfunction hip_function12; + HIPfunction hip_function2p; + HIPfunction hip_function2; + HIPfunction hip_function2e; + HIPfunction hip_function23; + HIPfunction hip_function3; + HIPfunction hip_function4; + HIPfunction hip_function_init2; + HIPfunction hip_function_loop2p; + HIPfunction hip_function_loop2; + HIPfunction hip_function_mp; + HIPfunction hip_function_mp_l; + HIPfunction hip_function_mp_r; + HIPfunction hip_function_amp; + HIPfunction hip_function_tm; + HIPfunction hip_function_memset; + HIPfunction hip_function_atinit; + HIPfunction hip_function_utf8toutf16le; + HIPfunction hip_function_decompress; + HIPfunction hip_function_aux1; + HIPfunction hip_function_aux2; + HIPfunction hip_function_aux3; + HIPfunction hip_function_aux4; - HIPdeviceptr hip_d_pws_buf; - HIPdeviceptr hip_d_pws_amp_buf; - HIPdeviceptr hip_d_pws_comp_buf; - HIPdeviceptr hip_d_pws_idx; - HIPdeviceptr hip_d_words_buf_l; - HIPdeviceptr hip_d_words_buf_r; - HIPdeviceptr hip_d_rules; - HIPdeviceptr hip_d_rules_c; - HIPdeviceptr hip_d_combs; - HIPdeviceptr hip_d_combs_c; - HIPdeviceptr hip_d_bfs; - HIPdeviceptr hip_d_bfs_c; - HIPdeviceptr hip_d_tm_c; - HIPdeviceptr hip_d_bitmap_s1_a; - HIPdeviceptr hip_d_bitmap_s1_b; - HIPdeviceptr hip_d_bitmap_s1_c; - HIPdeviceptr hip_d_bitmap_s1_d; - HIPdeviceptr hip_d_bitmap_s2_a; - HIPdeviceptr hip_d_bitmap_s2_b; - HIPdeviceptr hip_d_bitmap_s2_c; - HIPdeviceptr hip_d_bitmap_s2_d; - HIPdeviceptr hip_d_plain_bufs; - HIPdeviceptr hip_d_digests_buf; - HIPdeviceptr hip_d_digests_shown; - HIPdeviceptr hip_d_salt_bufs; - HIPdeviceptr hip_d_esalt_bufs; - HIPdeviceptr hip_d_tmps; - HIPdeviceptr hip_d_hooks; - HIPdeviceptr hip_d_result; - HIPdeviceptr hip_d_extra0_buf; - HIPdeviceptr hip_d_extra1_buf; - HIPdeviceptr hip_d_extra2_buf; - HIPdeviceptr hip_d_extra3_buf; - HIPdeviceptr hip_d_root_css_buf; - HIPdeviceptr hip_d_markov_css_buf; - HIPdeviceptr hip_d_st_digests_buf; - HIPdeviceptr hip_d_st_salts_buf; - HIPdeviceptr hip_d_st_esalts_buf; + HIPdeviceptr hip_d_pws_buf; + HIPdeviceptr hip_d_pws_amp_buf; + HIPdeviceptr hip_d_pws_comp_buf; + HIPdeviceptr hip_d_pws_idx; + HIPdeviceptr hip_d_rules; + HIPdeviceptr hip_d_rules_c; + HIPdeviceptr hip_d_combs; + HIPdeviceptr hip_d_combs_c; + HIPdeviceptr hip_d_bfs; + HIPdeviceptr hip_d_bfs_c; + HIPdeviceptr hip_d_tm_c; + HIPdeviceptr hip_d_bitmap_s1_a; + HIPdeviceptr hip_d_bitmap_s1_b; + HIPdeviceptr hip_d_bitmap_s1_c; + HIPdeviceptr hip_d_bitmap_s1_d; + HIPdeviceptr hip_d_bitmap_s2_a; + HIPdeviceptr hip_d_bitmap_s2_b; + HIPdeviceptr hip_d_bitmap_s2_c; + HIPdeviceptr hip_d_bitmap_s2_d; + HIPdeviceptr hip_d_plain_bufs; + HIPdeviceptr hip_d_digests_buf; + HIPdeviceptr hip_d_digests_shown; + HIPdeviceptr hip_d_salt_bufs; + HIPdeviceptr hip_d_esalt_bufs; + HIPdeviceptr hip_d_tmps; + HIPdeviceptr hip_d_hooks; + HIPdeviceptr hip_d_result; + HIPdeviceptr hip_d_extra0_buf; + HIPdeviceptr hip_d_extra1_buf; + HIPdeviceptr hip_d_extra2_buf; + HIPdeviceptr hip_d_extra3_buf; + HIPdeviceptr hip_d_root_css_buf; + HIPdeviceptr hip_d_markov_css_buf; + HIPdeviceptr hip_d_st_digests_buf; + HIPdeviceptr hip_d_st_salts_buf; + HIPdeviceptr hip_d_st_esalts_buf; // API: opencl @@ -1653,10 +1655,10 @@ typedef struct backend_ctx { bool enabled; - void *ocl; void *cuda; void *hip; - + void *ocl; + void *nvrtc; void *hiprtc; @@ -1667,6 +1669,7 @@ typedef struct backend_ctx int backend_devices_cnt; int backend_devices_active; + int cuda_devices_cnt; int cuda_devices_active; int hip_devices_cnt; @@ -1704,7 +1707,10 @@ typedef struct backend_ctx int nvrtc_driver_version; int cuda_driver_version; - // cuda + // hip + + int rc_hip_init; + int rc_hiprtc_init; int hiprtc_driver_version; int hip_driver_version; diff --git a/src/Makefile b/src/Makefile index acad8ddb4..e4832860c 100644 --- a/src/Makefile +++ b/src/Makefile @@ -4,7 +4,7 @@ ## SHARED ?= 0 -DEBUG := 1 +DEBUG := 0 PRODUCTION := 0 PRODUCTION_VERSION := v6.2.2 ENABLE_CUBIN ?= 1 diff --git a/src/autotune.c b/src/autotune.c index 04f8bc4c5..cbb1ff499 100644 --- a/src/autotune.c +++ b/src/autotune.c @@ -157,8 +157,9 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param const u32 kernel_power_max = device_param->hardware_power * kernel_accel_max; - int CL_rc; int CU_rc; + int HIP_rc; + int CL_rc; if (device_param->is_cuda == true) { @@ -167,6 +168,13 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (CU_rc == -1) return -1; } + if (device_param->is_hip == true) + { + HIP_rc = run_hip_kernel_atinit (hashcat_ctx, device_param, device_param->hip_d_pws_buf, kernel_power_max); + + if (HIP_rc == -1) return -1; + } + if (device_param->is_opencl == true) { CL_rc = run_opencl_kernel_atinit (hashcat_ctx, device_param, device_param->opencl_d_pws_buf, kernel_power_max); @@ -190,6 +198,13 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (CU_rc == -1) return -1; } + if (device_param->is_hip == true) + { + HIP_rc = hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t)); + + if (HIP_rc == -1) return -1; + } + if (device_param->is_opencl == true) { CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_rules, device_param->opencl_d_rules_c, 0, 0, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), 0, NULL, NULL); @@ -383,6 +398,27 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (CU_rc == -1) return -1; } + if (device_param->is_hip == true) + { + int HIP_rc; + + HIP_rc = run_hip_kernel_memset (hashcat_ctx, device_param, device_param->hip_d_pws_buf, 0, device_param->size_pws); + + if (HIP_rc == -1) return -1; + + HIP_rc = run_hip_kernel_memset (hashcat_ctx, device_param, device_param->hip_d_plain_bufs, 0, device_param->size_plains); + + if (HIP_rc == -1) return -1; + + HIP_rc = run_hip_kernel_memset (hashcat_ctx, device_param, device_param->hip_d_digests_shown, 0, device_param->size_shown); + + if (HIP_rc == -1) return -1; + + HIP_rc = run_hip_kernel_memset (hashcat_ctx, device_param, device_param->hip_d_result, 0, device_param->size_results); + + if (HIP_rc == -1) return -1; + } + if (device_param->is_opencl == true) { int CL_rc; @@ -456,6 +492,11 @@ HC_API_CALL void *thread_autotune (void *p) if (rc_cuCtxSetCurrent == -1) return NULL; } + if (device_param->is_hip == true) + { + if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return NULL; + } + const int rc_autotune = autotune (hashcat_ctx, device_param); if (rc_autotune == -1) @@ -463,5 +504,10 @@ HC_API_CALL void *thread_autotune (void *p) // we should do something here, tell hashcat main that autotune failed to abort } + if (device_param->is_hip == true) + { + if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) return NULL; + } + return NULL; } diff --git a/src/backend.c b/src/backend.c index 20ec98a0e..5b12cd2c5 100644 --- a/src/backend.c +++ b/src/backend.c @@ -55,6 +55,10 @@ static bool is_same_device (const hc_device_param_t *src, const hc_device_param_ if ((src->is_cuda == true) && (dst->is_cuda == true)) return false; + // HIP can't have aliases + + if ((src->is_hip == true) && (dst->is_hip == true)) return false; + // But OpenCL can have aliases if ((src->is_opencl == true) && (dst->is_opencl == true)) @@ -119,6 +123,10 @@ static int backend_ctx_find_alias_devices (hashcat_ctx_t *hashcat_ctx) if (alias_device->is_cuda == true) continue; + // this lets HIP devices survive over OpenCL + + if (alias_device->is_hip == true) continue; + // this lets native OpenCL runtime survive over generic OpenCL runtime if (alias_device->opencl_device_type & CL_DEVICE_TYPE_CPU) @@ -153,6 +161,7 @@ static bool is_same_device_type (const hc_device_param_t *src, const hc_device_p if (strcmp (src->device_name, dst->device_name) != 0) return false; if (src->is_cuda != dst->is_cuda) return false; + if (src->is_hip != dst->is_hip) return false; if (src->is_opencl != dst->is_opencl) return false; if (strcmp (src->device_name, dst->device_name) != 0) return false; @@ -960,6 +969,213 @@ 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 ("fixme.dll"); + #elif defined (__APPLE__) + hiprtc->lib = hc_dlopen ("fixme.dylib"); + #elif defined (__CYGWIN__) + hiprtc->lib = hc_dlopen ("fixme.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_HIPRTCGETPTX, HIPRTC, 1); + HC_LOAD_FUNC (hiprtc, hiprtcGetCodeSize, HIPRTC_HIPRTCGETPTXSIZE, 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); + HC_LOAD_FUNC (hiprtc, hiprtcVersion, HIPRTC_HIPRTCVERSION, 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; + + #if 0 + for(int i =0; i< numOptions; i++) + printf("Option_%d = %s\n", i, options[i]); + #endif + + 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 *ptxSizeRet) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; + + const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetCodeSize (prog, ptxSizeRet); + + 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 *ptx) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; + + const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetCode (prog, ptx); + + if (HIPRTC_err != HIPRTC_SUCCESS) + { + event_log_error (hashcat_ctx, "hiprtcGetCode(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); + + return -1; + } + + return 0; +} + +int hc_hiprtcVersion (hashcat_ctx_t *hashcat_ctx, int *major, int *minor) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; + + const hiprtcResult HIPRTC_err = hiprtc->hiprtcVersion (major, minor); + + if (HIPRTC_err != HIPRTC_SUCCESS) + { + event_log_error (hashcat_ctx, "hiprtcVersion(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); + + return -1; + } + + return 0; +} + // CUDA int cuda_init (hashcat_ctx_t *hashcat_ctx) @@ -2167,6 +2383,1215 @@ 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 ("fixme.dll"); + #elif defined (__APPLE__) + hip->lib = hc_dlopen ("fixme.dylib"); + #elif defined (__CYGWIN__) + hip->lib = hc_dlopen ("fixme.dll"); + #else + hip->lib = hc_dlopen ("libamdhip64.so"); + + //TODO: grab the 4 from the major RT version + if (hip->lib == NULL) hip->lib = hc_dlopen ("libamdhip64.so.4.2.40200"); + #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, hipCtxGetCacheConfig, hipCtxGetCacheConfig, HIP_HIPCTXGETCACHECONFIG, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipCtxGetCurrent, hipCtxGetCurrent, HIP_HIPCTXGETCURRENT, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipCtxGetSharedMemConfig, hipCtxGetSharedMemConfig, HIP_HIPCTXGETSHAREDMEMCONFIG, 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, hipCtxSetCacheConfig, hipCtxSetCacheConfig, HIP_HIPCTXSETCACHECONFIG, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipCtxSetCurrent, hipCtxSetCurrent, HIP_HIPCTXSETCURRENT, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipCtxSetSharedMemConfig, hipCtxSetSharedMemConfig, HIP_HIPCTXSETSHAREDMEMCONFIG, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipCtxSynchronize, hipCtxSynchronize, HIP_HIPCTXSYNCHRONIZE, 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, hipDeviceGet, hipDeviceGet, HIP_HIPDEVICEGET, 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, hipEventQuery, hipEventQuery, HIP_HIPEVENTQUERY, 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, hipFuncSetAttribute, hipFuncSetAttribute, HIP_HIPFUNCSETATTRIBUTE, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipFuncSetCacheConfig, hipFuncSetCacheConfig, HIP_HIPFUNCSETCACHECONFIG, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipFuncSetSharedMemConfig, hipFuncSetSharedMemConfig, HIP_HIPFUNCSETSHAREDMEMCONFIG, 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, hipMemAllocHost, hipMemAllocHost, HIP_HIPMEMALLOCHOST, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoD, hipMemcpyDtoD, HIP_HIPMEMCPYDTOD, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoH, hipMemcpyDtoH, HIP_HIPMEMCPYDTOH, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemcpyHtoD, hipMemcpyHtoD, HIP_HIPMEMCPYHTOD, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemFree, hipFree, HIP_HIPMEMFREE, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemFreeHost, hipFreeHost, HIP_HIPMEMFREEHOST, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemGetInfo, hipMemGetInfo, HIP_HIPMEMGETINFO, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemsetD32, hipMemsetD32, HIP_HIPMEMSETD32, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemsetD8, hipMemsetD8, HIP_HIPMEMSETD8, 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, hipModuleLoad, hipModuleLoad, HIP_HIPMODULELOAD, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipModuleLoadData, hipModuleLoadData, HIP_HIPMODULELOADDATA, 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, hipProfilerStart, hipProfilerStart, HIP_HIPPROFILERSTART, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipProfilerStop, hipProfilerStop, HIP_HIPPROFILERSTOP, 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); + HC_LOAD_FUNC_HIP (hip, hipStreamWaitEvent, hipStreamWaitEvent, HIP_HIPSTREAMWAITEVENT, HIP, 1); + #if defined (WITH_CUBINX) + HC_LOAD_FUNC_HIP (hip, hipLinkCreate, hipLinkCreate, HIP_HIPLINKCREATE, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipLinkAddData, hipLinkAddData, HIP_HIPLINKADDDATA, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipLinkDestroy, hipLinkDestroy, HIP_HIPLINKDESTROY, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipLinkComplete, hipLinkComplete, HIP_HIPLINKCOMPLETE, HIP, 1); + #endif + + 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_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 HIPresult HIP_err = hip->hipInit (Flags); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipInit(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipInit(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipDeviceGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPdevice_attribute attrib, HIPdevice dev) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipDeviceGetAttribute (pi, attrib, dev); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + 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 HIPresult HIP_err = hip->hipDeviceGetCount (count); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipDeviceGetCount(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipDeviceGetCount(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipDeviceGet (hashcat_ctx_t *hashcat_ctx, HIPdevice* device, int ordinal) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipDeviceGet (device, ordinal); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipDeviceGet(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipDeviceGet(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipDeviceGetName (hashcat_ctx_t *hashcat_ctx, char *name, int len, HIPdevice dev) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipDeviceGetName (name, len, dev); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + 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 dev) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipDeviceTotalMem (bytes, dev); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + 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 HIPresult HIP_err = hip->hipDriverGetVersion (driverVersion); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipDriverGetVersion(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipDriverGetVersion(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx, unsigned int flags, HIPdevice dev) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipCtxCreate (pctx, flags, dev); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + 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, HIPcontext ctx) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipCtxDestroy (ctx); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipCtxDestroy(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipCtxDestroy(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, HIPmodule *module, const void *image, unsigned int numOptions, HIPjit_option *options, void **optionValues) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipModuleLoadDataEx (module, image, numOptions, options, optionValues); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + 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 hmod) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipModuleUnload (hmod); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipModuleUnload(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipModuleUnload(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipCtxSetCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipCtxSetCurrent (ctx); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipCtxSetCurrent(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipCtxSetCurrent(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr *dptr, size_t bytesize) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipMemAlloc (dptr, bytesize); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + 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 dptr) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipMemFree (dptr); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipMemFree(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemFree(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, HIPdeviceptr srcDevice, size_t ByteCount) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipMemcpyDtoH (dstHost, srcDevice, ByteCount); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipMemcpyDtoH(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyDtoH(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, HIPdeviceptr srcDevice, size_t ByteCount) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipMemcpyDtoD (dstDevice, srcDevice, ByteCount); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, const void *srcHost, size_t ByteCount) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipMemcpyHtoD (dstDevice, srcHost, ByteCount); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipMemcpyHtoD(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyHtoD(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, HIPfunction *hfunc, HIPmodule hmod, const char *name) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipModuleGetFunction (hfunc, hmod, name); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + 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 *dptr, size_t *bytes, HIPmodule hmod, const char *name) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipModuleGetGlobal (dptr, bytes, hmod, name); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipModuleGetGlobal(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipModuleGetGlobal(): %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 HIPresult HIP_err = hip->hipMemGetInfo (free, total); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipMemGetInfo(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemGetInfo(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPfunction_attribute attrib, HIPfunction hfunc) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipFuncGetAttribute (pi, attrib, hfunc); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipFuncGetAttribute(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipFuncGetAttribute(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipFuncSetAttribute (hashcat_ctx_t *hashcat_ctx, HIPfunction hfunc, HIPfunction_attribute attrib, int value) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipFuncSetAttribute (hfunc, attrib, value); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipFuncSetAttribute(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipFuncSetAttribute(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, HIPstream *phStream, unsigned int Flags) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipStreamCreate (phStream, Flags); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + 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 hStream) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipStreamDestroy (hStream); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + 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 hStream) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipStreamSynchronize (hStream); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipStreamSynchronize(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipStreamSynchronize(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipLaunchKernel (hashcat_ctx_t *hashcat_ctx, HIPfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, HIPstream hStream, void **kernelParams, void **extra) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipLaunchKernel (f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipLaunchKernel(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipLaunchKernel(): %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 HIPresult HIP_err = hip->hipCtxSynchronize (); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipCtxSynchronize(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipCtxSynchronize(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipEventCreate (hashcat_ctx_t *hashcat_ctx, HIPevent *phEvent, unsigned int Flags) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipEventCreate (phEvent, Flags); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + 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 hEvent) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipEventDestroy (hEvent); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + 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 hStart, HIPevent hEnd) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipEventElapsedTime (pMilliseconds, hStart, hEnd); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipEventElapsedTime(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipEventElapsedTime(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipEventQuery (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipEventQuery (hEvent); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipEventQuery(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipEventQuery(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent, HIPstream hStream) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipEventRecord (hEvent, hStream); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + 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 hEvent) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipEventSynchronize (hEvent); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipEventSynchronize(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipEventSynchronize(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipCtxSetCacheConfig (hashcat_ctx_t *hashcat_ctx, HIPfunc_cache config) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipCtxSetCacheConfig (config); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipCtxSetCacheConfig(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipCtxSetCacheConfig(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipCtxPushCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipCtxPushCurrent (ctx); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipCtxPushCurrent(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipCtxPushCurrent(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipCtxPopCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipCtxPopCurrent (pctx); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipCtxPopCurrent(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipCtxPopCurrent(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipLinkCreate (hashcat_ctx_t *hashcat_ctx, unsigned int numOptions, HIPjit_option *options, void **optionValues, HIPlinkState *stateOut) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipLinkCreate (numOptions, options, optionValues, stateOut); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipLinkCreate(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipLinkCreate(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipLinkAddData (hashcat_ctx_t *hashcat_ctx, HIPlinkState state, HIPjitInputType type, void *data, size_t size, const char *name, unsigned int numOptions, HIPjit_option *options, void **optionValues) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipLinkAddData (state, type, data, size, name, numOptions, options, optionValues); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipLinkAddData(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipLinkAddData(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipLinkDestroy (hashcat_ctx_t *hashcat_ctx, HIPlinkState state) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipLinkDestroy (state); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipLinkDestroy(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipLinkDestroy(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipLinkComplete (hashcat_ctx_t *hashcat_ctx, HIPlinkState state, void **hipbinOut, size_t *sizeOut) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipLinkComplete (state, hipbinOut, sizeOut); + + if (HIP_err != HIP_SUCCESS) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + { + event_log_error (hashcat_ctx, "hipLinkComplete(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipLinkComplete(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + // OpenCL int ocl_init (hashcat_ctx_t *hashcat_ctx) @@ -2916,6 +4341,15 @@ int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, c if (hc_cuCtxPopCurrent (hashcat_ctx, &device_param->cuda_context) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return -1; + + if (hc_hipMemcpyDtoH (hashcat_ctx, &pw_idx, device_param->hip_d_pws_idx + (gidd * sizeof (pw_idx_t)), sizeof (pw_idx_t)) == -1) return -1; + + if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, gidd * sizeof (pw_idx_t), sizeof (pw_idx_t), &pw_idx, 0, NULL, NULL) == -1) return -1; @@ -2937,6 +4371,18 @@ int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, c } } + if (device_param->is_hip == true) + { + if (cnt > 0) + { + if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return -1; + + if (hc_hipMemcpyDtoH (hashcat_ctx,pw->i, device_param->hip_d_pws_comp_buf + (off * sizeof (u32)), cnt * sizeof (u32)) == -1) return -1; + + if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) return -1; + } + } + if (device_param->is_opencl == true) { if (cnt > 0) @@ -2986,6 +4432,11 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_tm_c, size_tm) == -1) return -1; } + if (device_param->is_hip == true) + { + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_tm_c, size_tm) == -1) return -1; + } + if (device_param->is_opencl == true) { if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_tm_c, size_tm) == -1) return -1; @@ -2998,6 +4449,11 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_bfs_c, device_param->cuda_d_tm_c, size_tm) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_bfs_c, device_param->hip_d_tm_c, size_tm) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tm_c, device_param->opencl_d_bfs_c, 0, 0, size_tm, 0, NULL, NULL) == -1) return -1; @@ -3059,6 +4515,11 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_pws_buf, device_param->cuda_d_pws_amp_buf, pws_cnt * sizeof (pw_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_pws_buf, device_param->hip_d_pws_amp_buf, pws_cnt * sizeof (pw_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_amp_buf, device_param->opencl_d_pws_buf, 0, 0, pws_cnt * sizeof (pw_t), 0, NULL, NULL) == -1) return -1; @@ -3079,6 +4540,11 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (run_cuda_kernel_utf8toutf16le (hashcat_ctx, device_param, device_param->cuda_d_pws_buf, pws_cnt) == -1) return -1; } + if (device_param->is_hip == true) + { + if (run_hip_kernel_utf8toutf16le (hashcat_ctx, device_param, device_param->hip_d_pws_buf, pws_cnt) == -1) return -1; + } + if (device_param->is_opencl == true) { if (run_opencl_kernel_utf8toutf16le (hashcat_ctx, device_param, device_param->opencl_d_pws_buf, pws_cnt) == -1) return -1; @@ -3096,6 +4562,11 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, pws_cnt * hashconfig->hook_size, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; @@ -3145,6 +4616,11 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, pws_cnt * hashconfig->hook_size, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; @@ -3234,6 +4710,11 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, pws_cnt * hashconfig->hook_size, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; @@ -3283,6 +4764,11 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, pws_cnt * hashconfig->hook_size) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, pws_cnt * hashconfig->hook_size, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; @@ -3390,6 +4876,11 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_tmps, device_param->size_tmps) == -1) return -1; } + if (device_param->is_hip == true) + { + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_tmps, device_param->size_tmps) == -1) return -1; + } + if (device_param->is_opencl == true) { if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_tmps, device_param->size_tmps) == -1) return -1; @@ -3403,6 +4894,11 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; } + if (device_param->is_hip == true) + { + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; + } + if (device_param->is_opencl == true) { if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_hooks, pws_cnt * hashconfig->hook_size) == -1) return -1; @@ -3557,6 +5053,99 @@ int run_cuda_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device return run_cuda_kernel_memset (hashcat_ctx, device_param, buf, 0, size); } +int run_hip_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 num) +{ + u64 num_elements = num; + + device_param->kernel_params_atinit[0] = (void *) &buf; + device_param->kernel_params_atinit_buf64[1] = num_elements; + + const u64 kernel_threads = device_param->kernel_wgs_atinit; + + num_elements = CEILDIV (num_elements, kernel_threads); + + HIPfunction function = device_param->hip_function_atinit; + + if (hc_hipLaunchKernel (hashcat_ctx, function, num_elements, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, device_param->kernel_params_atinit, NULL) == -1) return -1; + + if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; + + return 0; +} + +int run_hip_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 num) +{ + u64 num_elements = num; + + device_param->kernel_params_utf8toutf16le[0] = (void *) &buf; + device_param->kernel_params_utf8toutf16le_buf64[1] = num_elements; + + const u64 kernel_threads = device_param->kernel_wgs_utf8toutf16le; + + num_elements = CEILDIV (num_elements, kernel_threads); + + HIPfunction function = device_param->hip_function_utf8toutf16le; + + if (hc_hipLaunchKernel (hashcat_ctx, function, num_elements, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, device_param->kernel_params_utf8toutf16le, NULL) == -1) return -1; + + if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; + + return 0; +} + +int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u32 value, const u64 size) +{ + const u64 num16d = size / 16; + const u64 num16m = size % 16; + + if (num16d) + { + device_param->kernel_params_memset[0] = (void *) &buf; + device_param->kernel_params_memset_buf32[1] = value; + device_param->kernel_params_memset_buf64[2] = num16d; + + const u64 kernel_threads = device_param->kernel_wgs_memset; + + u64 num_elements = num16d; + + num_elements = CEILDIV (num_elements, kernel_threads); + + HIPfunction function = device_param->hip_function_memset; + + //HIP_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 0, sizeof (cl_mem), (void *) &buf); if (HIP_rc == -1) return -1; + //HIP_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 1, sizeof (cl_uint), device_param->kernel_params_memset[1]); if (HIP_rc == -1) return -1; + //HIP_rc = hc_clSetKernelArg (hashcat_ctx, kernel, 2, sizeof (cl_ulong), device_param->kernel_params_memset[2]); if (HIP_rc == -1) return -1; + + //const size_t global_work_size[3] = { num_elements, 1, 1 }; + //const size_t local_work_size[3] = { kernel_threads, 1, 1 }; + + if (hc_hipLaunchKernel (hashcat_ctx, function, num_elements, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, device_param->kernel_params_memset, NULL) == -1) return -1; + + if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; + } + + if (num16m) + { + u32 tmp[4]; + + tmp[0] = value; + tmp[1] = value; + tmp[2] = value; + tmp[3] = value; + + // Apparently are allowed to do this: https://devtalk.nvidia.com/default/topic/761515/how-to-copy-to-device-memory-with-offset-/ + + if (hc_hipMemcpyHtoD (hashcat_ctx, buf + (num16d * 16), tmp, num16m) == -1) return -1; + } + + return 0; +} + +int run_hip_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 size) +{ + return run_hip_kernel_memset (hashcat_ctx, device_param, buf, 0, size); +} + int run_opencl_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 num) { u64 num_elements = num; @@ -3860,6 +5449,105 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con } } + if (device_param->is_hip == true) + { + HIPfunction hip_function = NULL; + + if (device_param->is_hip == true) + { + switch (kern_run) + { + case KERN_RUN_1: hip_function = device_param->hip_function1; break; + case KERN_RUN_12: hip_function = device_param->hip_function12; break; + case KERN_RUN_2P: hip_function = device_param->hip_function2p; break; + case KERN_RUN_2: hip_function = device_param->hip_function2; break; + case KERN_RUN_2E: hip_function = device_param->hip_function2e; break; + case KERN_RUN_23: hip_function = device_param->hip_function23; break; + case KERN_RUN_3: hip_function = device_param->hip_function3; break; + case KERN_RUN_4: hip_function = device_param->hip_function4; break; + case KERN_RUN_INIT2: hip_function = device_param->hip_function_init2; break; + case KERN_RUN_LOOP2P: hip_function = device_param->hip_function_loop2p; break; + case KERN_RUN_LOOP2: hip_function = device_param->hip_function_loop2; break; + case KERN_RUN_AUX1: hip_function = device_param->hip_function_aux1; break; + case KERN_RUN_AUX2: hip_function = device_param->hip_function_aux2; break; + case KERN_RUN_AUX3: hip_function = device_param->hip_function_aux3; break; + case KERN_RUN_AUX4: hip_function = device_param->hip_function_aux4; break; + } + + if (hc_hipFuncSetAttribute (hashcat_ctx, hip_function, HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, dynamic_shared_mem) == -1) return -1; + } + + if (kernel_threads == 0) kernel_threads = 1; + + num_elements = CEILDIV (num_elements, kernel_threads); + + if (kern_run == KERN_RUN_1) + { + if (hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_SIMD_INIT) + { + num_elements = CEILDIV (num_elements, device_param->vector_width); + } + } + else if (kern_run == KERN_RUN_2) + { + if (hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_SIMD_LOOP) + { + num_elements = CEILDIV (num_elements, device_param->vector_width); + } + } + else if (kern_run == KERN_RUN_3) + { + if (hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_SIMD_COMP) + { + num_elements = CEILDIV (num_elements, device_param->vector_width); + } + } + else if (kern_run == KERN_RUN_INIT2) + { + if (hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_SIMD_INIT2) + { + num_elements = CEILDIV (num_elements, device_param->vector_width); + } + } + else if (kern_run == KERN_RUN_LOOP2) + { + if (hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_SIMD_LOOP2) + { + num_elements = CEILDIV (num_elements, device_param->vector_width); + } + } + + if (hc_hipEventRecord (hashcat_ctx, device_param->hip_event1, device_param->hip_stream) == -1) return -1; + + if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements, 1, 1, kernel_threads, 1, 1, dynamic_shared_mem, device_param->hip_stream, device_param->kernel_params, NULL) == -1) return -1; + + if (hc_hipEventRecord (hashcat_ctx, device_param->hip_event2, device_param->hip_stream) == -1) return -1; + + if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; + + if (hc_hipEventSynchronize (hashcat_ctx, device_param->hip_event2) == -1) return -1; + + float exec_ms; + + if (hc_hipEventElapsedTime (hashcat_ctx, &exec_ms, device_param->hip_event1, device_param->hip_event2) == -1) return -1; + + if (event_update) + { + u32 exec_pos = device_param->exec_pos; + + device_param->exec_msec[exec_pos] = exec_ms; + + exec_pos++; + + if (exec_pos == EXEC_CACHE) + { + exec_pos = 0; + } + + device_param->exec_pos = exec_pos; + } + } + if (device_param->is_opencl == true) { cl_kernel opencl_kernel = NULL; @@ -4089,6 +5777,32 @@ int run_kernel_mp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } + if (device_param->is_hip == true) + { + HIPfunction hip_function = NULL; + + void **hip_args = NULL; + + switch (kern_run) + { + case KERN_RUN_MP: hip_function = device_param->hip_function_mp; + hip_args = device_param->kernel_params_mp; + break; + case KERN_RUN_MP_R: hip_function = device_param->hip_function_mp_r; + hip_args = device_param->kernel_params_mp_r; + break; + case KERN_RUN_MP_L: hip_function = device_param->hip_function_mp_l; + hip_args = device_param->kernel_params_mp_l; + break; + } + + num_elements = CEILDIV (num_elements, kernel_threads); + + if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, hip_args, NULL) == -1) return -1; + + if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; + } + if (device_param->is_opencl == true) { cl_kernel opencl_kernel = NULL; @@ -4156,6 +5870,15 @@ int run_kernel_tm (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } + if (device_param->is_hip == true) + { + HIPfunction hip_function = device_param->hip_function_tm; + + if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements / kernel_threads, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, device_param->kernel_params_tm, NULL) == -1) return -1; + + if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; + } + if (device_param->is_opencl == true) { cl_kernel cuda_kernel = device_param->opencl_kernel_tm; @@ -4192,6 +5915,17 @@ int run_kernel_amp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } + if (device_param->is_hip == true) + { + num_elements = CEILDIV (num_elements, kernel_threads); + + HIPfunction hip_function = device_param->hip_function_amp; + + if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, device_param->kernel_params_amp, NULL) == -1) return -1; + + if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; + } + if (device_param->is_opencl == true) { num_elements = round_up_multiple_64 (num_elements, kernel_threads); @@ -4232,6 +5966,17 @@ int run_kernel_decompress (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; } + if (device_param->is_hip == true) + { + num_elements = CEILDIV (num_elements, kernel_threads); + + HIPfunction hip_function = device_param->hip_function_decompress; + + if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, device_param->kernel_params_decompress, NULL) == -1) return -1; + + if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; + } + if (device_param->is_opencl == true) { num_elements = round_up_multiple_64 (num_elements, kernel_threads); @@ -4290,6 +6035,20 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const } } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; + + const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; + + const u32 off = pw_idx->off; + + if (off) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; + } + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; @@ -4324,6 +6083,20 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const } } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; + + const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; + + const u32 off = pw_idx->off; + + if (off) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; + } + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; @@ -4392,6 +6165,20 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const } } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; + + const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; + + const u32 off = pw_idx->off; + + if (off) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; + } + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; @@ -4426,6 +6213,20 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const } } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; + + const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; + + const u32 off = pw_idx->off; + + if (off) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; + } + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; @@ -4458,6 +6259,20 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const } } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_idx, device_param->pws_idx, pws_cnt * sizeof (pw_idx_t)) == -1) return -1; + + const pw_idx_t *pw_idx = device_param->pws_idx + pws_cnt; + + const u32 off = pw_idx->off; + + if (off) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_comp_buf, device_param->pws_comp, off * sizeof (u32)) == -1) return -1; + } + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_idx, CL_TRUE, 0, pws_cnt * sizeof (pw_idx_t), device_param->pws_idx, 0, NULL, NULL) == -1) return -1; @@ -4679,6 +6494,11 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules + (innerloop_pos * sizeof (kernel_rule_t)), innerloop_left * sizeof (kernel_rule_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules + (innerloop_pos * sizeof (kernel_rule_t)), innerloop_left * sizeof (kernel_rule_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_rules, device_param->opencl_d_rules_c, innerloop_pos * sizeof (kernel_rule_t), 0, innerloop_left * sizeof (kernel_rule_t), 0, NULL, NULL) == -1) return -1; @@ -4799,6 +6619,11 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL) == -1) return -1; @@ -4817,6 +6642,11 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs, device_param->opencl_d_combs_c, 0, 0, innerloop_left * sizeof (pw_t), 0, NULL, NULL) == -1) return -1; @@ -4835,6 +6665,11 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs, device_param->opencl_d_combs_c, 0, 0, innerloop_left * sizeof (pw_t), 0, NULL, NULL) == -1) return -1; @@ -4956,6 +6791,11 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->combs_buf, innerloop_left * sizeof (pw_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_TRUE, 0, innerloop_left * sizeof (pw_t), device_param->combs_buf, 0, NULL, NULL) == -1) return -1; @@ -4974,6 +6814,11 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_combs_c, device_param->cuda_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_combs_c, device_param->hip_d_combs, innerloop_left * sizeof (pw_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs, device_param->opencl_d_combs_c, 0, 0, innerloop_left * sizeof (pw_t), 0, NULL, NULL) == -1) return -1; @@ -4994,6 +6839,11 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_bfs_c, device_param->cuda_d_bfs, innerloop_left * sizeof (bf_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_bfs_c, device_param->hip_d_bfs, innerloop_left * sizeof (bf_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bfs, device_param->opencl_d_bfs_c, 0, 0, innerloop_left * sizeof (bf_t), 0, NULL, NULL) == -1) return -1; @@ -5278,6 +7128,99 @@ int backend_ctx_init (hashcat_ctx_t *hashcat_ctx) } } + /** + * Load and map HIP library calls, then init HIP + */ + + int rc_hip_init = -1; + + if (user_options->backend_ignore_hip == false) + { + HIP_PTR *hip = (HIP_PTR *) hcmalloc (sizeof (HIP_PTR)); + + backend_ctx->hip = hip; + + rc_hip_init = hip_init (hashcat_ctx); + + if (rc_hip_init == -1) + { + backend_ctx->rc_hip_init = rc_hip_init; + + hip_close (hashcat_ctx); + } + + /** + * Load and map HIPRTC library calls + */ + + HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) hcmalloc (sizeof (HIPRTC_PTR)); + + backend_ctx->hiprtc = hiprtc; + + int rc_hiprtc_init = hiprtc_init (hashcat_ctx); + + if (rc_hiprtc_init == -1) + { + backend_ctx->rc_hiprtc_init = rc_hiprtc_init; + + hiprtc_close (hashcat_ctx); + } + + /** + * Check if both HIP and HIPRTC were load successful + */ + + if ((rc_hip_init == 0) && (rc_hiprtc_init == 0)) + { + // hiprtc version + + int hiprtc_major = 0; + int hiprtc_minor = 0; + + if (hc_hiprtcVersion (hashcat_ctx, &hiprtc_major, &hiprtc_minor) == -1) return -1; + + int hiprtc_driver_version = (hiprtc_major * 1000) + (hiprtc_minor * 10); + + backend_ctx->hiprtc_driver_version = hiprtc_driver_version; + + if (hiprtc_driver_version < 9000) + { + event_log_error (hashcat_ctx, "Outdated AMD HIPRTC driver version '%d' detected!", hiprtc_driver_version); + + event_log_warning (hashcat_ctx, "See hashcat.net for officially supported AMD HIP versions."); + event_log_warning (hashcat_ctx, NULL); + + return -1; + } + + // hip version + + int hip_driver_version = 10000; + + //if (hc_hipDriverGetVersion (hashcat_ctx, &hip_driver_version) == -1) return -1; + + backend_ctx->hip_driver_version = hip_driver_version; + + if (hip_driver_version < 9000) + { + event_log_error (hashcat_ctx, "Outdated AMD HIP driver version '%d' detected!", hip_driver_version); + + event_log_warning (hashcat_ctx, "See hashcat.net for officially supported AMD HIP versions."); + event_log_warning (hashcat_ctx, NULL); + + return -1; + } + } + else + { + rc_hip_init = -1; + rc_hiprtc_init = -1; + + hip_close (hashcat_ctx); + hiprtc_close (hashcat_ctx); + } + } + /** * Load and map OpenCL library calls */ @@ -5301,11 +7244,11 @@ int backend_ctx_init (hashcat_ctx_t *hashcat_ctx) * return if both CUDA and OpenCL initialization failed */ - if ((rc_cuda_init == -1) && (rc_ocl_init == -1)) + if ((rc_cuda_init == -1) && (rc_hip_init == -1) && (rc_ocl_init == -1)) { - event_log_error (hashcat_ctx, "ATTENTION! No OpenCL or CUDA installation found."); + event_log_error (hashcat_ctx, "ATTENTION! No OpenCL, HIP or CUDA installation found."); - event_log_warning (hashcat_ctx, "You are probably missing the CUDA or OpenCL runtime installation."); + event_log_warning (hashcat_ctx, "You are probably missing the CUDA, HIP or OpenCL runtime installation."); event_log_warning (hashcat_ctx, NULL); #if defined (__linux__) @@ -5366,6 +7309,18 @@ int backend_ctx_init (hashcat_ctx_t *hashcat_ctx) } } + /** + * HIP API: init + */ + + if (backend_ctx->hip) + { + if (hc_hipInit (hashcat_ctx, 0) == -1) + { + hip_close (hashcat_ctx); + } + } + /** * OpenCL API: init */ @@ -5614,11 +7569,11 @@ int backend_ctx_init (hashcat_ctx_t *hashcat_ctx) * Final checks */ - if ((backend_ctx->cuda == NULL) && (backend_ctx->ocl == NULL)) + if ((backend_ctx->cuda == NULL) && (backend_ctx->hip == NULL) && (backend_ctx->ocl == NULL)) { - event_log_error (hashcat_ctx, "ATTENTION! No OpenCL-compatible or CUDA-compatible platform found."); + event_log_error (hashcat_ctx, "ATTENTION! No OpenCL-compatible, HIP-compatible or CUDA-compatible platform found."); - event_log_warning (hashcat_ctx, "You are probably missing the OpenCL or CUDA runtime installation."); + event_log_warning (hashcat_ctx, "You are probably missing the OpenCL, CUDA or HIP runtime installation."); event_log_warning (hashcat_ctx, NULL); #if defined (__linux__) @@ -5664,9 +7619,12 @@ void backend_ctx_destroy (hashcat_ctx_t *hashcat_ctx) hcfree (backend_ctx->opencl_platforms_version); } - nvrtc_close (hashcat_ctx); - cuda_close (hashcat_ctx); - ocl_close (hashcat_ctx); + nvrtc_close (hashcat_ctx); + hiprtc_close (hashcat_ctx); + + cuda_close (hashcat_ctx); + hip_close (hashcat_ctx); + ocl_close (hashcat_ctx); memset (backend_ctx, 0, sizeof (backend_ctx_t)); } @@ -5689,6 +7647,8 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) int backend_devices_idx = 0; + // CUDA + int cuda_devices_cnt = 0; int cuda_devices_active = 0; @@ -5725,10 +7685,11 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) device_param->cuda_device = cuda_device; - device_param->is_cuda = true; - + device_param->is_cuda = true; + device_param->is_hip = false; device_param->is_opencl = false; + device_param->use_opencl12 = false; device_param->use_opencl20 = false; device_param->use_opencl21 = false; @@ -6047,6 +8008,390 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) backend_ctx->cuda_devices_cnt = cuda_devices_cnt; backend_ctx->cuda_devices_active = cuda_devices_active; + // HIP + + int hip_devices_cnt = 0; + int hip_devices_active = 0; + + if (backend_ctx->hip) + { + // device count + + if (hc_hipDeviceGetCount (hashcat_ctx, &hip_devices_cnt) == -1) + { + hip_close (hashcat_ctx); + } + + backend_ctx->hip_devices_cnt = hip_devices_cnt; + + // device specific + + for (int hip_devices_idx = 0; hip_devices_idx < hip_devices_cnt; hip_devices_idx++, backend_devices_idx++) + { + const u32 device_id = backend_devices_idx; + + hc_device_param_t *device_param = &devices_param[backend_devices_idx]; + + device_param->device_id = device_id; + + backend_ctx->backend_device_from_hip[hip_devices_idx] = backend_devices_idx; + + HIPdevice hip_device; + + if (hc_hipDeviceGet (hashcat_ctx, &hip_device, hip_devices_idx) == -1) + { + device_param->skipped = true; + continue; + } + + device_param->hip_device = hip_device; + + device_param->is_cuda = false; + device_param->is_hip = true; + device_param->is_opencl = false; + + device_param->use_opencl12 = false; + device_param->use_opencl20 = false; + device_param->use_opencl21 = false; + + // device_name + + char *device_name = (char *) hcmalloc (HCBUFSIZ_TINY); + + if (hc_hipDeviceGetName (hashcat_ctx, device_name, HCBUFSIZ_TINY, hip_device) == -1) + { + device_param->skipped = true; + hcfree (device_name); + continue; + } + + device_param->device_name = device_name; + + hc_string_trim_leading (device_name); + + hc_string_trim_trailing (device_name); + + // device_processors + + int device_processors = 0; + + if (hc_hipDeviceGetAttribute (hashcat_ctx, &device_processors, HIP_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + device_param->device_processors = device_processors; + + // device_global_mem, device_maxmem_alloc, device_available_mem + + size_t bytes = 0; + + if (hc_hipDeviceTotalMem (hashcat_ctx, &bytes, hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + device_param->device_global_mem = (u64) bytes; + + device_param->device_maxmem_alloc = (u64) bytes; + + device_param->device_available_mem = 0; + + // warp size + + int hip_warp_size = 0; + + if (hc_hipDeviceGetAttribute (hashcat_ctx, &hip_warp_size, HIP_DEVICE_ATTRIBUTE_WARP_SIZE, hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + device_param->hip_warp_size = hip_warp_size; + + // sm_minor, sm_major + + int sm_major = 0; + int sm_minor = 0; + + if (hc_hipDeviceGetAttribute (hashcat_ctx, &sm_major, HIP_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + if (hc_hipDeviceGetAttribute (hashcat_ctx, &sm_minor, HIP_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + device_param->sm_major = sm_major; + device_param->sm_minor = sm_minor; + + // device_maxworkgroup_size + + int device_maxworkgroup_size = 0; + + if (hc_hipDeviceGetAttribute (hashcat_ctx, &device_maxworkgroup_size, HIP_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + device_param->device_maxworkgroup_size = device_maxworkgroup_size; + + // max_clock_frequency + + int device_maxclock_frequency = 0; + + if (hc_hipDeviceGetAttribute (hashcat_ctx, &device_maxclock_frequency, HIP_DEVICE_ATTRIBUTE_CLOCK_RATE, hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + device_param->device_maxclock_frequency = device_maxclock_frequency / 1000; + + // pcie_bus, pcie_device, pcie_function + + int pci_domain_id_nv = 0; + int pci_bus_id_nv = 0; + int pci_slot_id_nv = 0; + + if (hc_hipDeviceGetAttribute (hashcat_ctx, &pci_domain_id_nv, HIP_DEVICE_ATTRIBUTE_PCI_DOMAIN_ID, hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + if (hc_hipDeviceGetAttribute (hashcat_ctx, &pci_bus_id_nv, HIP_DEVICE_ATTRIBUTE_PCI_BUS_ID, hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + if (hc_hipDeviceGetAttribute (hashcat_ctx, &pci_slot_id_nv, HIP_DEVICE_ATTRIBUTE_PCI_DEVICE_ID, hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + device_param->pcie_domain = (u8) (pci_domain_id_nv); + device_param->pcie_bus = (u8) (pci_bus_id_nv); + device_param->pcie_device = (u8) (pci_slot_id_nv >> 3); + device_param->pcie_function = (u8) (pci_slot_id_nv & 7); + + // kernel_exec_timeout + + int kernel_exec_timeout = 0; + + if (hc_hipDeviceGetAttribute (hashcat_ctx, &kernel_exec_timeout, HIP_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT, hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + device_param->kernel_exec_timeout = kernel_exec_timeout; + + // max_shared_memory_per_block + + int max_shared_memory_per_block = 0; + + if (hc_hipDeviceGetAttribute (hashcat_ctx, &max_shared_memory_per_block, HIP_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK, hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + if (max_shared_memory_per_block < 32768) + { + event_log_error (hashcat_ctx, "* Device #%u: This device's shared buffer size is too small.", device_id + 1); + + device_param->skipped = true; + } + + device_param->device_local_mem_size = max_shared_memory_per_block; + + // device_max_constant_buffer_size + + int device_max_constant_buffer_size = 0; + + if (hc_hipDeviceGetAttribute (hashcat_ctx, &device_max_constant_buffer_size, HIP_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY, hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + // TODO: broken on HIP? + + device_max_constant_buffer_size = 65536; + + if (device_max_constant_buffer_size < 65536) + { + event_log_error (hashcat_ctx, "* Device #%u: This device's local mem size is too small.", device_id + 1); + + device_param->skipped = true; + } + + // some attributes have to be hardcoded values because they are used for instance in the build options + + device_param->device_local_mem_type = CL_LOCAL; + device_param->opencl_device_type = CL_DEVICE_TYPE_GPU; + device_param->opencl_device_vendor_id = VENDOR_ID_AMD_USE_HIP; + device_param->opencl_platform_vendor_id = VENDOR_ID_AMD_USE_HIP; + + // or in the cached kernel checksum + + device_param->opencl_device_version = ""; + device_param->opencl_driver_version = ""; + + // or just to make sure they are not NULL + + device_param->opencl_device_vendor = ""; + device_param->opencl_device_c_version = ""; + + // skipped + + if ((backend_ctx->backend_devices_filter & (1ULL << device_id)) == 0) + { + device_param->skipped = true; + } + + #if !defined (__APPLE__) + if ((backend_ctx->opencl_device_types_filter & CL_DEVICE_TYPE_GPU) == 0) + { + device_param->skipped = true; + } + #endif + + if ((device_param->opencl_platform_vendor_id == VENDOR_ID_NV) && (device_param->opencl_device_vendor_id == VENDOR_ID_NV)) + { + need_nvml = true; + + #if defined (_WIN) || defined (__CYGWIN__) + need_nvapi = true; + #endif + } + + // CPU burning loop damper + // Value is given as number between 0-100 + // By default 8% + // in theory not needed with HIP + + device_param->spin_damp = (double) user_options->spin_damp / 100; + + // common driver check + + if (device_param->skipped == false) + { + if ((user_options->force == false) && (user_options->backend_info == false)) + { + // HIPDA does not support query nvidia driver version, therefore no driver checks here + // IF needed, could be retrieved using nvmlSystemGetDriverVersion() + + if (device_param->sm_major < 5) + { + if (user_options->quiet == false) event_log_warning (hashcat_ctx, "* Device #%u: This hardware has outdated HIPDA compute capability (%u.%u).", device_id + 1, device_param->sm_major, device_param->sm_minor); + if (user_options->quiet == false) event_log_warning (hashcat_ctx, " For modern OpenCL performance, upgrade to hardware that supports"); + if (user_options->quiet == false) event_log_warning (hashcat_ctx, " HIPDA compute capability version 5.0 (Maxwell) or higher."); + } + + if (device_param->kernel_exec_timeout != 0) + { + if (user_options->quiet == false) event_log_advice (hashcat_ctx, "* Device #%u: WARNING! Kernel exec timeout is not disabled.", device_id + 1); + if (user_options->quiet == false) event_log_advice (hashcat_ctx, " This may cause \"CL_OUT_OF_RESOURCES\" or related errors."); + if (user_options->quiet == false) event_log_advice (hashcat_ctx, " To disable the timeout, see: https://hashcat.net/q/timeoutpatch"); + } + } + + // activate device moved below, at end + } + + // instruction set + + // bcrypt optimization? + //const int rc_cuCtxSetCacheConfig = hc_hipCtxSetCacheConfig (hashcat_ctx, HIP_FUNC_CACHE_PREFER_SHARED); + // + //if (rc_cuCtxSetCacheConfig == -1) return -1; + + // const int sm = (device_param->sm_major * 10) + device_param->sm_minor; + + device_param->has_add = false; + device_param->has_addc = false; + device_param->has_sub = false; + device_param->has_subc = false; + device_param->has_bfe = false; + device_param->has_lop3 = false; + device_param->has_mov64 = false; + device_param->has_prmt = false; + + device_param->has_vadd = false; + device_param->has_vaddc = false; + device_param->has_vadd_co = false; + device_param->has_vaddc_co = false; + device_param->has_vsub = false; + device_param->has_vsubb = false; + device_param->has_vsub_co = false; + device_param->has_vsubb_co = false; + device_param->has_vadd3 = false; + device_param->has_vbfe = false; + device_param->has_vperm = false; + + // device_available_mem + + HIPcontext hip_context; + + if (hc_hipCtxCreate (hashcat_ctx, &hip_context, HIP_CTX_SCHED_BLOCKING_SYNC, device_param->hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + if (hc_hipCtxPushCurrent (hashcat_ctx, hip_context) == -1) + { + device_param->skipped = true; + continue; + } + + size_t free = 0; + size_t total = 0; + + if (hc_hipMemGetInfo (hashcat_ctx, &free, &total) == -1) + { + device_param->skipped = true; + continue; + } + + device_param->device_available_mem = (u64) free; + + if (hc_hipCtxPopCurrent (hashcat_ctx, &hip_context) == -1) + { + device_param->skipped = true; + continue; + } + + if (hc_hipCtxDestroy (hashcat_ctx, hip_context) == -1) + { + device_param->skipped = true; + continue; + } + + /** + * activate device + */ + + if (device_param->skipped == false) hip_devices_active++; + } + } + + backend_ctx->hip_devices_cnt = hip_devices_cnt; + backend_ctx->hip_devices_active = hip_devices_active; + + // OCL + int opencl_devices_cnt = 0; int opencl_devices_active = 0; @@ -6087,8 +8432,8 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) //device_param->opencl_platform = opencl_platform; - device_param->is_cuda = false; - + device_param->is_cuda = false; + device_param->is_hip = false; device_param->is_opencl = true; // store opencl platform i @@ -6829,7 +9174,6 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) device_param->spin_damp = (double) user_options->spin_damp / 100; - if (user_options->stdout_flag == false) { // recommend CUDA @@ -7065,12 +9409,12 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) // all devices combined go into backend_* variables - backend_ctx->backend_devices_cnt = cuda_devices_cnt + opencl_devices_cnt; - backend_ctx->backend_devices_active = cuda_devices_active + opencl_devices_active; + backend_ctx->backend_devices_cnt = cuda_devices_cnt + hip_devices_cnt + opencl_devices_cnt; + backend_ctx->backend_devices_active = cuda_devices_active + hip_devices_active + opencl_devices_active; // find duplicate devices - //if ((cuda_devices_cnt > 0) && (opencl_devices_cnt > 0)) + //if ((cuda_devices_cnt > 0) && (hip_devices_cnt > 0) && (opencl_devices_cnt > 0)) //{ // using force here enables both devices, which is the worst possible outcome // many users force by default, so this is not a good idea @@ -7171,6 +9515,11 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) */ } + if (backend_ctx->hip) + { + // TODO HIP + } + if (backend_ctx->ocl) { for (int backend_devices_cnt = 0; backend_devices_cnt < backend_ctx->backend_devices_cnt; backend_devices_cnt++) @@ -7449,6 +9798,8 @@ void backend_ctx_devices_destroy (hashcat_ctx_t *hashcat_ctx) backend_ctx->backend_devices_active = 0; backend_ctx->cuda_devices_cnt = 0; backend_ctx->cuda_devices_active = 0; + backend_ctx->hip_devices_cnt = 0; + backend_ctx->hip_devices_active = 0; backend_ctx->opencl_devices_cnt = 0; backend_ctx->opencl_devices_active = 0; @@ -7657,6 +10008,62 @@ static int get_cuda_kernel_dynamic_local_mem_size (hashcat_ctx_t *hashcat_ctx, C return 0; } +static int get_hip_kernel_wgs (hashcat_ctx_t *hashcat_ctx, HIPfunction function, u32 *result) +{ + int max_threads_per_block; + + if (hc_hipFuncGetAttribute (hashcat_ctx, &max_threads_per_block, HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, function) == -1) return -1; + + *result = (u32) max_threads_per_block; + + return 0; +} + +static int get_hip_kernel_local_mem_size (hashcat_ctx_t *hashcat_ctx, HIPfunction function, u64 *result) +{ + int shared_size_bytes; + + if (hc_hipFuncGetAttribute (hashcat_ctx, &shared_size_bytes, HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, function) == -1) return -1; + + *result = (u64) shared_size_bytes; + + return 0; +} + +static int get_hip_kernel_dynamic_local_mem_size (hashcat_ctx_t *hashcat_ctx, HIPfunction function, u64 *result) +{ + // AFAIK there's no way to query the maximum value for dynamic shared memory available (because it depends on kernel code). + // let's brute force it, therefore workaround the hashcat wrapper of cuFuncSetAttribute() + + #define MAX_ASSUMED_SHARED (1024 * 1024) + + u64 dynamic_shared_size_bytes = 0; + + for (int i = 1; i <= MAX_ASSUMED_SHARED; i++) + { + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const HIPresult HIP_err = hip->hipFuncSetAttribute (function, HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, i); + + if (HIP_err == HIP_SUCCESS) + { + dynamic_shared_size_bytes = i; + + continue; + } + + break; + } + + *result = dynamic_shared_size_bytes; + + if (hc_hipFuncSetAttribute (hashcat_ctx, function, HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, 0) == -1) return -1; + + return 0; +} + static int get_opencl_kernel_wgs (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_kernel kernel, u32 *result) { size_t work_group_size = 0; @@ -7774,7 +10181,7 @@ static u32 get_kernel_threads (const hc_device_param_t *device_param) return kernel_threads; } -static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const char *kernel_name, char *source_file, char *cached_file, const char *build_options_buf, const bool cache_disable, cl_program *opencl_program, CUmodule *cuda_module) +static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const char *kernel_name, char *source_file, char *cached_file, const char *build_options_buf, const bool cache_disable, cl_program *opencl_program, CUmodule *cuda_module, HIPmodule *hip_module) { const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; const folder_config_t *folder_config = hashcat_ctx->folder_config; @@ -8053,6 +10460,248 @@ static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_p hcfree (binary); } + if (device_param->is_hip == true) + { + hiprtcProgram program; + + if (hc_hiprtcCreateProgram (hashcat_ctx, &program, kernel_sources[0], kernel_name, 0, NULL, NULL) == -1) return false; + + char **hiprtc_options = (char **) hccalloc (6 + strlen (build_options_buf) + 1, sizeof (char *)); // ... + + //hiprtc_options[0] = "--restrict"; + //hiprtc_options[1] = "--device-as-default-execution-space"; + //hiprtc_options[2] = "--gpu-architecture"; + + //hc_asprintf (&hiprtc_options[3], "compute_%d%d", device_param->sm_major, device_param->sm_minor); + + // TODO HIP + + hiprtc_options[0] = ""; + hiprtc_options[1] = ""; + hiprtc_options[2] = ""; + hiprtc_options[3] = ""; + + hiprtc_options[4] = "-I"; + hiprtc_options[5] = folder_config->cpath_real; + + char *hiprtc_options_string = hcstrdup (build_options_buf); + + const int num_options = 6 + hiprtc_make_options_array_from_string (hiprtc_options_string, hiprtc_options + 6); + + const int rc_hiprtcCompileProgram = hc_hiprtcCompileProgram (hashcat_ctx, program, num_options, (const char * const *) hiprtc_options); + + hcfree (hiprtc_options_string); + hcfree (hiprtc_options); + + size_t build_log_size = 0; + + hc_hiprtcGetProgramLogSize (hashcat_ctx, program, &build_log_size); + + #if defined (DEBUG) + if ((build_log_size > 1) || (rc_hiprtcCompileProgram == -1)) + #else + if (rc_hiprtcCompileProgram == -1) + #endif + { + char *build_log = (char *) hcmalloc (build_log_size + 1); + + if (hc_hiprtcGetProgramLog (hashcat_ctx, program, build_log) == -1) + { + hcfree (build_log); + + return false; + } + + build_log[build_log_size] = 0; + + puts (build_log); + + hcfree (build_log); + } + + if (rc_hiprtcCompileProgram == -1) + { + event_log_error (hashcat_ctx, "* Device #%u: Kernel %s build failed.", device_param->device_id + 1, source_file); + + return false; + } + + size_t binary_size = 0; + + if (hc_hiprtcGetCodeSize (hashcat_ctx, program, &binary_size) == -1) return false; + + char *binary = (char *) hcmalloc (binary_size); + + if (hc_hiprtcGetCode (hashcat_ctx, program, binary) == -1) return false; + + if (hc_hiprtcDestroyProgram (hashcat_ctx, &program) == -1) return false; + + #define LOG_SIZE 8192 + + char *mod_info_log = (char *) hcmalloc (LOG_SIZE + 1); + char *mod_error_log = (char *) hcmalloc (LOG_SIZE + 1); + + int mod_cnt = 6; + + HIPjit_option mod_opts[7]; + void *mod_vals[7]; + + mod_opts[0] = HIP_JIT_TARGET_FROM_HIPCONTEXT; + mod_vals[0] = (void *) 0; + + mod_opts[1] = HIP_JIT_LOG_VERBOSE; + mod_vals[1] = (void *) 1; + + mod_opts[2] = HIP_JIT_INFO_LOG_BUFFER; + mod_vals[2] = (void *) mod_info_log; + + mod_opts[3] = HIP_JIT_INFO_LOG_BUFFER_SIZE_BYTES; + mod_vals[3] = (void *) LOG_SIZE; + + mod_opts[4] = HIP_JIT_ERROR_LOG_BUFFER; + mod_vals[4] = (void *) mod_error_log; + + mod_opts[5] = HIP_JIT_ERROR_LOG_BUFFER_SIZE_BYTES; + mod_vals[5] = (void *) LOG_SIZE; + + if (hashconfig->opti_type & OPTI_TYPE_REGISTER_LIMIT) + { + mod_opts[6] = HIP_JIT_MAX_REGISTERS; + mod_vals[6] = (void *) 128; + + mod_cnt++; + } + + #if defined (WITH_HIPBIN) + + char *jit_info_log = (char *) hcmalloc (LOG_SIZE + 1); + char *jit_error_log = (char *) hcmalloc (LOG_SIZE + 1); + + int jit_cnt = 6; + + HIPjit_option jit_opts[7]; + void *jit_vals[7]; + + jit_opts[0] = HIP_JIT_TARGET_FROM_HIPCONTEXT; + jit_vals[0] = (void *) 0; + + jit_opts[1] = HIP_JIT_LOG_VERBOSE; + jit_vals[1] = (void *) 1; + + jit_opts[2] = HIP_JIT_INFO_LOG_BUFFER; + jit_vals[2] = (void *) jit_info_log; + + jit_opts[3] = HIP_JIT_INFO_LOG_BUFFER_SIZE_BYTES; + jit_vals[3] = (void *) LOG_SIZE; + + jit_opts[4] = HIP_JIT_ERROR_LOG_BUFFER; + jit_vals[4] = (void *) jit_error_log; + + jit_opts[5] = HIP_JIT_ERROR_LOG_BUFFER_SIZE_BYTES; + jit_vals[5] = (void *) LOG_SIZE; + + if (hashconfig->opti_type & OPTI_TYPE_REGISTER_LIMIT) + { + jit_opts[6] = HIP_JIT_MAX_REGISTERS; + jit_vals[6] = (void *) 128; + + jit_cnt++; + } + + HIPlinkState state; + + if (hc_cuLinkCreate (hashcat_ctx, jit_cnt, jit_opts, jit_vals, &state) == -1) + { + event_log_error (hashcat_ctx, "* Device #%u: Kernel %s link failed. Error Log:", device_param->device_id + 1, source_file); + event_log_error (hashcat_ctx, "%s", jit_error_log); + event_log_error (hashcat_ctx, NULL); + + return false; + } + + if (hc_cuLinkAddData (hashcat_ctx, state, HIP_JIT_INPUT_CODE, binary, binary_size, kernel_name, 0, NULL, NULL) == -1) + { + event_log_error (hashcat_ctx, "* Device #%u: Kernel %s link failed. Error Log:", device_param->device_id + 1, source_file); + event_log_error (hashcat_ctx, "%s", jit_error_log); + event_log_error (hashcat_ctx, NULL); + + return false; + } + + void *cubin = NULL; + + size_t cubin_size = 0; + + if (hc_cuLinkComplete (hashcat_ctx, state, &cubin, &cubin_size) == -1) + { + event_log_error (hashcat_ctx, "* Device #%u: Kernel %s link failed. Error Log:", device_param->device_id + 1, source_file); + event_log_error (hashcat_ctx, "%s", jit_error_log); + event_log_error (hashcat_ctx, NULL); + + return false; + } + + #if defined (DEBUG) + event_log_info (hashcat_ctx, "* Device #%u: Kernel %s link successful. Info Log:", device_param->device_id + 1, source_file); + event_log_info (hashcat_ctx, "%s", jit_info_log); + event_log_info (hashcat_ctx, NULL); + #endif + + if (hc_cuModuleLoadDataEx (hashcat_ctx, hip_module, cubin, mod_cnt, mod_opts, mod_vals) == -1) + { + event_log_error (hashcat_ctx, "* Device #%u: Kernel %s load failed. Error Log:", device_param->device_id + 1, source_file); + event_log_error (hashcat_ctx, "%s", mod_error_log); + event_log_error (hashcat_ctx, NULL); + + return false; + } + + #if defined (DEBUG) + event_log_info (hashcat_ctx, "* Device #%u: Kernel %s load successful. Info Log:", device_param->device_id + 1, source_file); + event_log_info (hashcat_ctx, "%s", mod_info_log); + event_log_info (hashcat_ctx, NULL); + #endif + + if (cache_disable == false) + { + if (write_kernel_binary (hashcat_ctx, cached_file, cubin, cubin_size) == false) return false; + } + + if (hc_hipLinkDestroy (hashcat_ctx, state) == -1) return false; + + hcfree (jit_info_log); + hcfree (jit_error_log); + + #else + + if (hc_hipModuleLoadDataEx (hashcat_ctx, hip_module, binary, mod_cnt, mod_opts, mod_vals) == -1) + { + event_log_error (hashcat_ctx, "* Device #%u: Kernel %s load failed. Error Log:", device_param->device_id + 1, source_file); + event_log_error (hashcat_ctx, "%s", mod_error_log); + event_log_error (hashcat_ctx, NULL); + + return false; + } + + #if defined (DEBUG) + event_log_info (hashcat_ctx, "* Device #%u: Kernel %s load successful. Info Log:", device_param->device_id + 1, source_file); + event_log_info (hashcat_ctx, "%s", mod_info_log); + event_log_info (hashcat_ctx, NULL); + #endif + + if (cache_disable == false) + { + if (write_kernel_binary (hashcat_ctx, cached_file, binary, binary_size) == false) return false; + } + + #endif + + hcfree (mod_info_log); + hcfree (mod_error_log); + + hcfree (binary); + } + if (device_param->is_opencl == true) { size_t build_log_size = 0; @@ -8185,6 +10834,63 @@ static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_p hcfree (mod_error_log); } + if (device_param->is_hip == true) + { + #define LOG_SIZE 8192 + + char *mod_info_log = (char *) hcmalloc (LOG_SIZE + 1); + char *mod_error_log = (char *) hcmalloc (LOG_SIZE + 1); + + int mod_cnt = 6; + + HIPjit_option mod_opts[7]; + void *mod_vals[7]; + + mod_opts[0] = HIP_JIT_TARGET_FROM_HIPCONTEXT; + mod_vals[0] = (void *) 0; + + mod_opts[1] = HIP_JIT_LOG_VERBOSE; + mod_vals[1] = (void *) 1; + + mod_opts[2] = HIP_JIT_INFO_LOG_BUFFER; + mod_vals[2] = (void *) mod_info_log; + + mod_opts[3] = HIP_JIT_INFO_LOG_BUFFER_SIZE_BYTES; + mod_vals[3] = (void *) LOG_SIZE; + + mod_opts[4] = HIP_JIT_ERROR_LOG_BUFFER; + mod_vals[4] = (void *) mod_error_log; + + mod_opts[5] = HIP_JIT_ERROR_LOG_BUFFER_SIZE_BYTES; + mod_vals[5] = (void *) LOG_SIZE; + + if (hashconfig->opti_type & OPTI_TYPE_REGISTER_LIMIT) + { + mod_opts[6] = HIP_JIT_MAX_REGISTERS; + mod_vals[6] = (void *) 128; + + mod_cnt++; + } + + if (hc_hipModuleLoadDataEx (hashcat_ctx, hip_module, kernel_sources[0], mod_cnt, mod_opts, mod_vals) == -1) + { + event_log_error (hashcat_ctx, "* Device #%u: Kernel %s load failed. Error Log:", device_param->device_id + 1, source_file); + event_log_error (hashcat_ctx, "%s", mod_error_log); + event_log_error (hashcat_ctx, NULL); + + return false; + } + + #if defined (DEBUG) + event_log_info (hashcat_ctx, "* Device #%u: Kernel %s load successful. Info Log:", device_param->device_id + 1, source_file); + event_log_info (hashcat_ctx, "%s", mod_info_log); + event_log_info (hashcat_ctx, NULL); + #endif + + hcfree (mod_info_log); + hcfree (mod_error_log); + } + if (device_param->is_opencl == true) { if (hc_clCreateProgramWithBinary (hashcat_ctx, device_param->opencl_context, 1, &device_param->opencl_device, kernel_lengths, (const unsigned char **) kernel_sources, NULL, opencl_program) == -1) return false; @@ -8243,7 +10949,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if ((unstable_warning == true) && (user_options->force == false)) { event_log_warning (hashcat_ctx, "* Device #%u: Skipping hash-mode %u)", device_id + 1, hashconfig->hash_mode); - event_log_warning (hashcat_ctx, " This is due to a known CUDA/OpenCL runtime/driver issue (not a hashcat issue)"); + event_log_warning (hashcat_ctx, " This is due to a known CUDA/HIP/OpenCL runtime/driver issue (not a hashcat issue)"); event_log_warning (hashcat_ctx, " You can use --force to override, but do not report related errors."); device_param->skipped_warning = true; @@ -8282,6 +10988,13 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) vector_width = 1; } + if (device_param->is_hip == true) + { + // hip does not support this query + + vector_width = 1; + } + if (device_param->is_opencl == true) { if (hc_clGetDeviceInfo (hashcat_ctx, device_param->opencl_device, CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG, sizeof (vector_width), &vector_width, NULL) == -1) @@ -8300,6 +11013,13 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) vector_width = 1; } + if (device_param->is_hip == true) + { + // hip does not support this query + + vector_width = 1; + } + if (device_param->is_opencl == true) { if (hc_clGetDeviceInfo (hashcat_ctx, device_param->opencl_device, CL_DEVICE_NATIVE_VECTOR_WIDTH_INT, sizeof (vector_width), &vector_width, NULL) == -1) @@ -8547,6 +11267,21 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) } } + if (device_param->is_hip == true) + { + if (hc_hipCtxCreate (hashcat_ctx, &device_param->hip_context, HIP_CTX_SCHED_BLOCKING_SYNC, device_param->hip_device) == -1) + { + device_param->skipped = true; + continue; + } + + if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) + { + device_param->skipped = true; + continue; + } + } + if (device_param->is_opencl == true) { /* @@ -8592,6 +11327,19 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) } } + /** + * create stream for HIP devices + */ + + if (device_param->is_hip == true) + { + if (hc_hipStreamCreate (hashcat_ctx, &device_param->hip_stream, HIP_STREAM_DEFAULT) == -1) + { + device_param->skipped = true; + continue; + } + } + /** * create events for CUDA devices */ @@ -8611,6 +11359,25 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) } } + /** + * create events for HIP devices + */ + + if (device_param->is_hip == true) + { + if (hc_hipEventCreate (hashcat_ctx, &device_param->hip_event1, HIP_EVENT_BLOCKING_SYNC) == -1) + { + device_param->skipped = true; + continue; + } + + if (hc_hipEventCreate (hashcat_ctx, &device_param->hip_event2, HIP_EVENT_BLOCKING_SYNC) == -1) + { + device_param->skipped = true; + continue; + } + } + /** * create input buffers on device : calculate size of fixed memory buffers */ @@ -8726,7 +11493,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) int build_options_len = 0; - if (device_param->is_cuda == true) + if ((device_param->is_cuda == true) || (device_param->is_hip == true)) { // using a path with a space will break nvrtc_make_options_array_from_string() // we add it to options array in a clean way later @@ -8789,9 +11556,10 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) char device_name_chksum_amp_mp[HCBUFSIZ_TINY] = { 0 }; - const size_t dnclen_amp_mp = snprintf (device_name_chksum_amp_mp, HCBUFSIZ_TINY, "%d-%d-%d-%u-%s-%s-%s", + const size_t dnclen_amp_mp = snprintf (device_name_chksum_amp_mp, HCBUFSIZ_TINY, "%d-%d-%d-%d-%u-%s-%s-%s", backend_ctx->comptime, backend_ctx->cuda_driver_version, + backend_ctx->hip_driver_version, device_param->is_opencl, device_param->opencl_platform_vendor_id, device_param->device_name, @@ -8861,7 +11629,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) generate_cached_kernel_shared_filename (folder_config->cache_dir, device_name_chksum_amp_mp, cached_file); - const bool rc_load_kernel = load_kernel (hashcat_ctx, device_param, "shared_kernel", source_file, cached_file, build_options_buf, cache_disable, &device_param->opencl_program_shared, &device_param->cuda_module_shared); + const bool rc_load_kernel = load_kernel (hashcat_ctx, device_param, "shared_kernel", source_file, cached_file, build_options_buf, cache_disable, &device_param->opencl_program_shared, &device_param->cuda_module_shared, &device_param->hip_module_shared); if (rc_load_kernel == false) { @@ -8928,6 +11696,64 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) device_param->kernel_preferred_wgs_multiple_utf8toutf16le = device_param->cuda_warp_size; } + if (device_param->is_hip == true) + { + // GPU memset + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_memset, device_param->hip_module_shared, "gpu_memset") == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_memset, &device_param->kernel_wgs_memset) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_memset, &device_param->kernel_local_mem_size_memset) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_memset, &device_param->kernel_dynamic_local_mem_size_memset) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_memset = device_param->hip_warp_size; + + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_memset, 0, sizeof (cl_mem), device_param->kernel_params_memset[0]); if (CL_rc == -1) return -1; + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_memset, 1, sizeof (cl_uint), device_param->kernel_params_memset[1]); if (CL_rc == -1) return -1; + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_memset, 2, sizeof (cl_ulong), device_param->kernel_params_memset[2]); if (CL_rc == -1) return -1; + + // GPU autotune init + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_atinit, device_param->hip_module_shared, "gpu_atinit") == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_atinit, &device_param->kernel_wgs_atinit) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_atinit, &device_param->kernel_local_mem_size_atinit) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_atinit, &device_param->kernel_dynamic_local_mem_size_atinit) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_atinit = device_param->hip_warp_size; + + // CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_atinit, 0, sizeof (cl_mem), device_param->kernel_params_atinit[0]); if (CL_rc == -1) return -1; + // CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_atinit, 1, sizeof (cl_ulong), device_param->kernel_params_atinit[1]); if (CL_rc == -1) return -1; + + // GPU decompress + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_decompress, device_param->hip_module_shared, "gpu_decompress") == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_decompress, &device_param->kernel_wgs_decompress) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_decompress, &device_param->kernel_local_mem_size_decompress) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_decompress, &device_param->kernel_dynamic_local_mem_size_decompress) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_decompress = device_param->hip_warp_size; + + // GPU utf8 to utf16le conversion + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_utf8toutf16le, device_param->hip_module_shared, "gpu_utf8_to_utf16") == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_utf8toutf16le, &device_param->kernel_wgs_utf8toutf16le) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_utf8toutf16le, &device_param->kernel_local_mem_size_utf8toutf16le) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_utf8toutf16le, &device_param->kernel_dynamic_local_mem_size_utf8toutf16le) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_utf8toutf16le = device_param->hip_warp_size; + } + if (device_param->is_opencl == true) { // GPU memset @@ -9029,9 +11855,10 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) const u32 extra_value = (user_options->attack_mode == ATTACK_MODE_ASSOCIATION) ? ATTACK_MODE_ASSOCIATION : ATTACK_MODE_NONE; - const size_t dnclen = snprintf (device_name_chksum, HCBUFSIZ_TINY, "%d-%d-%d-%u-%s-%s-%s-%d-%u-%u-%s", + const size_t dnclen = snprintf (device_name_chksum, HCBUFSIZ_TINY, "%d-%d-%d-%d-%u-%s-%s-%s-%d-%u-%u-%s", backend_ctx->comptime, backend_ctx->cuda_driver_version, + backend_ctx->hip_driver_version, device_param->is_opencl, device_param->opencl_platform_vendor_id, device_param->device_name, @@ -9077,7 +11904,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) * load kernel */ - const bool rc_load_kernel = load_kernel (hashcat_ctx, device_param, "main_kernel", source_file, cached_file, build_options_module_buf, cache_disable, &device_param->opencl_program, &device_param->cuda_module); + const bool rc_load_kernel = load_kernel (hashcat_ctx, device_param, "main_kernel", source_file, cached_file, build_options_module_buf, cache_disable, &device_param->opencl_program, &device_param->cuda_module, &device_param->hip_module); if (rc_load_kernel == false) { @@ -9123,7 +11950,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) generate_cached_kernel_mp_filename (hashconfig->opti_type, hashconfig->opts_type, folder_config->cache_dir, device_name_chksum_amp_mp, cached_file); - const bool rc_load_kernel = load_kernel (hashcat_ctx, device_param, "mp_kernel", source_file, cached_file, build_options_buf, cache_disable, &device_param->opencl_program_mp, &device_param->cuda_module_mp); + const bool rc_load_kernel = load_kernel (hashcat_ctx, device_param, "mp_kernel", source_file, cached_file, build_options_buf, cache_disable, &device_param->opencl_program_mp, &device_param->cuda_module_mp, &device_param->hip_module_mp); if (rc_load_kernel == false) { @@ -9172,7 +11999,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) generate_cached_kernel_amp_filename (user_options_extra->attack_kern, folder_config->cache_dir, device_name_chksum_amp_mp, cached_file); - const bool rc_load_kernel = load_kernel (hashcat_ctx, device_param, "amp_kernel", source_file, cached_file, build_options_buf, cache_disable, &device_param->opencl_program_amp, &device_param->cuda_module_amp); + const bool rc_load_kernel = load_kernel (hashcat_ctx, device_param, "amp_kernel", source_file, cached_file, build_options_buf, cache_disable, &device_param->opencl_program_amp, &device_param->cuda_module_amp, &device_param->hip_module_amp); if (rc_load_kernel == false) { @@ -9354,6 +12181,116 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) } } + if (device_param->is_hip == true) + { + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_bitmap_s1_a, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_bitmap_s1_b, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_bitmap_s1_c, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_bitmap_s1_d, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_bitmap_s2_a, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_bitmap_s2_b, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_bitmap_s2_c, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_bitmap_s2_d, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_plain_bufs, size_plains) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_digests_buf, size_digests) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_digests_shown, size_shown) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_salt_bufs, size_salts) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_result, size_results) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_extra0_buf, size_extra_buffer / 4) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_extra1_buf, size_extra_buffer / 4) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_extra2_buf, size_extra_buffer / 4) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_extra3_buf, size_extra_buffer / 4) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_st_digests_buf, size_st_digests) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_st_salts_buf, size_st_salts) == -1) return -1; + + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s1_a, bitmap_ctx->bitmap_s1_a, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s1_b, bitmap_ctx->bitmap_s1_b, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s1_c, bitmap_ctx->bitmap_s1_c, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s1_d, bitmap_ctx->bitmap_s1_d, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s2_a, bitmap_ctx->bitmap_s2_a, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s2_b, bitmap_ctx->bitmap_s2_b, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s2_c, bitmap_ctx->bitmap_s2_c, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bitmap_s2_d, bitmap_ctx->bitmap_s2_d, bitmap_ctx->bitmap_size) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_digests_buf, hashes->digests_buf, size_digests) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_salt_bufs, hashes->salts_buf, size_salts) == -1) return -1; + + /** + * special buffers + */ + + if (user_options->slow_candidates == true) + { + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_rules_c, size_rules_c) == -1) return -1; + } + else + { + if (user_options_extra->attack_kern == ATTACK_KERN_STRAIGHT) + { + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_rules, size_rules) == -1) return -1; + + if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + size_t dummy = 0; + + if (hc_hipModuleGetGlobal (hashcat_ctx, &device_param->hip_d_rules_c, &dummy, device_param->hip_module, "generic_constant") == -1) return -1; + } + else + { + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_rules_c, size_rules_c) == -1) return -1; + } + + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_rules, straight_ctx->kernel_rules_buf, size_rules) == -1) return -1; + } + else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) + { + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_combs, size_combs) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_combs_c, size_combs) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_root_css_buf, size_root_css) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_markov_css_buf, size_markov_css) == -1) return -1; + } + else if (user_options_extra->attack_kern == ATTACK_KERN_BF) + { + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_bfs, size_bfs) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_root_css_buf, size_root_css) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_markov_css_buf, size_markov_css) == -1) return -1; + + if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + size_t dummy = 0; + + if (hc_hipModuleGetGlobal (hashcat_ctx, &device_param->hip_d_bfs_c, &dummy, device_param->hip_module, "generic_constant") == -1) return -1; + + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_tm_c, size_tm) == -1) return -1; + } + else + { + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_bfs_c, size_bfs) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_tm_c, size_tm) == -1) return -1; + } + } + } + + if (size_esalts) + { + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_esalt_bufs, size_esalts) == -1) return -1; + + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_esalt_bufs, hashes->esalts_buf, size_esalts) == -1) return -1; + } + + if (hashconfig->st_hash != NULL) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_st_digests_buf, hashes->st_digests_buf, size_st_digests) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_st_salts_buf, hashes->st_salts_buf, size_st_salts) == -1) return -1; + + if (size_esalts) + { + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_st_esalts_buf, size_st_esalts) == -1) return -1; + + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_st_esalts_buf, hashes->st_esalts_buf, size_st_esalts) == -1) return -1; + } + } + } + if (device_param->is_opencl == true) { if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_ONLY, bitmap_ctx->bitmap_size, NULL, &device_param->opencl_d_bitmap_s1_a) == -1) return -1; @@ -9488,6 +12425,34 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) device_param->kernel_params[23] = &device_param->cuda_d_extra3_buf; } + if (device_param->is_hip == true) + { + device_param->kernel_params[ 0] = NULL; // &device_param->hip_d_pws_buf; + device_param->kernel_params[ 1] = &device_param->hip_d_rules_c; + device_param->kernel_params[ 2] = &device_param->hip_d_combs_c; + device_param->kernel_params[ 3] = &device_param->hip_d_bfs_c; + device_param->kernel_params[ 4] = NULL; // &device_param->hip_d_tmps; + device_param->kernel_params[ 5] = NULL; // &device_param->hip_d_hooks; + device_param->kernel_params[ 6] = &device_param->hip_d_bitmap_s1_a; + device_param->kernel_params[ 7] = &device_param->hip_d_bitmap_s1_b; + device_param->kernel_params[ 8] = &device_param->hip_d_bitmap_s1_c; + device_param->kernel_params[ 9] = &device_param->hip_d_bitmap_s1_d; + device_param->kernel_params[10] = &device_param->hip_d_bitmap_s2_a; + device_param->kernel_params[11] = &device_param->hip_d_bitmap_s2_b; + device_param->kernel_params[12] = &device_param->hip_d_bitmap_s2_c; + device_param->kernel_params[13] = &device_param->hip_d_bitmap_s2_d; + device_param->kernel_params[14] = &device_param->hip_d_plain_bufs; + device_param->kernel_params[15] = &device_param->hip_d_digests_buf; + device_param->kernel_params[16] = &device_param->hip_d_digests_shown; + device_param->kernel_params[17] = &device_param->hip_d_salt_bufs; + device_param->kernel_params[18] = &device_param->hip_d_esalt_bufs; + device_param->kernel_params[19] = &device_param->hip_d_result; + device_param->kernel_params[20] = &device_param->hip_d_extra0_buf; + device_param->kernel_params[21] = &device_param->hip_d_extra1_buf; + device_param->kernel_params[22] = &device_param->hip_d_extra2_buf; + device_param->kernel_params[23] = &device_param->hip_d_extra3_buf; + } + if (device_param->is_opencl == true) { device_param->kernel_params[ 0] = NULL; // &device_param->opencl_d_pws_buf; @@ -9549,6 +12514,11 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) device_param->kernel_params_mp[0] = &device_param->cuda_d_combs; } + if (device_param->is_hip == true) + { + device_param->kernel_params_mp[0] = &device_param->hip_d_combs; + } + if (device_param->is_opencl == true) { device_param->kernel_params_mp[0] = &device_param->opencl_d_combs; @@ -9563,6 +12533,11 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) device_param->kernel_params_mp[0] = &device_param->cuda_d_combs; } + if (device_param->is_hip == true) + { + device_param->kernel_params_mp[0] = &device_param->hip_d_combs; + } + if (device_param->is_opencl == true) { device_param->kernel_params_mp[0] = &device_param->opencl_d_combs; @@ -9582,6 +12557,12 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) device_param->kernel_params_mp[2] = &device_param->cuda_d_markov_css_buf; } + if (device_param->is_hip == true) + { + device_param->kernel_params_mp[1] = &device_param->hip_d_root_css_buf; + device_param->kernel_params_mp[2] = &device_param->hip_d_markov_css_buf; + } + if (device_param->is_opencl == true) { device_param->kernel_params_mp[1] = &device_param->opencl_d_root_css_buf; @@ -9606,12 +12587,19 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) device_param->kernel_params_mp_l[0] = NULL; // (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) // ? &device_param->opencl_d_pws_buf // : &device_param->opencl_d_pws_amp_buf; + if (device_param->is_cuda == true) { device_param->kernel_params_mp_l[1] = &device_param->cuda_d_root_css_buf; device_param->kernel_params_mp_l[2] = &device_param->cuda_d_markov_css_buf; } + if (device_param->is_hip == true) + { + device_param->kernel_params_mp_l[1] = &device_param->hip_d_root_css_buf; + device_param->kernel_params_mp_l[2] = &device_param->hip_d_markov_css_buf; + } + if (device_param->is_opencl == true) { device_param->kernel_params_mp_l[1] = &device_param->opencl_d_root_css_buf; @@ -9640,6 +12628,13 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) device_param->kernel_params_mp_r[2] = &device_param->cuda_d_markov_css_buf; } + if (device_param->is_hip == true) + { + device_param->kernel_params_mp_r[0] = &device_param->hip_d_bfs; + device_param->kernel_params_mp_r[1] = &device_param->hip_d_root_css_buf; + device_param->kernel_params_mp_r[2] = &device_param->hip_d_markov_css_buf; + } + if (device_param->is_opencl == true) { device_param->kernel_params_mp_r[0] = &device_param->opencl_d_bfs; @@ -9666,6 +12661,15 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) device_param->kernel_params_amp[4] = &device_param->cuda_d_bfs_c; } + if (device_param->is_hip == true) + { + device_param->kernel_params_amp[0] = NULL; // &device_param->hip_d_pws_buf; + device_param->kernel_params_amp[1] = NULL; // &device_param->hip_d_pws_amp_buf; + device_param->kernel_params_amp[2] = &device_param->hip_d_rules_c; + device_param->kernel_params_amp[3] = &device_param->hip_d_combs_c; + device_param->kernel_params_amp[4] = &device_param->hip_d_bfs_c; + } + if (device_param->is_opencl == true) { device_param->kernel_params_amp[0] = NULL; // &device_param->opencl_d_pws_buf; @@ -9684,6 +12688,12 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) device_param->kernel_params_tm[1] = &device_param->cuda_d_tm_c; } + if (device_param->is_hip == true) + { + device_param->kernel_params_tm[0] = &device_param->hip_d_bfs_c; + device_param->kernel_params_tm[1] = &device_param->hip_d_tm_c; + } + if (device_param->is_opencl == true) { device_param->kernel_params_tm[0] = &device_param->opencl_d_bfs_c; @@ -9719,6 +12729,15 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) // : &device_param->cuda_d_pws_amp_buf; } + if (device_param->is_hip == true) + { + device_param->kernel_params_decompress[0] = NULL; // &device_param->hip_d_pws_idx; + device_param->kernel_params_decompress[1] = NULL; // &device_param->hip_d_pws_comp_buf; + device_param->kernel_params_decompress[2] = NULL; // (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + // ? &device_param->hip_d_pws_buf + // : &device_param->hip_d_pws_amp_buf; + } + if (device_param->is_opencl == true) { device_param->kernel_params_decompress[0] = NULL; // &device_param->opencl_d_pws_idx; @@ -10336,6 +13355,608 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) } } + if (device_param->is_hip == true) + { + char kernel_name[64] = { 0 }; + + if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + if (hashconfig->opti_type & OPTI_TYPE_SINGLE_HASH) + { + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) + { + // kernel1 + + snprintf (kernel_name, sizeof (kernel_name), "m%05u_s%02d", kern_type, 4); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function1, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function1, &device_param->kernel_wgs1) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_local_mem_size1) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_dynamic_local_mem_size1) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple1 = device_param->hip_warp_size; + + // kernel2 + + snprintf (kernel_name, sizeof (kernel_name), "m%05u_s%02d", kern_type, 8); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function2, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function2, &device_param->kernel_wgs2) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_local_mem_size2) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_dynamic_local_mem_size2) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple2 = device_param->hip_warp_size; + + // kernel3 + + snprintf (kernel_name, sizeof (kernel_name), "m%05u_s%02d", kern_type, 16); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function3, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function3, &device_param->kernel_wgs3) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_local_mem_size3) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_dynamic_local_mem_size3) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple3 = device_param->hip_warp_size; + } + else + { + snprintf (kernel_name, sizeof (kernel_name), "m%05u_sxx", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function4, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function4, &device_param->kernel_wgs4) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function4, &device_param->kernel_local_mem_size4) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function4, &device_param->kernel_dynamic_local_mem_size4) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple4 = device_param->hip_warp_size; + } + } + else + { + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) + { + // kernel1 + + snprintf (kernel_name, sizeof (kernel_name), "m%05u_m%02d", kern_type, 4); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function1, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function1, &device_param->kernel_wgs1) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_local_mem_size1) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_dynamic_local_mem_size1) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple1 = device_param->hip_warp_size; + + // kernel2 + + snprintf (kernel_name, sizeof (kernel_name), "m%05u_m%02d", kern_type, 8); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function2, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function2, &device_param->kernel_wgs2) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_local_mem_size2) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_dynamic_local_mem_size2) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple2 = device_param->hip_warp_size; + + // kernel3 + + snprintf (kernel_name, sizeof (kernel_name), "m%05u_m%02d", kern_type, 16); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function3, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function3, &device_param->kernel_wgs3) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_local_mem_size3) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_dynamic_local_mem_size3) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple3 = device_param->hip_warp_size; + } + else + { + snprintf (kernel_name, sizeof (kernel_name), "m%05u_mxx", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function4, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function4, &device_param->kernel_wgs4) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function4, &device_param->kernel_local_mem_size4) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function4, &device_param->kernel_dynamic_local_mem_size4) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple4 = device_param->hip_warp_size; + } + } + + if (user_options->slow_candidates == true) + { + } + else + { + if (user_options->attack_mode == ATTACK_MODE_BF) + { + if (hashconfig->opts_type & OPTS_TYPE_TM_KERNEL) + { + snprintf (kernel_name, sizeof (kernel_name), "m%05u_tm", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_tm, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_tm, &device_param->kernel_wgs_tm) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_tm, &device_param->kernel_local_mem_size_tm) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_tm, &device_param->kernel_dynamic_local_mem_size_tm) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_tm = device_param->hip_warp_size; + } + } + } + } + else + { + // kernel1 + + snprintf (kernel_name, sizeof (kernel_name), "m%05u_init", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function1, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function1, &device_param->kernel_wgs1) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_local_mem_size1) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function1, &device_param->kernel_dynamic_local_mem_size1) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple1 = device_param->hip_warp_size; + + // kernel2 + + snprintf (kernel_name, sizeof (kernel_name), "m%05u_loop", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function2, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function2, &device_param->kernel_wgs2) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_local_mem_size2) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2, &device_param->kernel_dynamic_local_mem_size2) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple2 = device_param->hip_warp_size; + + // kernel3 + + snprintf (kernel_name, sizeof (kernel_name), "m%05u_comp", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function3, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function3, &device_param->kernel_wgs3) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_local_mem_size3) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function3, &device_param->kernel_dynamic_local_mem_size3) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple3 = device_param->hip_warp_size; + + if (hashconfig->opts_type & OPTS_TYPE_LOOP_PREPARE) + { + // kernel2p + + snprintf (kernel_name, sizeof (kernel_name), "m%05u_loop_prepare", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function2p, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function2p, &device_param->kernel_wgs2p) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function2p, &device_param->kernel_local_mem_size2p) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2p, &device_param->kernel_dynamic_local_mem_size2p) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple2p = device_param->hip_warp_size; + } + + if (hashconfig->opts_type & OPTS_TYPE_LOOP_EXTENDED) + { + // kernel2e + + snprintf (kernel_name, sizeof (kernel_name), "m%05u_loop_extended", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function2e, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function2e, &device_param->kernel_wgs2e) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function2e, &device_param->kernel_local_mem_size2e) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function2e, &device_param->kernel_dynamic_local_mem_size2e) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple2e = device_param->hip_warp_size; + } + + // kernel12 + + if (hashconfig->opts_type & OPTS_TYPE_HOOK12) + { + snprintf (kernel_name, sizeof (kernel_name), "m%05u_hook12", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function12, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function12, &device_param->kernel_wgs12) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function12, &device_param->kernel_local_mem_size12) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function12, &device_param->kernel_dynamic_local_mem_size12) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple12 = device_param->hip_warp_size; + } + + // kernel23 + + if (hashconfig->opts_type & OPTS_TYPE_HOOK23) + { + snprintf (kernel_name, sizeof (kernel_name), "m%05u_hook23", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function23, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function23, &device_param->kernel_wgs23) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function23, &device_param->kernel_local_mem_size23) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function23, &device_param->kernel_dynamic_local_mem_size23) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple23 = device_param->hip_warp_size; + } + + // init2 + + if (hashconfig->opts_type & OPTS_TYPE_INIT2) + { + snprintf (kernel_name, sizeof (kernel_name), "m%05u_init2", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_init2, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_init2, &device_param->kernel_wgs_init2) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_init2, &device_param->kernel_local_mem_size_init2) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_init2, &device_param->kernel_dynamic_local_mem_size_init2) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_init2 = device_param->hip_warp_size; + } + + // loop2 prepare + + if (hashconfig->opts_type & OPTS_TYPE_LOOP2_PREPARE) + { + snprintf (kernel_name, sizeof (kernel_name), "m%05u_loop2_prepare", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_loop2p, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_loop2p, &device_param->kernel_wgs_loop2p) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_loop2p, &device_param->kernel_local_mem_size_loop2p) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_loop2p, &device_param->kernel_dynamic_local_mem_size_loop2p) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_loop2p = device_param->hip_warp_size; + } + + // loop2 + + if (hashconfig->opts_type & OPTS_TYPE_LOOP2) + { + snprintf (kernel_name, sizeof (kernel_name), "m%05u_loop2", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_loop2, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_loop2, &device_param->kernel_wgs_loop2) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_loop2, &device_param->kernel_local_mem_size_loop2) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_loop2, &device_param->kernel_dynamic_local_mem_size_loop2) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_loop2 = device_param->hip_warp_size; + } + + // aux1 + + if (hashconfig->opts_type & OPTS_TYPE_AUX1) + { + snprintf (kernel_name, sizeof (kernel_name), "m%05u_aux1", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_aux1, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_aux1, &device_param->kernel_wgs_aux1) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_aux1, &device_param->kernel_local_mem_size_aux1) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_aux1, &device_param->kernel_dynamic_local_mem_size_aux1) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_aux1 = device_param->hip_warp_size; + } + + // aux2 + + if (hashconfig->opts_type & OPTS_TYPE_AUX2) + { + snprintf (kernel_name, sizeof (kernel_name), "m%05u_aux2", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_aux2, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_aux2, &device_param->kernel_wgs_aux2) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_aux2, &device_param->kernel_local_mem_size_aux2) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_aux2, &device_param->kernel_dynamic_local_mem_size_aux2) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_aux2 = device_param->hip_warp_size; + } + + // aux3 + + if (hashconfig->opts_type & OPTS_TYPE_AUX3) + { + snprintf (kernel_name, sizeof (kernel_name), "m%05u_aux3", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_aux3, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_aux3, &device_param->kernel_wgs_aux3) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_aux3, &device_param->kernel_local_mem_size_aux3) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_aux3, &device_param->kernel_dynamic_local_mem_size_aux3) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_aux3 = device_param->hip_warp_size; + } + + // aux4 + + if (hashconfig->opts_type & OPTS_TYPE_AUX4) + { + snprintf (kernel_name, sizeof (kernel_name), "m%05u_aux4", kern_type); + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_aux4, device_param->hip_module, kernel_name) == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_aux4, &device_param->kernel_wgs_aux4) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_aux4, &device_param->kernel_local_mem_size_aux4) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_aux4, &device_param->kernel_dynamic_local_mem_size_aux4) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_aux4 = device_param->hip_warp_size; + } + } + + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_decompress, 0, sizeof (cl_mem), device_param->kernel_params_decompress[0]); if (CL_rc == -1) return -1; + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_decompress, 1, sizeof (cl_mem), device_param->kernel_params_decompress[1]); if (CL_rc == -1) return -1; + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_decompress, 2, sizeof (cl_mem), device_param->kernel_params_decompress[2]); if (CL_rc == -1) return -1; + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_decompress, 3, sizeof (cl_ulong), device_param->kernel_params_decompress[3]); if (CL_rc == -1) return -1; + + // MP start + + if (user_options->slow_candidates == true) + { + } + else + { + if (user_options->attack_mode == ATTACK_MODE_BF) + { + // mp_l + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_mp_l, device_param->hip_module_mp, "l_markov") == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_mp_l, &device_param->kernel_wgs_mp_l) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_mp_l, &device_param->kernel_local_mem_size_mp_l) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_mp_l, &device_param->kernel_dynamic_local_mem_size_mp_l) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_mp_l = device_param->hip_warp_size; + + // mp_r + + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_mp_r, device_param->hip_module_mp, "r_markov") == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_mp_r, &device_param->kernel_wgs_mp_r) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_mp_r, &device_param->kernel_local_mem_size_mp_r) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_mp_r, &device_param->kernel_dynamic_local_mem_size_mp_r) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_mp_r = device_param->hip_warp_size; + + if (user_options->attack_mode == ATTACK_MODE_BF) + { + if (hashconfig->opts_type & OPTS_TYPE_TM_KERNEL) + { + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_tm, 0, sizeof (cl_mem), device_param->kernel_params_tm[0]); if (CL_rc == -1) return -1; + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_tm, 1, sizeof (cl_mem), device_param->kernel_params_tm[1]); if (CL_rc == -1) return -1; + } + } + } + else if (user_options->attack_mode == ATTACK_MODE_HYBRID1) + { + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_mp, device_param->hip_module_mp, "C_markov") == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_mp, &device_param->kernel_wgs_mp) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_mp, &device_param->kernel_local_mem_size_mp) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_mp, &device_param->kernel_dynamic_local_mem_size_mp) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_mp = device_param->hip_warp_size; + } + else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) + { + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_mp, device_param->hip_module_mp, "C_markov") == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_mp, &device_param->kernel_wgs_mp) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_mp, &device_param->kernel_local_mem_size_mp) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_mp, &device_param->kernel_dynamic_local_mem_size_mp) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_mp = device_param->hip_warp_size; + } + } + + if (user_options->slow_candidates == true) + { + } + else + { + if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + // nothing to do + } + else + { + if (hc_hipModuleGetFunction (hashcat_ctx, &device_param->hip_function_amp, device_param->hip_module_amp, "amp") == -1) return -1; + + if (get_hip_kernel_wgs (hashcat_ctx, device_param->hip_function_amp, &device_param->kernel_wgs_amp) == -1) return -1; + + if (get_hip_kernel_local_mem_size (hashcat_ctx, device_param->hip_function_amp, &device_param->kernel_local_mem_size_amp) == -1) return -1; + + if (get_hip_kernel_dynamic_local_mem_size (hashcat_ctx, device_param->hip_function_amp, &device_param->kernel_dynamic_local_mem_size_amp) == -1) return -1; + + device_param->kernel_preferred_wgs_multiple_amp = device_param->hip_warp_size; + } + + /* + if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + // nothing to do + } + else + { + for (u32 i = 0; i < 5; i++) + { + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_amp, i, sizeof (cl_mem), device_param->kernel_params_amp[i]); + + //if (CL_rc == -1) return -1; + } + + for (u32 i = 5; i < 6; i++) + { + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_amp, i, sizeof (cl_uint), device_param->kernel_params_amp[i]); + + //if (CL_rc == -1) return -1; + } + + for (u32 i = 6; i < 7; i++) + { + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_amp, i, sizeof (cl_ulong), device_param->kernel_params_amp[i]); + + //if (CL_rc == -1) return -1; + } + } + */ + } + + // zero some data buffers + + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_plain_bufs, device_param->size_plains) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_digests_shown, device_param->size_shown) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_result, device_param->size_results) == -1) return -1; + + /** + * special buffers + */ + + if (user_options->slow_candidates == true) + { + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_rules_c, size_rules_c) == -1) return -1; + } + else + { + if (user_options_extra->attack_kern == ATTACK_KERN_STRAIGHT) + { + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_rules_c, size_rules_c) == -1) return -1; + } + else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) + { + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_combs, size_combs) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_combs_c, size_combs) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_root_css_buf, size_root_css) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_markov_css_buf, size_markov_css) == -1) return -1; + } + else if (user_options_extra->attack_kern == ATTACK_KERN_BF) + { + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_bfs, size_bfs) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_bfs_c, size_bfs) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_tm_c, size_tm) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_root_css_buf, size_root_css) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_markov_css_buf, size_markov_css) == -1) return -1; + } + } + + if (user_options->slow_candidates == true) + { + } + else + { + if ((user_options->attack_mode == ATTACK_MODE_HYBRID1) || (user_options->attack_mode == ATTACK_MODE_HYBRID2)) + { + /** + * prepare mp + */ + + if (user_options->attack_mode == ATTACK_MODE_HYBRID1) + { + device_param->kernel_params_mp_buf32[5] = 0; + device_param->kernel_params_mp_buf32[6] = 0; + device_param->kernel_params_mp_buf32[7] = 0; + + if (hashconfig->opts_type & OPTS_TYPE_PT_ADD01) device_param->kernel_params_mp_buf32[5] = full01; + if (hashconfig->opts_type & OPTS_TYPE_PT_ADD06) device_param->kernel_params_mp_buf32[5] = full06; + if (hashconfig->opts_type & OPTS_TYPE_PT_ADD80) device_param->kernel_params_mp_buf32[5] = full80; + if (hashconfig->opts_type & OPTS_TYPE_PT_ADDBITS14) device_param->kernel_params_mp_buf32[6] = 1; + if (hashconfig->opts_type & OPTS_TYPE_PT_ADDBITS15) device_param->kernel_params_mp_buf32[7] = 1; + } + else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) + { + device_param->kernel_params_mp_buf32[5] = 0; + device_param->kernel_params_mp_buf32[6] = 0; + device_param->kernel_params_mp_buf32[7] = 0; + } + + //for (u32 i = 0; i < 3; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp, i, sizeof (cl_mem), device_param->kernel_params_mp[i]); if (CL_rc == -1) return -1; } + } + else if (user_options->attack_mode == ATTACK_MODE_BF) + { + /** + * prepare mp_r and mp_l + */ + + device_param->kernel_params_mp_l_buf32[6] = 0; + device_param->kernel_params_mp_l_buf32[7] = 0; + device_param->kernel_params_mp_l_buf32[8] = 0; + + if (hashconfig->opts_type & OPTS_TYPE_PT_ADD01) device_param->kernel_params_mp_l_buf32[6] = full01; + if (hashconfig->opts_type & OPTS_TYPE_PT_ADD06) device_param->kernel_params_mp_l_buf32[6] = full06; + if (hashconfig->opts_type & OPTS_TYPE_PT_ADD80) device_param->kernel_params_mp_l_buf32[6] = full80; + if (hashconfig->opts_type & OPTS_TYPE_PT_ADDBITS14) device_param->kernel_params_mp_l_buf32[7] = 1; + if (hashconfig->opts_type & OPTS_TYPE_PT_ADDBITS15) device_param->kernel_params_mp_l_buf32[8] = 1; + + //for (u32 i = 0; i < 3; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_l, i, sizeof (cl_mem), device_param->kernel_params_mp_l[i]); if (CL_rc == -1) return -1; } + //for (u32 i = 0; i < 3; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_r, i, sizeof (cl_mem), device_param->kernel_params_mp_r[i]); if (CL_rc == -1) return -1; } + } + } + } + if (device_param->is_opencl == true) { // GPU memset @@ -11208,6 +14829,23 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_hooks, device_param->size_hooks) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_pws_buf, size_pws) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_pws_amp_buf, size_pws_amp) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_pws_comp_buf, size_pws_comp) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_pws_idx, size_pws_idx) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_tmps, size_tmps) == -1) return -1; + if (hc_hipMemAlloc (hashcat_ctx, &device_param->hip_d_hooks, size_hooks) == -1) return -1; + + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_pws_buf, device_param->size_pws) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_pws_amp_buf, device_param->size_pws_amp) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_pws_comp_buf, device_param->size_pws_comp) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_pws_idx, device_param->size_pws_idx) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_tmps, device_param->size_tmps) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_hooks, device_param->size_hooks) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clCreateBuffer (hashcat_ctx, device_param->opencl_context, CL_MEM_READ_WRITE, size_pws, NULL, &device_param->opencl_d_pws_buf) == -1) return -1; @@ -11279,6 +14917,13 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) device_param->kernel_params[ 5] = &device_param->cuda_d_hooks; } + if (device_param->is_hip == true) + { + device_param->kernel_params[ 0] = &device_param->hip_d_pws_buf; + device_param->kernel_params[ 4] = &device_param->hip_d_tmps; + device_param->kernel_params[ 5] = &device_param->hip_d_hooks; + } + if (device_param->is_opencl == true) { device_param->kernel_params[ 0] = &device_param->opencl_d_pws_buf; @@ -11308,6 +14953,15 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp, 0, sizeof (cl_mem), device_param->kernel_params_mp[0]); if (CL_rc == -1) return -1; } + if (device_param->is_hip == true) + { + device_param->kernel_params_mp[0] = (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + ? &device_param->hip_d_pws_buf + : &device_param->hip_d_pws_amp_buf; + + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp, 0, sizeof (cl_mem), device_param->kernel_params_mp[0]); if (CL_rc == -1) return -1; + } + if (device_param->is_opencl == true) { device_param->kernel_params_mp[0] = (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) @@ -11330,6 +14984,15 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_l, 0, sizeof (cl_mem), device_param->kernel_params_mp_l[0]); if (CL_rc == -1) return -1; } + if (device_param->is_hip == true) + { + device_param->kernel_params_mp_l[0] = (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + ? &device_param->hip_d_pws_buf + : &device_param->hip_d_pws_amp_buf; + + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_l, 0, sizeof (cl_mem), device_param->kernel_params_mp_l[0]); if (CL_rc == -1) return -1; + } + if (device_param->is_opencl == true) { device_param->kernel_params_mp_l[0] = (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) @@ -11355,6 +15018,15 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_amp, 1, sizeof (cl_mem), device_param->kernel_params_amp[1]); if (CL_rc == -1) return -1; } + if (device_param->is_hip == true) + { + device_param->kernel_params_amp[0] = &device_param->hip_d_pws_buf; + device_param->kernel_params_amp[1] = &device_param->hip_d_pws_amp_buf; + + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_amp, 0, sizeof (cl_mem), device_param->kernel_params_amp[0]); if (CL_rc == -1) return -1; + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_amp, 1, sizeof (cl_mem), device_param->kernel_params_amp[1]); if (CL_rc == -1) return -1; + } + if (device_param->is_opencl == true) { device_param->kernel_params_amp[0] = &device_param->opencl_d_pws_buf; @@ -11379,6 +15051,19 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_decompress, 2, sizeof (cl_mem), device_param->kernel_params_decompress[2]); if (CL_rc == -1) return -1; } + if (device_param->is_hip == true) + { + device_param->kernel_params_decompress[0] = &device_param->hip_d_pws_idx; + device_param->kernel_params_decompress[1] = &device_param->hip_d_pws_comp_buf; + device_param->kernel_params_decompress[2] = (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + ? &device_param->hip_d_pws_buf + : &device_param->hip_d_pws_amp_buf; + + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_decompress, 0, sizeof (cl_mem), device_param->kernel_params_decompress[0]); if (CL_rc == -1) return -1; + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_decompress, 1, sizeof (cl_mem), device_param->kernel_params_decompress[1]); if (CL_rc == -1) return -1; + //CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_decompress, 2, sizeof (cl_mem), device_param->kernel_params_decompress[2]); if (CL_rc == -1) return -1; + } + if (device_param->is_opencl == true) { device_param->kernel_params_decompress[0] = &device_param->opencl_d_pws_idx; @@ -11392,6 +15077,17 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_decompress, 2, sizeof (cl_mem), device_param->kernel_params_decompress[2]) == -1) return -1; } + // context + + if (device_param->is_hip == true) + { + if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) + { + device_param->skipped = true; + continue; + } + } + hardware_power_all += device_param->hardware_power; EVENT_DATA (EVENT_BACKEND_DEVICE_INIT_POST, &backend_devices_idx, sizeof (int)); @@ -11554,6 +15250,127 @@ void backend_session_destroy (hashcat_ctx_t *hashcat_ctx) device_param->cuda_context = NULL; } + if (device_param->is_hip == true) + { + if (device_param->hip_d_pws_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_pws_buf); + if (device_param->hip_d_pws_amp_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_pws_amp_buf); + if (device_param->hip_d_pws_comp_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_pws_comp_buf); + if (device_param->hip_d_pws_idx) hc_hipMemFree (hashcat_ctx, device_param->hip_d_pws_idx); + if (device_param->hip_d_rules) hc_hipMemFree (hashcat_ctx, device_param->hip_d_rules); + //if (device_param->hip_d_rules_c) hc_hipMemFree (hashcat_ctx, device_param->hip_d_rules_c); + if (device_param->hip_d_combs) hc_hipMemFree (hashcat_ctx, device_param->hip_d_combs); + if (device_param->hip_d_combs_c) hc_hipMemFree (hashcat_ctx, device_param->hip_d_combs_c); + if (device_param->hip_d_bfs) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bfs); + //if (device_param->hip_d_bfs_c) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bfs_c); + if (device_param->hip_d_bitmap_s1_a) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s1_a); + if (device_param->hip_d_bitmap_s1_b) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s1_b); + if (device_param->hip_d_bitmap_s1_c) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s1_c); + if (device_param->hip_d_bitmap_s1_d) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s1_d); + if (device_param->hip_d_bitmap_s2_a) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s2_a); + if (device_param->hip_d_bitmap_s2_b) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s2_b); + if (device_param->hip_d_bitmap_s2_c) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s2_c); + if (device_param->hip_d_bitmap_s2_d) hc_hipMemFree (hashcat_ctx, device_param->hip_d_bitmap_s2_d); + if (device_param->hip_d_plain_bufs) hc_hipMemFree (hashcat_ctx, device_param->hip_d_plain_bufs); + if (device_param->hip_d_digests_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_digests_buf); + if (device_param->hip_d_digests_shown) hc_hipMemFree (hashcat_ctx, device_param->hip_d_digests_shown); + if (device_param->hip_d_salt_bufs) hc_hipMemFree (hashcat_ctx, device_param->hip_d_salt_bufs); + if (device_param->hip_d_esalt_bufs) hc_hipMemFree (hashcat_ctx, device_param->hip_d_esalt_bufs); + if (device_param->hip_d_tmps) hc_hipMemFree (hashcat_ctx, device_param->hip_d_tmps); + if (device_param->hip_d_hooks) hc_hipMemFree (hashcat_ctx, device_param->hip_d_hooks); + if (device_param->hip_d_result) hc_hipMemFree (hashcat_ctx, device_param->hip_d_result); + if (device_param->hip_d_extra0_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_extra0_buf); + if (device_param->hip_d_extra1_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_extra1_buf); + if (device_param->hip_d_extra2_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_extra2_buf); + if (device_param->hip_d_extra3_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_extra3_buf); + if (device_param->hip_d_root_css_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_root_css_buf); + if (device_param->hip_d_markov_css_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_markov_css_buf); + if (device_param->hip_d_tm_c) hc_hipMemFree (hashcat_ctx, device_param->hip_d_tm_c); + if (device_param->hip_d_st_digests_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_st_digests_buf); + if (device_param->hip_d_st_salts_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_st_salts_buf); + if (device_param->hip_d_st_esalts_buf) hc_hipMemFree (hashcat_ctx, device_param->hip_d_st_esalts_buf); + + if (device_param->hip_event1) hc_hipEventDestroy (hashcat_ctx, device_param->hip_event1); + if (device_param->hip_event2) hc_hipEventDestroy (hashcat_ctx, device_param->hip_event2); + + if (device_param->hip_stream) hc_hipStreamDestroy (hashcat_ctx, device_param->hip_stream); + + if (device_param->hip_module) hc_hipModuleUnload (hashcat_ctx, device_param->hip_module); + if (device_param->hip_module_mp) hc_hipModuleUnload (hashcat_ctx, device_param->hip_module_mp); + if (device_param->hip_module_amp) hc_hipModuleUnload (hashcat_ctx, device_param->hip_module_amp); + if (device_param->hip_module_shared) hc_hipModuleUnload (hashcat_ctx, device_param->hip_module_shared); + + if (device_param->hip_context) hc_hipCtxDestroy (hashcat_ctx, device_param->hip_context); + + device_param->hip_d_pws_buf = 0; + device_param->hip_d_pws_amp_buf = 0; + device_param->hip_d_pws_comp_buf = 0; + device_param->hip_d_pws_idx = 0; + device_param->hip_d_rules = 0; + device_param->hip_d_rules_c = 0; + device_param->hip_d_combs = 0; + device_param->hip_d_combs_c = 0; + device_param->hip_d_bfs = 0; + device_param->hip_d_bfs_c = 0; + device_param->hip_d_bitmap_s1_a = 0; + device_param->hip_d_bitmap_s1_b = 0; + device_param->hip_d_bitmap_s1_c = 0; + device_param->hip_d_bitmap_s1_d = 0; + device_param->hip_d_bitmap_s2_a = 0; + device_param->hip_d_bitmap_s2_b = 0; + device_param->hip_d_bitmap_s2_c = 0; + device_param->hip_d_bitmap_s2_d = 0; + device_param->hip_d_plain_bufs = 0; + device_param->hip_d_digests_buf = 0; + device_param->hip_d_digests_shown = 0; + device_param->hip_d_salt_bufs = 0; + device_param->hip_d_esalt_bufs = 0; + device_param->hip_d_tmps = 0; + device_param->hip_d_hooks = 0; + device_param->hip_d_result = 0; + device_param->hip_d_extra0_buf = 0; + device_param->hip_d_extra1_buf = 0; + device_param->hip_d_extra2_buf = 0; + device_param->hip_d_extra3_buf = 0; + device_param->hip_d_root_css_buf = 0; + device_param->hip_d_markov_css_buf = 0; + device_param->hip_d_tm_c = 0; + device_param->hip_d_st_digests_buf = 0; + device_param->hip_d_st_salts_buf = 0; + device_param->hip_d_st_esalts_buf = 0; + + device_param->hip_function1 = NULL; + device_param->hip_function12 = NULL; + device_param->hip_function2p = NULL; + device_param->hip_function2 = NULL; + device_param->hip_function2e = NULL; + device_param->hip_function23 = NULL; + device_param->hip_function3 = NULL; + device_param->hip_function4 = NULL; + device_param->hip_function_init2 = NULL; + device_param->hip_function_loop2p = NULL; + device_param->hip_function_loop2 = NULL; + device_param->hip_function_mp = NULL; + device_param->hip_function_mp_l = NULL; + device_param->hip_function_mp_r = NULL; + device_param->hip_function_tm = NULL; + device_param->hip_function_amp = NULL; + device_param->hip_function_memset = NULL; + device_param->hip_function_atinit = NULL; + device_param->hip_function_utf8toutf16le = NULL; + device_param->hip_function_decompress = NULL; + device_param->hip_function_aux1 = NULL; + device_param->hip_function_aux2 = NULL; + device_param->hip_function_aux3 = NULL; + device_param->hip_function_aux4 = NULL; + + device_param->hip_module = NULL; + device_param->hip_module_mp = NULL; + device_param->hip_module_amp = NULL; + device_param->hip_module_shared = NULL; + + device_param->hip_context = NULL; + } + if (device_param->is_opencl == true) { if (device_param->opencl_d_pws_buf) hc_clReleaseMemObject (hashcat_ctx, device_param->opencl_d_pws_buf); @@ -11849,6 +15666,15 @@ int backend_session_update_mp (hashcat_ctx_t *hashcat_ctx) if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css) == -1) return -1; } + if (device_param->is_hip == true) + { + //for (u32 i = 3; i < 4; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp, i, sizeof (cl_ulong), device_param->kernel_params_mp[i]); if (CL_rc == -1) return -1; } + //for (u32 i = 4; i < 8; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp, i, sizeof (cl_uint), device_param->kernel_params_mp[i]); if (CL_rc == -1) return -1; } + + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css) == -1) return -1; + } + if (device_param->is_opencl == true) { for (u32 i = 3; i < 4; i++) { if (hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp, i, sizeof (cl_ulong), device_param->kernel_params_mp[i]) == -1) return -1; } @@ -11901,6 +15727,20 @@ int backend_session_update_mp_rl (hashcat_ctx_t *hashcat_ctx, const u32 css_cnt_ if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css) == -1) return -1; } + if (device_param->is_hip == true) + { + //for (u32 i = 3; i < 4; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_l, i, sizeof (cl_ulong), device_param->kernel_params_mp_l[i]); if (CL_rc == -1) return -1; } + //for (u32 i = 4; i < 8; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_l, i, sizeof (cl_uint), device_param->kernel_params_mp_l[i]); if (CL_rc == -1) return -1; } + //for (u32 i = 9; i < 9; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_l, i, sizeof (cl_ulong), device_param->kernel_params_mp_l[i]); if (CL_rc == -1) return -1; } + + //for (u32 i = 3; i < 4; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_r, i, sizeof (cl_ulong), device_param->kernel_params_mp_r[i]); if (CL_rc == -1) return -1; } + //for (u32 i = 4; i < 7; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_r, i, sizeof (cl_uint), device_param->kernel_params_mp_r[i]); if (CL_rc == -1) return -1; } + //for (u32 i = 8; i < 8; i++) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_r, i, sizeof (cl_ulong), device_param->kernel_params_mp_r[i]); if (CL_rc == -1) return -1; } + + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_root_css_buf, mask_ctx->root_css_buf, device_param->size_root_css) == -1) return -1; + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_markov_css_buf, mask_ctx->markov_css_buf, device_param->size_markov_css) == -1) return -1; + } + if (device_param->is_opencl == true) { for (u32 i = 3; i < 4; i++) { if (hc_clSetKernelArg (hashcat_ctx, device_param->opencl_kernel_mp_l, i, sizeof (cl_ulong), device_param->kernel_params_mp_l[i]) == -1) return -1; } diff --git a/src/dispatch.c b/src/dispatch.c index 25c40ea7f..d0cbfcfb6 100644 --- a/src/dispatch.c +++ b/src/dispatch.c @@ -350,6 +350,11 @@ HC_API_CALL void *thread_calc_stdin (void *p) if (hc_cuCtxSetCurrent (hashcat_ctx, device_param->cuda_context) == -1) return NULL; } + if (device_param->is_hip == true) + { + if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return NULL; + } + if (calc_stdin (hashcat_ctx, device_param) == -1) { status_ctx_t *status_ctx = hashcat_ctx->status_ctx; @@ -357,6 +362,11 @@ HC_API_CALL void *thread_calc_stdin (void *p) status_ctx->devices_status = STATUS_ERROR; } + if (device_param->is_hip == true) + { + if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) return NULL; + } + return NULL; } @@ -1584,6 +1594,11 @@ HC_API_CALL void *thread_calc (void *p) if (hc_cuCtxSetCurrent (hashcat_ctx, device_param->cuda_context) == -1) return NULL; } + if (device_param->is_hip == true) + { + if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return NULL; + } + if (calc (hashcat_ctx, device_param) == -1) { status_ctx_t *status_ctx = hashcat_ctx->status_ctx; @@ -1591,5 +1606,10 @@ HC_API_CALL void *thread_calc (void *p) status_ctx->devices_status = STATUS_ERROR; } + if (device_param->is_hip == true) + { + if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) return NULL; + } + return NULL; } diff --git a/src/hashes.c b/src/hashes.c index f1ee22334..27aa94370 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -322,6 +322,11 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl hc_cuMemcpyDtoH (hashcat_ctx, tmps, device_param->cuda_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size); } + if (device_param->is_hip == true) + { + hc_hipMemcpyDtoH (hashcat_ctx, tmps, device_param->hip_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size); + } + if (device_param->is_opencl == true) { hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tmps, CL_TRUE, plain->gidvid * hashconfig->tmp_size, hashconfig->tmp_size, tmps, 0, NULL, NULL); @@ -481,6 +486,7 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) u32 num_cracked = 0; int CU_rc; + int HIP_rc; int CL_rc; if (device_param->is_cuda == true) @@ -490,6 +496,13 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) if (CU_rc == -1) return -1; } + if (device_param->is_hip == true) + { + HIP_rc = hc_hipMemcpyDtoH (hashcat_ctx, &num_cracked, device_param->hip_d_result, sizeof (u32)); + + if (HIP_rc == -1) return -1; + } + if (device_param->is_opencl == true) { CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL); @@ -516,6 +529,13 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) if (CU_rc == -1) return -1; } + if (device_param->is_hip == true) + { + HIP_rc = hc_hipMemcpyDtoH (hashcat_ctx, cracked, device_param->hip_d_plain_bufs, num_cracked * sizeof (plain_t)); + + if (HIP_rc == -1) return -1; + } + if (device_param->is_opencl == true) { CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_plain_bufs, CL_TRUE, 0, num_cracked * sizeof (plain_t), cracked, 0, NULL, NULL); @@ -573,6 +593,13 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) if (CU_rc == -1) return -1; } + if (device_param->is_hip == true) + { + HIP_rc = hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_digests_shown + (salt_buf->digests_offset * sizeof (u32)), &hashes->digests_shown_tmp[salt_buf->digests_offset], salt_buf->digests_cnt * sizeof (u32)); + + if (HIP_rc == -1) return -1; + } + if (device_param->is_opencl == true) { CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (u32), salt_buf->digests_cnt * sizeof (u32), &hashes->digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL); @@ -611,6 +638,13 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) if (CU_rc == -1) return -1; } + if (device_param->is_hip == true) + { + HIP_rc = hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_result, &num_cracked, sizeof (u32)); + + if (HIP_rc == -1) return -1; + } + if (device_param->is_opencl == true) { CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL); diff --git a/src/selftest.c b/src/selftest.c index 219f6d771..ea08e4a43 100644 --- a/src/selftest.c +++ b/src/selftest.c @@ -32,6 +32,13 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param device_param->kernel_params[18] = &device_param->cuda_d_st_esalts_buf; } + if (device_param->is_hip == true) + { + device_param->kernel_params[15] = &device_param->hip_d_st_digests_buf; + device_param->kernel_params[17] = &device_param->hip_d_st_salts_buf; + device_param->kernel_params[18] = &device_param->hip_d_st_esalts_buf; + } + if (device_param->is_opencl == true) { device_param->kernel_params[15] = &device_param->opencl_d_st_digests_buf; @@ -91,6 +98,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; @@ -126,6 +138,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; @@ -190,6 +207,13 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_combs_c, &comb, 1 * sizeof (pw_t)) == -1) return -1; + + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_combs_c, CL_TRUE, 0, 1 * sizeof (pw_t), &comb, 0, NULL, NULL) == -1) return -1; @@ -225,6 +249,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; @@ -276,6 +305,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_bfs_c, &bf, 1 * sizeof (bf_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_bfs_c, &bf, 1 * sizeof (bf_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_bfs_c, CL_TRUE, 0, 1 * sizeof (bf_t), &bf, 0, NULL, NULL) == -1) return -1; @@ -372,6 +406,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; @@ -400,6 +439,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_pws_buf, &pw, 1 * sizeof (pw_t)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_pws_buf, CL_TRUE, 0, 1 * sizeof (pw_t), &pw, 0, NULL, NULL) == -1) return -1; @@ -446,6 +490,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (run_cuda_kernel_utf8toutf16le (hashcat_ctx, device_param, device_param->cuda_d_pws_buf, 1) == -1) return -1; } + if (device_param->is_hip == true) + { + if (run_hip_kernel_utf8toutf16le (hashcat_ctx, device_param, device_param->hip_d_pws_buf, 1) == -1) return -1; + } + if (device_param->is_opencl == true) { if (run_opencl_kernel_utf8toutf16le (hashcat_ctx, device_param, device_param->opencl_d_pws_buf, 1) == -1) return -1; @@ -463,6 +512,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, device_param->size_hooks) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, device_param->size_hooks) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; @@ -475,6 +529,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, device_param->size_hooks) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, device_param->size_hooks) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; @@ -526,6 +585,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_cuMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->cuda_d_hooks, device_param->size_hooks) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyDtoH (hashcat_ctx, device_param->hooks_buf, device_param->hip_d_hooks, device_param->size_hooks) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; @@ -538,6 +602,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_hooks, device_param->hooks_buf, device_param->size_hooks) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_hooks, device_param->hooks_buf, device_param->size_hooks) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_hooks, CL_TRUE, 0, device_param->size_hooks, device_param->hooks_buf, 0, NULL, NULL) == -1) return -1; @@ -617,6 +686,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hc_cuMemcpyDtoH (hashcat_ctx, &num_cracked, device_param->cuda_d_result, sizeof (u32)) == -1) return -1; } + if (device_param->is_hip == true) + { + if (hc_hipMemcpyDtoH (hashcat_ctx, &num_cracked, device_param->hip_d_result, sizeof (u32)) == -1) return -1; + } + if (device_param->is_opencl == true) { if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL) == -1) return -1; @@ -646,6 +720,20 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_result, device_param->size_results) == -1) return -1; } + if (device_param->is_hip == true) + { + device_param->kernel_params[15] = &device_param->hip_d_digests_buf; + device_param->kernel_params[17] = &device_param->hip_d_salt_bufs; + device_param->kernel_params[18] = &device_param->hip_d_esalt_bufs; + + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_pws_buf, device_param->size_pws) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_tmps, device_param->size_tmps) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_hooks, device_param->size_hooks) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_plain_bufs, device_param->size_plains) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_digests_shown, device_param->size_shown) == -1) return -1; + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_result, device_param->size_results) == -1) return -1; + } + if (device_param->is_opencl == true) { device_param->kernel_params[15] = &device_param->opencl_d_digests_buf; @@ -667,6 +755,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_rules_c, device_param->size_rules_c) == -1) return -1; } + if (device_param->is_hip == true) + { + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_rules_c, device_param->size_rules_c) == -1) return -1; + } + if (device_param->is_opencl == true) { if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_rules_c, device_param->size_rules_c) == -1) return -1; @@ -681,6 +774,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_rules_c, device_param->size_rules_c) == -1) return -1; } + if (device_param->is_hip == true) + { + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_rules_c, device_param->size_rules_c) == -1) return -1; + } + if (device_param->is_opencl == true) { if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_rules_c, device_param->size_rules_c) == -1) return -1; @@ -693,6 +791,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_combs_c, device_param->size_combs) == -1) return -1; } + if (device_param->is_hip == true) + { + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_combs_c, device_param->size_combs) == -1) return -1; + } + if (device_param->is_opencl == true) { if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_combs_c, device_param->size_combs) == -1) return -1; @@ -705,6 +808,11 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_bfs_c, device_param->size_bfs) == -1) return -1; } + if (device_param->is_hip == true) + { + if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_bfs_c, device_param->size_bfs) == -1) return -1; + } + if (device_param->is_opencl == true) { if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_bfs_c, device_param->size_bfs) == -1) return -1; @@ -713,19 +821,25 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param } // check return -//TODO: Add HIP in the above test. - if (num_cracked == 0 && false) + + if (num_cracked == 0) { hc_thread_mutex_lock (status_ctx->mux_display); + if (device_param->is_cuda == true) + { + event_log_error (hashcat_ctx, "* Device #%u: ATTENTION! CUDA kernel self-test failed.", device_param->device_id + 1); + } + + if (device_param->is_hip == true) + { + event_log_error (hashcat_ctx, "* Device #%u: ATTENTION! HIP kernel self-test failed.", device_param->device_id + 1); + } + if (device_param->is_opencl == true) { event_log_error (hashcat_ctx, "* Device #%u: ATTENTION! OpenCL kernel self-test failed.", device_param->device_id + 1); } - if (device_param->is_cuda == true) - { - event_log_error (hashcat_ctx, "* Device #%u: ATTENTION! CUDA kernel self-test failed.", device_param->device_id + 1); - } event_log_warning (hashcat_ctx, "Your device driver installation is probably broken."); event_log_warning (hashcat_ctx, "See also: https://hashcat.net/faq/wrongdriver"); @@ -735,6 +849,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param return -1; } + return 0; } @@ -763,6 +878,11 @@ HC_API_CALL void *thread_selftest (void *p) if (hc_cuCtxSetCurrent (hashcat_ctx, device_param->cuda_context) == -1) return NULL; } + if (device_param->is_hip == true) + { + if (hc_hipCtxPushCurrent (hashcat_ctx, device_param->hip_context) == -1) return NULL; + } + const int rc_selftest = selftest (hashcat_ctx, device_param); if (user_options->benchmark == true) @@ -781,5 +901,10 @@ HC_API_CALL void *thread_selftest (void *p) } } + if (device_param->is_hip == true) + { + if (hc_hipCtxPopCurrent (hashcat_ctx, &device_param->hip_context) == -1) return NULL; + } + return NULL; }