Jens Steube
61f87b2981
Try to load libOpenCL.so.1 if libOpenCL.so was not found
2016-02-23 07:57:52 +01:00
Jens Steube
f9834f3dfa
Zero pws_buf before reuse
2016-02-22 21:33:33 +01:00
Jens Steube
e6e5005a6b
Revert "Zero pws_buf before reuse"
...
This reverts commit b409e5e9e1
.
2016-02-22 21:32:38 +01:00
Jens Steube
b409e5e9e1
Zero pws_buf before reuse
2016-02-22 21:20:16 +01:00
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