Jens Steube
fb7bb04587
Do not use dynamic shared memory if dynamic_local_mem_size is a multiple of local_mem_size
2020-02-02 11:15:37 +01:00
Jens Steube
aef53f7e10
OpenCL Runtime: Allow the kernel to access post-48k shared memory region on CUDA. Requires both module and kernel preparation
2020-02-01 14:27:42 +01:00
Jens Steube
5e0eb288c9
Use __launch_bounds__ in CUDA as replacement for reqd_work_group_size() in OpenCL
2019-06-16 18:01:26 +02:00
Jens Steube
f689532e4c
Move P-box initializer values to constant memory in -m 3200
2019-06-03 14:40:51 +02:00
Jens Steube
03b2d3fb69
Fix use of LOCAL_VK and LOCAL_AS in -m 3200
2019-05-07 12:08:54 +02:00
Jens Steube
8ff8c5d536
Add LOCAL_VK to make use of __shared__
2019-05-07 09:01:32 +02:00
Jens Steube
d0bd33c9d1
Rename CONSTANT_AS to CONSTANT_VK
2019-05-06 14:34:16 +02:00
Jens Steube
89119bf24a
Add missing inc_platform.h include
2019-04-26 13:59:43 +02:00
jsteube
7c6970dbdd
Remove hard-coded static keyword from OpenCL kernels
2019-04-13 18:46:19 +02:00
jsteube
9ced13cc94
Get rid of CONSTSPEC macro in OpenCL kernels
2019-04-04 10:15:34 +02:00
Jens Steube
0fb3b3c83e
Declare internal functions in OpenCL kernels as static
2019-03-26 11:03:25 +01:00
jsteube
66d94b06e4
Get rid of src/rp_kernel_on_cpu.c and src/rp_kernel_on_cpu_optimized.c and use OpenCL emulated kernel version
2019-03-25 12:24:04 +01:00
jsteube
e80b1838e8
Rename some functions in inc_common.cl to avoid conflicts with bitops.c
2019-03-23 22:15:38 +01:00
jsteube
adeeaee84a
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 22:27:58 +01:00
jsteube
c9d60c079f
Prepare OpenCL kernels for non-static compilation
2019-03-22 15:16:25 +01:00
jsteube
7d4bea41a0
Get rid of OpenCL/inc_hash_constants.h and OpenCL/inc_hash_functions.cl
2019-03-21 23:00:38 +01:00
jsteube
5ef67a8ab7
Apply previous blowfish optimization for -m 3200 also on -m 9000 and -m 18600
2019-03-18 12:42:47 +01:00
jsteube
5ecbcde945
Cracking bcrypt: Use a feedback from the OpenCL runtime to dynamically find out optimal thread count
2019-03-17 14:17:35 +01:00
jsteube
26d7602bbd
No longer need reqd_work_group_size() in OpenCL kernels since modules taken care of this using threads_min and threads_max
2019-03-10 12:13:14 +01:00
Jens Steube
4bce25dd9d
Move remaining module specific typedefs and structures from inc_types.cl to kernels
2019-03-08 10:18:20 +01:00
jsteube
5da1e4b872
Fixed maximum password length limit which was announced as 256 but actually was 255
2019-02-26 21:20:07 +01:00
Jens Steube
2a6444c05a
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 15:26:46 +01:00
R. Yushaev
5de004103a
Replace kernel parameter lists with macros
...
Substitute long parameter lists in ~2900 kernel function declarations
with macros. This cleans up the code, reduces probability of copy-paste
errors and highlights the differences between kernel functions. Also
reduces the size of the OpenCL folder by ~3 MB.
2018-11-16 11:44:33 +01:00
Jens Steube
1c280e4a6e
Small performance boost for bcrypt on CPU
2018-08-02 14:20:04 +02:00
Jens Steube
a43d3ad176
Rename some hashcat specific OpenCL functions to avoid conflicts with existing OpenCL functions from OpenCL runtime
2018-07-22 12:20:20 +02:00
Jens Steube
02a2495349
Switched array pointer types in function declarations in order to be compatible with OpenCL 2.0
2018-07-22 11:47:42 +02:00
jsteube
32d6b3e10e
OpenCL kernels: Add '-pure' prefix to kernel filenames to avoid problems caused by reusing existing hashcat installation folder
2018-06-20 14:18:17 +02:00
