Jens Steube
9254603960
Fixed out-of-boundary read in PKZIP masterkey kernel if the password candidate has length zero.
...
Now that kernel threads are no longer fixed over the entire session, hardware_power and hardware_power_all
needs to be updated the same way as kernel_power and kernel_power_all for each call to inner2_loop().
3 years ago
Jukka Ojanen
5c0820b8d3
Assume hashconfig->tmp_size is the element size
3 years ago
Jens Steube
5e1d37c82e
Make unsupported AMD HIP runtime version error message a bit more human readable
3 years ago
Jens Steube
cb69e2d413
Added some HIP version checks, fall back to OpenCL automatically
...
Switched HIP version check from driverVersion to runtimeVersion
Fixed syntax check of HAS_VPERM macro in several kernel includes causing invalid error message for AMD GPUs on Windows
Updated AMD driver requirements
Updated docs/changes.txt with missing changes from previous commits
Fixed invalid vector data type in Murmur Hash in -a 3 mode
Fixed uninitialized variable warning in src/hashes.c
Fixed broken support for --generate-rules-func-min
3 years ago
Jens Steube
b2d1f42905
Fix self-test functionality if FIXED_LOCAL_SIZE_COMP is used
...
Fix -m 25700 datatype in -a 3 mode and maximum password length in pure kernel mode
Fix -m 12500, 23700 and 23800 if password is exactly length 128
3 years ago
Jens Steube
af5d346244
Merge pull request #2894 from jtojanen/master
...
Allow async execution
3 years ago
Jens Steube
20a7b9f992
Tuning-Database: Add new module function module_extra_tuningdb_block() to extend hashcat.hctune content from a plugin
...
See src/modules/module_08900.c as an example
3 years ago
Jukka Ojanen
3d7ce7162b
Merge branch 'master' of https://github.com/hashcat/hashcat
3 years ago
Jukka Ojanen
62a06f735f
kernel_memset32(): assume offset and size are bytes, not elements
3 years ago
Jukka Ojanen
fdbfae9a28
Modify OpenCL clEnqueueFillBuffer() workaround
3 years ago
Jukka Ojanen
dbe2bad098
Maintain code style
3 years ago
Jukka Ojanen
1ee222d43f
Replace free() with hcfree()
3 years ago
Jukka Ojanen
e352a79a05
Extend context in gidd_to_pw_t()
3 years ago
Jens Steube
d4997d1255
Added support for auto-tuning --kernel-threads (-T) on startup
3 years ago
Jukka Ojanen
e154f9e781
Few cleanups
3 years ago
Jukka Ojanen
c3195d0603
Merge branch 'master' of https://github.com/hashcat/hashcat
3 years ago
Jukka Ojanen
1064cce08c
Synchronize before hooks
3 years ago
Jukka Ojanen
81c2ec3caf
Small cleanup in gidd_to_pw_t()
3 years ago
Jukka Ojanen
7a8065d090
Do not call clWaitForEvents() after spin damper when we know that event status is CL_COMPLETE
3 years ago
Jukka Ojanen
c48e6a25a8
Enqueue several commands before clFlush()
3 years ago
Jens Steube
03ed06849a
Merge branch 'HIP44' into master
3 years ago
Jens Steube
d38d40c8ba
Unlock all GPU threads for AMD GPUs if WaveFront size is 32 (basically new models)
...
Add new hash-modes to tools/benchmark_deep.pl
Fix MINGW issue on 64 bit constant in refactored kernel-accel limiting section
3 years ago
Jukka Ojanen
9ed231c99c
Add comment to blocking OpenCL calls
3 years ago
Jens Steube
a4299b74af
Memory Management: Refactored the code responsible for limiting kernel accel in order to avoid out of -host- memory situations
3 years ago
Jukka Ojanen
d7cc8d7cd2
Revert due to module_03200.c, module_25600.c and module_25800.c using device_param->kernel_dynamic_local_mem_size_memset
3 years ago
Jukka Ojanen
8b590f651b
Remove unused gpu_memset and its references
3 years ago
Jens Steube
9c0a37accf
Update driver requirement
3 years ago
Jukka Ojanen
cdf27a1cb3
Implement async run_cuda_kernel_memset() and run_cuda_kernel_memset32()
3 years ago
Jukka Ojanen
d7de3550b1
Add fast event for CUDA and HIP, and use it to synchronize check_hash()
3 years ago
Jens Steube
f6de3e61e0
Merge pull request #2906 from hashcat/master
...
Backports
3 years ago
Jens Steube
fd2cb59d26
AMD GPUs: On Apple OpenCL platform, we ask for the preferred kernel thread size rather than hard-coding 32
...
ECC secp256k1: Removed the inline assembly code for AMD GPUs because the latest JIT compilers optimize it with the same efficiency
3 years ago
Jukka Ojanen
a642f7b233
Remove synchronous GPU memory copy functions
3 years ago
Jukka Ojanen
de5200cffc
Allow async execution of backend
3 years ago
Jens Steube
72e307fbce
Merge pull request #2901 from hashcat/master
...
Backport changes
3 years ago
Jens Steube
959a232828
Merge pull request #2885 from neheb/charfixes
...
const and char fixes
3 years ago
Jens Steube
84a4058edf
Merge pull request #2900 from hashcat/master
...
Backport changes
3 years ago
Jens Steube
640d95a00f
Vendor Detection: Add "Intel" as a valid vendor name for GPU on macOS
3 years ago
Jens Steube
5ffcaa980d
HIP Backend: Added support to support HIP 4.4 and later, but added check to rule out older versions because they are incompatible
3 years ago
Rosen Penev
adaf3f293b
make const char pointers actually const
...
const char* is a non const pointer that points to const data. Add
missing const.
Signed-off-by: Rosen Penev <rosenp@gmail.com>
3 years ago
Jukka Ojanen
9f9333f2ef
Allow async execution of run_opencl_kernel_bzero(), run_hip_kernel_bzero() and run_opencl_kernel_bzero()
3 years ago
Jukka Ojanen
fafd24237d
Define HC_ALIGN macro to control data alignment and use common constant zero buffer in run_cuda_kernel_bzero(), run_hip_kernel_bzero() and run_opencl_kernel_bzero().
3 years ago
Jukka Ojanen
2c2988518d
Remove all calls to clFlush()
3 years ago
Jens Steube
5024865d87
Kernel Threads: Use warp size / wavefront size query instead of hardcoded values as base for kernel threads
...
Kernel Cache: Add kernel threads into hash computation which is later used in the kernel cache filename
Remove some unused function symbol lookups in HIP library
3 years ago
Jens Steube
c990e252d3
Added option --multiply-accel-disable (short: -M) to disable multiply the kernel-accel with the multiprocessor count automatism
3 years ago
Jens Steube
a7a899e5a4
Backport changes from #2888 to HIP backend
3 years ago
Jukka Ojanen
cb923d6e46
Replace CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK with CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK_OPTIN
3 years ago
Jukka Ojanen
d23f2d6c2f
Calculation kernel dynamic memory size based on CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK
3 years ago
Jukka Ojanen
8674e23d79
Add async HIP memcpy functions: hc_hipMemcpyDtoDAsync(), hc_hipMemcpyDtoHAsync() and hc_hipMemcpyHtoDAsync(). Implement partially async HIP memset and bzero kernels.
3 years ago
Jukka Ojanen
4263cafdcf
Add async CUDA memcpy functions: hc_cuMemcpyDtoDAsync(), hc_cuMemcpyDtoHAsync() and hc_cuMemcpyHtoDAsync(). Implement partially async CUDA memset and bzero kernels.
3 years ago
Jukka Ojanen
4c0f6dd263
Merge branch 'hashcat:master' into master
3 years ago