arno
/
hashcat
mirror of https://github.com/hashcat/hashcat.git
1
0
Fork 0
Code Issues Releases Wiki Activity

HIP Backend: Added support to support HIP 4.4 and later, but added check to rule out older versions because they are incompatible

Browse Source
pull/2912/head^2
Jens Steube 3 years ago
parent edab99cdfa
commit 5ffcaa980d
12 changed files with 823 additions and 2052 deletions
Show Stats Download Patch File Download Diff File

25
OpenCL/inc_platform.cl
Unescape View File

@ -193,33 +193,21 @@ DECLSPEC u32 hc_atomic_dec (GLOBAL_AS u32 *p)
{ {
volatile const u32 val = 1; volatile const u32 val = 1;
return __atomic_fetch_sub (p, val, __ATOMIC_RELAXED); return atomicSub (p, val);
} }
DECLSPEC u32 hc_atomic_inc (GLOBAL_AS u32 *p) DECLSPEC u32 hc_atomic_inc (GLOBAL_AS u32 *p)
{ {
volatile const u32 val = 1; volatile const u32 val = 1;
return __atomic_fetch_add (p, val, __ATOMIC_RELAXED); return atomicAdd (p, val);
} }
DECLSPEC u32 hc_atomic_or (GLOBAL_AS u32 *p, volatile const u32 val) DECLSPEC u32 hc_atomic_or (GLOBAL_AS u32 *p, volatile const u32 val)
{ {
return __atomic_fetch_or (p, val, __ATOMIC_RELAXED); return atomicOr (p, val);
}
extern "C" __device__ __attribute__((pure)) double __ocml_log2_f64(double);
DECLSPEC double log2 (double x)
{
return __ocml_log2_f64 (x);
} }
extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_id(uint);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_group_id(uint);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_size(uint);
extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(uint);
DECLSPEC size_t get_global_id (const u32 dimindx) DECLSPEC size_t get_global_id (const u32 dimindx)
{ {
return (__ockl_get_group_id (dimindx) * __ockl_get_local_size (dimindx)) + __ockl_get_local_id (dimindx); return (__ockl_get_group_id (dimindx) * __ockl_get_local_size (dimindx)) + __ockl_get_local_id (dimindx);
@ -308,11 +296,8 @@ DECLSPEC u64 rotr64_S (const u64 a, const int n)
return out.v64; return out.v64;
} }
extern "C" __device__ int printf(const char *fmt, ...); #define FIXED_THREAD_COUNT(n) __launch_bounds__((n), 0)
//int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2))); #define SYNC_THREADS() __syncthreads ()
#define FIXED_THREAD_COUNT(n) __attribute__((amdgpu_flat_work_group_size (1, (n))))
#define SYNC_THREADS() __builtin_amdgcn_s_barrier ()
#endif #endif
#ifdef IS_OPENCL #ifdef IS_OPENCL

87
OpenCL/inc_types.h
Unescape View File

