Skip to content

[pull] master from hashcat:master #44

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
pull merged 1 commit into from
Feb 1, 2020
Merged

[pull] master from hashcat:master #44

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
25 changes: 25 additions & 0 deletions OpenCL/m03200-pure.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -356,6 +356,10 @@ CONSTANT_VK u32a c_pbox[18] =
L ^= P[17]; \
}

#ifdef DYNAMIC_LOCAL
extern __shared__ u32 lm[];
#endif

DECLSPEC void expand_key (u32 *E, u32 *W, const int len)
{
u8 *E_ptr = (u8 *) E;
Expand Down Expand Up @@ -456,6 +460,12 @@ KERNEL_FQ void FIXED_THREAD_COUNT(FIXED_LOCAL_SIZE) m03200_init (KERN_ATTR_TMPS
P[i] = c_pbox[i];
}

#ifdef DYNAMIC_LOCAL
u32 *S0 = lm + (lid * 1024) + 0;
u32 *S1 = lm + (lid * 1024) + 256;
u32 *S2 = lm + (lid * 1024) + 512;
u32 *S3 = lm + (lid * 1024) + 768;
#else
LOCAL_VK u32 S0_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S1_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S2_all[FIXED_LOCAL_SIZE][256];
Expand All @@ -465,6 +475,7 @@ KERNEL_FQ void FIXED_THREAD_COUNT(FIXED_LOCAL_SIZE) m03200_init (KERN_ATTR_TMPS
LOCAL_AS u32 *S1 = S1_all[lid];
LOCAL_AS u32 *S2 = S2_all[lid];
LOCAL_AS u32 *S3 = S3_all[lid];
#endif

for (u32 i = 0; i < 256; i++)
{
Expand Down Expand Up @@ -614,6 +625,12 @@ KERNEL_FQ void FIXED_THREAD_COUNT(FIXED_LOCAL_SIZE) m03200_loop (KERN_ATTR_TMPS
P[i] = tmps[gid].P[i];
}

#ifdef DYNAMIC_LOCAL
u32 *S0 = lm + (lid * 1024) + 0;
u32 *S1 = lm + (lid * 1024) + 256;
u32 *S2 = lm + (lid * 1024) + 512;
u32 *S3 = lm + (lid * 1024) + 768;
#else
LOCAL_VK u32 S0_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S1_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S2_all[FIXED_LOCAL_SIZE][256];
Expand All @@ -623,6 +640,7 @@ KERNEL_FQ void FIXED_THREAD_COUNT(FIXED_LOCAL_SIZE) m03200_loop (KERN_ATTR_TMPS
LOCAL_AS u32 *S1 = S1_all[lid];
LOCAL_AS u32 *S2 = S2_all[lid];
LOCAL_AS u32 *S3 = S3_all[lid];
#endif

for (u32 i = 0; i < 256; i++)
{
Expand Down Expand Up @@ -799,6 +817,12 @@ KERNEL_FQ void FIXED_THREAD_COUNT(FIXED_LOCAL_SIZE) m03200_comp (KERN_ATTR_TMPS
P[i] = tmps[gid].P[i];
}

#ifdef DYNAMIC_LOCAL
u32 *S0 = lm + (lid * 1024) + 0;
u32 *S1 = lm + (lid * 1024) + 256;
u32 *S2 = lm + (lid * 1024) + 512;
u32 *S3 = lm + (lid * 1024) + 768;
#else
LOCAL_VK u32 S0_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S1_all[FIXED_LOCAL_SIZE][256];
LOCAL_VK u32 S2_all[FIXED_LOCAL_SIZE][256];
Expand All @@ -808,6 +832,7 @@ KERNEL_FQ void FIXED_THREAD_COUNT(FIXED_LOCAL_SIZE) m03200_comp (KERN_ATTR_TMPS
LOCAL_AS u32 *S1 = S1_all[lid];
LOCAL_AS u32 *S2 = S2_all[lid];
LOCAL_AS u32 *S3 = S3_all[lid];
#endif

for (u32 i = 0; i < 256; i++)
{
Expand Down
1 change: 1 addition & 0 deletions docs/changes.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -115,6 +115,7 @@
- OpenCL Runtime: Workaround JiT compiler error on ROCm 2.3 driver if the 'inline' keyword is used in function declaration
- OpenCL Runtime: Workaround memory allocation error on AMD driver on Windows leading to CL_MEM_OBJECT_ALLOCATION_FAILURE
- OpenCL Runtime: Workaround ROCm OpenCL driver problem trying to write temporary file into readonly folder by setting TMPDIR
- OpenCL Runtime: Allow the kernel to access post-48k shared memory region on CUDA. Requires both module and kernel preparation
- Startup Checks: Improved the pidfile check: Do not just check for existing PID but also check executable filename
- Startup Checks: Prevent the user to modify options which are overwritten automatically in benchmark mode
- Startup Screen: Add extra warning when using --force
Expand Down
21 changes: 21 additions & 0 deletions include/types.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1132,6 +1132,27 @@ typedef struct hc_device_param
u64 kernel_local_mem_size_aux3;
u64 kernel_local_mem_size_aux4;

u64 kernel_dynamic_local_mem_size1;
u64 kernel_dynamic_local_mem_size12;
u64 kernel_dynamic_local_mem_size2;
u64 kernel_dynamic_local_mem_size23;
u64 kernel_dynamic_local_mem_size3;
u64 kernel_dynamic_local_mem_size4;
u64 kernel_dynamic_local_mem_size_init2;
u64 kernel_dynamic_local_mem_size_loop2;
u64 kernel_dynamic_local_mem_size_mp;
u64 kernel_dynamic_local_mem_size_mp_l;
u64 kernel_dynamic_local_mem_size_mp_r;
u64 kernel_dynamic_local_mem_size_amp;
u64 kernel_dynamic_local_mem_size_tm;
u64 kernel_dynamic_local_mem_size_memset;
u64 kernel_dynamic_local_mem_size_atinit;
u64 kernel_dynamic_local_mem_size_decompress;
u64 kernel_dynamic_local_mem_size_aux1;
u64 kernel_dynamic_local_mem_size_aux2;
u64 kernel_dynamic_local_mem_size_aux3;
u64 kernel_dynamic_local_mem_size_aux4;

u32 kernel_accel;
u32 kernel_accel_prev;
u32 kernel_accel_min;
Expand Down
Loading