mirror of
https://github.com/hashcat/hashcat.git
synced 2024-11-14 11:48:57 +00:00
Merge branch 'HIP44' into master
This commit is contained in:
Jens Steube
GitHub
commit
03ed06849a
@ -193,46 +193,35 @@ DECLSPEC u32 hc_atomic_dec (GLOBAL_AS u32 *p)
|
||||
{
|
||||
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)
|
||||
{
|
||||
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)
|
||||
{
|
||||
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)
|
||||
DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused)))
|
||||
{
|
||||
return __ocml_log2_f64 (x);
|
||||
return (blockIdx.x * blockDim.x) + threadIdx.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_local_id (const u32 dimindx __attribute__((unused)))
|
||||
{
|
||||
return (__ockl_get_group_id (dimindx) * __ockl_get_local_size (dimindx)) + __ockl_get_local_id (dimindx);
|
||||
return threadIdx.x;
|
||||
}
|
||||
|
||||
DECLSPEC size_t get_local_id (const u32 dimindx)
|
||||
DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
|
||||
{
|
||||
return __ockl_get_local_id (dimindx);
|
||||
}
|
||||
|
||||
DECLSPEC size_t get_local_size (const u32 dimindx)
|
||||
{
|
||||
return __ockl_get_local_size (dimindx);
|
||||
// verify
|
||||
return blockDim.x;
|
||||
}
|
||||
|
||||
DECLSPEC u32x rotl32 (const u32x a, const int n)
|
||||
@ -308,11 +297,8 @@ DECLSPEC u64 rotr64_S (const u64 a, const int n)
|
||||
return out.v64;
|
||||
}
|
||||
|
||||
extern "C" __device__ int printf(const char *fmt, ...);
|
||||
//int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2)));
|
||||
|
||||
#define FIXED_THREAD_COUNT(n) __attribute__((amdgpu_flat_work_group_size (1, (n))))
|
||||
#define SYNC_THREADS() __builtin_amdgcn_s_barrier ()
|
||||
#define FIXED_THREAD_COUNT(n) __launch_bounds__((n), 0)
|
||||
#define SYNC_THREADS() __syncthreads ()
|
||||
#endif
|
||||
|
||||
#ifdef IS_OPENCL
|
||||
|
||||
@ -21,96 +21,23 @@
|
||||
typedef unsigned char uchar;
|
||||
typedef unsigned short ushort;
|
||||
typedef unsigned int uint;
|
||||
typedef unsigned long long ulong;
|
||||
typedef unsigned long ulong;
|
||||
typedef unsigned long long ullong;
|
||||
#endif
|
||||
|
||||
#ifdef IS_HIP
|
||||
// https://github.com/llvm-mirror/clang/blob/master/lib/Headers/opencl-c-base.h
|
||||
|
||||
// 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)));
|
||||
|
||||
#ifdef IS_OPENCL
|
||||
typedef ulong ullong;
|
||||
typedef ulong2 ullong2;
|
||||
typedef ulong4 ullong4;
|
||||
typedef ulong8 ullong8;
|
||||
typedef ulong16 ullong16;
|
||||
#endif
|
||||
|
||||
#ifdef KERNEL_STATIC
|
||||
typedef uchar u8;
|
||||
typedef ushort u16;
|
||||
typedef uint u32;
|
||||
typedef ulong u64;
|
||||
typedef ullong u64;
|
||||
#else
|
||||
typedef uint8_t u8;
|
||||
typedef uint16_t u16;
|
||||
@ -910,7 +837,7 @@ typedef __device_builtin__ struct u64x u64x;
|
||||
typedef VTYPE(uchar, VECT_SIZE) u8x;
|
||||
typedef VTYPE(ushort, VECT_SIZE) u16x;
|
||||
typedef VTYPE(uint, VECT_SIZE) u32x;
|
||||
typedef VTYPE(ulong, VECT_SIZE) u64x;
|
||||
typedef VTYPE(ullong, VECT_SIZE) u64x;
|
||||
|
||||
#define make_u8x (u8x)
|
||||
#define make_u16x (u16x)
|
||||
|
||||
@ -32,10 +32,6 @@
|
||||
#define LOCAL_AS
|
||||
#define KERNEL_FQ extern "C" __global__
|
||||
#elif defined IS_HIP
|
||||
#define __device__ __attribute__((device))
|
||||
#define __constant__ __attribute__((constant))
|
||||
#define __shared__ __attribute__((shared))
|
||||
#define __global__ __attribute__((global))
|
||||
#define CONSTANT_VK __constant__
|
||||
#define CONSTANT_AS
|
||||
#define GLOBAL_AS
|
||||
|
||||
@ -73,18 +73,16 @@ enum{
|
||||
MZ_VERSION_ERROR = -6,
|
||||
MZ_PARAM_ERROR = -10000
|
||||
};
|
||||
typedef unsigned long mz_ulong;
|
||||
typedef ullong mz_ulong;
|
||||
|
||||
#ifndef MINIZ_NO_ZLIB_COMPATIBLE_NAMES
|
||||
typedef unsigned char Byte;
|
||||
typedef unsigned int uInt;
|
||||
typedef mz_ulong uLong;
|
||||
typedef Byte Bytef;
|
||||
typedef uInt uIntf;
|
||||
typedef char charf;
|
||||
typedef int intf;
|
||||
typedef void *voidpf;
|
||||
typedef uLong uLongf;
|
||||
typedef void *voidp;
|
||||
typedef void *const voidpc;
|
||||
#define Z_NULL 0
|
||||
@ -204,10 +202,6 @@ DECLSPEC void *memset(u8 *s, int c, u32 len){
|
||||
#define MZ_MIN(a, b) (((a) < (b)) ? (a) : (b))
|
||||
#define MZ_DEFAULT_WINDOW_BITS 15
|
||||
#define TINFL_LZ_DICT_SIZE 32768
|
||||
#define TINFL_MEMCPY(d, s, l) memcpy(d, s, l)
|
||||
#define TINFL_MEMCPY_G(d, s, l, p) memcpy_g(d, s, l, p)
|
||||
#define TINFL_MEMSET(p, c, l) memset(p, c, (u32)l)
|
||||
#define MZ_CLEAR_OBJ(obj) memset(&(obj), 0, sizeof(obj))
|
||||
|
||||
// hashcat-patched/hashcat-specific:
|
||||
#ifdef CRC32_IN_INFLATE
|
||||
@ -583,7 +577,7 @@ DECLSPEC tinfl_status tinfl_decompress(tinfl_decompressor *r, MAYBE_GLOBAL const
|
||||
TINFL_CR_RETURN(38, (decomp_flags & TINFL_FLAG_HAS_MORE_INPUT) ? TINFL_STATUS_NEEDS_MORE_INPUT : TINFL_STATUS_FAILED_CANNOT_MAKE_PROGRESS);
|
||||
}
|
||||
n = MZ_MIN(MZ_MIN((size_t)(pOut_buf_end - pOut_buf_cur), (size_t)(pIn_buf_end - pIn_buf_cur)), counter);
|
||||
TINFL_MEMCPY_G(pOut_buf_cur, pIn_buf_cur, n, pStream);
|
||||
memcpy_g(pOut_buf_cur, pIn_buf_cur, n, pStream);
|
||||
pIn_buf_cur += n;
|
||||
pOut_buf_cur += n;
|
||||
counter -= (mz_uint)n;
|
||||
@ -601,7 +595,7 @@ DECLSPEC tinfl_status tinfl_decompress(tinfl_decompressor *r, MAYBE_GLOBAL const
|
||||
mz_uint i;
|
||||
r->m_table_sizes[0] = 288;
|
||||
r->m_table_sizes[1] = 32;
|
||||
TINFL_MEMSET(r->m_tables[1].m_code_size, 5, 32);
|
||||
memset(r->m_tables[1].m_code_size, 5, 32);
|
||||
for (i = 0; i <= 143; ++i)
|
||||
*p++ = 8;
|
||||
for (; i <= 255; ++i)
|
||||
@ -618,7 +612,8 @@ DECLSPEC tinfl_status tinfl_decompress(tinfl_decompressor *r, MAYBE_GLOBAL const
|
||||
TINFL_GET_BITS(11, r->m_table_sizes[counter], "\05\05\04"[counter]);
|
||||
r->m_table_sizes[counter] += s_min_table_sizes[counter];
|
||||
}
|
||||
MZ_CLEAR_OBJ(r->m_tables[2].m_code_size);
|
||||
memset(r->m_tables[2].m_code_size, 0, TINFL_MAX_HUFF_SYMBOLS_0);
|
||||
|
||||
for (counter = 0; counter < r->m_table_sizes[2]; counter++)
|
||||
{
|
||||
mz_uint s;
|
||||
@ -633,9 +628,11 @@ DECLSPEC tinfl_status tinfl_decompress(tinfl_decompressor *r, MAYBE_GLOBAL const
|
||||
tinfl_huff_table *pTable;
|
||||
mz_uint i, j, used_syms, total, sym_index, next_code[17], total_syms[16];
|
||||
pTable = &r->m_tables[r->m_type];
|
||||
MZ_CLEAR_OBJ(total_syms);
|
||||
MZ_CLEAR_OBJ(pTable->m_look_up);
|
||||
MZ_CLEAR_OBJ(pTable->m_tree);
|
||||
|
||||
memset((u8 *) total_syms, 0, 64);
|
||||
memset((u8 *) pTable->m_look_up, 0, TINFL_FAST_LOOKUP_SIZE * 2);
|
||||
memset((u8 *) pTable->m_tree, 0, TINFL_MAX_HUFF_SYMBOLS_0 * 2 * 2);
|
||||
|
||||
for (i = 0; i < r->m_table_sizes[r->m_type]; ++i)
|
||||
total_syms[pTable->m_code_size[i]]++;
|
||||
used_syms = 0, total = 0;
|
||||
@ -707,15 +704,18 @@ DECLSPEC tinfl_status tinfl_decompress(tinfl_decompressor *r, MAYBE_GLOBAL const
|
||||
num_extra = "\02\03\07"[dist - 16];
|
||||
TINFL_GET_BITS(18, s, num_extra);
|
||||
s += "\03\03\013"[dist - 16];
|
||||
TINFL_MEMSET(r->m_len_codes + counter, (dist == 16) ? r->m_len_codes[counter - 1] : 0, s);
|
||||
|
||||
memset(r->m_len_codes + counter, (dist == 16) ? r->m_len_codes[counter - 1] : 0, s);
|
||||
|
||||
|
||||
counter += s;
|
||||
}
|
||||
if ((r->m_table_sizes[0] + r->m_table_sizes[1]) != counter)
|
||||
{
|
||||
TINFL_CR_RETURN_FOREVER(21, TINFL_STATUS_FAILED);
|
||||
}
|
||||
TINFL_MEMCPY(r->m_tables[0].m_code_size, r->m_len_codes, r->m_table_sizes[0]);
|
||||
TINFL_MEMCPY(r->m_tables[1].m_code_size, r->m_len_codes + r->m_table_sizes[0], r->m_table_sizes[1]);
|
||||
memcpy(r->m_tables[0].m_code_size, r->m_len_codes, r->m_table_sizes[0]);
|
||||
memcpy(r->m_tables[1].m_code_size, r->m_len_codes + r->m_table_sizes[0], r->m_table_sizes[1]);
|
||||
}
|
||||
}
|
||||
for (;;)
|
||||
|
||||
@ -24,7 +24,7 @@ typedef struct
|
||||
|
||||
} 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 )); }
|
||||
@ -41,15 +41,6 @@ inline __device__ uint4 rotate (const uint4 a, const int n)
|
||||
|
||||
#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)
|
||||
{
|
||||
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));
|
||||
|
||||
#if defined IS_CUDA
|
||||
#if defined IS_CUDA || defined IS_HIP
|
||||
|
||||
#define SALSA20_2R() \
|
||||
{ \
|
||||
|
||||
@ -31,7 +31,7 @@ typedef struct ethereum_scrypt
|
||||
|
||||
} 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 )); }
|
||||
@ -48,15 +48,6 @@ inline __device__ uint4 rotate (const uint4 a, const int n)
|
||||
|
||||
#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)
|
||||
{
|
||||
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));
|
||||
|
||||
#if defined IS_CUDA
|
||||
#if defined IS_CUDA || defined IS_HIP
|
||||
|
||||
#define SALSA20_2R() \
|
||||
{ \
|
||||
|
||||
@ -72,7 +72,7 @@ DECLSPEC int is_valid_bitcoinj (const u32 *w)
|
||||
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 )); }
|
||||
@ -89,15 +89,6 @@ inline __device__ uint4 rotate (const uint4 a, const int n)
|
||||
|
||||
#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)
|
||||
{
|
||||
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));
|
||||
|
||||
#if defined IS_CUDA
|
||||
#if defined IS_CUDA || defined IS_HIP
|
||||
|
||||
#define SALSA20_2R() \
|
||||
{ \
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
##
|
||||
|
||||
- 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
|
||||
|
||||
@ -10,8 +10,8 @@
|
||||
hashcat v6.2.3
|
||||
==============
|
||||
|
||||
AMD GPUs on Linux require "RadeonOpenCompute (ROCm)" Software Platform (3.1 or later)
|
||||
AMD GPUs on Windows require "AMD Radeon Adrenalin 2020 Edition" (20.2.2 or later)
|
||||
AMD GPUs on Linux require "AMD ROCm" (4.4 or later)
|
||||
AMD GPUs on Windows require "AMD Radeon Adrenalin 2020 Edition" (21.2.1 or later)
|
||||
Intel CPUs require "OpenCL Runtime for Intel Core and Intel Xeon Processors" (16.1.1 or later)
|
||||
NVIDIA GPUs require "NVIDIA Driver" (440.64 or later) and "CUDA Toolkit" (9.0 or later)
|
||||
|
||||
|
||||
@ -89,53 +89,50 @@ 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_hipCreateProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram *prog, const char *src, const char *name, int numHeaders, const char * const *headers, const char * const *includeNames);
|
||||
int hc_hipDestroyProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram *prog);
|
||||
int hc_hipCompileProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, int numOptions, const char * const *options);
|
||||
int hc_hipGetProgramLogSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *logSizeRet);
|
||||
int hc_hipGetProgramLog (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *log);
|
||||
int hc_hipGetCodeSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *codeSizeRet);
|
||||
int hc_hipGetCode (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *code);
|
||||
|
||||
int hc_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, 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_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, hipCtx_t *pctx, unsigned int flags, hipDevice_t dev);
|
||||
int hc_hipCtxDestroy (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx);
|
||||
int hc_hipCtxPopCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t *pctx);
|
||||
int hc_hipCtxPushCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx);
|
||||
int hc_hipCtxSetCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx);
|
||||
int hc_hipCtxSynchronize (hashcat_ctx_t *hashcat_ctx);
|
||||
int hc_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_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_hipDeviceGetName (hashcat_ctx_t *hashcat_ctx, char *name, int len, hipDevice_t dev);
|
||||
int hc_hipDeviceTotalMem (hashcat_ctx_t *hashcat_ctx, size_t *bytes, hipDevice_t dev);
|
||||
int hc_hipDriverGetVersion (hashcat_ctx_t *hashcat_ctx, int *driverVersion);
|
||||
int hc_hipEventCreate (hashcat_ctx_t *hashcat_ctx, HIPevent *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_hipEventCreate (hashcat_ctx_t *hashcat_ctx, hipEvent_t *phEvent, unsigned int Flags);
|
||||
int hc_hipEventDestroy (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent);
|
||||
int hc_hipEventElapsedTime (hashcat_ctx_t *hashcat_ctx, float *pMilliseconds, hipEvent_t hStart, hipEvent_t hEnd);
|
||||
int hc_hipEventQuery (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent);
|
||||
int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent, hipStream_t hStream);
|
||||
int hc_hipEventSynchronize (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent);
|
||||
int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, hipFunction_attribute attrib, hipFunction_t hfunc);
|
||||
int hc_hipInit (hashcat_ctx_t *hashcat_ctx, unsigned int Flags);
|
||||
int hc_hipLaunchKernel (hashcat_ctx_t *hashcat_ctx, HIPfunction 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_hipLaunchKernel (hashcat_ctx_t *hashcat_ctx, hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra);
|
||||
int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize);
|
||||
int hc_hipMemFree (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dptr);
|
||||
int hc_hipMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount);
|
||||
int hc_hipMemcpyDtoHAsync (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream);
|
||||
int hc_hipMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount);
|
||||
int hc_hipMemcpyDtoDAsync (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream);
|
||||
int hc_hipMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount);
|
||||
int hc_hipMemcpyHtoDAsync (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount, hipStream_t hStream);
|
||||
int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, hipFunction_t *hfunc, hipModule_t hmod, const char *name);
|
||||
int hc_hipModuleGetGlobal (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name);
|
||||
int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, hipModule_t *module, const void *image, unsigned int numOptions, hipJitOption *options, void **optionValues);
|
||||
int hc_hipModuleUnload (hashcat_ctx_t *hashcat_ctx, hipModule_t hmod);
|
||||
int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, hipStream_t *phStream, unsigned int Flags);
|
||||
int hc_hipStreamDestroy (hashcat_ctx_t *hashcat_ctx, hipStream_t hStream);
|
||||
int hc_hipStreamSynchronize (hashcat_ctx_t *hashcat_ctx, hipStream_t hStream);
|
||||
|
||||
int 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);
|
||||
@ -184,10 +181,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_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_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_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_t buf, const u64 num);
|
||||
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_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_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 num);
|
||||
|
||||
1437
include/ext_hip.h
1437
include/ext_hip.h
File diff suppressed because it is too large
Load Diff
@ -6,41 +6,26 @@
|
||||
#ifndef _EXT_HIPRTC_H
|
||||
#define _EXT_HIPRTC_H
|
||||
|
||||
/**
|
||||
* from hip_runtime.h (/opt/rocm/hip/include/hip/amd_detail/hiprtc.h)
|
||||
*/
|
||||
// start: amd_detail/hiprtc.h
|
||||
|
||||
/**
|
||||
* \ingroup error
|
||||
* \brief The enumerated type hiprtcResult defines API call result codes.
|
||||
* HIPRTC API functions return hiprtcResult to indicate the call
|
||||
* result.
|
||||
*/
|
||||
typedef enum {
|
||||
HIPRTC_SUCCESS = 0,
|
||||
HIPRTC_ERROR_OUT_OF_MEMORY = 1,
|
||||
HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2,
|
||||
HIPRTC_ERROR_INVALID_INPUT = 3,
|
||||
HIPRTC_ERROR_INVALID_PROGRAM = 4,
|
||||
HIPRTC_ERROR_INVALID_OPTION = 5,
|
||||
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
|
||||
typedef enum hiprtcResult {
|
||||
HIPRTC_SUCCESS = 0,
|
||||
HIPRTC_ERROR_OUT_OF_MEMORY = 1,
|
||||
HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2,
|
||||
HIPRTC_ERROR_INVALID_INPUT = 3,
|
||||
HIPRTC_ERROR_INVALID_PROGRAM = 4,
|
||||
HIPRTC_ERROR_INVALID_OPTION = 5,
|
||||
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;
|
||||
|
||||
/**
|
||||
* \ingroup compilation
|
||||
* \brief hiprtcProgram is the unit of compilation, and an opaque handle for
|
||||
* 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;
|
||||
typedef struct _hiprtcProgram* hiprtcProgram;
|
||||
|
||||
// stop: amd_detail/hiprtc.h
|
||||
|
||||
#ifdef _WIN32
|
||||
#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_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_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_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_HIPRTCGETPROGRAMLOGSIZE) (hiprtcProgram, size_t *);
|
||||
typedef const char * (HIPRTC_API_CALL *HIPRTC_HIPRTCGETERRORSTRING) (hiprtcResult);
|
||||
typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCVERSION) (int *, int *);
|
||||
typedef const char * (HIPRTC_API_CALL *HIPRTC_HIPRTCGETERRORSTRING) (hiprtcResult);
|
||||
|
||||
typedef struct hc_hiprtc_lib
|
||||
{
|
||||
@ -70,13 +54,12 @@ typedef struct hc_hiprtc_lib
|
||||
HIPRTC_HIPRTCCOMPILEPROGRAM hiprtcCompileProgram;
|
||||
HIPRTC_HIPRTCCREATEPROGRAM hiprtcCreateProgram;
|
||||
HIPRTC_HIPRTCDESTROYPROGRAM hiprtcDestroyProgram;
|
||||
HIPRTC_HIPRTCGETCODE hiprtcGetCode;
|
||||
HIPRTC_HIPRTCGETCODESIZE hiprtcGetCodeSize;
|
||||
HIPRTC_HIPRTCGETLOWEREDNAME hiprtcGetLoweredName;
|
||||
HIPRTC_HIPRTCGETPTX hiprtcGetCode;
|
||||
HIPRTC_HIPRTCGETPTXSIZE hiprtcGetCodeSize;
|
||||
HIPRTC_HIPRTCGETPROGRAMLOG hiprtcGetProgramLog;
|
||||
HIPRTC_HIPRTCGETPROGRAMLOGSIZE hiprtcGetProgramLogSize;
|
||||
HIPRTC_HIPRTCGETERRORSTRING hiprtcGetErrorString;
|
||||
HIPRTC_HIPRTCVERSION hiprtcVersion;
|
||||
|
||||
} hc_hiprtc_lib_t;
|
||||
|
||||
|
||||
143
include/types.h
143
include/types.h
@ -1502,80 +1502,80 @@ typedef struct hc_device_param
|
||||
|
||||
int hip_warp_size;
|
||||
|
||||
HIPdevice hip_device;
|
||||
HIPcontext hip_context;
|
||||
HIPstream hip_stream;
|
||||
hipDevice_t hip_device;
|
||||
hipCtx_t hip_context;
|
||||
hipStream_t hip_stream;
|
||||
|
||||
HIPevent hip_event1;
|
||||
HIPevent hip_event2;
|
||||
hipEvent_t hip_event1;
|
||||
hipEvent_t hip_event2;
|
||||
|
||||
HIPmodule hip_module;
|
||||
HIPmodule hip_module_shared;
|
||||
HIPmodule hip_module_mp;
|
||||
HIPmodule hip_module_amp;
|
||||
hipModule_t hip_module;
|
||||
hipModule_t hip_module_shared;
|
||||
hipModule_t hip_module_mp;
|
||||
hipModule_t hip_module_amp;
|
||||
|
||||
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_bzero;
|
||||
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;
|
||||
hipFunction_t hip_function1;
|
||||
hipFunction_t hip_function12;
|
||||
hipFunction_t hip_function2p;
|
||||
hipFunction_t hip_function2;
|
||||
hipFunction_t hip_function2e;
|
||||
hipFunction_t hip_function23;
|
||||
hipFunction_t hip_function3;
|
||||
hipFunction_t hip_function4;
|
||||
hipFunction_t hip_function_init2;
|
||||
hipFunction_t hip_function_loop2p;
|
||||
hipFunction_t hip_function_loop2;
|
||||
hipFunction_t hip_function_mp;
|
||||
hipFunction_t hip_function_mp_l;
|
||||
hipFunction_t hip_function_mp_r;
|
||||
hipFunction_t hip_function_amp;
|
||||
hipFunction_t hip_function_tm;
|
||||
hipFunction_t hip_function_memset;
|
||||
hipFunction_t hip_function_bzero;
|
||||
hipFunction_t hip_function_atinit;
|
||||
hipFunction_t hip_function_utf8toutf16le;
|
||||
hipFunction_t hip_function_decompress;
|
||||
hipFunction_t hip_function_aux1;
|
||||
hipFunction_t hip_function_aux2;
|
||||
hipFunction_t hip_function_aux3;
|
||||
hipFunction_t 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_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_t hip_d_pws_buf;
|
||||
hipDeviceptr_t hip_d_pws_amp_buf;
|
||||
hipDeviceptr_t hip_d_pws_comp_buf;
|
||||
hipDeviceptr_t hip_d_pws_idx;
|
||||
hipDeviceptr_t hip_d_rules;
|
||||
hipDeviceptr_t hip_d_rules_c;
|
||||
hipDeviceptr_t hip_d_combs;
|
||||
hipDeviceptr_t hip_d_combs_c;
|
||||
hipDeviceptr_t hip_d_bfs;
|
||||
hipDeviceptr_t hip_d_bfs_c;
|
||||
hipDeviceptr_t hip_d_tm_c;
|
||||
hipDeviceptr_t hip_d_bitmap_s1_a;
|
||||
hipDeviceptr_t hip_d_bitmap_s1_b;
|
||||
hipDeviceptr_t hip_d_bitmap_s1_c;
|
||||
hipDeviceptr_t hip_d_bitmap_s1_d;
|
||||
hipDeviceptr_t hip_d_bitmap_s2_a;
|
||||
hipDeviceptr_t hip_d_bitmap_s2_b;
|
||||
hipDeviceptr_t hip_d_bitmap_s2_c;
|
||||
hipDeviceptr_t hip_d_bitmap_s2_d;
|
||||
hipDeviceptr_t hip_d_plain_bufs;
|
||||
hipDeviceptr_t hip_d_digests_buf;
|
||||
hipDeviceptr_t hip_d_digests_shown;
|
||||
hipDeviceptr_t hip_d_salt_bufs;
|
||||
hipDeviceptr_t hip_d_esalt_bufs;
|
||||
hipDeviceptr_t hip_d_tmps;
|
||||
hipDeviceptr_t hip_d_hooks;
|
||||
hipDeviceptr_t hip_d_result;
|
||||
hipDeviceptr_t hip_d_extra0_buf;
|
||||
hipDeviceptr_t hip_d_extra1_buf;
|
||||
hipDeviceptr_t hip_d_extra2_buf;
|
||||
hipDeviceptr_t hip_d_extra3_buf;
|
||||
hipDeviceptr_t hip_d_root_css_buf;
|
||||
hipDeviceptr_t hip_d_markov_css_buf;
|
||||
hipDeviceptr_t hip_d_st_digests_buf;
|
||||
hipDeviceptr_t hip_d_st_salts_buf;
|
||||
hipDeviceptr_t hip_d_st_esalts_buf;
|
||||
|
||||
// API: opencl
|
||||
|
||||
@ -1726,8 +1726,7 @@ typedef struct backend_ctx
|
||||
int rc_hip_init;
|
||||
int rc_hiprtc_init;
|
||||
|
||||
int hiprtc_driver_version;
|
||||
int hip_driver_version;
|
||||
int hip_driverVersion;
|
||||
|
||||
// opencl
|
||||
|
||||
|
||||
1703
src/backend.c
1703
src/backend.c
File diff suppressed because it is too large
Load Diff
@ -170,11 +170,6 @@ bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE
|
||||
// it leads to CL_KERNEL_WORK_GROUP_SIZE to return 0 and later we will divide with 0
|
||||
// workaround would be to rewrite kernel to use global memory
|
||||
|
||||
if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
|
||||
if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD)
|
||||
{
|
||||
return true;
|
||||
|
||||
@ -170,11 +170,6 @@ bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE
|
||||
// it leads to CL_KERNEL_WORK_GROUP_SIZE to return 0 and later we will divide with 0
|
||||
// workaround would be to rewrite kernel to use global memory
|
||||
|
||||
if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
|
||||
if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD)
|
||||
{
|
||||
return true;
|
||||
|
||||
@ -170,11 +170,6 @@ bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE
|
||||
// it leads to CL_KERNEL_WORK_GROUP_SIZE to return 0 and later we will divide with 0
|
||||
// workaround would be to rewrite kernel to use global memory
|
||||
|
||||
if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
|
||||
if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD)
|
||||
{
|
||||
return true;
|
||||
|
||||
@ -818,9 +818,9 @@ void backend_info (hashcat_ctx_t *hashcat_ctx)
|
||||
event_log_info (hashcat_ctx, NULL);
|
||||
|
||||
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);
|
||||
|
||||
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)
|
||||
{
|
||||
int hip_devices_cnt = backend_ctx->hip_devices_cnt;
|
||||
int hip_driver_version = backend_ctx->hip_driver_version;
|
||||
int hip_devices_cnt = backend_ctx->hip_devices_cnt;
|
||||
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 };
|
||||
|
||||
|
||||
Loading…
Reference in New Issue
Block a user