@ -21,96 +21,19 @@
typedef unsigned char uchar; typedef unsigned char uchar;
typedef unsigned short ushort; typedef unsigned short ushort;
typedef unsigned int uint; typedef unsigned int uint;
typedef unsigned long long ulong; typedef unsigned long ulong;
typedef unsigned long long ullong;
#endif #endif
#ifdef IS_HIP #ifdef IS_OPENCL
// https://github.com/llvm-mirror/clang/blob/master/lib/Headers/opencl-c-base.h typedef ulong ullong;
// built-in scalar data types:
/**
* An unsigned 8-bit integer.
*/
typedef unsigned char uchar;
/**
* An unsigned 16-bit integer.
*/
typedef unsigned short ushort;
/**
* An unsigned 32-bit integer.
*/
typedef unsigned int uint;
/**
* An unsigned 64-bit integer.
*/
typedef unsigned long ulong;
/**
* The unsigned integer type of the result of the sizeof operator. This
* is a 32-bit unsigned integer if CL_DEVICE_ADDRESS_BITS
* defined in table 4.3 is 32-bits and is a 64-bit unsigned integer if
* CL_DEVICE_ADDRESS_BITS is 64-bits.
*/
typedef __SIZE_TYPE__ size_t;
// built-in vector data types:
typedef char char2 __attribute__((ext_vector_type(2)));
typedef char char3 __attribute__((ext_vector_type(3)));
typedef char char4 __attribute__((ext_vector_type(4)));
typedef char char8 __attribute__((ext_vector_type(8)));
typedef char char16 __attribute__((ext_vector_type(16)));
typedef uchar uchar2 __attribute__((ext_vector_type(2)));
typedef uchar uchar3 __attribute__((ext_vector_type(3)));
typedef uchar uchar4 __attribute__((ext_vector_type(4)));
typedef uchar uchar8 __attribute__((ext_vector_type(8)));
typedef uchar uchar16 __attribute__((ext_vector_type(16)));
typedef short short2 __attribute__((ext_vector_type(2)));
typedef short short3 __attribute__((ext_vector_type(3)));
typedef short short4 __attribute__((ext_vector_type(4)));
typedef short short8 __attribute__((ext_vector_type(8)));
typedef short short16 __attribute__((ext_vector_type(16)));
typedef ushort ushort2 __attribute__((ext_vector_type(2)));
typedef ushort ushort3 __attribute__((ext_vector_type(3)));
typedef ushort ushort4 __attribute__((ext_vector_type(4)));
typedef ushort ushort8 __attribute__((ext_vector_type(8)));
typedef ushort ushort16 __attribute__((ext_vector_type(16)));
typedef int int2 __attribute__((ext_vector_type(2)));
typedef int int3 __attribute__((ext_vector_type(3)));
typedef int int4 __attribute__((ext_vector_type(4)));
typedef int int8 __attribute__((ext_vector_type(8)));
typedef int int16 __attribute__((ext_vector_type(16)));
typedef uint uint2 __attribute__((ext_vector_type(2)));
typedef uint uint3 __attribute__((ext_vector_type(3)));
typedef uint uint4 __attribute__((ext_vector_type(4)));
typedef uint uint8 __attribute__((ext_vector_type(8)));
typedef uint uint16 __attribute__((ext_vector_type(16)));
typedef long long2 __attribute__((ext_vector_type(2)));
typedef long long3 __attribute__((ext_vector_type(3)));
typedef long long4 __attribute__((ext_vector_type(4)));
typedef long long8 __attribute__((ext_vector_type(8)));
typedef long long16 __attribute__((ext_vector_type(16)));
typedef ulong ulong2 __attribute__((ext_vector_type(2)));
typedef ulong ulong3 __attribute__((ext_vector_type(3)));
typedef ulong ulong4 __attribute__((ext_vector_type(4)));
typedef ulong ulong8 __attribute__((ext_vector_type(8)));
typedef ulong ulong16 __attribute__((ext_vector_type(16)));
typedef float float2 __attribute__((ext_vector_type(2)));
typedef float float3 __attribute__((ext_vector_type(3)));
typedef float float4 __attribute__((ext_vector_type(4)));
typedef float float8 __attribute__((ext_vector_type(8)));
typedef float float16 __attribute__((ext_vector_type(16)));
#endif #endif
#ifdef KERNEL_STATIC #ifdef KERNEL_STATIC
typedef uchar u8; typedef uchar u8;
typedef ushort u16; typedef ushort u16;
typedef uint u32; typedef uint u32;
typedef ulong u64; typedef ullong u64;
#else #else
typedef uint8_t u8; typedef uint8_t u8;
typedef uint16_t u16; typedef uint16_t u16;

13
OpenCL/m08900-pure.cl
Unescape View File

