DeforaNetworks/hashcat
3
0
Fork 0
Code Releases Activity

Add missing backports from code base v6.2.2

Browse Source 
Fix context to thread management
Fix missing code in selftest.c, autotune.c, hashes.c, dispatch.c and backend.c
Use IS_HIP depending code makes it easier for future optimization related to inline assembly calls - instead of using IS_CUDA || IS_HIP
See TODO markers for more optimizations / next steps
This commit is contained in:
Jens Steube
2021-07-11 12:38:59 +02:00
parent 5fd51268ca
commit 1b84a9e53b
15 changed files with 4485 additions and 277 deletions
Download Patch File Download Diff File Expand all files Collapse all files
OpenCL/inc_common.cl
-4
View File
@@ -3,10 +3,6 @@
* License.....: MIT
*/
#ifdef IS_HIP
#include <hip/hip_runtime.h>
#endif
#include "inc_vendor.h"
#include "inc_types.h"
#include "inc_platform.h"
OpenCL/inc_common.h
-1
View File
@@ -105,7 +105,6 @@
MAYBE_UNUSED const u64 pws_pos, \
MAYBE_UNUSED const u64 gid_max
#endif
/*
* Shortcut macros for usage in the actual kernels
*
OpenCL/inc_platform.cl
+105 -4
View File
@@ -2,9 +2,6 @@
* Author......: See docs/credits.txt
* License.....: MIT
*/
#ifdef IS_HIP
#include <hip_runtime.h>
#endif
#include "inc_vendor.h"
#include "inc_types.h"
@@ -63,7 +60,111 @@ DECLSPEC u64 rotr64_S (const u64 a, const int n)
#endif
#if defined IS_CUDA || defined IS_HIP
#if defined IS_CUDA
#if ATTACK_EXEC == 11
CONSTANT_VK u32 generic_constant[8192]; // 32k
#if ATTACK_KERN == 0
#define bfs_buf g_bfs_buf
#define rules_buf ((const kernel_rule_t *) generic_constant)
#define words_buf_s g_words_buf_s
#define words_buf_r g_words_buf_r
#elif ATTACK_KERN == 1
#define bfs_buf g_bfs_buf
#define rules_buf g_rules_buf
#define words_buf_s g_words_buf_s
#define words_buf_r g_words_buf_r
#elif ATTACK_KERN == 3
#define rules_buf g_rules_buf
#define bfs_buf ((const bf_t *) generic_constant)
#define words_buf_s ((const bs_word_t *) generic_constant)
#define words_buf_r ((const u32x *) generic_constant)
#endif
#endif
DECLSPEC u32 hc_atomic_dec (GLOBAL_AS u32 *p)
{
volatile const u32 val = 1;
return atomicSub (p, val);
}
DECLSPEC u32 hc_atomic_inc (GLOBAL_AS u32 *p)
{
volatile const u32 val = 1;
return atomicAdd (p, val);
}
DECLSPEC u32 hc_atomic_or (GLOBAL_AS u32 *p, volatile const u32 val)
{
return atomicOr (p, val);
}
DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused)))
{
return (blockIdx.x * blockDim.x) + threadIdx.x;
}
DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused)))
{
return threadIdx.x;
}
DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
{
// verify
return blockDim.x;
}
DECLSPEC u32x rotl32 (const u32x a, const int n)
{
return ((a << n) | ((a >> (32 - n))));
}
DECLSPEC u32x rotr32 (const u32x a, const int n)
{
return ((a >> n) | ((a << (32 - n))));
}
DECLSPEC u32 rotl32_S (const u32 a, const int n)
{
return ((a << n) | ((a >> (32 - n))));
}
DECLSPEC u32 rotr32_S (const u32 a, const int n)
{
return ((a >> n) | ((a << (32 - n))));
}
DECLSPEC u64x rotl64 (const u64x a, const int n)
{
return ((a << n) | ((a >> (64 - n))));
}
DECLSPEC u64x rotr64 (const u64x a, const int n)
{
return ((a >> n) | ((a << (64 - n))));
}
DECLSPEC u64 rotl64_S (const u64 a, const int n)
{
return ((a << n) | ((a >> (64 - n))));
}
DECLSPEC u64 rotr64_S (const u64 a, const int n)
{
return ((a >> n) | ((a << (64 - n))));
}
#define FIXED_THREAD_COUNT(n) __launch_bounds__((n), 0)
#define SYNC_THREADS() __syncthreads ()
#endif
#if defined IS_HIP
#if ATTACK_EXEC == 11
OpenCL/inc_platform.h
+22 -3
View File
@@ -21,7 +21,7 @@ DECLSPEC u64 rotl64_S (const u64 a, const int n);
DECLSPEC u64 rotr64_S (const u64 a, const int n);
#endif
#if defined IS_CUDA || defined IS_HIP
#ifdef IS_CUDA
DECLSPEC u32 hc_atomic_dec (volatile GLOBAL_AS u32 *p);
DECLSPEC u32 hc_atomic_inc (volatile GLOBAL_AS u32 *p);
DECLSPEC u32 hc_atomic_or (volatile GLOBAL_AS u32 *p, volatile const u32 val);
@@ -39,10 +39,29 @@ DECLSPEC u64x rotr64 (const u64x a, const int n);
DECLSPEC u64 rotl64_S (const u64 a, const int n);
DECLSPEC u64 rotr64_S (const u64 a, const int n);
#ifdef IS_HIP
#define rotate(a,n) (((a) << (n)) | ((a) >> (32 - (n))))
//#define rotate(a,n) (((a) << (n)) | ((a) >> (32 - (n))))
#define bitselect(a,b,c) ((a) ^ ((c) & ((b) ^ (a))))
#endif
#ifdef IS_HIP
DECLSPEC u32 hc_atomic_dec (volatile GLOBAL_AS u32 *p);
DECLSPEC u32 hc_atomic_inc (volatile GLOBAL_AS u32 *p);
DECLSPEC u32 hc_atomic_or (volatile GLOBAL_AS u32 *p, volatile const u32 val);
DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused)));
DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused)));
DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)));
DECLSPEC u32x rotl32 (const u32x a, const int n);
DECLSPEC u32x rotr32 (const u32x a, const int n);
DECLSPEC u32 rotl32_S (const u32 a, const int n);
DECLSPEC u32 rotr32_S (const u32 a, const int n);
DECLSPEC u64x rotl64 (const u64x a, const int n);
DECLSPEC u64x rotr64 (const u64x a, const int n);
DECLSPEC u64 rotl64_S (const u64 a, const int n);
DECLSPEC u64 rotr64_S (const u64 a, const int n);
//#define rotate(a,n) (((a) << (n)) | ((a) >> (32 - (n))))
#define bitselect(a,b,c) ((a) ^ ((c) & ((b) ^ (a))))
#endif
OpenCL/inc_types.h
+3 -3
View File
@@ -16,12 +16,12 @@
#define DIGESTS_OFFSET digests_offset_host
#endif
#if defined IS_CUDA || defined IS_HIP
#ifdef IS_CUDA
//https://docs.nvidia.com/cuda/nvrtc/index.html#integer-size
typedef unsigned char uchar;
typedef unsigned short ushort;
typedef unsigned int uint;
typedef unsigned long long xulong;
typedef unsigned long long ulong;
#endif
#ifdef KERNEL_STATIC
@@ -68,7 +68,7 @@ typedef u64 u64x;
#define make_u64x (u64)
#else
#if defined IS_CUDA || defined IS_HIP
#if defined IS_CUDA
#if VECT_SIZE == 2
OpenCL/inc_vendor.h
+28 -10
View File
@@ -16,6 +16,10 @@
#define IS_OPENCL
#endif
#ifdef IS_HIP
#include <hip/hip_runtime.h>
#endif
#if defined IS_NATIVE
#define CONSTANT_VK
#define CONSTANT_AS
@@ -23,7 +27,14 @@
#define LOCAL_VK
#define LOCAL_AS
#define KERNEL_FQ
#elif (defined IS_CUDA) || (defined IS_HIP)
#elif defined IS_CUDA
#define CONSTANT_VK __constant__
#define CONSTANT_AS
#define GLOBAL_AS
#define LOCAL_VK __shared__
#define LOCAL_AS
#define KERNEL_FQ extern "C" __global__
#elif defined IS_HIP
#define CONSTANT_VK __constant__
#define CONSTANT_AS
#define GLOBAL_AS
@@ -78,12 +89,14 @@
#define IS_MESA
#define IS_GENERIC
#elif VENDOR_ID == (1 << 5)
//#define IS_NV //TODO: FIX ME HIP
#define IS_POCL
#define IS_GENERIC
#define IS_NV
#elif VENDOR_ID == (1 << 6)
#define IS_POCL
#define IS_GENERIC
#elif VENDOR_ID == (1 << 8)
#define IS_AMD_USE_HIP
// TODO HIP optimization potential
#define IS_GENERIC
#else
#define IS_GENERIC
#endif
@@ -116,14 +129,12 @@
*/
#if defined IS_AMD && defined IS_GPU
#define DECLSPEC inline static __device__
#else
#ifdef IS_HIP
#define DECLSPEC __device__
#define DECLSPEC inline static
#elif defined IS_HIP
#define DECLSPEC __device__
#else
#define DECLSPEC
#endif
#endif
/**
* AMD specific
@@ -141,11 +152,18 @@
// Whitelist some OpenCL specific functions
// This could create more stable kernels on systems with bad OpenCL drivers
#if defined IS_CUDA || defined IS_HIP
#ifdef IS_CUDA
#define USE_BITSELECT
#define USE_ROTATE
#endif
#ifdef IS_HIP
//TODO HIP
//#define USE_BITSELECT
//#define USE_ROTATE
//#define USE_SWIZZLE
#endif
#ifdef IS_ROCM
#define USE_BITSELECT
#define USE_ROTATE
OpenCL/shared.cl
+5
View File
@@ -126,6 +126,11 @@ KERNEL_FQ void gpu_memset (GLOBAL_AS uint4 *buf, const u32 value, const u64 gid_
r.y = value;
r.z = value;
r.w = value;
#elif defined IS_HIP
r.x = value;
r.y = value;
r.z = value;
r.w = value;
#endif
buf[gid] = r;