Commit Graph

3014 Commits

Author SHA1 Message Date
Jens Steube a6bc1d3cc0 Experimental kernel-thread autotuner 2019-05-11 11:58:18 +02:00
Jens Steube 2b0f657564 Limit some modules to 64 threads in order to get additional L1 cache 2019-05-10 23:51:09 +02:00
Jens Steube d59474fded Testwise unlock full thread count on NVidia 2019-05-10 17:27:15 +02:00
Jens Steube d378aa7ab9 Show host memory requirement on startup 2019-05-10 16:37:49 +02:00
Jens Steube 46f737c5af Use real constant memory on CUDA 2019-05-10 13:22:26 +02:00
Jens Steube 5d14a59304 Need 3.x nvrtc minimum 2019-05-10 10:11:12 +02:00
Jens Steube 54feb62e94 brute-force nvrtc .dll name 2019-05-09 22:17:13 +02:00
Jens Steube a2b5981303 Fix some library names 2019-05-09 21:20:50 +02:00
Jens Steube be8f29ca39 Only warn about broken NVIDIA driver 2019-05-09 16:30:08 +02:00
Jens Steube 39e150fc1e Use xxx_v2 CUDA symbols 2019-05-09 14:37:14 +02:00
Jens Steube 33028314f0 Add hc_cuCtxSetCacheConfig() 2019-05-09 00:04:05 +02:00
Jens Steube fb82bfc169 Improve thread handling based on FIXED_LOCAL_SIZE 2019-05-08 23:30:07 +02:00
Jens Steube 3a3df091c7 Fix CUDA num_elements 2019-05-08 22:42:52 +02:00
Jens Steube a01cc28ec6 Merge pull request #2014 from hashcat/master
Backports
2019-05-07 21:48:47 +02:00
Jens Steube 363e789b89 Assume local nvrtc.dll and cuda.dll on windows 2019-05-07 16:52:08 +02:00
philsmd f45a726376 fixes #2012: add Deflate support for 7-Zip using zlib 2019-05-07 15:31:11 +02:00
Jens Steube a7d04adba3 Fix opencl_devices_active and backend_devices_active 2019-05-07 14:17:29 +02:00
Jens Steube 8ff8c5d536 Add LOCAL_VK to make use of __shared__ 2019-05-07 09:01:32 +02:00
Jens Steube bbed0cd67a Fix test.sh and bitsliced algos 2019-05-06 15:06:02 +02:00
Jens Steube d0bd33c9d1 Rename CONSTANT_AS to CONSTANT_VK 2019-05-06 14:34:16 +02:00
Jens Steube 64c495dfa5 Use CUDA stream for all cuLaunchKernel() invocations 2019-05-06 11:23:34 +02:00
Jens Steube d94f582097 Replace CEILDIV() with round_up_multiple_64() 2019-05-06 09:36:07 +02:00
Jens Steube e9c04c2446 More CUDA implementation 2019-05-05 21:15:46 +02:00
Jens Steube 08dc1acc02 More CUDA rewrites 2019-05-05 11:57:54 +02:00
Jens Steube ec9925f3b1 Warnings self-check and autotune with CUDA 2019-05-04 21:52:00 +02:00
Jens Steube 4df00033d7 Prepare CUDA events 2019-05-04 10:44:03 +02:00
Jens Steube f2948460c9 Some first kernel invocations 2019-05-04 10:13:43 +02:00
Jens Steube 5ee033673c Disable name mangling in NVRTC's PTX output and more 2019-05-03 15:50:07 +02:00
Jens Steube 503304f36a Add some first CUDA device memory allocations and host buffer copies 2019-05-03 12:07:06 +02:00
Jens Steube 50a6e720ca More OpenCL variables rename 2019-05-02 17:30:46 +02:00
Jens Steube af8e317cf4 Begin renaming some OpenCL only variables 2019-05-02 17:12:59 +02:00
Jens Steube a6fa7a2749 Add support for some first CUDA module loader 2019-05-02 14:58:52 +02:00
Jens Steube 456c57a6d0 Set vector width size for CUDA 2019-05-01 18:20:19 +02:00
Jens Steube 3c4f4df771 Rename some more variables 2019-05-01 15:52:56 +02:00
Jens Steube 495d89f831 Find alias devices across different backend API's 2019-05-01 07:27:10 +02:00
Jens Steube 6fd936b43a Removed --opencl-platforms filter in order to force backend device numbers to stay constant 2019-04-30 16:24:13 +02:00
Jens Steube e3500ff4aa Add CUDA device attributes to -I 2019-04-30 13:38:44 +02:00
Jens Steube c80e516c3c Merge pull request #2007 from hashcat/master
Backport some fixes
2019-04-29 12:27:27 +02:00
Jens Steube 6caa78695f Backport -m 19500 fix to -m 19300 2019-04-29 12:26:22 +02:00
Jens Steube 922fa2e351 Use md5 for generatic unique salt for salt sorter in -m 19500 2019-04-29 12:01:53 +02:00
Jens Steube d862458ab5 Begin renaming API specific variables in backend section 2019-04-29 10:21:59 +02:00
Jens Steube d73c0ac8a9 More CUDA attribute queries 2019-04-28 18:54:26 +02:00
Jens Steube a415422123 Initialize CUDA devices and some first attribute queries 2019-04-28 14:45:50 +02:00
Jens Steube 222be0b0dc Merge pull request #2005 from hashcat/master
bcrypt changes
2019-04-27 16:35:38 +02:00
Jens Steube 200e72dba3 Limit -T maximum on -m 3200 to what's possible based on device specific shared memory available 2019-04-27 16:15:18 +02:00
Jens Steube d67de66453 Disable kernel cache on -m 3200 2019-04-27 16:00:29 +02:00
Jens Steube 58213c81d6 Add vector datatypes operators 2019-04-26 22:07:56 +02:00
Jens Steube 052e42ccef Fix CUDA_ARCH value 2019-04-26 15:14:48 +02:00
Jens Steube 06171958ee Add --gpu-architecture to NVRTC build options 2019-04-26 15:10:02 +02:00
Jens Steube 9faba41848 Use nvrtc to compile PTX (resulting PTX not yet used) 2019-04-26 13:28:44 +02:00