@ -24,7 +24,7 @@ typedef struct
} scrypt_tmp_t; } scrypt_tmp_t;
#if defined IS_CUDA #if defined IS_CUDA || defined IS_HIP
inline __device__ uint4 operator & (const uint4 a, const u32 b) { return make_uint4 ((a.x & b ), (a.y & b ), (a.z & b ), (a.w & b )); } inline __device__ uint4 operator & (const uint4 a, const u32 b) { return make_uint4 ((a.x & b ), (a.y & b ), (a.z & b ), (a.w & b )); }
inline __device__ uint4 operator << (const uint4 a, const u32 b) { return make_uint4 ((a.x << b ), (a.y << b ), (a.z << b ), (a.w << b )); } inline __device__ uint4 operator << (const uint4 a, const u32 b) { return make_uint4 ((a.x << b ), (a.y << b ), (a.z << b ), (a.w << b )); }
@ -41,15 +41,6 @@ inline __device__ uint4 rotate (const uint4 a, const int n)
#endif #endif
#if defined IS_HIP
inline __device__ uint4 rotate (const uint4 a, const int n)
{
return ((a << n) | ((a >> (32 - n))));
}
#endif
DECLSPEC uint4 hc_swap32_4 (uint4 v) DECLSPEC uint4 hc_swap32_4 (uint4 v)
{ {
return (rotate ((v & 0x00FF00FF), 24u) | rotate ((v & 0xFF00FF00), 8u)); return (rotate ((v & 0x00FF00FF), 24u) | rotate ((v & 0xFF00FF00), 8u));
@ -66,7 +57,7 @@ DECLSPEC uint4 hc_swap32_4 (uint4 v)
#define ADD_ROTATE_XOR(r,i1,i2,s) (r) ^= rotate ((i1) + (i2), (s)); #define ADD_ROTATE_XOR(r,i1,i2,s) (r) ^= rotate ((i1) + (i2), (s));
#if defined IS_CUDA #if defined IS_CUDA || defined IS_HIP
#define SALSA20_2R() \ #define SALSA20_2R() \
{ \ { \

13
OpenCL/m15700-pure.cl
Unescape View File

@ -31,7 +31,7 @@ typedef struct ethereum_scrypt
} ethereum_scrypt_t; } ethereum_scrypt_t;
#if defined IS_CUDA #if defined IS_CUDA || defined IS_HIP
inline __device__ uint4 operator & (const uint4 a, const u32 b) { return make_uint4 ((a.x & b ), (a.y & b ), (a.z & b ), (a.w & b )); } inline __device__ uint4 operator & (const uint4 a, const u32 b) { return make_uint4 ((a.x & b ), (a.y & b ), (a.z & b ), (a.w & b )); }
inline __device__ uint4 operator << (const uint4 a, const u32 b) { return make_uint4 ((a.x << b ), (a.y << b ), (a.z << b ), (a.w << b )); } inline __device__ uint4 operator << (const uint4 a, const u32 b) { return make_uint4 ((a.x << b ), (a.y << b ), (a.z << b ), (a.w << b )); }
@ -48,15 +48,6 @@ inline __device__ uint4 rotate (const uint4 a, const int n)
#endif #endif
#if defined IS_HIP
inline __device__ uint4 rotate (const uint4 a, const int n)
{
return ((a << n) | ((a >> (32 - n))));
}
#endif
DECLSPEC uint4 hc_swap32_4 (uint4 v) DECLSPEC uint4 hc_swap32_4 (uint4 v)
{ {
return (rotate ((v & 0x00FF00FF), 24u) | rotate ((v & 0xFF00FF00), 8u)); return (rotate ((v & 0x00FF00FF), 24u) | rotate ((v & 0xFF00FF00), 8u));
@ -73,7 +64,7 @@ DECLSPEC uint4 hc_swap32_4 (uint4 v)
#define ADD_ROTATE_XOR(r,i1,i2,s) (r) ^= rotate ((i1) + (i2), (s)); #define ADD_ROTATE_XOR(r,i1,i2,s) (r) ^= rotate ((i1) + (i2), (s));
#if defined IS_CUDA #if defined IS_CUDA || defined IS_HIP
#define SALSA20_2R() \ #define SALSA20_2R() \
{ \ { \

13
OpenCL/m22700-pure.cl
Unescape View File

@ -72,7 +72,7 @@ DECLSPEC int is_valid_bitcoinj (const u32 *w)
return 1; return 1;
} }
#if defined IS_CUDA #if defined IS_CUDA || defined IS_HIP
inline __device__ uint4 operator & (const uint4 a, const u32 b) { return make_uint4 ((a.x & b ), (a.y & b ), (a.z & b ), (a.w & b )); } inline __device__ uint4 operator & (const uint4 a, const u32 b) { return make_uint4 ((a.x & b ), (a.y & b ), (a.z & b ), (a.w & b )); }
inline __device__ uint4 operator << (const uint4 a, const u32 b) { return make_uint4 ((a.x << b ), (a.y << b ), (a.z << b ), (a.w << b )); } inline __device__ uint4 operator << (const uint4 a, const u32 b) { return make_uint4 ((a.x << b ), (a.y << b ), (a.z << b ), (a.w << b )); }
@ -89,15 +89,6 @@ inline __device__ uint4 rotate (const uint4 a, const int n)
#endif #endif
#if defined IS_HIP
inline __device__ uint4 rotate (const uint4 a, const int n)
{
return ((a << n) | ((a >> (32 - n))));
}
#endif
DECLSPEC uint4 hc_swap32_4 (uint4 v) DECLSPEC uint4 hc_swap32_4 (uint4 v)
{ {
return (rotate ((v & 0x00FF00FF), 24u) | rotate ((v & 0xFF00FF00), 8u)); return (rotate ((v & 0x00FF00FF), 24u) | rotate ((v & 0xFF00FF00), 8u));
@ -114,7 +105,7 @@ DECLSPEC uint4 hc_swap32_4 (uint4 v)
#define ADD_ROTATE_XOR(r,i1,i2,s) (r) ^= rotate ((i1) + (i2), (s)); #define ADD_ROTATE_XOR(r,i1,i2,s) (r) ^= rotate ((i1) + (i2), (s));
#if defined IS_CUDA #if defined IS_CUDA || defined IS_HIP
#define SALSA20_2R() \ #define SALSA20_2R() \
{ \ { \

1
docs/changes.txt
Unescape View File

@ -5,6 +5,7 @@
## ##
- Added option --multiply-accel-disable (short: -M) to disable multiply the kernel-accel with the multiprocessor count automatism - Added option --multiply-accel-disable (short: -M) to disable multiply the kernel-accel with the multiprocessor count automatism
- HIP Backend: Added support to support HIP 4.4 and later, but added check to rule out older versions because they are incompatible
## ##
## Bugs ## Bugs

91
include/backend.h
Unescape View File

@ -88,53 +88,50 @@ int hc_cuLinkAddData (hashcat_ctx_t *hashcat_ctx, CUlinkState state,
int hc_cuLinkDestroy (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_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_hipCreateProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram *prog, const char *src, const char *name, int numHeaders, const char * const *headers, const char * const *includeNames);
int hc_nvrtcDestroyProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram *prog); int hc_hipDestroyProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram *prog);
int hc_nvrtcCompileProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, int numOptions, const char * const *options); int hc_hipCompileProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, int numOptions, const char * const *options);
int hc_nvrtcGetProgramLogSize (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *logSizeRet); int hc_hipGetProgramLogSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *logSizeRet);
int hc_nvrtcGetProgramLog (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *log); int hc_hipGetProgramLog (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *log);
int hc_nvrtcGetPTXSize (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *ptxSizeRet); int hc_hipGetCodeSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *codeSizeRet);
int hc_nvrtcGetPTX (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *ptx); int hc_hipGetCode (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *code);
int hc_nvrtcVersion (hashcat_ctx_t *hashcat_ctx, int *major, int *minor);
int hc_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, hipCtx_t *pctx, unsigned int flags, hipDevice_t dev);
int hc_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx, unsigned int flags, HIPdevice dev); int hc_hipCtxDestroy (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx);
int hc_hipCtxDestroy (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx); int hc_hipCtxPopCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t *pctx);
int hc_hipCtxSetCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx); int hc_hipCtxPushCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx);
int hc_hipCtxSetCacheConfig (hashcat_ctx_t *hashcat_ctx, HIPfunc_cache config); int hc_hipCtxSetCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx);
int hc_hipCtxSynchronize (hashcat_ctx_t *hashcat_ctx); 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_hipDeviceGet (hashcat_ctx_t *hashcat_ctx, hipDevice_t *device, int ordinal);
int hc_hipDeviceGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, hipDeviceAttribute_t attrib, hipDevice_t dev);
int hc_hipDeviceGetCount (hashcat_ctx_t *hashcat_ctx, int *count); int hc_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_t dev);
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_t 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_hipDriverGetVersion (hashcat_ctx_t *hashcat_ctx, int *driverVersion);
int hc_hipEventCreate (hashcat_ctx_t *hashcat_ctx, HIPevent *phEvent, unsigned int Flags); int hc_hipEventCreate (hashcat_ctx_t *hashcat_ctx, hipEvent_t *phEvent, unsigned int Flags);
int hc_hipEventDestroy (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); int hc_hipEventDestroy (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent);
int hc_hipEventElapsedTime (hashcat_ctx_t *hashcat_ctx, float *pMilliseconds, HIPevent hStart, HIPevent hEnd); int hc_hipEventElapsedTime (hashcat_ctx_t *hashcat_ctx, float *pMilliseconds, hipEvent_t hStart, hipEvent_t hEnd);
int hc_hipEventQuery (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); int hc_hipEventQuery (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent);
int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent, HIPstream hStream); int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent, hipStream_t hStream);
int hc_hipEventSynchronize (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); int hc_hipEventSynchronize (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent);
int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPfunction_attribute attrib, HIPfunction hfunc); int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, hipFunction_attribute attrib, hipFunction_t 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_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_hipLaunchKernel (hashcat_ctx_t *hashcat_ctx, hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra);
int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr *dptr, size_t bytesize); int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize);
int hc_hipMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, HIPdeviceptr srcDevice, size_t ByteCount); int hc_hipMemFree (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dptr);
int hc_hipMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, HIPdeviceptr srcDevice, size_t ByteCount); int hc_hipMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount);
int hc_hipMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, const void *srcHost, size_t ByteCount); int hc_hipMemcpyDtoHAsync (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream);
int hc_hipMemFree (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dptr); int hc_hipMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount);
int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, HIPfunction *hfunc, HIPmodule hmod, const char *name); int hc_hipMemcpyDtoDAsync (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream);
int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, HIPmodule *module, const void *image, unsigned int numOptions, HIPjit_option *options, void **optionValues); int hc_hipMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount);
int hc_hipModuleUnload (hashcat_ctx_t *hashcat_ctx, HIPmodule hmod); int hc_hipMemcpyHtoDAsync (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount, hipStream_t hStream);
int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, HIPstream *phStream, unsigned int Flags); int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, hipFunction_t *hfunc, hipModule_t hmod, const char *name);
int hc_hipStreamDestroy (hashcat_ctx_t *hashcat_ctx, HIPstream hStream); int hc_hipModuleGetGlobal (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name);
int hc_hipStreamSynchronize (hashcat_ctx_t *hashcat_ctx, HIPstream hStream); int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, hipModule_t *module, const void *image, unsigned int numOptions, hipJitOption *options, void **optionValues);
int hc_hipCtxPushCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx); int hc_hipModuleUnload (hashcat_ctx_t *hashcat_ctx, hipModule_t hmod);
int hc_hipCtxPopCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx); int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, hipStream_t *phStream, unsigned int Flags);
int hc_hipLinkCreate (hashcat_ctx_t *hashcat_ctx, unsigned int numOptions, HIPjit_option *options, void **optionValues, HIPlinkState *stateOut); int hc_hipStreamDestroy (hashcat_ctx_t *hashcat_ctx, hipStream_t hStream);
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_hipStreamSynchronize (hashcat_ctx_t *hashcat_ctx, hipStream_t hStream);
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_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); 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);
@ -183,10 +180,10 @@ int run_cuda_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param
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_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_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_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t 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_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t 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_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t 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_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t 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_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 num);

