Gabriele Gristina
b3d3b31c3e
Metal: added support for vectors up to 4
2022-02-10 21:53:08 +01:00
Jens Steube
be75e4b4ea
Rename STR() to M2S() to avoid future collisions and move from kernel source to command line parameter
2022-02-07 09:57:08 +01:00
Gabriele Gristina
9d36245d51
Kernels: Set the default Address Space Qualifier for any pointer, refactored / updated KERN_ATTR macros and rc4 cipher functions, in order to support Apple Metal runtime
2022-02-04 19:54:00 +01:00
Gabriele Gristina
2e4a136758
Refactored standard kernel includes in order to support Apple Metal runtime, updated backend, test units and status code
2022-01-18 22:52:14 +01:00
Jens Steube
0abdcb1ae5
Rename GID_MAX to GID_CNT to avoid naming conflict with existing macro
2022-01-04 22:57:26 +01:00
Jens Steube
668d2179cd
Kernels: Refactored standard kernel declaration to use a structure holding u32/u64 attributes to reduce the number of attributes
2022-01-04 18:07:18 +01:00
Jens Steube
04d5e5a119
New Attack-Mode: Association Attack. Like JtR's single mode. Very early
...
stage. See hashcat Forum for detailed writeup.
2020-09-29 15:56:32 +02:00
Jens Steube
b1d5f92c2d
Do not use __local keyword in -m 5500 for devices without real shared memory
2020-03-06 15:48:01 +01:00
Jens Steube
6b7d064118
Replace (u32x) (...) with make_u32x (...)
2019-05-08 15:21:22 +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
Jens Steube
00e1e32492
Replace barrier() with SYNC_THREADS()
2019-04-26 13:34:07 +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
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
dc9279c95c
New Strategy: Instead of using volatile, mark the mode as unstable. Remove all volatiles
2019-03-03 19:18:56 +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
Jens Steube
6d39fb1feb
Make all kernel parameter macro helper a function
2018-11-16 14:17:01 +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
jsteube
68bff94980
Workaround rocm OpenCL runtime bug when copy data from constant to local memory
2018-08-12 18:04:33 +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
e4e1c1d515
We can't mix inline functions with static constants
2018-02-07 14:16:27 +01:00
jsteube
3e08750900
OpenCL Kernels: Add general function declaration keyword (static inline)
2018-02-06 19:12:24 +01:00
jsteube
53f3da9f63
OpenCL Kernels: Use static declaraction for uXXa variables used in __constant space
2018-01-18 23:19:31 +01:00
jsteube
1d04de3a8e
Limit kernel-loops in straight-mode to 256, therefore allow rules to be stored in constant memory
2017-08-23 12:43:59 +02:00
jsteube
319799bbbf
Switch the datatypes of the variables responsible for work-item count and work-item size from u32 to u64
2017-08-19 16:39:22 +02:00
jsteube
34d882a116
Rename inc_rp.X to inc_rp_optimized.X
2017-08-11 11:25:47 +02:00
jsteube
beb6ee2061
Add OPTI_TYPE_OPTIMIZED_KERNEL
...
Rename unconverted fast hash kernels to optimized kernels
Finalize some converted fast hashes to default kernels
2017-07-18 14:45:15 +02:00
