Jens Steube
8ebf4b9858
Fixed 'E' rule in pure kernel mode which was ignoring letters that are in positions that are multiples of 4
3 years ago
Jens Steube
9dffc69089
Merge pull request #2448 from philsmd/refactor_cosmetic_fix
...
cosmetic: minor code style fixes
4 years ago
philsmd
e59f61e8cf
cosmetic: minor code style fixes
4 years ago
philsmd
1e2bc78fd0
rule engine: add zero-length check for rule 'z'
4 years ago
Jens Steube
27df7429ce
Fix current password length check in y/Y rules in pure kernel mode
4 years ago
Jens Steube
f5527bb937
Fix mangle_dupeword_times() in OpenCL/inc_rp.cl
5 years ago
Jens Steube
a74cbe3461
Fixed out-of-boundary read in pure kernel rule engine rule 'p' if parameter is set to 2 or higher
5 years ago
Jens Steube
d9a92afecc
Change out-of-boundary fix in order to re-enable password length 256 with rules in pure kernel mode
5 years ago
Jens Steube
270210a8ab
Fix out-of-boundary read in rule engines
5 years ago
Jens Steube
8ff8c5d536
Add LOCAL_VK to make use of __shared__
5 years ago
Jens Steube
89119bf24a
Add missing inc_platform.h include
6 years ago
jsteube
ffd8ec9001
Fixed output password of 'e' rule in pure and cpu rule engine if separator character is also the first letter
6 years ago
jsteube
ea7d9c50b6
Fix missing const keyword in some includes
6 years ago
jsteube
fb8a9d7c40
Get rid of cpu_sha1.c
6 years ago
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
6 years ago
jsteube
194fd7e6d1
Fix some invalid code sections caused from conversion
6 years ago
jsteube
977199698f
Move some macros from .h to .cl sources
6 years ago
jsteube
e80b1838e8
Rename some functions in inc_common.cl to avoid conflicts with bitops.c
6 years ago
jsteube
adeeaee84a
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
6 years ago
jsteube
5f5468be6f
Add missing DECLSPEC in OpenCL rule functions
6 years ago
jsteube
305a044ec6
Remove some old function headers
6 years ago
jsteube
84d6b8ecc1
Add function prototypes in OpenCL kernels to make some compilers happy
6 years ago
jsteube
874635cc49
Do not use a vector function to write into a scalar variable even if vectorize support is disabled
7 years ago
jsteube
512fb5f6fb
No inline keyword for rules
7 years ago
jsteube
3e08750900
OpenCL Kernels: Add general function declaration keyword (static inline)
7 years ago
jsteube
503f676cb0
Remove static keywords from OpenCL functions, they can cause old NV JiT compiler to fail
7 years ago
mhasbini
de7ccd88ef
Fix overflow in mangle_dupechar_last function
7 years ago
jsteube
1d04de3a8e
Limit kernel-loops in straight-mode to 256, therefore allow rules to be stored in constant memory
7 years ago
jsteube
8853884f2a
Fix append_four_byte() in case sm8 is 0
7 years ago
jsteube
f32e113942
Add missing case in append_block() in pure kernel rule engine
7 years ago
jsteube
ad1ce462d1
Get rid of ceil() in OpenCL kernels
7 years ago
jsteube
5b5bdf3889
Replace some SIMD related function calls
7 years ago
jsteube
967e96728d
Make all the OpenCL kernel function includes static
7 years ago
jsteube
ec874c1d59
Optimized the following pure kernel rule engine functions:
...
- mangle_lrest()
- mangle_lrest_ufirst()
- mangle_urest()
- mangle_urest_lfirst()
- mangle_trest()
- mangle_toggle_at()
- mangle_reverse()
- mangle_dupeword()
- mangle_reflect()
- mangle_rotate_left()
- mangle_rotate_right()
- mangle_switch_first()
- mangle_switch_last()
- mangle_switch_at()
- mangle_title_sep()
- mangle_title_sep()
Added some helper functions:
- generate_cmask()
- append_four_byte()
- append_three_byte()
- append_two_byte()
- append_one_byte()
- append_block()
- exchange_byte()
Removed some helper functions:
- upper_at()
- lower_at()
- toggle_at()
- mangle_switch()
NOTE: Changes need to be backported to CPU when finished
7 years ago
jsteube
9a57c4b20e
Fix missing boundary check in pure kernel rule engine
7 years ago
jsteube
0b0abb5c12
Prepare pure kernel rule engine for performance optimization
7 years ago
jsteube
cb1fe251de
Add missing functions to pure kernel rule engine
7 years ago
jsteube
8a0d21360b
Prepare for pure kernel rule engine support
7 years ago
jsteube
34d882a116
Rename inc_rp.X to inc_rp_optimized.X
7 years ago
jsteube
c5c12f89c1
Rewrite code around amd_bytealign to be of type BE to save a branch afterwards
7 years ago
jsteube
a53d9e09de
Fix some issue with offset_minus_4
7 years ago
jsteube
5e34ec348e
Optimize kernels for ROCm 1.6
...
- Remove inline keywords
- Remove volatile keywords where it causes ROCm to slow down
- Replace DES functions (looks like bitselect somehow is no longer mapped to BFI_INT)
7 years ago
mhasbini
5734741392
Add support for rule: eX
8 years ago
Jens Steube
7fe575e204
Add const qualifier to variable declaration of matching global memory objects
8 years ago
jsteube
d081ac2ba1
Fix some uninitialized variables in rp_kernel_on_cpu.c
8 years ago
jsteube
1ed89eb261
Fix rule-engine optimizations for vector datatypes
8 years ago
jsteube
faf6f1932d
Optimize append_block8()
8 years ago
jsteube
bd12a1eb0d
Optimize rule_op_mangle_dupechar_last()
8 years ago
jsteube
03fde2cdc3
Optimize rule_op_mangle_delete_last()
8 years ago
jsteube
a4f7ba2c60
Optimize rule_op_mangle_rotate_right()
8 years ago