1437
include/ext_hip.h
View File

File diff suppressed because it is too large Load Diff

61
include/ext_hiprtc.h
Unescape View File

@ -6,41 +6,26 @@
#ifndef _EXT_HIPRTC_H #ifndef _EXT_HIPRTC_H
#define _EXT_HIPRTC_H #define _EXT_HIPRTC_H
/** // start: amd_detail/hiprtc.h
* from hip_runtime.h (/opt/rocm/hip/include/hip/amd_detail/hiprtc.h)
*/
/** typedef enum hiprtcResult {
* \ingroup error HIPRTC_SUCCESS = 0,
* \brief The enumerated type hiprtcResult defines API call result codes. HIPRTC_ERROR_OUT_OF_MEMORY = 1,
* HIPRTC API functions return hiprtcResult to indicate the call HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2,
* result. HIPRTC_ERROR_INVALID_INPUT = 3,
*/ HIPRTC_ERROR_INVALID_PROGRAM = 4,
typedef enum { HIPRTC_ERROR_INVALID_OPTION = 5,
HIPRTC_SUCCESS = 0, HIPRTC_ERROR_COMPILATION = 6,
HIPRTC_ERROR_OUT_OF_MEMORY = 1, HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7,
HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2, HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8,
HIPRTC_ERROR_INVALID_INPUT = 3, HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9,
HIPRTC_ERROR_INVALID_PROGRAM = 4, HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10,
HIPRTC_ERROR_INVALID_OPTION = 5, HIPRTC_ERROR_INTERNAL_ERROR = 11
HIPRTC_ERROR_COMPILATION = 6,
HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7,
HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8,
HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9,
HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10,
HIPRTC_ERROR_INTERNAL_ERROR = 11
} hiprtcResult; } hiprtcResult;
/** typedef struct _hiprtcProgram* hiprtcProgram;
* \ingroup compilation
* \brief hiprtcProgram is the unit of compilation, and an opaque handle for // stop: amd_detail/hiprtc.h
* a program.
*
* To compile a CUDA program string, an instance of hiprtcProgram must be
* created first with ::hiprtcCreateProgram, then compiled with
* ::hiprtcCompileProgram.
*/
typedef struct _hiprtcProgram *hiprtcProgram;
#ifdef _WIN32 #ifdef _WIN32
#define HIPRTCAPI __stdcall #define HIPRTCAPI __stdcall
@ -54,13 +39,12 @@ typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCADDNAMEEXPRESSION) (hiprtc
typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCCOMPILEPROGRAM) (hiprtcProgram, int, const char * const *); typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCCOMPILEPROGRAM) (hiprtcProgram, int, const char * const *);
typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCCREATEPROGRAM) (hiprtcProgram *, const char *, const char *, int, const char * const *, const char * const *); typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCCREATEPROGRAM) (hiprtcProgram *, const char *, const char *, int, const char * const *, const char * const *);
typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCDESTROYPROGRAM) (hiprtcProgram *); typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCDESTROYPROGRAM) (hiprtcProgram *);
typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETCODE) (hiprtcProgram, char *);
typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETCODESIZE) (hiprtcProgram, size_t *);
typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETLOWEREDNAME) (hiprtcProgram, const char * const, const char **); typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETLOWEREDNAME) (hiprtcProgram, const char * const, const char **);
typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETPTX) (hiprtcProgram, char *);
typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETPTXSIZE) (hiprtcProgram, size_t *);
typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETPROGRAMLOG) (hiprtcProgram, char *); typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETPROGRAMLOG) (hiprtcProgram, char *);
typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETPROGRAMLOGSIZE) (hiprtcProgram, size_t *); typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETPROGRAMLOGSIZE) (hiprtcProgram, size_t *);
typedef const char * (HIPRTC_API_CALL *HIPRTC_HIPRTCGETERRORSTRING) (hiprtcResult); typedef const char * (HIPRTC_API_CALL *HIPRTC_HIPRTCGETERRORSTRING) (hiprtcResult);
typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCVERSION) (int *, int *);
typedef struct hc_hiprtc_lib typedef struct hc_hiprtc_lib
{ {
@ -70,13 +54,12 @@ typedef struct hc_hiprtc_lib
HIPRTC_HIPRTCCOMPILEPROGRAM hiprtcCompileProgram; HIPRTC_HIPRTCCOMPILEPROGRAM hiprtcCompileProgram;
HIPRTC_HIPRTCCREATEPROGRAM hiprtcCreateProgram; HIPRTC_HIPRTCCREATEPROGRAM hiprtcCreateProgram;
HIPRTC_HIPRTCDESTROYPROGRAM hiprtcDestroyProgram; HIPRTC_HIPRTCDESTROYPROGRAM hiprtcDestroyProgram;
HIPRTC_HIPRTCGETCODE hiprtcGetCode;
HIPRTC_HIPRTCGETCODESIZE hiprtcGetCodeSize;
HIPRTC_HIPRTCGETLOWEREDNAME hiprtcGetLoweredName; HIPRTC_HIPRTCGETLOWEREDNAME hiprtcGetLoweredName;
HIPRTC_HIPRTCGETPTX hiprtcGetCode;
HIPRTC_HIPRTCGETPTXSIZE hiprtcGetCodeSize;
HIPRTC_HIPRTCGETPROGRAMLOG hiprtcGetProgramLog; HIPRTC_HIPRTCGETPROGRAMLOG hiprtcGetProgramLog;
HIPRTC_HIPRTCGETPROGRAMLOGSIZE hiprtcGetProgramLogSize; HIPRTC_HIPRTCGETPROGRAMLOGSIZE hiprtcGetProgramLogSize;
HIPRTC_HIPRTCGETERRORSTRING hiprtcGetErrorString; HIPRTC_HIPRTCGETERRORSTRING hiprtcGetErrorString;
HIPRTC_HIPRTCVERSION hiprtcVersion;
} hc_hiprtc_lib_t; } hc_hiprtc_lib_t;

