hashcat/OpenCL/inc_vendor.h

/**
 * Author......: See docs/credits.txt
 * License.....: MIT
 */

#ifndef _INC_VENDOR_H
#define _INC_VENDOR_H

#if defined _CPU_OPENCL_EMU_H
#define IS_NATIVE
#elif defined __CUDACC__
#define IS_CUDA
#elif defined __HIPCC__
#define IS_HIP
#else
#define IS_OPENCL
#endif

#if defined IS_NATIVE
#define CONSTANT_VK
#define CONSTANT_AS
#define GLOBAL_AS
#define LOCAL_VK
#define LOCAL_AS
#define KERNEL_FQ
#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
#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_OPENCL
#define CONSTANT_VK __constant
#define CONSTANT_AS __constant
#define GLOBAL_AS   __global
#define LOCAL_VK    __local
#define LOCAL_AS    __local
#define KERNEL_FQ   __kernel
#endif

#ifndef MAYBE_UNUSED
#define MAYBE_UNUSED
#endif

/**
 * 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
#endif

/**
 * vendor specific
 */

#if   VENDOR_ID == (1 << 0)
#define IS_AMD
#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)
#define IS_NV
#elif VENDOR_ID == (1 << 6)
#define IS_POCL
#define IS_GENERIC
#elif VENDOR_ID == (1 << 8)
#define IS_AMD_USE_HIP
#else
#define IS_GENERIC
#endif

#if defined IS_AMD && HAS_VPERM == 1
#define IS_ROCM
#endif

#define LOCAL_MEM_TYPE_LOCAL  1
#define LOCAL_MEM_TYPE_GLOBAL 2

#if LOCAL_MEM_TYPE == LOCAL_MEM_TYPE_LOCAL
#define REAL_SHM
#endif

// 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

#ifdef REAL_SHM
#define SHM_TYPE LOCAL_AS
#else
#define SHM_TYPE CONSTANT_AS
#endif

/**
 * function declarations can have a large influence depending on the opencl runtime
 * fast but pure kernels on rocm is a good example
 */

#if defined IS_AMD && defined IS_GPU
#define DECLSPEC inline static
#elif defined IS_HIP
#define DECLSPEC __device__
#else
#define DECLSPEC
#endif

#define HC_INLINE0 __attribute__ ((noinline))
#define HC_INLINE1 __attribute__ ((inline))

/**
 * AMD specific
 */

#ifdef IS_AMD
#if defined(cl_amd_media_ops)
#pragma OPENCL EXTENSION cl_amd_media_ops  : enable
#endif
#if defined(cl_amd_media_ops2)
#pragma OPENCL EXTENSION cl_amd_media_ops2 : enable
#endif
#endif

// Whitelist some OpenCL specific functions
// This could create more stable kernels on systems with bad OpenCL drivers

#ifdef IS_CUDA
#define USE_BITSELECT
#define USE_ROTATE
#endif

#ifdef IS_HIP
#define USE_BITSELECT
#define USE_ROTATE
#endif

#ifdef IS_ROCM
#define USE_BITSELECT
#define USE_ROTATE
#endif

#ifdef IS_INTEL_SDK
#ifdef IS_CPU
//#define USE_BITSELECT
//#define USE_ROTATE
#endif
#endif

#ifdef IS_OPENCL
//#define USE_BITSELECT
//#define USE_ROTATE
//#define USE_SWIZZLE
#endif

#endif
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`
Remove inline static keyword in inc_vendor.h for HIP platform since it's the default setting with HIP 4.4 2021-11-02 07:12:13 +00:00			`#define DECLSPEC __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
Renamed macro INLINE to HC_INLINE to avoid naming conflict with INLINE on MacOS 2021-12-24 15:40:43 +00:00			`#define HC_INLINE0 __attribute__ ((noinline))`
			`#define HC_INLINE1 __attribute__ ((inline))`
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 09:05:58 +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`