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
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
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
magnum
a5be8a75ed
Allow and support vector-width 16, which is current maximum for
...
OpenCL. Closes #226 .
2016-02-18 08:51:45 +01:00
Fist0urs
c3dabdd69e
Initial commit
2016-02-16 16:34:46 +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
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
Jens Steube
2dfe9200d6
Remove Wordload-Profiles from tuningdb, we can simply calculate it
2016-02-11 14:16:08 +01:00
jsteube
5c01349ba6
Update event handling to workaround event handling error in nvidia opencl runtime
2016-02-11 09:54:50 +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
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
0362df07ed
Merge pull request #210 from gm4tr1x/hd4000v0
...
Fix SHA-3(Keccak) issue with HD4000 (issue #191 ) and reorder macro
2016-02-06 22:07:38 +01:00
Gabriele 'matrix' Gristina
225fc28dff
Fix SHA-3(Keccak) issue with HD4000 (issue #191 ) and reorder macro
2016-02-06 18:18:21 +01:00
Gabriele 'matrix' Gristina
e2db8afdf8
Add missing pthread_setaffinity_np for osx
2016-02-06 18:00:04 +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
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
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
Gabriele 'matrix' Gristina
0f159d42e6
Fixed HC_LOAD_FUNC macro
2016-02-01 13:07:14 +01:00
Jens Steube
b8285cbce2
Fix broken ADL on windows
2016-02-01 12:18:05 +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
Gabriele 'matrix' Gristina
d120ede655
Update ADL Library Loader
2016-01-31 23:48:38 +01:00
Gabriele 'matrix' Gristina
b542d4a431
Fixed loops/accel for 6800, 11300 and 11600 (osx)
2016-01-31 21:48:08 +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
Jens Steube
6846348602
Use a different workaround for a catalyst bug which takes effect in sha512crypt
2016-01-31 17:24:12 +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
Jens Steube
060af27dd9
Merge branch 'master' of https://github.com/hashcat/oclHashcat
2016-01-29 18:38:46 +01:00
Jens Steube
eb60d6bb23
Remove MD4/MD5 *H1/*H2 functions and use original H functions. Modern compilers will find this easy optimization automatically
2016-01-29 18:38:34 +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
3952fa7e06
Updated gpu accel for -m 6211 and default runtime value to 8
2016-01-27 23:39:58 +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
Jens Steube
9ec2c86cd0
Merge pull request #165 from gm4tr1x/gcc-warnings
...
Fixed gcc warnings
2016-01-27 18:19:50 +01:00
Jens Steube
9d5a757ae5
Merge pull request #157 from gm4tr1x/loops-accel
...
Improved osx support with custom kernel loops and gpu accel (issue #63 )
2016-01-27 18:19:11 +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
6cbe3cee50
Cosmetix fix for types.h
2016-01-25 12:28:06 +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
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
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
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
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
49d0767aa8
Run weak-hash checks only in straight-attack mode, this greatly reduces code complexity
2016-01-13 22:22:46 +01:00
Jens Steube
52d71de0d4
Added support to utilize multiple different OpenCL platforms in parallel, ex: AMD + NV
2016-01-13 17:10:40 +01:00
Jens Steube
6b4e4c060a
Replace typedef for bool with stdbool.h
2016-01-12 18:56:26 +01:00
Jens Steube
0ff49c5b44
Fix segfault in related to weak-hash-check on CPU devices
...
Move typedef for bool datatype to common.h
2016-01-12 17:18:01 +01:00
Jens Steube
730044e26e
Added option --opencl-device-types to filter for specific OpenCL device types
2016-01-12 00:44:28 +01:00
philsmd
02cecce247
fixed hardware monitor: mapping between supported ADL devices and available GPU (and CPUs) fixed
2016-01-12 00:30:28 +01:00
Jens Steube
1d628ca6e0
The weak-hash-check feature didn't work for algorithms that made use of some automatic optimizer flags
...
We'll enforce a weak-hash-check on an _a0 kernel for them
2016-01-10 21:02:03 +01:00
jsteube
3cdee2e06e
We need an esalt to hold the unused IV otherwise the unique salt count becomes incorrect in multihash sitations
2016-01-09 21:35:29 +01:00
jsteube
f0a84a2410
Added new hash mode -m 13000 = RAR5
2016-01-09 20:34:12 +01:00
jsteube
332b3c35e5
Added new hash mode -m 12900 = Android FDE (Samsung DEK)
2016-01-09 00:49:54 +01:00
jsteube
e8229af09b
Fix for -m 1500, -m 3000: Missing "defined" keyword
...
Fix for -m 1000: MD4_H1 copy/paste error
Fix for -m 8900, -m 9300: Invalid value for device_processor_cores for CPU devices
Fix for -m 9100: Variable salt2 initializer needed some clauses for clearness
Temporary limit gpu_accel for CPU devices to 1 for development phase
Mark pocl as too bleeding edge for production use, recommend native drivers
Remove workarounds for pocl
Rename VENDOR_ID_UNKNOWN to VENDOR_ID_GENERIC in host code
Rename IS_UNKNOWN to IS_GENERIC in kernel code
2016-01-07 20:14:34 +01:00
jsteube
6e680aa31c
Workaround for either pocl or llvm fails and produces invalid optimized code
2016-01-06 11:48:18 +01:00
jsteube
331188167c
Replace the substring GPU to a more appropriate "device" or "kernel" substring depending on the context
2016-01-05 08:26:44 +01:00
jsteube
894140b816
Generate a more easy to read filename for cached kernels
2016-01-04 20:56:15 +01:00
jsteube
8baf705f6a
Preparation for distinguish between OpenCL device types at runtime (mostly for HMS)
2016-01-04 20:12:34 +01:00
jsteube
c4b1c8e7ca
small fix for unknown opencl platforms
2016-01-04 16:13:10 +01:00
jsteube
3b589e3aac
Prepare for allow other OpenCL platforms thans AMD and NVidia
2016-01-04 13:17:20 +01:00
jsteube
0428514f61
Extended support from 14 to 255 functions calls per rule on GPU
2016-01-03 19:49:47 +01:00
philsmd
d378319ea6
this patch makes it much clearer where the sessions under ~/.hashcat are located
2016-01-03 17:11:36 +01:00
jsteube
d008a45f3c
Activate reordering or files to help integration into linux distributions
...
Details can be found here: https://github.com/hashcat/oclHashcat/issues/20
2016-01-03 00:40:31 +01:00
jsteube
31292946e2
Fix some scrypt default values
2015-12-30 23:06:45 +01:00
jsteube
c17bf5e865
- Fixed PHY memory handling for scrypt based algorithms
...
- Bring back kernel exec timeout checking for NV
2015-12-23 15:51:55 +01:00
jsteube
84568e5b3d
Increase bcrypt speed for NV
2015-12-22 23:43:39 +01:00
jsteube
7be2c2fd8d
Fixed some speeds:
...
1800
3200
7100
7200
7400
8200
9300
9000
9600
11300
11600
2015-12-21 14:29:04 +01:00
jsteube
378258d789
Fix caching system for use with AMD and NV, drop BINARY_KERNEL define
2015-12-21 12:01:38 +01:00
jsteube
9115547601
Revert "Fix -m 10700"
...
This reverts commit 15da53da38
.
2015-12-19 18:18:28 +01:00
jsteube
15da53da38
Fix -m 10700
2015-12-19 18:17:19 +01:00
jsteube
91249942ab
Fix -m 1800 for NV
2015-12-17 19:19:31 +01:00
jsteube
9c392b472e
Removed useless comments
2015-12-16 11:18:57 +01:00
jsteube
c29c7c093f
Fix use of LOP3
2015-12-16 11:11:31 +01:00
Jens Steube
0f1ae86be9
Fix invalid VENDOR_ID for AMD
2015-12-16 10:00:33 +01:00
philsmd
cff683eac1
change in library loading: libnvidia-ml now does load dynamically only on NVidia systems
2015-12-15 20:34:07 +01:00
philsmd
3ab7a23cd5
fixed hardware monitor for amd/nv (one problem that is still there since last commit: dynamic loading of libnvidia-ml)
2015-12-15 18:41:11 +01:00
jsteube
0bf4e3c34a
- Dropped all vector code since new GPU's are all scalar, makes the code much easier
...
- Some performance on low-end GPU may drop because of that, but only for a few hash-modes
- Dropped scalar code (aka warp) since we do not have any vector datatypes anymore
- Renamed C++ overloading functions memcat32_9 -> memcat_c32_w4x4_a3x4
- Still need to fix kernels to new function names, needs to be done manually
- Temperature Management needs to be rewritten partially because of conflicting datatypes names
- Added code to create different codepaths for NV on AMD in runtime in host (see data.vendor_id)
- Added code to create different codepaths for NV on AMD in runtime in kernels (see IS_NV and IS_AMD)
- First tests working for -m 0, for example
- Great performance increases in general for NV so far
- Tested amp_* and markov_* kernel
- Migrated special NV optimizations for rule processor
2015-12-15 12:04:22 +01:00
Jens Steube
586441fa25
Revert "Fixed a bug where oclHashcat rejected to load a rule which calls 15 functions although it is supported"
...
This reverts commit f230ed73dc
.
Some kernel require too much constant memory and do not compile anymore
2015-12-13 18:46:21 +01:00
Jens Steube
f230ed73dc
Fixed a bug where oclHashcat rejected to load a rule which calls 15 functions although it is supported
2015-12-13 14:54:39 +01:00
Jens Steube
1537390fae
Prepare reordering or files to help integration into linux distributions as discussed in https://github.com/hashcat/oclHashcat/issues/20
...
TODOS:
- Let oclHashcat actually use the new paths
- Find a better way for native compilation
- Replace /bin/cp with /usr/bin/install where it has to copy files recursive
2015-12-13 12:21:36 +01:00
jsteube
968265fffb
- Prepared for JIT use of hash-mode 1500, 8900 and 9300, works already on OpenCL (AMD)
...
- Changed PROMPT
2015-12-07 21:37:12 +01:00
philsmd
946b4a37d0
implements the enhancement mentioned in issue #10 (cancel on next checkpoint)
2015-12-07 18:31:45 +01:00