From 2f4003cded144c7264605850e5227d63e77c74be Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Fri, 15 Apr 2022 20:33:11 +0200 Subject: [PATCH] Synchronize SCRYPT code for -m 28200 with -m 8900 as source --- OpenCL/m28200-pure.cl | 4 ++-- src/modules/module_28200.c | 16 ++++++++++------ 2 files changed, 12 insertions(+), 8 deletions(-) diff --git a/OpenCL/m28200-pure.cl b/OpenCL/m28200-pure.cl index a34757aba..5e8df55ea 100644 --- a/OpenCL/m28200-pure.cl +++ b/OpenCL/m28200-pure.cl @@ -34,7 +34,7 @@ typedef struct exodus } exodus_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 )); } @@ -67,7 +67,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() \ { \ diff --git a/src/modules/module_28200.c b/src/modules/module_28200.c index de43be20a..f28730534 100644 --- a/src/modules/module_28200.c +++ b/src/modules/module_28200.c @@ -25,7 +25,8 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_MP_MULTI_DISABLE | OPTS_TYPE_NATIVE_THREADS - | OPTS_TYPE_LOOP_PREPARE; + | OPTS_TYPE_LOOP_PREPARE + | OPTS_TYPE_SELF_TEST_DISABLE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "EXODUS:16384:8:1:IYkXZgFETRmFp4wQXyP8XMe3LtuOw8wMdLcBVQ+9YWE=:lq0W9ekN5sC0O7Xw:UD4a6mUUhkTbQtGWitXHZUg0pQ4RHI6W/KUyYE95m3k=:ZuNQckXOtr4r21x+DT1zpQ=="; @@ -503,13 +504,16 @@ const char *module_extra_tuningdb_block (MAYBE_UNUSED const hashconfig_t *hashco const char *extra_tuningdb_block = "DEVICE_TYPE_CPU * 28200 1 N A\n" "DEVICE_TYPE_GPU * 28200 1 N A\n" - "GeForce_GTX_980 * 28200 1 16 A\n" - "GeForce_GTX_1080 * 28200 1 45 A\n" + "GeForce_GTX_980 * 28200 1 29 A\n" + "GeForce_GTX_1080 * 28200 1 15 A\n" "GeForce_RTX_2080_Ti * 28200 1 68 A\n" - "GeForce_RTX_3080 * 28200 1 68 A\n" + "GeForce_RTX_3060_Ti * 28200 1 51 A\n" + "GeForce_RTX_3070 * 28200 1 46 A\n" "GeForce_RTX_3090 * 28200 1 82 A\n" - "ALIAS_AMD_Vega64 * 28200 1 60 A\n" - "ALIAS_AMD_RX6900XT * 28200 1 60 A\n" + "ALIAS_AMD_RX480 * 28200 1 15 A\n" + "ALIAS_AMD_Vega64 * 28200 1 28 A\n" + "ALIAS_AMD_MI100 * 28200 1 79 A\n" + "ALIAS_AMD_RX6900XT * 28200 1 59 A\n" ; return extra_tuningdb_block;