151
include/types.h
Unescape View File

@ -1502,80 +1502,80 @@ typedef struct hc_device_param
int hip_warp_size; int hip_warp_size;
HIPdevice hip_device; hipDevice_t hip_device;
HIPcontext hip_context; hipCtx_t hip_context;
HIPstream hip_stream; hipStream_t hip_stream;
HIPevent hip_event1; hipEvent_t hip_event1;
HIPevent hip_event2; hipEvent_t hip_event2;
HIPmodule hip_module; hipModule_t hip_module;
HIPmodule hip_module_shared; hipModule_t hip_module_shared;
HIPmodule hip_module_mp; hipModule_t hip_module_mp;
HIPmodule hip_module_amp; hipModule_t hip_module_amp;
HIPfunction hip_function1; hipFunction_t hip_function1;
HIPfunction hip_function12; hipFunction_t hip_function12;
HIPfunction hip_function2p; hipFunction_t hip_function2p;
HIPfunction hip_function2; hipFunction_t hip_function2;
HIPfunction hip_function2e; hipFunction_t hip_function2e;
HIPfunction hip_function23; hipFunction_t hip_function23;
HIPfunction hip_function3; hipFunction_t hip_function3;
HIPfunction hip_function4; hipFunction_t hip_function4;
HIPfunction hip_function_init2; hipFunction_t hip_function_init2;
HIPfunction hip_function_loop2p; hipFunction_t hip_function_loop2p;
HIPfunction hip_function_loop2; hipFunction_t hip_function_loop2;
HIPfunction hip_function_mp; hipFunction_t hip_function_mp;
HIPfunction hip_function_mp_l; hipFunction_t hip_function_mp_l;
HIPfunction hip_function_mp_r; hipFunction_t hip_function_mp_r;
HIPfunction hip_function_amp; hipFunction_t hip_function_amp;
HIPfunction hip_function_tm; hipFunction_t hip_function_tm;
HIPfunction hip_function_memset; hipFunction_t hip_function_memset;
HIPfunction hip_function_bzero; hipFunction_t hip_function_bzero;
HIPfunction hip_function_atinit; hipFunction_t hip_function_atinit;
HIPfunction hip_function_utf8toutf16le; hipFunction_t hip_function_utf8toutf16le;
HIPfunction hip_function_decompress; hipFunction_t hip_function_decompress;
HIPfunction hip_function_aux1; hipFunction_t hip_function_aux1;
HIPfunction hip_function_aux2; hipFunction_t hip_function_aux2;
HIPfunction hip_function_aux3; hipFunction_t hip_function_aux3;
HIPfunction hip_function_aux4; hipFunction_t hip_function_aux4;
HIPdeviceptr hip_d_pws_buf; hipDeviceptr_t hip_d_pws_buf;
HIPdeviceptr hip_d_pws_amp_buf; hipDeviceptr_t hip_d_pws_amp_buf;
HIPdeviceptr hip_d_pws_comp_buf; hipDeviceptr_t hip_d_pws_comp_buf;
HIPdeviceptr hip_d_pws_idx; hipDeviceptr_t hip_d_pws_idx;
HIPdeviceptr hip_d_rules; hipDeviceptr_t hip_d_rules;
HIPdeviceptr hip_d_rules_c; hipDeviceptr_t hip_d_rules_c;
HIPdeviceptr hip_d_combs; hipDeviceptr_t hip_d_combs;
HIPdeviceptr hip_d_combs_c; hipDeviceptr_t hip_d_combs_c;
HIPdeviceptr hip_d_bfs; hipDeviceptr_t hip_d_bfs;
HIPdeviceptr hip_d_bfs_c; hipDeviceptr_t hip_d_bfs_c;
HIPdeviceptr hip_d_tm_c; hipDeviceptr_t hip_d_tm_c;
HIPdeviceptr hip_d_bitmap_s1_a; hipDeviceptr_t hip_d_bitmap_s1_a;
HIPdeviceptr hip_d_bitmap_s1_b; hipDeviceptr_t hip_d_bitmap_s1_b;
HIPdeviceptr hip_d_bitmap_s1_c; hipDeviceptr_t hip_d_bitmap_s1_c;
HIPdeviceptr hip_d_bitmap_s1_d; hipDeviceptr_t hip_d_bitmap_s1_d;
HIPdeviceptr hip_d_bitmap_s2_a; hipDeviceptr_t hip_d_bitmap_s2_a;
HIPdeviceptr hip_d_bitmap_s2_b; hipDeviceptr_t hip_d_bitmap_s2_b;
HIPdeviceptr hip_d_bitmap_s2_c; hipDeviceptr_t hip_d_bitmap_s2_c;
HIPdeviceptr hip_d_bitmap_s2_d; hipDeviceptr_t hip_d_bitmap_s2_d;
HIPdeviceptr hip_d_plain_bufs; hipDeviceptr_t hip_d_plain_bufs;
HIPdeviceptr hip_d_digests_buf; hipDeviceptr_t hip_d_digests_buf;
HIPdeviceptr hip_d_digests_shown; hipDeviceptr_t hip_d_digests_shown;
HIPdeviceptr hip_d_salt_bufs; hipDeviceptr_t hip_d_salt_bufs;
HIPdeviceptr hip_d_esalt_bufs; hipDeviceptr_t hip_d_esalt_bufs;
HIPdeviceptr hip_d_tmps; hipDeviceptr_t hip_d_tmps;
HIPdeviceptr hip_d_hooks; hipDeviceptr_t hip_d_hooks;
HIPdeviceptr hip_d_result; hipDeviceptr_t hip_d_result;
HIPdeviceptr hip_d_extra0_buf; hipDeviceptr_t hip_d_extra0_buf;
HIPdeviceptr hip_d_extra1_buf; hipDeviceptr_t hip_d_extra1_buf;
HIPdeviceptr hip_d_extra2_buf; hipDeviceptr_t hip_d_extra2_buf;
HIPdeviceptr hip_d_extra3_buf; hipDeviceptr_t hip_d_extra3_buf;
HIPdeviceptr hip_d_root_css_buf; hipDeviceptr_t hip_d_root_css_buf;
HIPdeviceptr hip_d_markov_css_buf; hipDeviceptr_t hip_d_markov_css_buf;
HIPdeviceptr hip_d_st_digests_buf; hipDeviceptr_t hip_d_st_digests_buf;
HIPdeviceptr hip_d_st_salts_buf; hipDeviceptr_t hip_d_st_salts_buf;
HIPdeviceptr hip_d_st_esalts_buf; hipDeviceptr_t hip_d_st_esalts_buf;
// API: opencl // API: opencl
@ -1726,8 +1726,7 @@ typedef struct backend_ctx
int rc_hip_init; int rc_hip_init;
int rc_hiprtc_init; int rc_hiprtc_init;
int hiprtc_driver_version; int hip_driverVersion;
int hip_driver_version;
// opencl // opencl

