Gabriele Gristina
b3d3b31c3e
Metal: added support for vectors up to 4
3 years ago
Gabriele Gristina
9d36245d51
Kernels: Set the default Address Space Qualifier for any pointer, refactored / updated KERN_ATTR macros and rc4 cipher functions, in order to support Apple Metal runtime
3 years ago
Jens Steube
cb7f99ef79
Renamed macro INLINE to HC_INLINE to avoid naming conflict with INLINE on MacOS
3 years ago
Jens Steube
89cd5bd78b
Remove inline static keyword in inc_vendor.h for HIP platform since it's the default setting with HIP 4.4
3 years ago
Jens Steube
aee8e559c4
PDF Kernel (10700): Improved performance on AMD GPU by using shared memory for the scratch buffer and disable inlining to save spilling
...
Inspired by https://github.com/reger-men/hashcat/blob/6.2.4/OpenCL/m10700-optimized.cl
3 years ago
Jens Steube
3d4e2aec43
Work around segmentation fault in Intel JiT 2021.12.6.0.19_160000 compiling hc_enc_next()/hc_enc_next_global()
3 years ago
Jens Steube
3f6c5a0042
Update module_unstable_warning() for -m 172xx on HIP
3 years ago
Jens Steube
257098a301
Get rid of hip/hip_runtime.h dependancy
3 years ago
Jens Steube
45e65dd05a
Backport more ROCm based optimizations to HIP
3 years ago
Jens Steube
cf512faa53
Update large switch() cases in inc_common.cl and some inline assembly common functions for devices managed with HIP backend
3 years ago
Jens Steube
7faf6859d6
Backport hand-optimized compiler settings in modules from ROCM to HIP
...
Backport DECLSPEC settings from ROCM to HIP
3 years ago
Jens Steube
1b84a9e53b
Add missing backports from code base v6.2.2
...
Fix context to thread management
Fix missing code in selftest.c, autotune.c, hashes.c, dispatch.c and backend.c
Use IS_HIP depending code makes it easier for future optimization related to inline assembly calls - instead of using IS_CUDA || IS_HIP
See TODO markers for more optimizations / next steps
3 years ago
Jens Steube
a22f8149fc
Merge branch 'HIP' into hip
3 years ago
reger-men
ea7b74389f
First draft HIP Version
3 years ago
Jens Steube
9bf0f36d0a
Get rid of MAYBE_VOLATILE for context position by replacing it with zero length check
3 years ago
Jens Steube
ddb641b843
Add option to force disable real SHM access to be used from within the module
5 years ago
Jens Steube
61fe90bacb
Use oldschool SHA1 kernel for CPU it's slightly faster
5 years ago
Jens Steube
b4bac70bd6
Remove inline keyword in DECLSPEC for CPU
5 years ago
Jens Steube
e53bff0fb0
Reenable bitselect() and rotate() on Intel SDK
5 years ago
Jens Steube
c90d83c3eb
Prepare for UNROLL whitelisting
5 years ago
Jens Steube
3561e7b8d7
Add special ROCM detection in OpenCL/inc_vendor.h
5 years ago
Jens Steube
3a5544a554
Help some compiler with 64 bit constants
5 years ago
Jens Steube
cf4cee2f2f
Update selection of API to make use of bitselect and rotate
5 years ago
Jens Steube
89f9ef45b6
Whitelist some OpenCL specific functions
5 years ago
Jens Steube
8ff8c5d536
Add LOCAL_VK to make use of __shared__
5 years ago
Jens Steube
d0bd33c9d1
Rename CONSTANT_AS to CONSTANT_VK
5 years ago
Jens Steube
5ee033673c
Disable name mangling in NVRTC's PTX output and more
5 years ago
Jens Steube
9faba41848
Use nvrtc to compile PTX (resulting PTX not yet used)
6 years ago
Jens Steube
4b986de5fb
Prepare native CUDA hybrid integration
6 years ago
Jens Steube
38c1029f2e
Need volatile for IRIS GPU on Mac OSX for -m 2500 and -m 2501
6 years ago
jsteube
7c6970dbdd
Remove hard-coded static keyword from OpenCL kernels
6 years ago
jsteube
b7cdca09c4
OpenCL Runtime: Workaround JiT compiler error on ROCM 2.3 driver if the 'inline' keyword is used in function declaration
6 years ago
jsteube
d7d716f3ab
Make it easier to include OpenCL kernels into modules
6 years ago
jsteube
9ced13cc94
Get rid of CONSTSPEC macro in OpenCL kernels
6 years ago
Jens Steube
0fb3b3c83e
Declare internal functions in OpenCL kernels as static
6 years ago
jsteube
66d94b06e4
Get rid of src/rp_kernel_on_cpu.c and src/rp_kernel_on_cpu_optimized.c and use OpenCL emulated kernel version
6 years ago
jsteube
e80b1838e8
Rename some functions in inc_common.cl to avoid conflicts with bitops.c
6 years ago
jsteube
adeeaee84a
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
6 years ago
jsteube
7d4bea41a0
Get rid of OpenCL/inc_hash_constants.h and OpenCL/inc_hash_functions.cl
6 years ago