Gabriele Gristina
b3d3b31c3e
Metal: added support for vectors up to 4
2022-02-10 21:53:08 +01:00
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
2022-02-04 19:54:00 +01:00
Jens Steube
cb7f99ef79
Renamed macro INLINE to HC_INLINE to avoid naming conflict with INLINE on MacOS
2021-12-24 16:40:43 +01:00
Jens Steube
89cd5bd78b
Remove inline static keyword in inc_vendor.h for HIP platform since it's the default setting with HIP 4.4
2021-11-02 08:12:13 +01:00
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
2021-10-31 10:05:58 +01:00
Jens Steube
3d4e2aec43
Work around segmentation fault in Intel JiT 2021.12.6.0.19_160000 compiling hc_enc_next()/hc_enc_next_global()
2021-08-03 08:34:37 +02:00
Jens Steube
3f6c5a0042
Update module_unstable_warning() for -m 172xx on HIP
2021-07-23 21:09:55 +02:00
Jens Steube
257098a301
Get rid of hip/hip_runtime.h dependancy
2021-07-18 21:14:45 +02:00
Jens Steube
45e65dd05a
Backport more ROCm based optimizations to HIP
2021-07-15 23:34:27 +02:00
Jens Steube
cf512faa53
Update large switch() cases in inc_common.cl and some inline assembly common functions for devices managed with HIP backend
2021-07-14 17:06:20 +02:00
Jens Steube
7faf6859d6
Backport hand-optimized compiler settings in modules from ROCM to HIP
...
Backport DECLSPEC settings from ROCM to HIP
2021-07-13 20:45:01 +02:00
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
2021-07-11 12:38:59 +02:00
Jens Steube
a22f8149fc
Merge branch 'HIP' into hip
2021-07-10 21:34:09 +02:00
reger-men
ea7b74389f
First draft HIP Version
2021-07-09 03:50:40 +00:00
Jens Steube
9bf0f36d0a
Get rid of MAYBE_VOLATILE for context position by replacing it with zero length check
2021-05-09 11:43:32 +02:00
Jens Steube
ddb641b843
Add option to force disable real SHM access to be used from within the module
2020-03-20 16:20:22 +01:00
Jens Steube
61fe90bacb
Use oldschool SHA1 kernel for CPU it's slightly faster
2020-03-03 12:36:55 +01:00
Jens Steube
b4bac70bd6
Remove inline keyword in DECLSPEC for CPU
2020-03-03 08:52:26 +01:00
Jens Steube
e53bff0fb0
Reenable bitselect() and rotate() on Intel SDK
2020-03-02 16:07:13 +01:00
Jens Steube
c90d83c3eb
Prepare for UNROLL whitelisting
2020-02-15 12:44:12 +01:00
Jens Steube
3561e7b8d7
Add special ROCM detection in OpenCL/inc_vendor.h
2020-01-25 12:09:39 +01:00
Jens Steube
3a5544a554
Help some compiler with 64 bit constants
2020-01-21 22:09:56 +01:00
Jens Steube
cf4cee2f2f
Update selection of API to make use of bitselect and rotate
2020-01-20 09:20:12 +01:00
Jens Steube
89f9ef45b6
Whitelist some OpenCL specific functions
2020-01-12 13:32:02 +01:00
Jens Steube
8ff8c5d536
Add LOCAL_VK to make use of __shared__
2019-05-07 09:01:32 +02:00
Jens Steube
d0bd33c9d1
Rename CONSTANT_AS to CONSTANT_VK
2019-05-06 14:34:16 +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
9faba41848
Use nvrtc to compile PTX (resulting PTX not yet used)
2019-04-26 13:28:44 +02:00
Jens Steube
4b986de5fb
Prepare native CUDA hybrid integration
2019-04-25 14:45:17 +02:00
Jens Steube
38c1029f2e
Need volatile for IRIS GPU on Mac OSX for -m 2500 and -m 2501
2019-04-17 13:21:35 +02:00
jsteube
7c6970dbdd
Remove hard-coded static keyword from OpenCL kernels
2019-04-13 18:46:19 +02:00
jsteube
b7cdca09c4
OpenCL Runtime: Workaround JiT compiler error on ROCM 2.3 driver if the 'inline' keyword is used in function declaration
2019-04-13 13:46:55 +02:00
jsteube
d7d716f3ab
Make it easier to include OpenCL kernels into modules
2019-04-04 20:01:37 +02:00
jsteube
9ced13cc94
Get rid of CONSTSPEC macro in OpenCL kernels
2019-04-04 10:15:34 +02:00
Jens Steube
0fb3b3c83e
Declare internal functions in OpenCL kernels as static
2019-03-26 11:03:25 +01:00
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
2019-03-25 12:24:04 +01:00
jsteube
e80b1838e8
Rename some functions in inc_common.cl to avoid conflicts with bitops.c
2019-03-23 22:15:38 +01:00
jsteube
adeeaee84a
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 22:27:58 +01:00
jsteube
7d4bea41a0
Get rid of OpenCL/inc_hash_constants.h and OpenCL/inc_hash_functions.cl
2019-03-21 23:00:38 +01:00
