Commit Graph

169 Commits (9f5a22a3ab2f8c0a6ec46e86177ce1a3fe4f2600)

Author SHA1 Message Date
Jens Steube 8c3808bad5 Fix NUL filename on windows
4 years ago
Jens Steube 3e4d110fd2 Add stderr redirection the regular way
4 years ago
Jens Steube 125e9ec863 Do not redirect stderr to /dev/null to prevent rocm 3.1 from crashing on debian
4 years ago
Jens Steube f381e1bbf8 Remove force_recompile functionality, doesn't work with cubin anymore
4 years ago
Jens Steube f96e35649d Change bitsliced kernels from 3d to 2d invocation mode for slightly better performance
4 years ago
Jens Steube d9473358ef Add support for OPTS_TYPE_LOOP_EXTENDED kernel for special cases like VeraCrypt
4 years ago
Jens Steube c90d83c3eb Prepare for UNROLL whitelisting
4 years ago
Jens Steube 4788c61dd2 Add OPTI_TYPE_REGISTER_LIMIT flag to enable register limiting in CUDA
4 years ago
Jens Steube 17a64f5019 Set a fixed register count maximumfor CUDA kernel. This prevents kernels going out of control and to have negative effects on other kernels from the same source code (For instance 16600)
4 years ago
Jens Steube c40f474c2e Add special module option to indicate the kernel is using dynamic shared memory
4 years ago
Jens Steube fb7bb04587 Do not use dynamic shared memory if dynamic_local_mem_size is a multiple of local_mem_size
4 years ago
Jens Steube 96a2c36f53 Reduce CUDA Toolkit minimum version to 9.0 (even 8.0 should be sufficient)
4 years ago
Jens Steube aef53f7e10 OpenCL Runtime: Allow the kernel to access post-48k shared memory region on CUDA. Requires both module and kernel preparation
4 years ago
Jens Steube 1fc37c25f9 OpenCL Kernels: Moved "gpu_decompress", "gpu_memset" and "gpu_atinit" into new OpenCL/shared.cl in order to reduce compile time
4 years ago
Jens Steube 08163501cf Add option to disable cubin cache binaries and moved some redundant kernel load code into specific function
4 years ago
Jens Steube 01085cdab2 Move cujit_opts allocation closer to the calling functions because CUDA library needs it reinitialized after each use
4 years ago
Jens Steube 346637ec43 Improve cujit logging
4 years ago
Jens Steube 66ae5125ce Cache cubin instead of PTX to decrease startup time
4 years ago
Jens Steube cc4fd48ace Optimize hook buffer size to be copied
4 years ago
Jens Steube 041a777025 OpenCL Runtime: Unlocked maximum thread count for NVIDIA GPU
4 years ago
Jens Steube ccacc508cb Reenabled support for Intel GPU OpenCL runtime (Beignet and NEO) because a workaround was found (force -cl-std=CL2.0)
4 years ago
Jens Steube fe372dffb7 Add RDNA ISA instructions test for ADD/ADDC/SUB/SUBB
4 years ago
Jens Steube df5e2361d3 Disable inline assembly instruction tests for CUDA and refer to documented requirements
4 years ago
Jens Steube d0fb171da9 Added new options --backend-ignore-cuda and --backend-ingore-opencl, to ignore CUDA and/or OpenCL interface from being load on startup
4 years ago
Jens Steube b3690fcd05 Backport instruction test cache from CUDA to OpenCL
4 years ago
Jens Steube 2b4d0656d5 Cache inline assembly instruction check results for same devices types
4 years ago
Jens Steube 5d1d48f5d7 Do not check for COPY_PW limits in outside kernels
4 years ago
Jens Steube 53254b45aa Backport inc_ecc_secp256k1 inline assembly code for AMD ISA
5 years ago
Jens Steube bfd95d42f6 - OpenCL Runtime: Reenabled support for Intel GPU OpenCL runtime
5 years ago
Jens Steube 2884bded32 Initialize some variable to make scan-build happy
5 years ago
Jens Steube 00b9f4c557 Add kernel accel minimum limit check
5 years ago
Jens Steube 424777ae28 Add kernel accel limiter based on kernel threads to reduce host memory requirements
5 years ago
Jens Steube f7c3ced548 Fix use of calloc() in backend.c
5 years ago
Jens Steube c4dd020685 Add support for NVIDIA Jetson AGX Xavier developer kit
5 years ago
Jens Steube 53e96a12a0 Improve automatic calculation of hook threads value
5 years ago
Jens Steube fe8c17f4c7 Support pause/abort in hooks
5 years ago
Jens Steube 9c2c73c6cc Clear hook buffers after full kernel chain is finished
5 years ago
Jens Steube 7458e4f487 Add per-device available memory test of static data (hashlist, ruleset) before test of dynamic data (-n based)
5 years ago
Rosen Penev a6edb84157
Fix extra semicolon warnings
5 years ago
Jens Steube c12470b978
Merge pull request #2188 from neheb/cast
5 years ago
Jens Steube a8555fa048 Support use of all available CPU cores for hash-mode specific hooks
5 years ago
Rosen Penev fd8150769d Add casts where needed in C++ mode
5 years ago
Jens Steube 57a149276c Do alias check only in case both CUDA and OpenCL devices were detected
5 years ago
Jens Steube 97c9e86d15 Filehandling: Print a truncation warning in case an oversized line was detected
5 years ago
Rosen Penev dca1a86315
Run through Clang's bugprone-macro-parentheses
5 years ago
Rosen Penev 6dc72ebcc5
Run through Clang's readability-else-after-return
5 years ago
Rosen Penev fb75164126
Run through Clang's google-readability-casting
5 years ago
Rosen Penev 2f76326c37
Run through Clang's android-cloexec checkers
5 years ago
Rosen Penev 98e17d5774
Run through clang-tidy's readability-uppercase-literal-suffix
5 years ago
Gabriele Gristina ae62e597ce (backend) remove unused *rc* vars and cleanup
5 years ago
Jens Steube a7fd1e40f8
Merge pull request #2075 from matrix/zlib_support_2
5 years ago
Gabriele Gristina 2db6dfcd4e fix HCFILE with potfile BUG and something else related to HCFILE wrong usage
5 years ago
Gabriele Gristina ea786f715f avoid logical negation operator
5 years ago
Gabriele Gristina 3161aec3da fix the comments :)
5 years ago
Gabriele Gristina 5679ca3344 Rewrite hc_fopen to better handling file descriptor locking/unlocking functions, saving kernels binary from plain to gzip format
5 years ago
Gabriele Gristina caf34e0e83 Fix some *print* format arguments
5 years ago
Gabriele Gristina 5d3ed3e754 Remove union from HCFILE, using std file ops in ocl_check_dri, remove debug comments
5 years ago
Gabriele Gristina c2e634c426 switch is_gzip from short to bool
5 years ago
Gabriele Gristina 481c752456 No more compress functions, update example.dict.gz, remove some comments
5 years ago
Gabriele Gristina 398c89c75c switch almost all FILE ops, potfile is the only missing
5 years ago
Jens Steube 2cda236a18 OpenCL Runtime: Do not run a shared- and constant-memory size check if their memory type is of type global memory (typically CPU)
5 years ago
Jens Steube 6dfb474adf OpenCL Runtime: Do not run a shared- and constant-memory size check if their memory type is of type global memory (typically CPU)
5 years ago
Gabriele Gristina b2529af172 remove original commented code
5 years ago
Gabriele Gristina 6cb4abd526 Add zlib support v2
5 years ago
Jens Steube 955bfeaa14 Improve performance of bitsliced algorithms on ROCm
5 years ago
Jens Steube 5e0eb288c9 Use __launch_bounds__ in CUDA as replacement for reqd_work_group_size() in OpenCL
5 years ago
Jens Steube c2fc849e2c Fix minimum threads_per_block check
5 years ago
Jens Steube 0568c0746a Emulate effect of reqd_work_group_size() in CUDA
5 years ago
Jens Steube 44ecc83d82 Do some CUDA and NVRTC version checks on startup
5 years ago
Jens Steube 03ed89684e Use --restrict nvrtc option by default
5 years ago
Jens Steube 87c336e822 Fix format warning in backend.c
5 years ago
Jens Steube 1f6c82b6d1 Add hc_cuModuleLoadDataExLog wrapper function for more detailed error logging from CUDA
5 years ago
Jens Steube ce8a6fde0a Fix status screen current password query
5 years ago
Jens Steube f84eaa2e4d Fix bitsliced algorithm brute-force with CUDA
5 years ago
Jens Steube 523e0f7151 Fix free unallocated memory in case OpenCL initialization failed
5 years ago
Jens Steube bca03bb7ed CUDA offers a nice way to query available device memory, no need to brute force
5 years ago
Jens Steube a6bc1d3cc0 Experimental kernel-thread autotuner
5 years ago
Jens Steube d59474fded Testwise unlock full thread count on NVidia
5 years ago
Jens Steube d378aa7ab9 Show host memory requirement on startup
5 years ago
Jens Steube 46f737c5af Use real constant memory on CUDA
5 years ago
Jens Steube 5d14a59304 Need 3.x nvrtc minimum
5 years ago
Jens Steube 54feb62e94 brute-force nvrtc .dll name
5 years ago
Jens Steube a2b5981303 Fix some library names
5 years ago
Jens Steube be8f29ca39 Only warn about broken NVIDIA driver
5 years ago
Jens Steube 39e150fc1e Use xxx_v2 CUDA symbols
5 years ago
Jens Steube 33028314f0 Add hc_cuCtxSetCacheConfig()
5 years ago
Jens Steube fb82bfc169 Improve thread handling based on FIXED_LOCAL_SIZE
5 years ago
Jens Steube 3a3df091c7 Fix CUDA num_elements
5 years ago
Jens Steube 363e789b89 Assume local nvrtc.dll and cuda.dll on windows
5 years ago
Jens Steube a7d04adba3 Fix opencl_devices_active and backend_devices_active
5 years ago
Jens Steube 8ff8c5d536 Add LOCAL_VK to make use of __shared__
5 years ago
Jens Steube bbed0cd67a Fix test.sh and bitsliced algos
5 years ago
Jens Steube d0bd33c9d1 Rename CONSTANT_AS to CONSTANT_VK
5 years ago
Jens Steube 64c495dfa5 Use CUDA stream for all cuLaunchKernel() invocations
5 years ago
Jens Steube d94f582097 Replace CEILDIV() with round_up_multiple_64()
5 years ago
Jens Steube e9c04c2446 More CUDA implementation
5 years ago
Jens Steube 08dc1acc02 More CUDA rewrites
5 years ago
Jens Steube ec9925f3b1 Warnings self-check and autotune with CUDA
5 years ago
Jens Steube 4df00033d7 Prepare CUDA events
5 years ago
Jens Steube f2948460c9 Some first kernel invocations
5 years ago