Jens Steube
d9a92afecc
Change out-of-boundary fix in order to re-enable password length 256 with rules in pure kernel mode
2019-11-26 11:26:56 +01:00
Jens Steube
270210a8ab
Fix out-of-boundary read in rule engines
2019-11-20 14:35:47 +01:00
Jens Steube
8ff8c5d536
Add LOCAL_VK to make use of __shared__
2019-05-07 09:01:32 +02:00
Jens Steube
89119bf24a
Add missing inc_platform.h include
2019-04-26 13:59:43 +02:00
jsteube
ffd8ec9001
Fixed output password of 'e' rule in pure and cpu rule engine if separator character is also the first letter
2019-03-28 13:07:39 +01:00
jsteube
ea7d9c50b6
Fix missing const keyword in some includes
2019-03-27 11:38:49 +01:00
jsteube
fb8a9d7c40
Get rid of cpu_sha1.c
2019-03-25 14:08:59 +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
194fd7e6d1
Fix some invalid code sections caused from conversion
2019-03-24 13:46:06 +01:00
jsteube
977199698f
Move some macros from .h to .cl sources
2019-03-24 00:26:10 +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
5f5468be6f
Add missing DECLSPEC in OpenCL rule functions
2019-03-22 13:14:25 +01:00
jsteube
305a044ec6
Remove some old function headers
2019-03-22 09:52:54 +01:00
jsteube
84d6b8ecc1
Add function prototypes in OpenCL kernels to make some compilers happy
2019-03-09 09:05:44 +01:00
jsteube
874635cc49
Do not use a vector function to write into a scalar variable even if vectorize support is disabled
2018-02-16 10:48:49 +01:00
jsteube
512fb5f6fb
No inline keyword for rules
2018-02-07 15:02:58 +01:00
jsteube
3e08750900
OpenCL Kernels: Add general function declaration keyword (static inline)
2018-02-06 19:12:24 +01:00
jsteube
503f676cb0
Remove static keywords from OpenCL functions, they can cause old NV JiT compiler to fail
2017-10-20 13:23:43 +02:00
mhasbini
de7ccd88ef
Fix overflow in mangle_dupechar_last function
2017-09-16 20:43:38 +03: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
8853884f2a
Fix append_four_byte() in case sm8 is 0
2017-08-21 16:04:43 +02:00
jsteube
f32e113942
Add missing case in append_block() in pure kernel rule engine
2017-08-20 15:08:51 +02:00
jsteube
ad1ce462d1
Get rid of ceil() in OpenCL kernels
2017-08-17 13:43:35 +02:00
jsteube
5b5bdf3889
Replace some SIMD related function calls
2017-08-17 10:18:17 +02:00
jsteube
967e96728d
Make all the OpenCL kernel function includes static
2017-08-16 20:27:17 +02:00
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
2017-08-13 16:43:46 +02:00
jsteube
9a57c4b20e
Fix missing boundary check in pure kernel rule engine
2017-08-12 13:32:05 +02:00
jsteube
0b0abb5c12
Prepare pure kernel rule engine for performance optimization
2017-08-12 13:04:52 +02:00
jsteube
cb1fe251de
Add missing functions to pure kernel rule engine
2017-08-11 22:58:37 +02:00
jsteube
8a0d21360b
Prepare for pure kernel rule engine support
2017-08-11 16:09:12 +02:00
jsteube
34d882a116
Rename inc_rp.X to inc_rp_optimized.X
2017-08-11 11:25:47 +02:00
jsteube
c5c12f89c1
Rewrite code around amd_bytealign to be of type BE to save a branch afterwards
2017-08-05 19:46:56 +02:00
jsteube
a53d9e09de
Fix some issue with offset_minus_4
2017-08-04 14:12:58 +02:00
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)
2017-07-22 18:05:18 +02:00
mhasbini
5734741392
Add support for rule: eX
2017-04-30 16:23:39 +03:00
Jens Steube
7fe575e204
Add const qualifier to variable declaration of matching global memory objects
2016-11-22 20:20:34 +01:00
jsteube
d081ac2ba1
Fix some uninitialized variables in rp_kernel_on_cpu.c
2016-11-12 16:30:21 +01:00
jsteube
1ed89eb261
Fix rule-engine optimizations for vector datatypes
2016-11-06 12:25:38 +01:00
jsteube
faf6f1932d
Optimize append_block8()
2016-11-05 16:32:59 +01:00
jsteube
bd12a1eb0d
Optimize rule_op_mangle_dupechar_last()
2016-11-01 14:54:45 +01:00
jsteube
03fde2cdc3
Optimize rule_op_mangle_delete_last()
2016-11-01 14:50:54 +01:00
jsteube
a4f7ba2c60
Optimize rule_op_mangle_rotate_right()
2016-11-01 14:44:35 +01:00
jsteube
9038955bc6
Optimize rule_op_mangle_rotate_left()
2016-11-01 14:35:51 +01:00
jsteube
f3dd6fe741
Some cleanup before going to optimize more code in inc_rp.c
2016-11-01 13:01:15 +01:00
jsteube
ae5109a5a2
Optimize $ (append) rule performance
2016-11-01 03:22:55 +01:00
jsteube
3daf0af480
Added docs/credits.txt
...
Added docs/team.txt
2016-09-11 22:20:15 +02:00
jsteube
8acf5b3879
Add support for @ rule (RULE_OP_MANGLE_PURGECHAR) to use on GPU
2016-07-25 22:54:07 +02:00
jsteube
0154d636bf
Slightly increased NVidias rule-processing performance by using generic instructions instead of byte_perm()
2016-07-25 21:51:04 +02:00
jsteube
5caf32447e
Some experiment. This should increase rule-based cracking on nvidia slightly
2016-07-19 21:07:30 +02:00
