1
0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-17 20:28:13 +00:00
Commit Graph

342 Commits

Author SHA1 Message Date
Jens Steube
6c10ca5853 Fix use of find_kernel_power_div() 2016-02-22 12:50:04 +01:00
Jens Steube
6bc98368ba Get rid of old pw_cache mechanism to control host-based vector data-types 2016-02-22 11:57:37 +01:00
Jens Steube
ec869595d1 Add auto-append for 0x01 and 0x80 in ATTACK_MODE_HYBRID2 2016-02-22 10:36:04 +01:00
Jens Steube
bb732686b1 Autotune should respect --quiet parameter 2016-02-21 18:23:04 +01:00
Jens Steube
9a128ce667 Fix force_jit_compilation usage 2016-02-21 13:30:28 +01:00
Jens Steube
7190dcf855 Prepare NEW_SIMD_MODE for -a 1 kernels 2016-02-20 16:13:06 +01:00
Jens Steube
990f973c48 Add hard-wired Device-Name for Tuning-Database which matches all Device-Types:
- DEVICE_TYPE_CPU
- DEVICE_TYPE_GPU
- DEVICE_TYPE_ACCELERATOR
Optimized IRIS GPU workaround settings
2016-02-20 00:58:22 +01:00
Fist0urs
62bed36638 Add verification of checksum for -m 13100
This avoid collisions by arc4'ing all data then hmac-md5
when valid ASN1 structures headers are found.
Performance should not be impacted.
2016-02-19 23:12:46 +01:00
Jens Steube
96bde85378 Allow mixed settings auto and fixed values for parameter and inside tuningdb for loops and accel 2016-02-19 18:58:03 +01:00
Jens Steube
14fe45b1f0 Simplify steps for autotune 2016-02-18 14:44:44 +01:00
Jens Steube
a83622faaf Fix final autotune balancing 2016-02-18 13:04:25 +01:00
Jens Steube
f05944395a Move fake word generation to try_run() to avoid unwanted caching 2016-02-18 10:49:23 +01:00
magnum
a5be8a75ed Allow and support vector-width 16, which is current maximum for
OpenCL. Closes #226.
2016-02-18 08:51:45 +01:00
magnum
998605ef89 Simplify code. No change in functionality at all. 2016-02-17 18:24:35 +01:00
philsmd
8298afa462 minor: some formatting fixes re: #225 (-m 13100 = Kerberos 5 TGS-REP etype 23) 2016-02-17 11:02:40 +01:00
Fist0urs
c3dabdd69e Initial commit 2016-02-16 16:34:46 +01:00
Jens Steube
4bc02e5ac8 Replace OpenCL event's with own timers, this saves us a thread-pair spawned on nvidia's OpenCL runtime and doesn't affect others 2016-02-15 22:27:30 +01:00
Jens Steube
c44b50c2e2 Use CL_DEVICE_NATIVE_VECTOR_WIDTH_* instead of CL_DEVICE_PREFERRED_VECTOR_WIDTH_* 2016-02-15 21:19:41 +01:00
Jens Steube
1af0f9c67b Fix autotune in -i mode 2016-02-15 21:08:16 +01:00
jsteube
bbadabe278 Fix another source of 0H/s in benchmark mode 2016-02-15 20:32:01 +01:00
Jens Steube
0cbe2347cd Merge pull request #223 from gm4tr1x/clEnqueueNDRangeKernelMod
Get rid of hc_clEnqueueNDRangeKernel double call
2016-02-15 18:56:03 +01:00
Jens Steube
709f79a7da Add option to repeat the kernel on the device NUM times to increase benchmark accuracy 2016-02-15 17:42:09 +01:00
Jens Steube
eaf7a7cb2f Final balance meassurements for autotune 2016-02-15 14:57:03 +01:00
Jens Steube
a9e3ef0576 In benchmark-mode, do not depend any longer on a fixed time, better use a single iteration instead 2016-02-15 12:38:54 +01:00
Jens Steube
8af57d55f4 Disable GPU-Temp monitoring during benchmark 2016-02-14 20:15:14 +01:00
jsteube
f1e0cd6d56 Autotune: 200 loops should be fine for a start 2016-02-14 20:07:16 +01:00
Jens Steube
515385c57d Add dedicated steps for loops and accel 2016-02-14 18:23:21 +01:00
Jens Steube
72e0553e44 Fixed a missing condition 2016-02-14 15:50:29 +01:00
Jens Steube
c09bc848f7 Autotuning engine prototype 2016-02-14 15:45:52 +01:00
Jens Steube
c0293928dd Prepare for better autotuning capabilities 2016-02-13 16:07:58 +01:00
Gabriele 'matrix' Gristina
2a1d0d21ff Get rid of hc_clEnqueueNDRangeKernel double call 2016-02-12 17:11:23 +01:00
Gabriele 'matrix' Gristina
c97dbf56cb Fix bug in weak_hash_check(), wrong kernel_loops assigned 2016-02-12 14:04:13 +01:00
Jens Steube
2dfe9200d6 Remove Wordload-Profiles from tuningdb, we can simply calculate it 2016-02-11 14:16:08 +01:00
Jens Steube
10dc25f807 Merge pull request #217 from magnumripper/master
Use device's preferred vector width, not its native one.
2016-02-11 10:08:57 +01:00
jsteube
5c01349ba6 Update event handling to workaround event handling error in nvidia opencl runtime 2016-02-11 09:54:50 +01:00
magnum
162bc25a3d Use device's preferred vector width, not the native one. Also, don't
assume vector width for 'long' is half of that for 'int'.
2016-02-10 23:43:08 +01:00
Gabriele 'Matrix' Gristina
20ee62e196 Revert "Use per-device timer resolution in total_time calculations" 2016-02-10 22:31:43 +01:00
Jens Steube
02ad834c7f Added new concept of a tuning database; tryout phase 2016-02-10 20:40:21 +01:00
Jens Steube
8d85b7539d Fixed a bug in line counter: Conditional jump or move depends on uninitialised value 2016-02-10 16:56:51 +01:00
Jens Steube
e80541d6dd Check for allocatable device-memory depending on kernel_accel amplifier before trying to allocate 2016-02-10 13:31:04 +01:00
Gabriele 'matrix' Gristina
523a7bfe93 Use CL_DEVICE_PROFILING_TIMER_RESOLUTION per-device value in total_time calculations 2016-02-09 21:48:18 +01:00
Jens Steube
9fc360e39d Added the execution time of the running kernel to the status display 2016-02-09 20:01:50 +01:00
Gabriele 'matrix' Gristina
74be8e2e25 skip device if buildProgram() fail 2016-02-07 21:20:10 +01:00
Jens Steube
6d37c123a5 Merge branch 'master' of https://github.com/hashcat/oclHashcat 2016-02-06 23:22:38 +01:00
Jens Steube
ab5ecd8eaa Fix some updated buffer sizes and new optimizer descriptions 2016-02-06 23:22:26 +01:00
Jens Steube
952c20e0f8 Merge pull request #211 from gm4tr1x/issue150
Fixed "VENDOR_ID" format (probably relate with issue #150)
2016-02-06 22:08:34 +01:00
Gabriele 'matrix' Gristina
91fae7e49a Fixed bug in 8900 (probably relate with issue #150) 2016-02-06 19:36:00 +01:00
Gabriele 'matrix' Gristina
e2db8afdf8 Add missing pthread_setaffinity_np for osx 2016-02-06 18:00:04 +01:00
Gabriele 'matrix' Gristina
5953130062 Fix clEnqueueNDRangeKernel() error -54 2016-02-06 15:40:22 +01:00
Jens Steube
8650212b80 Remove Workload display in benchmark, remove unused variable 2016-02-05 17:26:51 +01:00
Jens Steube
fa0e6fb78e Merge pull request #204 from gm4tr1x/LoopAccelMod
gpu loops/accel per-device
2016-02-05 17:17:04 +01:00
Gabriele 'matrix' Gristina
fa7465aa86 gpu loops/accel per-device 2016-02-05 15:27:09 +01:00
Jens Steube
22834781f5 Hack in NvAPI_GPU_GetCoolerSettings(); Enables query fanspeed in percentage on windows 2016-02-05 11:28:31 +01:00
Jens Steube
6c5938201b Revert "Merge pull request #202 from magnumripper/master"
This reverts commit 28edfbd654, reversing
changes made to b6f10a2a81.
2016-02-04 23:32:18 +01:00
magnum
947a927ed2 Do not silently ignore -w2 for -b. Closes #201. 2016-02-04 23:20:37 +01:00
magnum
68dbaa6f8a Bugfix: putenv(3) was used with a stack variable, that's a no-no. 2016-02-04 22:49:34 +01:00
Gabriele 'matrix' Gristina
89e5ed2574 handling possible clEnqueueFillBuffer not found in opencl lib 2016-02-04 21:06:19 +01:00
Jens Steube
2381af313d Some cleanups 2016-02-04 15:47:52 +01:00
Jens Steube
4d2be9073e Consolidate hardware monitoring interface status on startup 2016-02-04 15:07:38 +01:00
Gabriele 'matrix' Gristina
74db620d51 Add missing memset and update ext_nvapi.c header 2016-02-04 13:36:52 +01:00
Jens Steube
e01038083c Update Makefile and BUILD.md 2016-02-04 09:44:52 +01:00
Gabriele 'matrix' Gristina
e67eec1ae0 get rid of nvidia nvapi static linking 2016-02-03 20:53:00 +01:00
Gabriele 'matrix' Gristina
f3407cd7f4 HWMon mod 2016-02-02 01:14:33 +01:00
Gabriele 'matrix' Gristina
88087295b3 Removed deprecated and unused ADL function (ADL_DisplayEnable_Set) 2016-02-01 15:39:44 +01:00
Jens Steube
f256363448 Merge pull request #190 from gm4tr1x/fixLoopsAccelV3
Fixed loops/accel for 6800, 11300 and 11600 (osx)
2016-02-01 08:08:20 +01:00
Jens Steube
102dee00eb Merge pull request #192 from gm4tr1x/ADL_LibraryLoader
Update ADL/OpenCL Library Loader
2016-02-01 08:08:13 +01:00
Gabriele 'matrix' Gristina
d120ede655 Update ADL Library Loader 2016-01-31 23:48:38 +01:00
jsteube
d0d3507dc2 Fix TI series detection 2016-01-31 22:53:28 +01:00
jsteube
21525b74e1 Remove more unnecessary libOpenCL*.a related stuff 2016-01-31 22:01:32 +01:00
Gabriele 'matrix' Gristina
b542d4a431 Fixed loops/accel for 6800, 11300 and 11600 (osx) 2016-01-31 21:48:08 +01:00
Gabriele 'matrix' Gristina
dcb4888dd5 Remove unnecessary libOpenCL*.a from Makefile 2016-01-31 19:48:40 +01:00
Jens Steube
55d4904dd3 Merge pull request #185 from gm4tr1x/fixLoopsAccelv2
Fixed some gpu loops/accel settings
2016-01-31 18:27:21 +01:00
Gabriele 'matrix' Gristina
53faf6a24a Fixed some gpu loops/accel settings 2016-01-31 17:38:05 +01:00
Gabriele 'matrix' Gristina
0a907fc210 Fixed typo in help 2016-01-31 13:59:36 +01:00
Gabriele 'matrix' Gristina
0f0984fe86 Fixed all gpu code (see PR #179 for details) 2016-01-30 23:02:15 +01:00
Jens Steube
9aea991424 Use a rc for first result, easier to read 2016-01-30 22:47:13 +01:00
Gabriele 'matrix' Gristina
f9598f691d Fixed CL_INVALID_WORK_GROUP_SIZE error with Apple CPU 2016-01-30 22:28:41 +01:00
Gabriele 'matrix' Gristina
d147d89f5e cleanup unnecessary clGetProgramBuildInfo code 2016-01-30 20:05:39 +01:00
Gabriele 'matrix' Gristina
5d4160dc8c Fixed gcc warnings about missing field initializers 2016-01-30 14:59:26 +01:00
Jens Steube
2c4ad77275 Some systems don't like using the same buffer for both input and output at the same time with snprintf() 2016-01-29 22:47:42 +01:00
magnum
ce170ea980 Acquire an exclusive lock before writing to any file. Note that in
some cases we never explicitly unlock a file because fclose will do
it implicitly. Closes #172.
2016-01-29 16:48:19 +01:00
Gabriele 'matrix' Gristina
ebd28553a5 Handling clBuildProgram failure (show build log on errors) 2016-01-28 21:02:36 +01:00
Gabriele 'matrix' Gristina
3952fa7e06 Updated gpu accel for -m 6211 and default runtime value to 8 2016-01-27 23:39:58 +01:00
Jens Steube
bfe3f148d7 As long as we depend on NVAPI for windows we also need to add a path to its includes 2016-01-27 20:09:17 +01:00
Gabriele 'matrix' Gristina
9eefb4bcd7 Fixed get_profile_dir and get_session_dir byte missing bug 2016-01-27 19:29:51 +01:00
Jens Steube
2232b1632e Merge pull request #169 from magnumripper/onetwenty
Drop dependencies on non-distributable ADL/NVML headers.
2016-01-27 19:15:31 +01:00
Jens Steube
af7ff57b3d Merge pull request #168 from magnumripper/onesixfour
Do not create a kernel cache file if build failed (actually if size ended up as zero).
2016-01-27 18:52:10 +01:00
Jens Steube
98ec1bcee7 Merge pull request #156 from gm4tr1x/master
Updated test.sh (support for osx and some fixes)
2016-01-27 18:51:28 +01:00
magnum
115d2b6a5a Drop dependencies on non-distributable ADL/NVML headers. The needed glue
is copied into our respective local headers. Should close #120.
2016-01-27 18:48:54 +01:00
Jens Steube
cae457df0c Merge pull request #167 from gm4tr1x/issue84
Issue84
2016-01-27 18:43:18 +01:00
magnum
f5d6f9b6d4 Do not create a kernel cache file if build failed. Also disregard any
existing cache files with size of zero. Should close #164.
2016-01-27 18:39:45 +01:00
Jens Steube
9ec2c86cd0 Merge pull request #165 from gm4tr1x/gcc-warnings
Fixed gcc warnings
2016-01-27 18:19:50 +01:00
Gabriele 'matrix' Gristina
b7d3a7c5bc Fixed gcc warnings 2016-01-27 01:08:35 +01:00
Gabriele 'matrix' Gristina
5bae9de3a3 Implemented OpenCL library loader 2016-01-26 21:40:49 +01:00
Gabriele 'matrix' Gristina
6ba392c0df Modified as @jsteube request 2016-01-26 12:07:19 +01:00
Gabriele 'matrix' Gristina
1124687270 OpenCL runtime environment is no longer required with --keyspace option 2016-01-25 15:44:04 +01:00
Gabriele 'matrix' Gristina
dcc4b5a60f Added custom OSX kernel loops and gpu accel default values
Disabled some not yet worked hash types in OSX
2016-01-24 22:48:01 +01:00
Gabriele 'matrix' Gristina
479327fc20 Re-enable gpu-temp-disable argument for osx build 2016-01-24 17:49:15 +01:00
Gabriele 'matrix' Gristina
58359f3b7e Fixed 'buf' initialization for windows 2016-01-24 17:20:01 +01:00
Jens Steube
0997ae029e Remove NO-BREAK SPACE character (utf8) from sources 2016-01-24 17:08:49 +01:00
Jens Steube
9c89b58f76 Merge pull request #153 from gm4tr1x/missingcheck
Added missing check for in_len in _old_apply_rule function
2016-01-24 16:24:57 +01:00
Jens Steube
d275748c7c Merge pull request #152 from gm4tr1x/buf
Optimized memset calls and also some initializations
2016-01-24 16:24:10 +01:00
Gabriele 'matrix' Gristina
17d885ffa4 Added missing check for in_len in _old_apply_rule function 2016-01-24 15:13:43 +01:00
Gabriele 'matrix' Gristina
8674959a93 @philsmd suggestion (PR #152) 2016-01-24 14:33:19 +01:00
Gabriele 'matrix' Gristina
3abacf515c Optimized memset calls and also some initializations 2016-01-24 13:25:47 +01:00
Gabriele 'matrix' Gristina
8d891ef0ba Avoid sprintf 2016-01-24 13:05:02 +01:00
Jens Steube
20b0c23af1 Merge pull request #146 from gm4tr1x/master
Fixed some memory allocations and other small things
2016-01-23 15:43:47 +01:00
Jens Steube
1d3795a3ab Converted _a3 kernels, use SIMD for CPU and GPU 2016-01-23 15:32:31 +01:00
Gabriele 'matrix' Gristina
076597fe01 Fixed some memory allocations and other small things 2016-01-21 17:20:02 +01:00
Jens Steube
4c0e520fd8 Test convert for -m 1000 with -a 0 for SIMD, speed is now on par or faster than hashcat 2016-01-21 16:47:38 +01:00
Jens Steube
add18eaa6d Fix for last commit 2016-01-21 10:22:57 +01:00
Jens Steube
bfc4495bd3 Expanded version information as discussed in https://github.com/hashcat/oclHashcat/issues/138 2016-01-21 10:17:22 +01:00
Gabriele 'matrix' Gristina
5da79e4411 Added support for build without ADL/NVML/NVAPI (issue #120)
Added support for build OSX native binaries (issue #63)
2016-01-20 20:55:09 +01:00
philsmd
751a364125 #137: fixes padding for base64_encode () and base64_decode() 2016-01-19 22:32:20 +01:00
philsmd
e1ae2538c6 cosmetic fix: some fixed typos e.g. to -> too 2016-01-19 20:47:18 +01:00
philsmd
5e92020180 buffer overflow fix for -m 8900 = scrypt 2016-01-19 19:04:52 +01:00
Jens Steube
a62b7ed06e Upgrade kernel to support dynamic local work sizes 2016-01-19 16:06:03 +01:00
Jens Steube
44b0cb4e65 Merge pull request #135 from gm4tr1x/master
Fixed a buffer overflow in ascii_digest for hash type 8300
2016-01-17 22:29:12 +01:00
jsteube
e3c0c80b6f Prepare new SIMD code for kernel, -m 0, 10, 20, 1000 should work in -a 3 mode and other hopefully stay unaffected 2016-01-17 22:17:50 +01:00
Gabriele 'matrix' Gristina
29636ca0a3 Re-fix ascii_digest 2016-01-17 22:10:19 +01:00
Gabriele 'matrix' Gristina
e9bc4caed3 Fixed a buffer overflow in ascii_digest for hash type 8300 2016-01-17 18:52:25 +01:00
Jens Steube
45431d9201 Fixed a buffer overflow in potfile handling 2016-01-17 02:18:07 +01:00
jsteube
471c10c4f7 Make sure to spawn threads for all devices, even skipped, and then instantly return 2016-01-16 15:12:42 +01:00
jsteube
954be23a12 Fix some code formating 2016-01-16 13:47:43 +01:00
Jens Steube
b9c61eb66f Merge pull request #134 from yhfudev/add-cl-error-messages
human-readable error message for the OpenCL error codes.
2016-01-16 13:42:34 +01:00
jsteube
0524e78617 Enable forced JIT compilation for MD5, if requested 2016-01-16 12:36:56 +01:00
jsteube
09de56b720 Make sure to select an active device for weak_hash_check 2016-01-16 12:24:08 +01:00
Jens Steube
36251e1fe5 Merge pull request #132 from neheb/master
Remove some memsets
2016-01-16 11:47:17 +01:00
yhfudev
edc6c920a9 human-readable error message for the OpenCL error codes. 2016-01-16 00:41:31 -05:00
jsteube
9fb506f15f Fixed out of bounds access in -m 11300 2016-01-15 23:25:24 +01:00
Mangix
54c7c23ad4 Remove some memsets 2016-01-15 10:51:47 -08:00
Jens Steube
bc24e3d79b Add bitness to filename for cached kernels, otherwise a user would load a 32 bit kernel with a 64 bit host binary in case he used the 32 bit binary before which leads to a segfault 2016-01-15 17:48:27 +01:00
Jens Steube
98b1a9370c Fix AMD OpenCL runtime bug in clGetProgramInfo() 2016-01-15 17:32:43 +01:00
Jens Steube
b35d2dd2eb Fix native compiler 2016-01-15 17:24:51 +01:00
Jens Steube
05505030b8 Fix CFLAGS 2016-01-15 17:24:20 +01:00
jsteube
5ae5a4bc25 Reorder device mapping 2016-01-15 17:23:07 +01:00
Jens Steube
63e06f582b Prepare fix for datatypes, not final 2016-01-15 17:16:43 +01:00
Jens Steube
2e61685f48 Fix outputting wrong password for weak-hash-check in case user defined -r or -g 2016-01-14 23:03:23 +01:00
Jens Steube
ed91e69ae3 Fix out of bounds access in -m 7500 2016-01-14 22:13:06 +01:00
Jens Steube
f3b85a6363 Fix some clang compiler warnings 2016-01-14 21:45:11 +01:00
jsteube
76612ac031 Fix more missing casts to uint in rule-engine 2016-01-14 20:54:41 +01:00
jsteube
497b41fee4 Fix more missing casts to uint 2016-01-14 20:44:11 +01:00
jsteube
d3981ee66a Fix missing cast to uint 2016-01-14 20:42:07 +01:00
Jens Steube
7bfee268d4 Fixed some heap buffer overflow 2016-01-14 20:30:38 +01:00
Jens Steube
245301c9b4 Started optimizing some of the OpenCL kernel for latest AMD Catalyst 15.12:
- Replaced SBOX for DES:

replaced JtR's * Bitslice DES S-boxes making use of a vector conditional select operation (e.g., vsel on PowerPC with AltiVec).
with     JtR's * Bitslice DES S-boxes for x86 with MMX/SSE2/AVX and for typical RISC architectures.

Performance increased for DEScrypt from 355MH/s to 405MH/s and for LM from 11100MH/s to 12000MH/s

BTW, the same effect can be seen with non-maxwell GPU's

- Remove some volatile keywords no longer needed thanks to fixed catalyst bugs

- Fix weak-hash-check parameter for use with tools/test.sh
2016-01-14 19:44:47 +01:00
Jens Steube
24b5aa6226 Merge pull request #128 from philsmd/pr-set-iter-for-dcc2
set default iteration count for -m 2100 = DCC2 to 10240
2016-01-14 09:02:13 +01:00
philsmd
cbe6be9246 set default iteration count for -m 2100 = DCC2 to 10240 2016-01-14 00:14:43 +01:00
philsmd
88cbb45ca5 add skipped devices to the --benchmark output 2016-01-13 23:38:07 +01:00
philsmd
7b53ccab70 minor: remove quiet_sav, was just intended for debugging 2016-01-13 23:06:03 +01:00
Jens Steube
49d0767aa8 Run weak-hash checks only in straight-attack mode, this greatly reduces code complexity 2016-01-13 22:22:46 +01:00