Gabriele Gristina
409c81f16c
remove inc_rp_common.cl include from kernels
2025-07-19 13:46:00 +02:00
Gabriele Gristina
360f91f218
Merge branch 'master' into rule_purgeclass
2025-07-19 11:44:16 +02:00
Gabriele Gristina
0df324c7d8
Merge remote-tracking branch 'upstream/master' into rule_purgeclass
2025-07-15 02:42:40 +02:00
Gabriele Gristina
8f389bb205
Fix the same Apple Intel bug with Metal on inc_rp_optimized, this time on Apple Silicon
2025-07-15 02:06:35 +02:00
Gabriele Gristina
f39f7cc877
cleanup
2025-07-14 16:09:09 +02:00
Gabriele Gristina
867b3d4df1
Fixed bug in inc_rp_optimized.cl on Apple Intel with Metal
...
Fixed old/critical bug on Apple Intel with Metal by patching inc_rp_optimized.cl.
Tested on Apple Intel and Silicon with Metal/OpenCL and on Linux with CUDA, HIP, OpenCL GPU/CPU
Metal Backend: parallelize pipeline state object (PSO) compilation internally
Set unexported setting, setShouldMaximizeConcurrentCompilation, to boost kernel build process on Apple Metal (only >= 3)
2025-07-14 00:25:58 +02:00
Jens Steube
06344910a4
Refactored HIP kernel code for improved performance and cleanup
...
- Replaced inline asm in hc_byte_perm() with __builtin_amdgcn_perm()
- Replaced inline asm in hc_bytealign() with __builtin_amdgcn_alignbyte()
- Defined HC_INLINE as default for HIP, significantly boosting kernel performance of pure kernels
- Removed IS_ROCM from inc_vendor.h as it's no longer needed
- Removed backend-specific code from several hash-modes and inc_rp_optimized.cl, as hc_bytealign_S() is now available on all backends
2025-07-10 13:31:00 +02:00
Gabriele Gristina
477c5df690
mv common parts between inc_rp.cl/h inc_rp_optimized.cl/h to inc_rp_common.cl/h
2025-06-27 01:23:16 +02:00
Gabriele Gristina
06810bdbee
fix HIPRTC_ERROR_COMPILATION: candidate function not viable: call to __host__ function from __device__ function
2025-04-23 00:13:17 +02:00
Gabriele Gristina
25af3a1c73
Add RULE_OP_MANGLE_TITLE_SEP_CLASS
2024-11-03 01:53:00 +01:00
Gabriele Gristina
53f57fb7ae
Add support to character class rules
2024-11-02 22:02:18 +01:00
Gabriele Gristina
cedaef797a
Rename purgeclass to purgechar_class, fix lookup tables
2024-11-02 03:55:22 +01:00
Gabriele Gristina
d37d26de74
Update s_lookup/s_lookup_optimized tables
2024-11-02 00:21:33 +01:00
Gabriele Gristina
caf1f6a478
Improve Purge rule backwards compatibility and simplify code
2024-11-01 18:27:51 +01:00
Gabriele Gristina
33092b13c3
Extended the Purge (@) rule operation to support Classes, Fixed d3ad0ne.rule to support the new Purge (@) rule handling
2024-10-31 15:23:18 +01:00
Jens Steube
ac98130fef
Fix double use of macro HC_INLINE
2022-11-08 19:05:18 +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
861e644057
OpenCL Backend: added workaround to make optimized kernels work on Apple Silicon
2021-12-24 17:30:49 +01:00
Jens Steube
f4e52ca533
Add new rule function '3' to switch the case of the first letter after occurrence N of char X
2021-08-01 00:04:10 +02:00
Jens Steube
cf512faa53
Update large switch() cases in inc_common.cl and some inline assembly common functions for devices managed with HIP backend
2021-07-14 17:06:20 +02:00
Jens Steube
4658e470a2
OpenCL Kernels: Added datatypes to literals of some 64 bit kernel constants
2020-07-22 14:06:58 +02:00
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
3ca3d1cc60
Fix kernel_rules variable name
2019-05-11 14:34:10 +02:00
Jens Steube
89119bf24a
Add missing inc_platform.h include
2019-04-26 13:59:43 +02:00
jsteube
ace765bf96
Fix 'E' rule in optimized mode if password is longer than 16 characters
2019-04-05 12:24:53 +02:00
jsteube
5a1d929628
Fix some missing code changes after hc_bytealign() was changed
2019-03-28 12:26:24 +01:00
jsteube
ea7d9c50b6
Fix missing const keyword in some includes
2019-03-27 11:38:49 +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
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
540b405e3a
Replace IS_ROCM with HAS_VPERM and HAS_VADD3
2019-02-24 10:12:48 +01:00
jsteube
fbf434146d
Add set_mark_1x4() and set_mark_1x4_S()
2018-07-28 18:03:18 +02:00
Jens Steube
14c444fd47
Replace c_append_helper_mini[] table with on-the-fly calculation in order to workaround compiler bugs in AMD OpenCL runtime
2018-07-28 13:59:55 +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
Jens Steube
81b229c08a
Make new c_append_helper a bit more AMD friendly
2018-02-17 15:18:19 +01:00
jsteube
3a23b275e5
Improved c_append_helper[] handling
2018-02-17 14:24:29 +01:00
jsteube
5951207365
Get rid of some old volatiles
2018-02-09 19:18:30 +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
jsteube
b169653b8f
Fix missing return value in rule_op_mangle_toggle_at()
2017-09-08 22:49:49 +02:00
jsteube
9125062ffc
Move volatiles for AMD closer to the problem
2017-09-08 13:32:19 +02:00
jsteube
ac9f1da747
Add fine-tuned AMD GCN control macros
2017-09-07 20:33:43 +02:00
jsteube
7bfd343ec9
Optimized rule_op_mangle_dupechar_last(), rule_op_mangle_rotate_right(), rule_op_mangle_rotate_left() and append_block1() in rule engine
2017-08-27 16:47:21 +02:00
jsteube
9b73c464d2
Fix typo in macro
2017-08-24 17:19:16 +02:00