973
src/backend.c
View File

File diff suppressed because it is too large Load Diff

10
src/terminal.c
Unescape View File

@ -818,9 +818,9 @@ void backend_info (hashcat_ctx_t *hashcat_ctx)
event_log_info (hashcat_ctx, NULL); event_log_info (hashcat_ctx, NULL);
int hip_devices_cnt = backend_ctx->hip_devices_cnt; int hip_devices_cnt = backend_ctx->hip_devices_cnt;
int hip_driver_version = backend_ctx->hip_driver_version; int hip_driverVersion = backend_ctx->hip_driverVersion;
event_log_info (hashcat_ctx, "HIP.Version.: %d.%d", hip_driver_version / 1000, (hip_driver_version % 100) / 10); event_log_info (hashcat_ctx, "HIP.Version.: %d.%d", hip_driverVersion / 100, hip_driverVersion % 10);
event_log_info (hashcat_ctx, NULL); event_log_info (hashcat_ctx, NULL);
for (int hip_devices_idx = 0; hip_devices_idx < hip_devices_cnt; hip_devices_idx++) for (int hip_devices_idx = 0; hip_devices_idx < hip_devices_cnt; hip_devices_idx++)
@ -1014,10 +1014,10 @@ void backend_info_compact (hashcat_ctx_t *hashcat_ctx)
if (backend_ctx->hip) if (backend_ctx->hip)
{ {
int hip_devices_cnt = backend_ctx->hip_devices_cnt; int hip_devices_cnt = backend_ctx->hip_devices_cnt;
int hip_driver_version = backend_ctx->hip_driver_version; int hip_driverVersion = backend_ctx->hip_driverVersion;
const size_t len = event_log_info (hashcat_ctx, "HIP API (HIP %d.%d)", hip_driver_version / 1000, (hip_driver_version % 100) / 10); const size_t len = event_log_info (hashcat_ctx, "HIP API (HIP %d.%d)", hip_driverVersion / 100, hip_driverVersion % 10);
char line[HCBUFSIZ_TINY] = { 0 }; char line[HCBUFSIZ_TINY] = { 0 };

Write Preview
Loading…
Cancel
Save