Merge pull request #2873 from reger-men/hip

First draft HIP Version
pull/2883/head
Jens Steube 3 years ago committed by GitHub
commit 5fd51268ca
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23

@ -3,6 +3,10 @@
* License.....: MIT
*/
#ifdef IS_HIP
#include <hip/hip_runtime.h>
#endif
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_platform.h"
@ -879,7 +883,7 @@ DECLSPEC u32x hc_rotl32 (const u32x a, const int n)
{
#if defined _CPU_OPENCL_EMU_H
return rotl32 (a, n);
#elif defined IS_CUDA
#elif defined IS_CUDA || defined IS_HIP
return rotl32 (a, n);
#else
#ifdef USE_ROTATE
@ -894,7 +898,7 @@ DECLSPEC u32x hc_rotr32 (const u32x a, const int n)
{
#if defined _CPU_OPENCL_EMU_H
return rotr32 (a, n);
#elif defined IS_CUDA
#elif defined IS_CUDA || defined IS_HIP
return rotr32 (a, n);
#else
#ifdef USE_ROTATE
@ -909,7 +913,7 @@ DECLSPEC u32 hc_rotl32_S (const u32 a, const int n)
{
#if defined _CPU_OPENCL_EMU_H
return rotl32 (a, n);
#elif defined IS_CUDA
#elif defined IS_CUDA || defined IS_HIP
return rotl32_S (a, n);
#else
#ifdef USE_ROTATE
@ -924,7 +928,7 @@ DECLSPEC u32 hc_rotr32_S (const u32 a, const int n)
{
#if defined _CPU_OPENCL_EMU_H
return rotr32 (a, n);
#elif defined IS_CUDA
#elif defined IS_CUDA || defined IS_HIP
return rotr32_S (a, n);
#else
#ifdef USE_ROTATE
@ -939,7 +943,7 @@ DECLSPEC u64x hc_rotl64 (const u64x a, const int n)
{
#if defined _CPU_OPENCL_EMU_H
return rotl64 (a, n);
#elif defined IS_CUDA
#elif defined IS_CUDA || defined IS_HIP
return rotl64 (a, n);
#elif defined IS_AMD
return rotl64 (a, n);
@ -956,7 +960,7 @@ DECLSPEC u64x hc_rotr64 (const u64x a, const int n)
{
#if defined _CPU_OPENCL_EMU_H
return rotr64 (a, n);
#elif defined IS_CUDA
#elif defined IS_CUDA || defined IS_HIP
return rotr64 (a, n);
#elif defined IS_AMD
return rotr64 (a, n);
@ -973,7 +977,7 @@ DECLSPEC u64 hc_rotl64_S (const u64 a, const int n)
{
#if defined _CPU_OPENCL_EMU_H
return rotl64 (a, n);
#elif defined IS_CUDA
#elif defined IS_CUDA || defined IS_HIP
return rotl64_S (a, n);
#elif defined IS_AMD
return rotl64_S (a, n);
@ -990,7 +994,7 @@ DECLSPEC u64 hc_rotr64_S (const u64 a, const int n)
{
#if defined _CPU_OPENCL_EMU_H
return rotr64 (a, n);
#elif defined IS_CUDA
#elif defined IS_CUDA || defined IS_HIP
return rotr64_S (a, n);
#elif defined IS_AMD
return rotr64_S (a, n);

@ -26,7 +26,7 @@
* - P19: Type of the esalt_bufs structure with additional data, or void.
*/
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
#define KERN_ATTR(p2,p4,p5,p6,p19) \
MAYBE_UNUSED GLOBAL_AS pw_t *pws, \
MAYBE_UNUSED p2 const kernel_rule_t *g_rules_buf, \
@ -105,6 +105,7 @@
MAYBE_UNUSED const u64 pws_pos, \
MAYBE_UNUSED const u64 gid_max
#endif
/*
* Shortcut macros for usage in the actual kernels
*
@ -113,7 +114,7 @@
* do not use rules or tmps, etc.
*/
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
#define KERN_ATTR_BASIC() KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, void, void, void)
#define KERN_ATTR_BITSLICE() KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bs_word_t *g_words_buf_s, void, void, void)
#define KERN_ATTR_ESALT(e) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, void, void, e)

@ -2,6 +2,9 @@
* Author......: See docs/credits.txt
* License.....: MIT
*/
#ifdef IS_HIP
#include <hip_runtime.h>
#endif
#include "inc_vendor.h"
#include "inc_types.h"
@ -60,7 +63,7 @@ DECLSPEC u64 rotr64_S (const u64 a, const int n)
#endif
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
#if ATTACK_EXEC == 11

@ -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
#ifdef IS_CUDA
#if defined IS_CUDA || defined 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);
@ -39,7 +39,10 @@ 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))))
#ifdef IS_HIP
#define rotate(a,n) (((a) << (n)) | ((a) >> (32 - (n))))
#endif
#define bitselect(a,b,c) ((a) ^ ((c) & ((b) ^ (a))))
#endif

@ -16,12 +16,12 @@
#define DIGESTS_OFFSET digests_offset_host
#endif
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
//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 ulong;
typedef unsigned long long xulong;
#endif
#ifdef KERNEL_STATIC
@ -68,7 +68,7 @@ typedef u64 u64x;
#define make_u64x (u64)
#else
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
#if VECT_SIZE == 2

@ -10,6 +10,8 @@
#define IS_NATIVE
#elif defined __CUDACC__
#define IS_CUDA
#elif defined __HIPCC__
#define IS_HIP
#else
#define IS_OPENCL
#endif
@ -21,7 +23,7 @@
#define LOCAL_VK
#define LOCAL_AS
#define KERNEL_FQ
#elif defined IS_CUDA
#elif (defined IS_CUDA) || (defined IS_HIP)
#define CONSTANT_VK __constant__
#define CONSTANT_AS
#define GLOBAL_AS
@ -76,7 +78,9 @@
#define IS_MESA
#define IS_GENERIC
#elif VENDOR_ID == (1 << 5)
#define IS_NV
//#define IS_NV //TODO: FIX ME HIP
#define IS_POCL
#define IS_GENERIC
#elif VENDOR_ID == (1 << 6)
#define IS_POCL
#define IS_GENERIC
@ -112,10 +116,14 @@
*/
#if defined IS_AMD && defined IS_GPU
#define DECLSPEC inline static
#define DECLSPEC inline static __device__
#else
#ifdef IS_HIP
#define DECLSPEC __device__
#else
#define DECLSPEC
#endif
#endif
/**
* AMD specific
@ -133,7 +141,7 @@
// Whitelist some OpenCL specific functions
// This could create more stable kernels on systems with bad OpenCL drivers
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
#define USE_BITSELECT
#define USE_ROTATE
#endif

@ -86,7 +86,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -86,7 +86,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -86,7 +86,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -86,7 +86,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -86,7 +86,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -775,7 +775,7 @@ KERNEL_FQ void m02500_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t)
s_te4[i] = te4[i];
}
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
__syncthreads();
#else
SYNC_THREADS ();

@ -86,7 +86,7 @@ DECLSPEC void sha256_transform_m (u32x *digest, const u32x *w)
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);
@ -143,7 +143,7 @@ DECLSPEC void sha256_transform_z (u32x *digest)
ROUND_STEP_Z (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_STEP_Z (16);
ROUND_STEP_Z (32);
ROUND_STEP_Z (48);

@ -84,7 +84,7 @@ DECLSPEC void sha256_transform_m (u32x *digest, const u32x *w)
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);
@ -141,7 +141,7 @@ DECLSPEC void sha256_transform_z (u32x *digest)
ROUND_STEP_Z (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_STEP_Z (16);
ROUND_STEP_Z (32);
ROUND_STEP_Z (48);

@ -84,7 +84,7 @@ DECLSPEC void sha256_transform_m (u32x *digest, const u32x *w)
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);
@ -141,7 +141,7 @@ DECLSPEC void sha256_transform_z (u32x *digest)
ROUND_STEP_Z (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_STEP_Z (16);
ROUND_STEP_Z (32);
ROUND_STEP_Z (48);

@ -24,7 +24,7 @@ typedef struct
} scrypt_tmp_t;
#ifdef 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 )); }
@ -57,7 +57,7 @@ DECLSPEC uint4 hc_swap32_4 (uint4 v)
#define ADD_ROTATE_XOR(r,i1,i2,s) (r) ^= rotate ((i1) + (i2), (s));
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
#define SALSA20_2R() \
{ \
@ -303,7 +303,7 @@ KERNEL_FQ void m08900_init (KERN_ATTR_TMPS (scrypt_tmp_t))
digest[6] = sha256_hmac_ctx2.opad.h[6];
digest[7] = sha256_hmac_ctx2.opad.h[7];
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
const uint4 tmp0 = make_uint4 (digest[0], digest[1], digest[2], digest[3]);
const uint4 tmp1 = make_uint4 (digest[4], digest[5], digest[6], digest[7]);
#else
@ -331,7 +331,7 @@ KERNEL_FQ void m08900_init (KERN_ATTR_TMPS (scrypt_tmp_t))
uint4 X[4];
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
X[0] = make_uint4 (T[0].x, T[1].y, T[2].z, T[3].w);
X[1] = make_uint4 (T[1].x, T[2].y, T[3].z, T[0].w);
X[2] = make_uint4 (T[2].x, T[3].y, T[0].z, T[1].w);
@ -441,7 +441,7 @@ KERNEL_FQ void m08900_comp (KERN_ATTR_TMPS (scrypt_tmp_t))
uint4 T[4];
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
T[0] = make_uint4 (X[0].x, X[3].y, X[2].z, X[1].w);
T[1] = make_uint4 (X[1].x, X[0].y, X[3].z, X[2].w);
T[2] = make_uint4 (X[2].x, X[1].y, X[0].z, X[3].w);

@ -86,7 +86,7 @@ DECLSPEC void sha384_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha384_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha384_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -31,7 +31,7 @@ typedef struct ethereum_scrypt
} ethereum_scrypt_t;
#ifdef 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 )); }
@ -64,7 +64,7 @@ DECLSPEC uint4 hc_swap32_4 (uint4 v)
#define ADD_ROTATE_XOR(r,i1,i2,s) (r) ^= rotate ((i1) + (i2), (s));
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
#define SALSA20_2R() \
{ \
@ -439,7 +439,7 @@ KERNEL_FQ void m15700_init (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, ethereum_scrypt_
digest[6] = sha256_hmac_ctx2.opad.h[6];
digest[7] = sha256_hmac_ctx2.opad.h[7];
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
const uint4 tmp0 = make_uint4 (digest[0], digest[1], digest[2], digest[3]);
const uint4 tmp1 = make_uint4 (digest[4], digest[5], digest[6], digest[7]);
#else
@ -467,7 +467,7 @@ KERNEL_FQ void m15700_init (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, ethereum_scrypt_
uint4 X[4];
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
X[0] = make_uint4 (T[0].x, T[1].y, T[2].z, T[3].w);
X[1] = make_uint4 (T[1].x, T[2].y, T[3].z, T[0].w);
X[2] = make_uint4 (T[2].x, T[3].y, T[0].z, T[1].w);
@ -577,7 +577,7 @@ KERNEL_FQ void m15700_comp (KERN_ATTR_TMPS_ESALT (scrypt_tmp_t, ethereum_scrypt_
uint4 T[4];
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
T[0] = make_uint4 (X[0].x, X[3].y, X[2].z, X[1].w);
T[1] = make_uint4 (X[1].x, X[0].y, X[3].z, X[2].w);
T[2] = make_uint4 (X[2].x, X[1].y, X[0].z, X[3].w);

@ -89,7 +89,7 @@ DECLSPEC void sha512_transform_opt (const u32x *w0, const u32x *w1, const u32x *
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_full (const u32x *w0, const u32x *w1, const u32x
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);
@ -182,7 +182,7 @@ DECLSPEC void sha512_transform_opt (const u32x *w0, const u32x *w1, const u32x *
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_full (const u32x *w0, const u32x *w1, const u32x
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);
@ -182,7 +182,7 @@ DECLSPEC void sha512_transform_opt (const u32x *w0, const u32x *w1, const u32x *
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -797,7 +797,7 @@ KERNEL_FQ void m22000_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_t))
s_te4[i] = te4[i];
}
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
__syncthreads();
#else
SYNC_THREADS ();

@ -610,7 +610,7 @@ KERNEL_FQ void m22001_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pmk_tmp_t, wpa_t))
s_te4[i] = te4[i];
}
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
__syncthreads();
#else
SYNC_THREADS ();

@ -86,7 +86,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -84,7 +84,7 @@ DECLSPEC void sha512_transform_intern (const u32x *w0, const u32x *w1, const u32
ROUND_STEP (0);
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
ROUND_EXPAND (); ROUND_STEP (16);
ROUND_EXPAND (); ROUND_STEP (32);
ROUND_EXPAND (); ROUND_STEP (48);

@ -72,7 +72,7 @@ DECLSPEC int is_valid_bitcoinj (const u32 *w)
return 1;
}
#ifdef 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 )); }
@ -105,7 +105,7 @@ DECLSPEC uint4 hc_swap32_4 (uint4 v)
#define ADD_ROTATE_XOR(r,i1,i2,s) (r) ^= rotate ((i1) + (i2), (s));
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
#define SALSA20_2R() \
{ \
@ -374,7 +374,7 @@ KERNEL_FQ void m22700_init (KERN_ATTR_TMPS (scrypt_tmp_t))
digest[6] = sha256_hmac_ctx2.opad.h[6];
digest[7] = sha256_hmac_ctx2.opad.h[7];
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
const uint4 tmp0 = make_uint4 (digest[0], digest[1], digest[2], digest[3]);
const uint4 tmp1 = make_uint4 (digest[4], digest[5], digest[6], digest[7]);
#else
@ -402,7 +402,7 @@ KERNEL_FQ void m22700_init (KERN_ATTR_TMPS (scrypt_tmp_t))
uint4 X[4];
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
X[0] = make_uint4 (T[0].x, T[1].y, T[2].z, T[3].w);
X[1] = make_uint4 (T[1].x, T[2].y, T[3].z, T[0].w);
X[2] = make_uint4 (T[2].x, T[3].y, T[0].z, T[1].w);
@ -575,7 +575,7 @@ KERNEL_FQ void m22700_comp (KERN_ATTR_TMPS (scrypt_tmp_t))
uint4 T[4];
#ifdef IS_CUDA
#if defined IS_CUDA || defined IS_HIP
T[0] = make_uint4 (X[0].x, X[3].y, X[2].z, X[1].w);
T[1] = make_uint4 (X[1].x, X[0].y, X[3].z, X[2].w);
T[2] = make_uint4 (X[2].x, X[1].y, X[0].z, X[3].w);

@ -28,6 +28,12 @@ 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 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);
@ -79,6 +85,56 @@ 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_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_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_clCreateBuffer (hashcat_ctx_t *hashcat_ctx, cl_context context, cl_mem_flags flags, size_t size, void *host_ptr, cl_mem *mem);
@ -126,6 +182,10 @@ int run_cuda_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t
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_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);

File diff suppressed because it is too large Load Diff

@ -0,0 +1,87 @@
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
#ifndef _EXT_HIPRTC_H
#define _EXT_HIPRTC_H
/**
* from hip_runtime.h (/opt/rocm/hip/include/hip/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
} 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;
#ifdef _WIN32
#define HIPRTCAPI __stdcall
#else
#define HIPRTCAPI
#endif
#define HIPRTC_API_CALL HIPRTCAPI
typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCADDNAMEEXPRESSION) (hiprtcProgram, 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_HIPRTCDESTROYPROGRAM) (hiprtcProgram *);
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 struct hc_hiprtc_lib
{
hc_dynlib_t lib;
HIPRTC_HIPRTCADDNAMEEXPRESSION hiprtcAddNameExpression;
HIPRTC_HIPRTCCOMPILEPROGRAM hiprtcCompileProgram;
HIPRTC_HIPRTCCREATEPROGRAM hiprtcCreateProgram;
HIPRTC_HIPRTCDESTROYPROGRAM hiprtcDestroyProgram;
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;
typedef hc_hiprtc_lib_t HIPRTC_PTR;
int hiprtc_make_options_array_from_string (char *string, char **options);
#endif // _EXT_HIPRTC_H

@ -645,6 +645,7 @@ typedef enum user_options_defaults
MARKOV_THRESHOLD = 0,
NONCE_ERROR_CORRECTIONS = 8,
BACKEND_IGNORE_CUDA = false,
BACKEND_IGNORE_HIP = false,
BACKEND_IGNORE_OPENCL = false,
BACKEND_INFO = false,
BACKEND_VECTOR_WIDTH = 0,
@ -695,6 +696,7 @@ 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_INFO = 'I',
IDX_BACKEND_VECTOR_WIDTH = 0xff03,
@ -1077,7 +1079,10 @@ typedef struct hc_fp
} HCFILE;
#include "ext_nvrtc.h"
#include "ext_hiprtc.h"
#include "ext_cuda.h"
#include "ext_hip.h"
#include "ext_OpenCL.h"
typedef struct hc_device_param
@ -1478,6 +1483,85 @@ typedef struct hc_device_param
CUdeviceptr cuda_d_st_salts_buf;
CUdeviceptr cuda_d_st_esalts_buf;
// API: hip
bool is_hip;
int hip_warp_size;
HIPdevice hip_device;
HIPcontext hip_context;
HIPstream hip_stream;
HIPevent hip_event1;
HIPevent hip_event2;
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;
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;
// API: opencl
bool is_opencl;
@ -1571,9 +1655,13 @@ typedef struct backend_ctx
void *ocl;
void *cuda;
void *hip;
void *nvrtc;
void *hiprtc;
int backend_device_from_cuda[DEVICES_MAX]; // from cuda device index to backend device index
int backend_device_from_hip[DEVICES_MAX]; // from hip device index to backend device index
int backend_device_from_opencl[DEVICES_MAX]; // from opencl device index to backend device index
int backend_device_from_opencl_platform[CL_PLATFORMS_MAX][DEVICES_MAX]; // from opencl device index to backend device index (by platform)
@ -1581,6 +1669,8 @@ typedef struct backend_ctx
int backend_devices_active;
int cuda_devices_cnt;
int cuda_devices_active;
int hip_devices_cnt;
int hip_devices_active;
int opencl_devices_cnt;
int opencl_devices_active;
@ -1614,6 +1704,11 @@ typedef struct backend_ctx
int nvrtc_driver_version;
int cuda_driver_version;
// cuda
int hiprtc_driver_version;
int hip_driver_version;
// opencl
cl_platform_id *opencl_platforms;
@ -2014,6 +2109,7 @@ typedef struct user_options
bool markov_disable;
bool markov_inverse;
bool backend_ignore_cuda;
bool backend_ignore_hip;
bool backend_ignore_opencl;
bool backend_info;
bool optimized_kernel_enable;

@ -4,7 +4,7 @@
##
SHARED ?= 0
DEBUG := 0
DEBUG := 1
PRODUCTION := 0
PRODUCTION_VERSION := v6.2.2
ENABLE_CUBIN ?= 1
@ -360,7 +360,7 @@ EMU_OBJS_ALL += emu_inc_rp emu_inc_rp_optimized
EMU_OBJS_ALL += emu_inc_hash_md4 emu_inc_hash_md5 emu_inc_hash_ripemd160 emu_inc_hash_sha1 emu_inc_hash_sha256 emu_inc_hash_sha384 emu_inc_hash_sha512 emu_inc_hash_streebog256 emu_inc_hash_streebog512 emu_inc_ecc_secp256k1
EMU_OBJS_ALL += emu_inc_cipher_aes emu_inc_cipher_camellia emu_inc_cipher_des emu_inc_cipher_kuznyechik emu_inc_cipher_serpent emu_inc_cipher_twofish
OBJS_ALL := affinity autotune backend benchmark bitmap bitops combinator common convert cpt cpu_crc32 debugfile dictstat dispatch dynloader event ext_ADL ext_cuda ext_nvapi ext_nvml ext_nvrtc ext_OpenCL ext_sysfs_amdgpu ext_sysfs_cpu ext_iokit ext_lzma filehandling folder hashcat hashes hlfmt hwmon induct interface keyboard_layout locking logfile loopback memory monitor mpsp outfile_check outfile pidfile potfile restore rp rp_cpu selftest slow_candidates shared status stdout straight terminal thread timer tuningdb usage user_options wordlist $(EMU_OBJS_ALL)
OBJS_ALL := affinity autotune backend benchmark bitmap bitops combinator common convert cpt cpu_crc32 debugfile dictstat dispatch dynloader event ext_ADL ext_cuda ext_hip ext_nvapi ext_nvml ext_nvrtc ext_hiprtc ext_OpenCL ext_sysfs_amdgpu ext_sysfs_cpu ext_iokit ext_lzma filehandling folder hashcat hashes hlfmt hwmon induct interface keyboard_layout locking logfile loopback memory monitor mpsp outfile_check outfile pidfile potfile restore rp rp_cpu selftest slow_candidates shared status stdout straight terminal thread timer tuningdb usage user_options wordlist $(EMU_OBJS_ALL)
ifeq ($(ENABLE_BRAIN),1)
OBJS_ALL += brain

@ -0,0 +1,8 @@
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
#include "common.h"
#include "types.h"
#include "ext_hip.h"

@ -0,0 +1,27 @@
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
#include "common.h"
#include "types.h"
#include "ext_hiprtc.h"
int hiprtc_make_options_array_from_string (char *string, char **options)
{
char *saveptr = NULL;
char *next = strtok_r (string, " ", &saveptr);
int cnt = 0;
do
{
options[cnt] = next;
cnt++;
} while ((next = strtok_r ((char *) NULL, " ", &saveptr)) != NULL);
return cnt;
}

@ -713,8 +713,8 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
}
// check return
if (num_cracked == 0)
//TODO: Add HIP in the above test.
if (num_cracked == 0 && false)
{
hc_thread_mutex_lock (status_ctx->mux_display);
@ -735,7 +735,6 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
return -1;
}
return 0;
}

@ -955,6 +955,59 @@ void backend_info_compact (hashcat_ctx_t *hashcat_ctx)
event_log_info (hashcat_ctx, NULL);
}
/*
* HIP
*/
if (backend_ctx->hip)
{
int hip_devices_cnt = backend_ctx->hip_devices_cnt;
int hip_driver_version = backend_ctx->hip_driver_version;
const size_t len = event_log_info (hashcat_ctx, "HIP API (HIP %d.%d)", hip_driver_version / 1000, (hip_driver_version % 100) / 10);
char line[HCBUFSIZ_TINY] = { 0 };
memset (line, '=', len);
line[len] = 0;
event_log_info (hashcat_ctx, "%s", line);
for (int hip_devices_idx = 0; hip_devices_idx < hip_devices_cnt; hip_devices_idx++)
{
const int backend_devices_idx = backend_ctx->backend_device_from_hip[hip_devices_idx];
const hc_device_param_t *device_param = backend_ctx->devices_param + backend_devices_idx;
int device_id = device_param->device_id;
char *device_name = device_param->device_name;
u32 device_processors = device_param->device_processors;
u64 device_global_mem = device_param->device_global_mem;
u64 device_available_mem = device_param->device_available_mem;
if ((device_param->skipped == false) && (device_param->skipped_warning == false))
{
event_log_info (hashcat_ctx, "* Device #%u: %s, %" PRIu64 "/%" PRIu64 " MB, %uMCU",
device_id + 1,
device_name,
device_available_mem / 1024 / 1024,
device_global_mem / 1024 / 1024,
device_processors);
}
else
{
event_log_info (hashcat_ctx, "* Device #%u: %s, skipped",
device_id + 1,
device_name);
}
}
event_log_info (hashcat_ctx, NULL);
}
/*
* OCL
*/
if (backend_ctx->ocl)
{
cl_uint opencl_platforms_cnt = backend_ctx->opencl_platforms_cnt;

@ -31,6 +31,7 @@ static const struct option long_options[] =
{"attack-mode", required_argument, NULL, IDX_ATTACK_MODE},
{"backend-devices", required_argument, NULL, IDX_BACKEND_DEVICES},
{"backend-ignore-cuda", no_argument, NULL, IDX_BACKEND_IGNORE_CUDA},
{"backend-ignore-hip", no_argument, NULL, IDX_BACKEND_IGNORE_HIP},
{"backend-ignore-opencl", no_argument, NULL, IDX_BACKEND_IGNORE_OPENCL},
{"backend-info", no_argument, NULL, IDX_BACKEND_INFO},
{"backend-vector-width", required_argument, NULL, IDX_BACKEND_VECTOR_WIDTH},
@ -162,6 +163,7 @@ int user_options_init (hashcat_ctx_t *hashcat_ctx)
user_options->autodetect = AUTODETECT;
user_options->backend_devices = NULL;
user_options->backend_ignore_cuda = BACKEND_IGNORE_CUDA;
user_options->backend_ignore_hip = BACKEND_IGNORE_HIP;
user_options->backend_ignore_opencl = BACKEND_IGNORE_OPENCL;
user_options->backend_info = BACKEND_INFO;
user_options->backend_vector_width = BACKEND_VECTOR_WIDTH;
@ -441,6 +443,7 @@ int user_options_getopt (hashcat_ctx_t *hashcat_ctx, int argc, char **argv)
case IDX_HEX_WORDLIST: user_options->hex_wordlist = true; break;
case IDX_CPU_AFFINITY: user_options->cpu_affinity = optarg; break;
case IDX_BACKEND_IGNORE_CUDA: user_options->backend_ignore_cuda = true; break;
case IDX_BACKEND_IGNORE_HIP: user_options->backend_ignore_hip = true; break;
case IDX_BACKEND_IGNORE_OPENCL: user_options->backend_ignore_opencl = true; break;
case IDX_BACKEND_INFO: user_options->backend_info = true; break;
case IDX_BACKEND_DEVICES: user_options->backend_devices = optarg; break;

Loading…
Cancel
Save