arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-18 22:28:14 +00:00
Code Issues Releases Wiki Activity
fbcd7bebe9
hashcat/OpenCL/inc_vendor.h

178 lines
3.2 KiB
C
Raw Normal View History

Initial commit
2015-12-04 14:47:52 +00:00
/**
Added docs/credits.txt Added docs/team.txt
2016-09-11 20:20:15 +00:00
* Author......: See docs/credits.txt
Initial commit
2015-12-04 14:47:52 +00:00
* License.....: MIT
*/
Rename some functions in inc_common.cl to avoid conflicts with bitops.c
2019-03-23 21:15:38 +00:00
#ifndef _INC_VENDOR_H
#define _INC_VENDOR_H
Prepare native CUDA hybrid integration
2019-04-25 12:45:17 +00:00
#if defined _CPU_OPENCL_EMU_H
#define IS_NATIVE
#elif defined __CUDACC__
#define IS_CUDA
First draft HIP Version
2021-07-09 03:50:40 +00:00
#elif defined __HIPCC__
#define IS_HIP
Prepare native CUDA hybrid integration
2019-04-25 12:45:17 +00:00
#else
#define IS_OPENCL
#endif
#if defined IS_NATIVE
Rename CONSTANT_AS to CONSTANT_VK
2019-05-06 12:34:16 +00:00
#define CONSTANT_VK
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 21:27:58 +00:00
#define CONSTANT_AS
#define GLOBAL_AS
Add LOCAL_VK to make use of __shared__
2019-05-07 07:01:32 +00:00
#define LOCAL_VK
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 21:27:58 +00:00
#define LOCAL_AS
#define KERNEL_FQ
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 10:38:59 +00:00
#elif defined IS_CUDA
#define CONSTANT_VK __constant__
#define CONSTANT_AS
#define GLOBAL_AS
#define LOCAL_VK __shared__
#define LOCAL_AS
#define KERNEL_FQ extern "C" __global__
#elif defined IS_HIP
Rename CONSTANT_AS to CONSTANT_VK
2019-05-06 12:34:16 +00:00
#define CONSTANT_VK __constant__
Prepare native CUDA hybrid integration
2019-04-25 12:45:17 +00:00
#define CONSTANT_AS
#define GLOBAL_AS
Add LOCAL_VK to make use of __shared__
2019-05-07 07:01:32 +00:00
#define LOCAL_VK __shared__
Prepare native CUDA hybrid integration
2019-04-25 12:45:17 +00:00
#define LOCAL_AS
Disable name mangling in NVRTC's PTX output and more
2019-05-03 13:50:07 +00:00
#define KERNEL_FQ extern "C" __global__
Prepare native CUDA hybrid integration
2019-04-25 12:45:17 +00:00
#elif defined IS_OPENCL
Rename CONSTANT_AS to CONSTANT_VK
2019-05-06 12:34:16 +00:00
#define CONSTANT_VK __constant
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 21:27:58 +00:00
#define CONSTANT_AS __constant
#define GLOBAL_AS __global
Add LOCAL_VK to make use of __shared__
2019-05-07 07:01:32 +00:00
#define LOCAL_VK __local
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 21:27:58 +00:00
#define LOCAL_AS __local
#define KERNEL_FQ __kernel
#endif
Make it easier to include OpenCL kernels into modules
2019-04-04 18:01:37 +00:00
#ifndef MAYBE_UNUSED
#define MAYBE_UNUSED
#endif
Prepare for a more dynamic #pragma unroll use
2016-05-09 19:32:12 +00:00
/**
* device type
*/
#define DEVICE_TYPE_CPU 2
#define DEVICE_TYPE_GPU 4
#define DEVICE_TYPE_ACCEL 8
#if DEVICE_TYPE == DEVICE_TYPE_CPU
#define IS_CPU
#elif DEVICE_TYPE == DEVICE_TYPE_GPU
#define IS_GPU
#elif DEVICE_TYPE == DEVICE_TYPE_ACCEL
#define IS_ACCEL
Test fix for broken AMD support after https://github.com/hashcat/oclHashcat/pull/303
2016-04-20 19:19:15 +00:00
#endif
Initial commit
2015-12-04 14:47:52 +00:00
/**
* vendor specific
*/
Prepare for a more dynamic #pragma unroll use
2016-05-09 19:32:12 +00:00
#if VENDOR_ID == (1 << 0)
Initial commit
2015-12-04 14:47:52 +00:00
#define IS_AMD
Added SIMD code for WPA/WPA2
2016-05-14 17:45:51 +00:00
#elif VENDOR_ID == (1 << 1)
#define IS_APPLE
#define IS_GENERIC
#elif VENDOR_ID == (1 << 2)
#define IS_INTEL_BEIGNET
#define IS_GENERIC
#elif VENDOR_ID == (1 << 3)
#define IS_INTEL_SDK
#define IS_GENERIC
#elif VENDOR_ID == (1 << 4)
#define IS_MESA
#define IS_GENERIC
#elif VENDOR_ID == (1 << 5)
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 10:38:59 +00:00
#define IS_NV
Added SIMD code for WPA/WPA2
2016-05-14 17:45:51 +00:00
#elif VENDOR_ID == (1 << 6)
#define IS_POCL
#define IS_GENERIC
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 10:38:59 +00:00
#elif VENDOR_ID == (1 << 8)
#define IS_AMD_USE_HIP
Added support to utilize multiple different OpenCL platforms in parallel, ex: AMD + NV
2016-01-13 16:10:40 +00:00
#else
Fix for -m 1500, -m 3000: Missing "defined" keyword Fix for -m 1000: MD4_H1 copy/paste error Fix for -m 8900, -m 9300: Invalid value for device_processor_cores for CPU devices Fix for -m 9100: Variable salt2 initializer needed some clauses for clearness Temporary limit gpu_accel for CPU devices to 1 for development phase Mark pocl as too bleeding edge for production use, recommend native drivers Remove workarounds for pocl Rename VENDOR_ID_UNKNOWN to VENDOR_ID_GENERIC in host code Rename IS_UNKNOWN to IS_GENERIC in kernel code
2016-01-07 19:14:34 +00:00
#define IS_GENERIC
Prepare for allow other OpenCL platforms thans AMD and NVidia
2016-01-04 12:17:20 +00:00
#endif
Add special ROCM detection in OpenCL/inc_vendor.h
2020-01-25 11:09:39 +00:00
#if defined IS_AMD && HAS_VPERM == 1
#define IS_ROCM
#endif
OpenCL Device: Do a real query on OpenCL local memory type instead of just assuming it
2018-11-20 09:06:34 +00:00
#define LOCAL_MEM_TYPE_LOCAL 1
#define LOCAL_MEM_TYPE_GLOBAL 2
#if LOCAL_MEM_TYPE == LOCAL_MEM_TYPE_LOCAL
Remove SCR_TYPE macro from OpenCL code Disable REAL_SHM access to AMD platform devices
2018-08-13 10:10:03 +00:00
#define REAL_SHM
#endif
Add option to force disable real SHM access to be used from within the module
2020-03-20 15:20:22 +00:00
// So far, only used by -m 22100 and only affects NVIDIA on OpenCL. CUDA seems to work fine.
#ifdef FORCE_DISABLE_SHM
#undef REAL_SHM
#endif
Remove SCR_TYPE macro from OpenCL code Disable REAL_SHM access to AMD platform devices
2018-08-13 10:10:03 +00:00
#ifdef REAL_SHM
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 21:27:58 +00:00
#define SHM_TYPE LOCAL_AS
Remove SCR_TYPE macro from OpenCL code Disable REAL_SHM access to AMD platform devices
2018-08-13 10:10:03 +00:00
#else
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 21:27:58 +00:00
#define SHM_TYPE CONSTANT_AS
Remove SCR_TYPE macro from OpenCL code Disable REAL_SHM access to AMD platform devices
2018-08-13 10:10:03 +00:00
#endif
OpenCL Kernels: Add general function declaration keyword (static inline)
2018-02-06 18:12:24 +00:00
/**
* function declarations can have a large influence depending on the opencl runtime
Remove hard-coded static keyword from OpenCL kernels
2019-04-13 16:46:19 +00:00
* fast but pure kernels on rocm is a good example
OpenCL Kernels: Add general function declaration keyword (static inline)
2018-02-06 18:12:24 +00:00
*/
Use oldschool SHA1 kernel for CPU it's slightly faster
2020-03-03 11:36:55 +00:00
#if defined IS_AMD && defined IS_GPU
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 10:38:59 +00:00
#define DECLSPEC inline static
#elif defined IS_HIP
Backport hand-optimized compiler settings in modules from ROCM to HIP Backport DECLSPEC settings from ROCM to HIP
2021-07-13 18:45:01 +00:00
#define DECLSPEC inline static __device__
Use oldschool SHA1 kernel for CPU it's slightly faster
2020-03-03 11:36:55 +00:00
#else
#define DECLSPEC
NV JiT doesn't like static inline keywords
2018-02-06 21:05:15 +00:00
#endif
OpenCL Kernels: Add general function declaration keyword (static inline)
2018-02-06 18:12:24 +00:00
Initial commit
2015-12-04 14:47:52 +00:00
/**
* AMD specific
*/
#ifdef IS_AMD
Rename some hashcat specific OpenCL functions to avoid conflicts with existing OpenCL functions from OpenCL runtime
2018-07-22 10:20:20 +00:00
#if defined(cl_amd_media_ops)
Prepare for allow other OpenCL platforms thans AMD and NVidia
2016-01-04 12:17:20 +00:00
#pragma OPENCL EXTENSION cl_amd_media_ops : enable
Rename some hashcat specific OpenCL functions to avoid conflicts with existing OpenCL functions from OpenCL runtime
2018-07-22 10:20:20 +00:00
#endif
#if defined(cl_amd_media_ops2)
Initial commit
2015-12-04 14:47:52 +00:00
#pragma OPENCL EXTENSION cl_amd_media_ops2 : enable
#endif
Rename some hashcat specific OpenCL functions to avoid conflicts with existing OpenCL functions from OpenCL runtime
2018-07-22 10:20:20 +00:00
#endif
Initial commit
2015-12-04 14:47:52 +00:00
Whitelist some OpenCL specific functions
2020-01-12 12:32:02 +00:00
// Whitelist some OpenCL specific functions
// This could create more stable kernels on systems with bad OpenCL drivers
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 10:38:59 +00:00
#ifdef IS_CUDA
Update selection of API to make use of bitselect and rotate
2020-01-20 08:20:12 +00:00
#define USE_BITSELECT
#define USE_ROTATE
#endif
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 10:38:59 +00:00
#ifdef IS_HIP
Backport more ROCm based optimizations to HIP
2021-07-15 21:34:27 +00:00
#define USE_BITSELECT
#define USE_ROTATE
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 10:38:59 +00:00
#endif
Add special ROCM detection in OpenCL/inc_vendor.h
2020-01-25 11:09:39 +00:00
#ifdef IS_ROCM
#define USE_BITSELECT
#define USE_ROTATE
#endif
Reenable bitselect() and rotate() on Intel SDK
2020-03-02 15:07:13 +00:00
#ifdef IS_INTEL_SDK
Remove inline keyword in DECLSPEC for CPU
2020-03-03 07:52:26 +00:00
#ifdef IS_CPU
//#define USE_BITSELECT
//#define USE_ROTATE
#endif
Reenable bitselect() and rotate() on Intel SDK
2020-03-02 15:07:13 +00:00
#endif
Update selection of API to make use of bitselect and rotate
2020-01-20 08:20:12 +00:00
#ifdef IS_OPENCL
Help some compiler with 64 bit constants
2020-01-21 21:09:56 +00:00
//#define USE_BITSELECT
//#define USE_ROTATE
//#define USE_SWIZZLE
Whitelist some OpenCL specific functions
2020-01-12 12:32:02 +00:00
#endif
Rename some functions in inc_common.cl to avoid conflicts with bitops.c
2019-03-23 21:15:38 +00:00
#endif
Reference in New Issue Copy Permalink