From 5ffcaa980d0e4d95cd573bebffa35e66883b7432 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Fri, 23 Jul 2021 16:04:34 +0200 Subject: [PATCH 1/4] HIP Backend: Added support to support HIP 4.4 and later, but added check to rule out older versions because they are incompatible --- OpenCL/inc_platform.cl | 25 +- OpenCL/inc_types.h | 87 +- OpenCL/m08900-pure.cl | 13 +- OpenCL/m15700-pure.cl | 13 +- OpenCL/m22700-pure.cl | 13 +- docs/changes.txt | 1 + include/backend.h | 89 +-- include/ext_hip.h | 1437 +++++++++------------------------ include/ext_hiprtc.h | 61 +- include/types.h | 143 ++-- src/backend.c | 1699 +++++++++++++++------------------------- src/terminal.c | 10 +- 12 files changed, 1181 insertions(+), 2410 deletions(-) diff --git a/OpenCL/inc_platform.cl b/OpenCL/inc_platform.cl index df0e210a4..47cfc9a84 100644 --- a/OpenCL/inc_platform.cl +++ b/OpenCL/inc_platform.cl @@ -193,33 +193,21 @@ DECLSPEC u32 hc_atomic_dec (GLOBAL_AS u32 *p) { volatile const u32 val = 1; - return __atomic_fetch_sub (p, val, __ATOMIC_RELAXED); + return atomicSub (p, val); } DECLSPEC u32 hc_atomic_inc (GLOBAL_AS u32 *p) { volatile const u32 val = 1; - return __atomic_fetch_add (p, val, __ATOMIC_RELAXED); + return atomicAdd (p, val); } DECLSPEC u32 hc_atomic_or (GLOBAL_AS u32 *p, volatile const u32 val) { - return __atomic_fetch_or (p, val, __ATOMIC_RELAXED); + return atomicOr (p, val); } -extern "C" __device__ __attribute__((pure)) double __ocml_log2_f64(double); - -DECLSPEC double log2 (double x) -{ - return __ocml_log2_f64 (x); -} - -extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_id(uint); -extern "C" __device__ __attribute__((const)) size_t __ockl_get_group_id(uint); -extern "C" __device__ __attribute__((const)) size_t __ockl_get_local_size(uint); -extern "C" __device__ __attribute__((const)) size_t __ockl_get_num_groups(uint); - DECLSPEC size_t get_global_id (const u32 dimindx) { return (__ockl_get_group_id (dimindx) * __ockl_get_local_size (dimindx)) + __ockl_get_local_id (dimindx); @@ -308,11 +296,8 @@ DECLSPEC u64 rotr64_S (const u64 a, const int n) return out.v64; } -extern "C" __device__ int printf(const char *fmt, ...); -//int printf(__constant const char* st, ...) __attribute__((format(printf, 1, 2))); - -#define FIXED_THREAD_COUNT(n) __attribute__((amdgpu_flat_work_group_size (1, (n)))) -#define SYNC_THREADS() __builtin_amdgcn_s_barrier () +#define FIXED_THREAD_COUNT(n) __launch_bounds__((n), 0) +#define SYNC_THREADS() __syncthreads () #endif #ifdef IS_OPENCL diff --git a/OpenCL/inc_types.h b/OpenCL/inc_types.h index 0c715da66..4caf595de 100644 --- a/OpenCL/inc_types.h +++ b/OpenCL/inc_types.h @@ -21,96 +21,19 @@ typedef unsigned char uchar; typedef unsigned short ushort; typedef unsigned int uint; -typedef unsigned long long ulong; +typedef unsigned long ulong; +typedef unsigned long long ullong; #endif -#ifdef IS_HIP -// https://github.com/llvm-mirror/clang/blob/master/lib/Headers/opencl-c-base.h - -// built-in scalar data types: - -/** - * An unsigned 8-bit integer. - */ -typedef unsigned char uchar; - -/** - * An unsigned 16-bit integer. - */ -typedef unsigned short ushort; - -/** - * An unsigned 32-bit integer. - */ -typedef unsigned int uint; - -/** - * An unsigned 64-bit integer. - */ -typedef unsigned long ulong; - -/** - * The unsigned integer type of the result of the sizeof operator. This - * is a 32-bit unsigned integer if CL_DEVICE_ADDRESS_BITS - * defined in table 4.3 is 32-bits and is a 64-bit unsigned integer if - * CL_DEVICE_ADDRESS_BITS is 64-bits. - */ -typedef __SIZE_TYPE__ size_t; - -// built-in vector data types: -typedef char char2 __attribute__((ext_vector_type(2))); -typedef char char3 __attribute__((ext_vector_type(3))); -typedef char char4 __attribute__((ext_vector_type(4))); -typedef char char8 __attribute__((ext_vector_type(8))); -typedef char char16 __attribute__((ext_vector_type(16))); -typedef uchar uchar2 __attribute__((ext_vector_type(2))); -typedef uchar uchar3 __attribute__((ext_vector_type(3))); -typedef uchar uchar4 __attribute__((ext_vector_type(4))); -typedef uchar uchar8 __attribute__((ext_vector_type(8))); -typedef uchar uchar16 __attribute__((ext_vector_type(16))); -typedef short short2 __attribute__((ext_vector_type(2))); -typedef short short3 __attribute__((ext_vector_type(3))); -typedef short short4 __attribute__((ext_vector_type(4))); -typedef short short8 __attribute__((ext_vector_type(8))); -typedef short short16 __attribute__((ext_vector_type(16))); -typedef ushort ushort2 __attribute__((ext_vector_type(2))); -typedef ushort ushort3 __attribute__((ext_vector_type(3))); -typedef ushort ushort4 __attribute__((ext_vector_type(4))); -typedef ushort ushort8 __attribute__((ext_vector_type(8))); -typedef ushort ushort16 __attribute__((ext_vector_type(16))); -typedef int int2 __attribute__((ext_vector_type(2))); -typedef int int3 __attribute__((ext_vector_type(3))); -typedef int int4 __attribute__((ext_vector_type(4))); -typedef int int8 __attribute__((ext_vector_type(8))); -typedef int int16 __attribute__((ext_vector_type(16))); -typedef uint uint2 __attribute__((ext_vector_type(2))); -typedef uint uint3 __attribute__((ext_vector_type(3))); -typedef uint uint4 __attribute__((ext_vector_type(4))); -typedef uint uint8 __attribute__((ext_vector_type(8))); -typedef uint uint16 __attribute__((ext_vector_type(16))); -typedef long long2 __attribute__((ext_vector_type(2))); -typedef long long3 __attribute__((ext_vector_type(3))); -typedef long long4 __attribute__((ext_vector_type(4))); -typedef long long8 __attribute__((ext_vector_type(8))); -typedef long long16 __attribute__((ext_vector_type(16))); -typedef ulong ulong2 __attribute__((ext_vector_type(2))); -typedef ulong ulong3 __attribute__((ext_vector_type(3))); -typedef ulong ulong4 __attribute__((ext_vector_type(4))); -typedef ulong ulong8 __attribute__((ext_vector_type(8))); -typedef ulong ulong16 __attribute__((ext_vector_type(16))); -typedef float float2 __attribute__((ext_vector_type(2))); -typedef float float3 __attribute__((ext_vector_type(3))); -typedef float float4 __attribute__((ext_vector_type(4))); -typedef float float8 __attribute__((ext_vector_type(8))); -typedef float float16 __attribute__((ext_vector_type(16))); - +#ifdef IS_OPENCL +typedef ulong ullong; #endif #ifdef KERNEL_STATIC typedef uchar u8; typedef ushort u16; typedef uint u32; -typedef ulong u64; +typedef ullong u64; #else typedef uint8_t u8; typedef uint16_t u16; diff --git a/OpenCL/m08900-pure.cl b/OpenCL/m08900-pure.cl index 2bd1de39a..74f8a0e66 100644 --- a/OpenCL/m08900-pure.cl +++ b/OpenCL/m08900-pure.cl @@ -24,7 +24,7 @@ typedef struct } scrypt_tmp_t; -#if defined IS_CUDA +#if defined IS_CUDA || defined IS_HIP inline __device__ uint4 operator & (const uint4 a, const u32 b) { return make_uint4 ((a.x & b ), (a.y & b ), (a.z & b ), (a.w & b )); } inline __device__ uint4 operator << (const uint4 a, const u32 b) { return make_uint4 ((a.x << b ), (a.y << b ), (a.z << b ), (a.w << b )); } @@ -41,15 +41,6 @@ inline __device__ uint4 rotate (const uint4 a, const int n) #endif -#if defined IS_HIP - -inline __device__ uint4 rotate (const uint4 a, const int n) -{ - return ((a << n) | ((a >> (32 - n)))); -} - -#endif - DECLSPEC uint4 hc_swap32_4 (uint4 v) { return (rotate ((v & 0x00FF00FF), 24u) | rotate ((v & 0xFF00FF00), 8u)); @@ -66,7 +57,7 @@ DECLSPEC uint4 hc_swap32_4 (uint4 v) #define ADD_ROTATE_XOR(r,i1,i2,s) (r) ^= rotate ((i1) + (i2), (s)); -#if defined IS_CUDA +#if defined IS_CUDA || defined IS_HIP #define SALSA20_2R() \ { \ diff --git a/OpenCL/m15700-pure.cl b/OpenCL/m15700-pure.cl index 10a7aaa14..09819b085 100644 --- a/OpenCL/m15700-pure.cl +++ b/OpenCL/m15700-pure.cl @@ -31,7 +31,7 @@ typedef struct ethereum_scrypt } ethereum_scrypt_t; -#if defined IS_CUDA +#if defined IS_CUDA || defined IS_HIP inline __device__ uint4 operator & (const uint4 a, const u32 b) { return make_uint4 ((a.x & b ), (a.y & b ), (a.z & b ), (a.w & b )); } inline __device__ uint4 operator << (const uint4 a, const u32 b) { return make_uint4 ((a.x << b ), (a.y << b ), (a.z << b ), (a.w << b )); } @@ -48,15 +48,6 @@ inline __device__ uint4 rotate (const uint4 a, const int n) #endif -#if defined IS_HIP - -inline __device__ uint4 rotate (const uint4 a, const int n) -{ - return ((a << n) | ((a >> (32 - n)))); -} - -#endif - DECLSPEC uint4 hc_swap32_4 (uint4 v) { return (rotate ((v & 0x00FF00FF), 24u) | rotate ((v & 0xFF00FF00), 8u)); @@ -73,7 +64,7 @@ DECLSPEC uint4 hc_swap32_4 (uint4 v) #define ADD_ROTATE_XOR(r,i1,i2,s) (r) ^= rotate ((i1) + (i2), (s)); -#if defined IS_CUDA +#if defined IS_CUDA || defined IS_HIP #define SALSA20_2R() \ { \ diff --git a/OpenCL/m22700-pure.cl b/OpenCL/m22700-pure.cl index 4ecc345ce..a28b458c2 100644 --- a/OpenCL/m22700-pure.cl +++ b/OpenCL/m22700-pure.cl @@ -72,7 +72,7 @@ DECLSPEC int is_valid_bitcoinj (const u32 *w) return 1; } -#if defined IS_CUDA +#if defined IS_CUDA || defined IS_HIP inline __device__ uint4 operator & (const uint4 a, const u32 b) { return make_uint4 ((a.x & b ), (a.y & b ), (a.z & b ), (a.w & b )); } inline __device__ uint4 operator << (const uint4 a, const u32 b) { return make_uint4 ((a.x << b ), (a.y << b ), (a.z << b ), (a.w << b )); } @@ -89,15 +89,6 @@ inline __device__ uint4 rotate (const uint4 a, const int n) #endif -#if defined IS_HIP - -inline __device__ uint4 rotate (const uint4 a, const int n) -{ - return ((a << n) | ((a >> (32 - n)))); -} - -#endif - DECLSPEC uint4 hc_swap32_4 (uint4 v) { return (rotate ((v & 0x00FF00FF), 24u) | rotate ((v & 0xFF00FF00), 8u)); @@ -114,7 +105,7 @@ DECLSPEC uint4 hc_swap32_4 (uint4 v) #define ADD_ROTATE_XOR(r,i1,i2,s) (r) ^= rotate ((i1) + (i2), (s)); -#if defined IS_CUDA +#if defined IS_CUDA || defined IS_HIP #define SALSA20_2R() \ { \ diff --git a/docs/changes.txt b/docs/changes.txt index 92cebc836..50bbfb5ae 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -5,6 +5,7 @@ ## - Added option --multiply-accel-disable (short: -M) to disable multiply the kernel-accel with the multiprocessor count automatism +- HIP Backend: Added support to support HIP 4.4 and later, but added check to rule out older versions because they are incompatible ## ## Bugs diff --git a/include/backend.h b/include/backend.h index 957ac229d..30e79bc89 100644 --- a/include/backend.h +++ b/include/backend.h @@ -88,53 +88,50 @@ int hc_cuLinkAddData (hashcat_ctx_t *hashcat_ctx, CUlinkState state, int hc_cuLinkDestroy (hashcat_ctx_t *hashcat_ctx, CUlinkState state); int hc_cuLinkComplete (hashcat_ctx_t *hashcat_ctx, CUlinkState state, void **cubinOut, size_t *sizeOut); -int hc_nvrtcCreateProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram *prog, const char *src, const char *name, int numHeaders, const char * const *headers, const char * const *includeNames); -int hc_nvrtcDestroyProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram *prog); -int hc_nvrtcCompileProgram (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, int numOptions, const char * const *options); -int hc_nvrtcGetProgramLogSize (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *logSizeRet); -int hc_nvrtcGetProgramLog (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *log); -int hc_nvrtcGetPTXSize (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, size_t *ptxSizeRet); -int hc_nvrtcGetPTX (hashcat_ctx_t *hashcat_ctx, nvrtcProgram prog, char *ptx); -int hc_nvrtcVersion (hashcat_ctx_t *hashcat_ctx, int *major, int *minor); +int hc_hipCreateProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram *prog, const char *src, const char *name, int numHeaders, const char * const *headers, const char * const *includeNames); +int hc_hipDestroyProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram *prog); +int hc_hipCompileProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, int numOptions, const char * const *options); +int hc_hipGetProgramLogSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *logSizeRet); +int hc_hipGetProgramLog (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *log); +int hc_hipGetCodeSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *codeSizeRet); +int hc_hipGetCode (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *code); -int hc_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx, unsigned int flags, HIPdevice dev); -int hc_hipCtxDestroy (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx); -int hc_hipCtxSetCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx); -int hc_hipCtxSetCacheConfig (hashcat_ctx_t *hashcat_ctx, HIPfunc_cache config); +int hc_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, hipCtx_t *pctx, unsigned int flags, hipDevice_t dev); +int hc_hipCtxDestroy (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx); +int hc_hipCtxPopCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t *pctx); +int hc_hipCtxPushCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx); +int hc_hipCtxSetCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx); int hc_hipCtxSynchronize (hashcat_ctx_t *hashcat_ctx); -int hc_hipDeviceGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPdevice_attribute attrib, HIPdevice dev); +int hc_hipDeviceGet (hashcat_ctx_t *hashcat_ctx, hipDevice_t *device, int ordinal); +int hc_hipDeviceGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, hipDeviceAttribute_t attrib, hipDevice_t dev); int hc_hipDeviceGetCount (hashcat_ctx_t *hashcat_ctx, int *count); -int hc_hipDeviceGet (hashcat_ctx_t *hashcat_ctx, HIPdevice *device, int ordinal); -int hc_hipDeviceGetName (hashcat_ctx_t *hashcat_ctx, char *name, int len, HIPdevice dev); -int hc_hipDeviceTotalMem (hashcat_ctx_t *hashcat_ctx, size_t *bytes, HIPdevice dev); +int hc_hipDeviceGetName (hashcat_ctx_t *hashcat_ctx, char *name, int len, hipDevice_t dev); +int hc_hipDeviceTotalMem (hashcat_ctx_t *hashcat_ctx, size_t *bytes, hipDevice_t dev); int hc_hipDriverGetVersion (hashcat_ctx_t *hashcat_ctx, int *driverVersion); -int hc_hipEventCreate (hashcat_ctx_t *hashcat_ctx, HIPevent *phEvent, unsigned int Flags); -int hc_hipEventDestroy (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); -int hc_hipEventElapsedTime (hashcat_ctx_t *hashcat_ctx, float *pMilliseconds, HIPevent hStart, HIPevent hEnd); -int hc_hipEventQuery (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); -int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent, HIPstream hStream); -int hc_hipEventSynchronize (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent); -int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPfunction_attribute attrib, HIPfunction hfunc); -//int hc_hipFuncSetAttribute (hashcat_ctx_t *hashcat_ctx, HIPfunction hfunc, HIPfunction_attribute attrib, int value); +int hc_hipEventCreate (hashcat_ctx_t *hashcat_ctx, hipEvent_t *phEvent, unsigned int Flags); +int hc_hipEventDestroy (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent); +int hc_hipEventElapsedTime (hashcat_ctx_t *hashcat_ctx, float *pMilliseconds, hipEvent_t hStart, hipEvent_t hEnd); +int hc_hipEventQuery (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent); +int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent, hipStream_t hStream); +int hc_hipEventSynchronize (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent); +int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, hipFunction_attribute attrib, hipFunction_t hfunc); int hc_hipInit (hashcat_ctx_t *hashcat_ctx, unsigned int Flags); -int hc_hipLaunchKernel (hashcat_ctx_t *hashcat_ctx, HIPfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, HIPstream hStream, void **kernelParams, void **extra); -int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr *dptr, size_t bytesize); -int hc_hipMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, HIPdeviceptr srcDevice, size_t ByteCount); -int hc_hipMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, HIPdeviceptr srcDevice, size_t ByteCount); -int hc_hipMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, const void *srcHost, size_t ByteCount); -int hc_hipMemFree (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dptr); -int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, HIPfunction *hfunc, HIPmodule hmod, const char *name); -int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, HIPmodule *module, const void *image, unsigned int numOptions, HIPjit_option *options, void **optionValues); -int hc_hipModuleUnload (hashcat_ctx_t *hashcat_ctx, HIPmodule hmod); -int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, HIPstream *phStream, unsigned int Flags); -int hc_hipStreamDestroy (hashcat_ctx_t *hashcat_ctx, HIPstream hStream); -int hc_hipStreamSynchronize (hashcat_ctx_t *hashcat_ctx, HIPstream hStream); -int hc_hipCtxPushCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx); -int hc_hipCtxPopCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx); -int hc_hipLinkCreate (hashcat_ctx_t *hashcat_ctx, unsigned int numOptions, HIPjit_option *options, void **optionValues, HIPlinkState *stateOut); -int hc_hipLinkAddData (hashcat_ctx_t *hashcat_ctx, HIPlinkState state, HIPjitInputType type, void *data, size_t size, const char *name, unsigned int numOptions, HIPjit_option *options, void **optionValues); -int hc_hipLinkDestroy (hashcat_ctx_t *hashcat_ctx, HIPlinkState state); -int hc_hipLinkComplete (hashcat_ctx_t *hashcat_ctx, HIPlinkState state, void **cubinOut, size_t *sizeOut); +int hc_hipLaunchKernel (hashcat_ctx_t *hashcat_ctx, hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra); +int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize); +int hc_hipMemFree (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dptr); +int hc_hipMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount); +int hc_hipMemcpyDtoHAsync (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream); +int hc_hipMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount); +int hc_hipMemcpyDtoDAsync (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream); +int hc_hipMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount); +int hc_hipMemcpyHtoDAsync (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount, hipStream_t hStream); +int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, hipFunction_t *hfunc, hipModule_t hmod, const char *name); +int hc_hipModuleGetGlobal (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name); +int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, hipModule_t *module, const void *image, unsigned int numOptions, hipJitOption *options, void **optionValues); +int hc_hipModuleUnload (hashcat_ctx_t *hashcat_ctx, hipModule_t hmod); +int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, hipStream_t *phStream, unsigned int Flags); +int hc_hipStreamDestroy (hashcat_ctx_t *hashcat_ctx, hipStream_t hStream); +int hc_hipStreamSynchronize (hashcat_ctx_t *hashcat_ctx, hipStream_t hStream); int hc_clBuildProgram (hashcat_ctx_t *hashcat_ctx, cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data); int hc_clCompileProgram (hashcat_ctx_t *hashcat_ctx, cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, cl_uint num_input_headers, const cl_program *input_headers, const char **header_include_names, void (CL_CALLBACK *pfn_notify) (cl_program program, void *user_data), void *user_data); @@ -183,10 +180,10 @@ int run_cuda_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param int run_cuda_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u32 value, const u64 size); int run_cuda_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 size); -int run_hip_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 num); -int run_hip_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 num); -int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u32 value, const u64 size); -int run_hip_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 size); +int run_hip_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 num); +int run_hip_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 num); +int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u32 value, const u64 size); +int run_hip_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 size); int run_opencl_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 num); int run_opencl_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 num); diff --git a/include/ext_hip.h b/include/ext_hip.h index a99d1e5a9..1477c20c4 100644 --- a/include/ext_hip.h +++ b/include/ext_hip.h @@ -6,995 +6,344 @@ #ifndef _EXT_HIP_H #define _EXT_HIP_H -/** - * TODO: FIX ME - */ +// The general Idea with HIP is to use it for AMD GPU since we use CUDA for NV +// Therefore, we need to take certain items, such as hipDeviceptr_t from driver specific paths like amd_driver_types.h +// We just need to keep this in mind in case we need to update these constants from future SDK versions -#define __HIP_API_VERSION 4221131 +// start: amd_driver_types.h -/** - * HIP device pointer - * HIPdeviceptr is defined as an unsigned integer type whose size matches the size of a pointer on the target platform. - */ -#if __HIP_API_VERSION >= 3020 +typedef void* hipDeviceptr_t; -#if defined(_WIN64) || defined(__LP64__) -typedef unsigned long long HIPdeviceptr; +typedef enum hipFunction_attribute { + HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, + HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, + HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES, + HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES, + HIP_FUNC_ATTRIBUTE_NUM_REGS, + HIP_FUNC_ATTRIBUTE_PTX_VERSION, + HIP_FUNC_ATTRIBUTE_BINARY_VERSION, + HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA, + HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, + HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT, + HIP_FUNC_ATTRIBUTE_MAX +}hipFunction_attribute; + +// stop: amd_driver_types.h + +// start: hip_runtime_api.h + +typedef int hipDevice_t; +typedef struct ihipCtx_t* hipCtx_t; +typedef struct ihipEvent_t* hipEvent_t; +typedef struct ihipStream_t* hipStream_t; +typedef struct ihipModule_t* hipModule_t; +typedef struct ihipModuleSymbol_t* hipFunction_t; + +// Ignoring error-code return values from hip APIs is discouraged. On C++17, +// we can make that yield a warning +#if __cplusplus >= 201703L +#define __HIP_NODISCARD [[nodiscard]] #else -typedef unsigned int HIPdeviceptr; +#define __HIP_NODISCARD #endif -#endif /* __HIP_API_VERSION >= 3020 */ +typedef enum __HIP_NODISCARD hipError_t { + hipSuccess = 0, ///< Successful completion. + hipErrorInvalidValue = 1, ///< One or more of the parameters passed to the API call is NULL + ///< or not in an acceptable range. + hipErrorOutOfMemory = 2, + // Deprecated + hipErrorMemoryAllocation = 2, ///< Memory allocation error. + hipErrorNotInitialized = 3, + // Deprecated + hipErrorInitializationError = 3, + hipErrorDeinitialized = 4, + hipErrorProfilerDisabled = 5, + hipErrorProfilerNotInitialized = 6, + hipErrorProfilerAlreadyStarted = 7, + hipErrorProfilerAlreadyStopped = 8, + hipErrorInvalidConfiguration = 9, + hipErrorInvalidPitchValue = 12, + hipErrorInvalidSymbol = 13, + hipErrorInvalidDevicePointer = 17, ///< Invalid Device Pointer + hipErrorInvalidMemcpyDirection = 21, ///< Invalid memory copy direction + hipErrorInsufficientDriver = 35, + hipErrorMissingConfiguration = 52, + hipErrorPriorLaunchFailure = 53, + hipErrorInvalidDeviceFunction = 98, + hipErrorNoDevice = 100, ///< Call to hipGetDeviceCount returned 0 devices + hipErrorInvalidDevice = 101, ///< DeviceID must be in range 0...#compute-devices. + hipErrorInvalidImage = 200, + hipErrorInvalidContext = 201, ///< Produced when input context is invalid. + hipErrorContextAlreadyCurrent = 202, + hipErrorMapFailed = 205, + // Deprecated + hipErrorMapBufferObjectFailed = 205, ///< Produced when the IPC memory attach failed from ROCr. + hipErrorUnmapFailed = 206, + hipErrorArrayIsMapped = 207, + hipErrorAlreadyMapped = 208, + hipErrorNoBinaryForGpu = 209, + hipErrorAlreadyAcquired = 210, + hipErrorNotMapped = 211, + hipErrorNotMappedAsArray = 212, + hipErrorNotMappedAsPointer = 213, + hipErrorECCNotCorrectable = 214, + hipErrorUnsupportedLimit = 215, + hipErrorContextAlreadyInUse = 216, + hipErrorPeerAccessUnsupported = 217, + hipErrorInvalidKernelFile = 218, ///< In CUDA DRV, it is CUDA_ERROR_INVALID_PTX + hipErrorInvalidGraphicsContext = 219, + hipErrorInvalidSource = 300, + hipErrorFileNotFound = 301, + hipErrorSharedObjectSymbolNotFound = 302, + hipErrorSharedObjectInitFailed = 303, + hipErrorOperatingSystem = 304, + hipErrorInvalidHandle = 400, + // Deprecated + hipErrorInvalidResourceHandle = 400, ///< Resource handle (hipEvent_t or hipStream_t) invalid. + hipErrorNotFound = 500, + hipErrorNotReady = 600, ///< Indicates that asynchronous operations enqueued earlier are not + ///< ready. This is not actually an error, but is used to distinguish + ///< from hipSuccess (which indicates completion). APIs that return + ///< this error include hipEventQuery and hipStreamQuery. + hipErrorIllegalAddress = 700, + hipErrorLaunchOutOfResources = 701, ///< Out of resources error. + hipErrorLaunchTimeOut = 702, + hipErrorPeerAccessAlreadyEnabled = + 704, ///< Peer access was already enabled from the current device. + hipErrorPeerAccessNotEnabled = + 705, ///< Peer access was never enabled from the current device. + hipErrorSetOnActiveProcess = 708, + hipErrorContextIsDestroyed = 709, + hipErrorAssert = 710, ///< Produced when the kernel calls assert. + hipErrorHostMemoryAlreadyRegistered = + 712, ///< Produced when trying to lock a page-locked memory. + hipErrorHostMemoryNotRegistered = + 713, ///< Produced when trying to unlock a non-page-locked memory. + hipErrorLaunchFailure = + 719, ///< An exception occurred on the device while executing a kernel. + hipErrorCooperativeLaunchTooLarge = + 720, ///< This error indicates that the number of blocks launched per grid for a kernel + ///< that was launched via cooperative launch APIs exceeds the maximum number of + ///< allowed blocks for the current device + hipErrorNotSupported = 801, ///< Produced when the hip API is not supported/implemented + hipErrorStreamCaptureUnsupported = 900, ///< The operation is not permitted when the stream + ///< is capturing. + hipErrorStreamCaptureInvalidated = 901, ///< The current capture sequence on the stream + ///< has been invalidated due to a previous error. + hipErrorStreamCaptureMerge = 902, ///< The operation would have resulted in a merge of + ///< two independent capture sequences. + hipErrorStreamCaptureUnmatched = 903, ///< The capture was not initiated in this stream. + hipErrorStreamCaptureUnjoined = 904, ///< The capture sequence contains a fork that was not + ///< joined to the primary stream. + hipErrorStreamCaptureIsolation = 905, ///< A dependency would have been created which crosses + ///< the capture sequence boundary. Only implicit + ///< in-stream ordering dependencies are allowed + ///< to cross the boundary + hipErrorStreamCaptureImplicit = 906, ///< The operation would have resulted in a disallowed + ///< implicit dependency on a current capture sequence + ///< from hipStreamLegacy. + hipErrorCapturedEvent = 907, ///< The operation is not permitted on an event which was last + ///< recorded in a capturing stream. + hipErrorStreamCaptureWrongThread = 908, ///< A stream capture sequence not initiated with + ///< the hipStreamCaptureModeRelaxed argument to + ///< hipStreamBeginCapture was passed to + ///< hipStreamEndCapture in a different thread. + hipErrorUnknown = 999, //< Unknown error. + // HSA Runtime Error Codes start here. + hipErrorRuntimeMemory = 1052, ///< HSA runtime memory call returned error. Typically not seen + ///< in production systems. + hipErrorRuntimeOther = 1053, ///< HSA runtime call other than memory returned error. Typically + ///< not seen in production systems. + hipErrorTbd ///< Marker that more error codes are needed. +} hipError_t; + +#undef __HIP_NODISCARD + +typedef enum hipDeviceAttribute_t { + hipDeviceAttributeCudaCompatibleBegin = 0, + + hipDeviceAttributeEccEnabled = hipDeviceAttributeCudaCompatibleBegin, ///< Whether ECC support is enabled. + hipDeviceAttributeAccessPolicyMaxWindowSize, ///< Cuda only. The maximum size of the window policy in bytes. + hipDeviceAttributeAsyncEngineCount, ///< Cuda only. Asynchronous engines number. + hipDeviceAttributeCanMapHostMemory, ///< Whether host memory can be mapped into device address space + hipDeviceAttributeCanUseHostPointerForRegisteredMem,///< Cuda only. Device can access host registered memory + ///< at the same virtual address as the CPU + hipDeviceAttributeClockRate, ///< Peak clock frequency in kilohertz. + hipDeviceAttributeComputeMode, ///< Compute mode that device is currently in. + hipDeviceAttributeComputePreemptionSupported, ///< Cuda only. Device supports Compute Preemption. + hipDeviceAttributeConcurrentKernels, ///< Device can possibly execute multiple kernels concurrently. + hipDeviceAttributeConcurrentManagedAccess, ///< Device can coherently access managed memory concurrently with the CPU + hipDeviceAttributeCooperativeLaunch, ///< Support cooperative launch + hipDeviceAttributeCooperativeMultiDeviceLaunch, ///< Support cooperative launch on multiple devices + hipDeviceAttributeDeviceOverlap, ///< Cuda only. Device can concurrently copy memory and execute a kernel. + ///< Deprecated. Use instead asyncEngineCount. + hipDeviceAttributeDirectManagedMemAccessFromHost, ///< Host can directly access managed memory on + ///< the device without migration + hipDeviceAttributeGlobalL1CacheSupported, ///< Cuda only. Device supports caching globals in L1 + hipDeviceAttributeHostNativeAtomicSupported, ///< Cuda only. Link between the device and the host supports native atomic operations + hipDeviceAttributeIntegrated, ///< Device is integrated GPU + hipDeviceAttributeIsMultiGpuBoard, ///< Multiple GPU devices. + hipDeviceAttributeKernelExecTimeout, ///< Run time limit for kernels executed on the device + hipDeviceAttributeL2CacheSize, ///< Size of L2 cache in bytes. 0 if the device doesn't have L2 cache. + hipDeviceAttributeLocalL1CacheSupported, ///< caching locals in L1 is supported + hipDeviceAttributeLuid, ///< Cuda only. 8-byte locally unique identifier in 8 bytes. Undefined on TCC and non-Windows platforms + hipDeviceAttributeLuidDeviceNodeMask, ///< Cuda only. Luid device node mask. Undefined on TCC and non-Windows platforms + hipDeviceAttributeComputeCapabilityMajor, ///< Major compute capability version number. + hipDeviceAttributeManagedMemory, ///< Device supports allocating managed memory on this system + hipDeviceAttributeMaxBlocksPerMultiProcessor, ///< Cuda only. Max block size per multiprocessor + hipDeviceAttributeMaxBlockDimX, ///< Max block size in width. + hipDeviceAttributeMaxBlockDimY, ///< Max block size in height. + hipDeviceAttributeMaxBlockDimZ, ///< Max block size in depth. + hipDeviceAttributeMaxGridDimX, ///< Max grid size in width. + hipDeviceAttributeMaxGridDimY, ///< Max grid size in height. + hipDeviceAttributeMaxGridDimZ, ///< Max grid size in depth. + hipDeviceAttributeMaxSurface1D, ///< Maximum size of 1D surface. + hipDeviceAttributeMaxSurface1DLayered, ///< Cuda only. Maximum dimensions of 1D layered surface. + hipDeviceAttributeMaxSurface2D, ///< Maximum dimension (width, height) of 2D surface. + hipDeviceAttributeMaxSurface2DLayered, ///< Cuda only. Maximum dimensions of 2D layered surface. + hipDeviceAttributeMaxSurface3D, ///< Maximum dimension (width, height, depth) of 3D surface. + hipDeviceAttributeMaxSurfaceCubemap, ///< Cuda only. Maximum dimensions of Cubemap surface. + hipDeviceAttributeMaxSurfaceCubemapLayered, ///< Cuda only. Maximum dimension of Cubemap layered surface. + hipDeviceAttributeMaxTexture1DWidth, ///< Maximum size of 1D texture. + hipDeviceAttributeMaxTexture1DLayered, ///< Cuda only. Maximum dimensions of 1D layered texture. + hipDeviceAttributeMaxTexture1DLinear, ///< Maximum number of elements allocatable in a 1D linear texture. + ///< Use cudaDeviceGetTexture1DLinearMaxWidth() instead on Cuda. + hipDeviceAttributeMaxTexture1DMipmap, ///< Cuda only. Maximum size of 1D mipmapped texture. + hipDeviceAttributeMaxTexture2DWidth, ///< Maximum dimension width of 2D texture. + hipDeviceAttributeMaxTexture2DHeight, ///< Maximum dimension hight of 2D texture. + hipDeviceAttributeMaxTexture2DGather, ///< Cuda only. Maximum dimensions of 2D texture if gather operations performed. + hipDeviceAttributeMaxTexture2DLayered, ///< Cuda only. Maximum dimensions of 2D layered texture. + hipDeviceAttributeMaxTexture2DLinear, ///< Cuda only. Maximum dimensions (width, height, pitch) of 2D textures bound to pitched memory. + hipDeviceAttributeMaxTexture2DMipmap, ///< Cuda only. Maximum dimensions of 2D mipmapped texture. + hipDeviceAttributeMaxTexture3DWidth, ///< Maximum dimension width of 3D texture. + hipDeviceAttributeMaxTexture3DHeight, ///< Maximum dimension height of 3D texture. + hipDeviceAttributeMaxTexture3DDepth, ///< Maximum dimension depth of 3D texture. + hipDeviceAttributeMaxTexture3DAlt, ///< Cuda only. Maximum dimensions of alternate 3D texture. + hipDeviceAttributeMaxTextureCubemap, ///< Cuda only. Maximum dimensions of Cubemap texture + hipDeviceAttributeMaxTextureCubemapLayered, ///< Cuda only. Maximum dimensions of Cubemap layered texture. + hipDeviceAttributeMaxThreadsDim, ///< Maximum dimension of a block + hipDeviceAttributeMaxThreadsPerBlock, ///< Maximum number of threads per block. + hipDeviceAttributeMaxThreadsPerMultiProcessor, ///< Maximum resident threads per multiprocessor. + hipDeviceAttributeMaxPitch, ///< Maximum pitch in bytes allowed by memory copies + hipDeviceAttributeMemoryBusWidth, ///< Global memory bus width in bits. + hipDeviceAttributeMemoryClockRate, ///< Peak memory clock frequency in kilohertz. + hipDeviceAttributeComputeCapabilityMinor, ///< Minor compute capability version number. + hipDeviceAttributeMultiGpuBoardGroupID, ///< Cuda only. Unique ID of device group on the same multi-GPU board + hipDeviceAttributeMultiprocessorCount, ///< Number of multiprocessors on the device. + hipDeviceAttributeName, ///< Device name. + hipDeviceAttributePageableMemoryAccess, ///< Device supports coherently accessing pageable memory + ///< without calling hipHostRegister on it + hipDeviceAttributePageableMemoryAccessUsesHostPageTables, ///< Device accesses pageable memory via the host's page tables + hipDeviceAttributePciBusId, ///< PCI Bus ID. + hipDeviceAttributePciDeviceId, ///< PCI Device ID. + hipDeviceAttributePciDomainID, ///< PCI Domain ID. + hipDeviceAttributePersistingL2CacheMaxSize, ///< Cuda11 only. Maximum l2 persisting lines capacity in bytes + hipDeviceAttributeMaxRegistersPerBlock, ///< 32-bit registers available to a thread block. This number is shared + ///< by all thread blocks simultaneously resident on a multiprocessor. + hipDeviceAttributeMaxRegistersPerMultiprocessor, ///< 32-bit registers available per block. + hipDeviceAttributeReservedSharedMemPerBlock, ///< Cuda11 only. Shared memory reserved by CUDA driver per block. + hipDeviceAttributeMaxSharedMemoryPerBlock, ///< Maximum shared memory available per block in bytes. + hipDeviceAttributeSharedMemPerBlockOptin, ///< Cuda only. Maximum shared memory per block usable by special opt in. + hipDeviceAttributeSharedMemPerMultiprocessor, ///< Cuda only. Shared memory available per multiprocessor. + hipDeviceAttributeSingleToDoublePrecisionPerfRatio, ///< Cuda only. Performance ratio of single precision to double precision. + hipDeviceAttributeStreamPrioritiesSupported, ///< Cuda only. Whether to support stream priorities. + hipDeviceAttributeSurfaceAlignment, ///< Cuda only. Alignment requirement for surfaces + hipDeviceAttributeTccDriver, ///< Cuda only. Whether device is a Tesla device using TCC driver + hipDeviceAttributeTextureAlignment, ///< Alignment requirement for textures + hipDeviceAttributeTexturePitchAlignment, ///< Pitch alignment requirement for 2D texture references bound to pitched memory; + hipDeviceAttributeTotalConstantMemory, ///< Constant memory size in bytes. + hipDeviceAttributeTotalGlobalMem, ///< Global memory available on devicice. + hipDeviceAttributeUnifiedAddressing, ///< Cuda only. An unified address space shared with the host. + hipDeviceAttributeUuid, ///< Cuda only. Unique ID in 16 byte. + hipDeviceAttributeWarpSize, ///< Warp size in threads. + + hipDeviceAttributeCudaCompatibleEnd = 9999, + hipDeviceAttributeAmdSpecificBegin = 10000, + + hipDeviceAttributeClockInstructionRate = hipDeviceAttributeAmdSpecificBegin, ///< Frequency in khz of the timer used by the device-side "clock*" + hipDeviceAttributeArch, ///< Device architecture + hipDeviceAttributeMaxSharedMemoryPerMultiprocessor, ///< Maximum Shared Memory PerMultiprocessor. + hipDeviceAttributeGcnArch, ///< Device gcn architecture + hipDeviceAttributeGcnArchName, ///< Device gcnArch name in 256 bytes + hipDeviceAttributeHdpMemFlushCntl, ///< Address of the HDP_MEM_COHERENCY_FLUSH_CNTL register + hipDeviceAttributeHdpRegFlushCntl, ///< Address of the HDP_REG_COHERENCY_FLUSH_CNTL register + hipDeviceAttributeCooperativeMultiDeviceUnmatchedFunc, ///< Supports cooperative launch on multiple + ///< devices with unmatched functions + hipDeviceAttributeCooperativeMultiDeviceUnmatchedGridDim, ///< Supports cooperative launch on multiple + ///< devices with unmatched grid dimensions + hipDeviceAttributeCooperativeMultiDeviceUnmatchedBlockDim, ///< Supports cooperative launch on multiple + ///< devices with unmatched block dimensions + hipDeviceAttributeCooperativeMultiDeviceUnmatchedSharedMem, ///< Supports cooperative launch on multiple + ///< devices with unmatched shared memories + hipDeviceAttributeIsLargeBar, ///< Whether it is LargeBar + hipDeviceAttributeAsicRevision, ///< Revision of the GPU in this device + hipDeviceAttributeCanUseStreamWaitValue, ///< '1' if Device supports hipStreamWaitValue32() and + ///< hipStreamWaitValue64() , '0' otherwise. + + hipDeviceAttributeAmdSpecificEnd = 19999, + hipDeviceAttributeVendorSpecificBegin = 20000, + // Extended attributes for vendors +} hipDeviceAttribute_t; + +//! Flags that can be used with hipStreamCreateWithFlags +#define hipStreamDefault \ + 0x00 ///< Default stream creation flags. These are used with hipStreamCreate(). +#define hipStreamNonBlocking 0x01 ///< Stream does not implicitly synchronize with null stream + + +//! Flags that can be used with hipEventCreateWithFlags: +#define hipEventDefault 0x0 ///< Default flags +#define hipEventBlockingSync \ + 0x1 ///< Waiting will yield CPU. Power-friendly and usage-friendly but may increase latency. +#define hipEventDisableTiming \ + 0x2 ///< Disable event's capability to record timing information. May improve performance. +#define hipEventInterprocess 0x4 ///< Event can support IPC. @warning - not supported in HIP. +#define hipEventReleaseToDevice \ + 0x40000000 /// < Use a device-scope release when recording this event. This flag is useful to + /// obtain more precise timings of commands between events. The flag is a no-op on + /// CUDA platforms. +#define hipEventReleaseToSystem \ + 0x80000000 /// < Use a system-scope release when recording this event. This flag is + /// useful to make non-coherent host memory visible to the host. The flag is a + /// no-op on CUDA platforms. + + +#define hipDeviceScheduleAuto 0x0 ///< Automatically select between Spin and Yield +#define hipDeviceScheduleSpin \ + 0x1 ///< Dedicate a CPU core to spin-wait. Provides lowest latency, but burns a CPU core and + ///< may consume more power. +#define hipDeviceScheduleYield \ + 0x2 ///< Yield the CPU to the operating system when waiting. May increase latency, but lowers + ///< power and is friendlier to other threads in the system. +#define hipDeviceScheduleBlockingSync 0x4 +#define hipDeviceScheduleMask 0x7 +#define hipDeviceMapHost 0x8 +#define hipDeviceLmemResizeToMax 0x16 + +typedef enum hipJitOption { + hipJitOptionMaxRegisters = 0, + hipJitOptionThreadsPerBlock, + hipJitOptionWallTime, + hipJitOptionInfoLogBuffer, + hipJitOptionInfoLogBufferSizeBytes, + hipJitOptionErrorLogBuffer, + hipJitOptionErrorLogBufferSizeBytes, + hipJitOptionOptimizationLevel, + hipJitOptionTargetFromContext, + hipJitOptionTarget, + hipJitOptionFallbackStrategy, + hipJitOptionGenerateDebugInfo, + hipJitOptionLogVerbose, + hipJitOptionGenerateLineInfo, + hipJitOptionCacheMode, + hipJitOptionSm3xOpt, + hipJitOptionFastCompile, + hipJitOptionNumOptions +} hipJitOption; + +// stop: hip_runtime_api.h -typedef int HIPdevice; /**< HIP device */ -typedef struct HIPctx_st *HIPcontext; /**< HIP context */ -typedef struct HIPevent_st *HIPevent; /**< HIP event */ -typedef struct HIPfunc_st *HIPfunction; /**< HIP function */ -typedef struct HIPmod_st *HIPmodule; /**< HIP module */ -typedef struct HIPstream_st *HIPstream; /**< HIP stream */ -typedef struct HIPlinkState_st *HIPlinkState; - - -typedef enum hipError_enum { - /** - * The API call returned with no errors. In the case of query calls, this - * also means that the operation being queried is complete (see - * ::hipEventQuery() and ::hipStreamQuery()). - */ - HIP_SUCCESS = 0, - - /** - * This indicates that one or more of the parameters passed to the API call - * is not within an acceptable range of values. - */ - HIP_ERROR_INVALID_VALUE = 1, - - /** - * The API call failed because it was unable to allocate enough memory to - * perform the requested operation. - */ - HIP_ERROR_OUT_OF_MEMORY = 2, - - /** - * This indicates that the HIP driver has not been initialized with - * ::hipInit() or that initialization has failed. - */ - HIP_ERROR_NOT_INITIALIZED = 3, - - /** - * This indicates that the HIP driver is in the process of shutting down. - */ - HIP_ERROR_DEINITIALIZED = 4, - - /** - * This indicates profiler is not initialized for this run. This can - * happen when the application is running with external profiling tools - * like visual profiler. - */ - HIP_ERROR_PROFILER_DISABLED = 5, - - /** - * \deprecated - * This error return is deprecated as of HIP 5.0. It is no longer an error - * to attempt to enable/disable the profiling via ::hipProfilerStart or - * ::hipProfilerStop without initialization. - */ - HIP_ERROR_PROFILER_NOT_INITIALIZED = 6, - - /** - * \deprecated - * This error return is deprecated as of HIP 5.0. It is no longer an error - * to call hipProfilerStart() when profiling is already enabled. - */ - HIP_ERROR_PROFILER_ALREADY_STARTED = 7, - - /** - * \deprecated - * This error return is deprecated as of HIP 5.0. It is no longer an error - * to call hipProfilerStop() when profiling is already disabled. - */ - HIP_ERROR_PROFILER_ALREADY_STOPPED = 8, - - /** - * This indicates that no HIP-capable devices were detected by the installed - * HIP driver. - */ - HIP_ERROR_NO_DEVICE = 100, - - /** - * This indicates that the device ordinal supplied by the user does not - * correspond to a valid HIP device. - */ - HIP_ERROR_INVALID_DEVICE = 101, - - - /** - * This indicates that the device kernel image is invalid. This can also - * indicate an invalid HIP module. - */ - HIP_ERROR_INVALID_IMAGE = 200, - - /** - * This most frequently indicates that there is no context bound to the - * hiprrent thread. This can also be returned if the context passed to an - * API call is not a valid handle (such as a context that has had - * ::hipCtxDestroy() invoked on it). This can also be returned if a user - * mixes different API versions (i.e. 3010 context with 3020 API calls). - * See ::hipCtxGetApiVersion() for more details. - */ - HIP_ERROR_INVALID_CONTEXT = 201, - - /** - * This indicated that the context being supplied as a parameter to the - * API call was already the active context. - * \deprecated - * This error return is deprecated as of HIP 3.2. It is no longer an - * error to attempt to push the active context via ::hipCtxPushCurrent(). - */ - HIP_ERROR_CONTEXT_ALREADY_CURRENT = 202, - - /** - * This indicates that a map or register operation has failed. - */ - HIP_ERROR_MAP_FAILED = 205, - - /** - * This indicates that an unmap or unregister operation has failed. - */ - HIP_ERROR_UNMAP_FAILED = 206, - - /** - * This indicates that the specified array is currently mapped and thus - * cannot be destroyed. - */ - HIP_ERROR_ARRAY_IS_MAPPED = 207, - - /** - * This indicates that the resource is already mapped. - */ - HIP_ERROR_ALREADY_MAPPED = 208, - - /** - * This indicates that there is no kernel image available that is suitable - * for the device. This can occur when a user specifies code generation - * options for a particular HIP source file that do not include the - * corresponding device configuration. - */ - HIP_ERROR_NO_BINARY_FOR_GPU = 209, - - /** - * This indicates that a resource has already been acquired. - */ - HIP_ERROR_ALREADY_ACQUIRED = 210, - - /** - * This indicates that a resource is not mapped. - */ - HIP_ERROR_NOT_MAPPED = 211, - - /** - * This indicates that a mapped resource is not available for access as an - * array. - */ - HIP_ERROR_NOT_MAPPED_AS_ARRAY = 212, - - /** - * This indicates that a mapped resource is not available for access as a - * pointer. - */ - HIP_ERROR_NOT_MAPPED_AS_POINTER = 213, - - /** - * This indicates that an uncorrectable ECC error was detected during - * execution. - */ - HIP_ERROR_ECC_UNCORRECTABLE = 214, - - /** - * This indicates that the ::HIPlimit passed to the API call is not - * supported by the active device. - */ - HIP_ERROR_UNSUPPORTED_LIMIT = 215, - - /** - * This indicates that the ::HIPcontext passed to the API call can - * only be bound to a single CPU thread at a time but is already - * bound to a CPU thread. - */ - HIP_ERROR_CONTEXT_ALREADY_IN_USE = 216, - - /** - * This indicates that peer access is not supported across the given - * devices. - */ - HIP_ERROR_PEER_ACCESS_UNSUPPORTED = 217, - - /** - * This indicates that a PTX JIT compilation failed. - */ - HIP_ERROR_INVALID_PTX = 218, - - /** - * This indicates an error with OpenGL or DirectX context. - */ - HIP_ERROR_INVALID_GRAPHICS_CONTEXT = 219, - - /** - * This indicates that an uncorrectable NVLink error was detected during the - * execution. - */ - HIP_ERROR_NVLINK_UNCORRECTABLE = 220, - - /** - * This indicates that the PTX JIT compiler library was not found. - */ - HIP_ERROR_JIT_COMPILER_NOT_FOUND = 221, - - /** - * This indicates that the device kernel source is invalid. - */ - HIP_ERROR_INVALID_SOURCE = 300, - - /** - * This indicates that the file specified was not found. - */ - HIP_ERROR_FILE_NOT_FOUND = 301, - - /** - * This indicates that a link to a shared object failed to resolve. - */ - HIP_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND = 302, - - /** - * This indicates that initialization of a shared object failed. - */ - HIP_ERROR_SHARED_OBJECT_INIT_FAILED = 303, - - /** - * This indicates that an OS call failed. - */ - HIP_ERROR_OPERATING_SYSTEM = 304, - - /** - * This indicates that a resource handle passed to the API call was not - * valid. Resource handles are opaque types like ::HIPstream and ::HIPevent. - */ - HIP_ERROR_INVALID_HANDLE = 400, - - /** - * This indicates that a resource required by the API call is not in a - * valid state to perform the requested operation. - */ - HIP_ERROR_ILLEGAL_STATE = 401, - - /** - * This indicates that a named symbol was not found. Examples of symbols - * are global/constant variable names, texture names, and surface names. - */ - HIP_ERROR_NOT_FOUND = 500, - - /** - * This indicates that asynchronous operations issued previously have not - * completed yet. This result is not actually an error, but must be indicated - * differently than ::HIP_SUCCESS (which indicates completion). Calls that - * may return this value include ::hipEventQuery() and ::hipStreamQuery(). - */ - HIP_ERROR_NOT_READY = 600, - - /** - * While executing a kernel, the device encountered a - * load or store instruction on an invalid memory address. - * This leaves the process in an inconsistent state and any further HIP work - * will return the same error. To continue using HIP, the process must be terminated - * and relaunched. - */ - HIP_ERROR_ILLEGAL_ADDRESS = 700, - - /** - * This indicates that a launch did not occur because it did not have - * appropriate resources. This error usually indicates that the user has - * attempted to pass too many arguments to the device kernel, or the - * kernel launch specifies too many threads for the kernel's register - * count. Passing arguments of the wrong size (i.e. a 64-bit pointer - * when a 32-bit int is expected) is equivalent to passing too many - * arguments and can also result in this error. - */ - HIP_ERROR_LAUNCH_OUT_OF_RESOURCES = 701, - - /** - * This indicates that the device kernel took too long to execute. This can - * only occur if timeouts are enabled - see the device attribute - * ::HIP_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT for more information. - * This leaves the process in an inconsistent state and any further HIP work - * will return the same error. To continue using HIP, the process must be terminated - * and relaunched. - */ - HIP_ERROR_LAUNCH_TIMEOUT = 702, - - /** - * This error indicates a kernel launch that uses an incompatible texturing - * mode. - */ - HIP_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING = 703, - - /** - * This error indicates that a call to ::hipCtxEnablePeerAccess() is - * trying to re-enable peer access to a context which has already - * had peer access to it enabled. - */ - HIP_ERROR_PEER_ACCESS_ALREADY_ENABLED = 704, - - /** - * This error indicates that ::hipCtxDisablePeerAccess() is - * trying to disable peer access which has not been enabled yet - * via ::hipCtxEnablePeerAccess(). - */ - HIP_ERROR_PEER_ACCESS_NOT_ENABLED = 705, - - /** - * This error indicates that the primary context for the specified device - * has already been initialized. - */ - HIP_ERROR_PRIMARY_CONTEXT_ACTIVE = 708, - - /** - * This error indicates that the context hiprrent to the calling thread - * has been destroyed using ::hipCtxDestroy, or is a primary context which - * has not yet been initialized. - */ - HIP_ERROR_CONTEXT_IS_DESTROYED = 709, - - /** - * A device-side assert triggered during kernel execution. The context - * cannot be used anymore, and must be destroyed. All existing device - * memory allocations from this context are invalid and must be - * reconstructed if the program is to continue using HIP. - */ - HIP_ERROR_ASSERT = 710, - - /** - * This error indicates that the hardware resources required to enable - * peer access have been exhausted for one or more of the devices - * passed to ::hipCtxEnablePeerAccess(). - */ - HIP_ERROR_TOO_MANY_PEERS = 711, - - /** - * This error indicates that the memory range passed to ::hipMemHostRegister() - * has already been registered. - */ - HIP_ERROR_HOST_MEMORY_ALREADY_REGISTERED = 712, - - /** - * This error indicates that the pointer passed to ::hipMemHostUnregister() - * does not correspond to any currently registered memory region. - */ - HIP_ERROR_HOST_MEMORY_NOT_REGISTERED = 713, - - /** - * While executing a kernel, the device encountered a stack error. - * This can be due to stack corruption or exceeding the stack size limit. - * This leaves the process in an inconsistent state and any further HIP work - * will return the same error. To continue using HIP, the process must be terminated - * and relaunched. - */ - HIP_ERROR_HARDWARE_STACK_ERROR = 714, - - /** - * While executing a kernel, the device encountered an illegal instruction. - * This leaves the process in an inconsistent state and any further HIP work - * will return the same error. To continue using HIP, the process must be terminated - * and relaunched. - */ - HIP_ERROR_ILLEGAL_INSTRUCTION = 715, - - /** - * While executing a kernel, the device encountered a load or store instruction - * on a memory address which is not aligned. - * This leaves the process in an inconsistent state and any further HIP work - * will return the same error. To continue using HIP, the process must be terminated - * and relaunched. - */ - HIP_ERROR_MISALIGNED_ADDRESS = 716, - - /** - * While executing a kernel, the device encountered an instruction - * which can only operate on memory locations in certain address spaces - * (global, shared, or local), but was supplied a memory address not - * belonging to an allowed address space. - * This leaves the process in an inconsistent state and any further HIP work - * will return the same error. To continue using HIP, the process must be terminated - * and relaunched. - */ - HIP_ERROR_INVALID_ADDRESS_SPACE = 717, - - /** - * While executing a kernel, the device program counter wrapped its address space. - * This leaves the process in an inconsistent state and any further HIP work - * will return the same error. To continue using HIP, the process must be terminated - * and relaunched. - */ - HIP_ERROR_INVALID_PC = 718, - - /** - * An exception occurred on the device while executing a kernel. Common - * causes include dereferencing an invalid device pointer and accessing - * out of bounds shared memory. Less common cases can be system specific - more - * information about these cases can be found in the system specific user guide. - * This leaves the process in an inconsistent state and any further HIP work - * will return the same error. To continue using HIP, the process must be terminated - * and relaunched. - */ - HIP_ERROR_LAUNCH_FAILED = 719, - - /** - * This error indicates that the number of blocks launched per grid for a kernel that was - * launched via either ::hipLaunchCooperativeKernel or ::hipLaunchCooperativeKernelMultiDevice - * exceeds the maximum number of blocks as allowed by ::hipOccupancyMaxActiveBlocksPerMultiprocessor - * or ::hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags times the number of multiprocessors - * as specified by the device attribute ::HIP_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT. - */ - HIP_ERROR_COOPERATIVE_LAUNCH_TOO_LARGE = 720, - - /** - * This error indicates that the attempted operation is not permitted. - */ - HIP_ERROR_NOT_PERMITTED = 800, - - /** - * This error indicates that the attempted operation is not supported - * on the current system or device. - */ - HIP_ERROR_NOT_SUPPORTED = 801, - - /** - * This error indicates that the system is not yet ready to start any HIP - * work. To continue using HIP, verify the system configuration is in a - * valid state and all required driver daemons are actively running. - * More information about this error can be found in the system specific - * user guide. - */ - HIP_ERROR_SYSTEM_NOT_READY = 802, - - /** - * This error indicates that there is a mismatch between the versions of - * the display driver and the HIP driver. Refer to the compatibility documentation - * for supported versions. - */ - HIP_ERROR_SYSTEM_DRIVER_MISMATCH = 803, - - /** - * This error indicates that the system was upgraded to run with forward compatibility - * but the visible hardware detected by HIP does not support this configuration. - * Refer to the compatibility documentation for the supported hardware matrix or ensure - * that only supported hardware is visible during initialization via the HIP_VISIBLE_DEVICES - * environment variable. - */ - HIP_ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE = 804, - - /** - * This error indicates that the operation is not permitted when - * the stream is capturing. - */ - HIP_ERROR_STREAM_CAPTURE_UNSUPPORTED = 900, - - /** - * This error indicates that the current capture sequence on the stream - * has been invalidated due to a previous error. - */ - HIP_ERROR_STREAM_CAPTURE_INVALIDATED = 901, - - /** - * This error indicates that the operation would have resulted in a merge - * of two independent capture sequences. - */ - HIP_ERROR_STREAM_CAPTURE_MERGE = 902, - - /** - * This error indicates that the capture was not initiated in this stream. - */ - HIP_ERROR_STREAM_CAPTURE_UNMATCHED = 903, - - /** - * This error indicates that the capture sequence contains a fork that was - * not joined to the primary stream. - */ - HIP_ERROR_STREAM_CAPTURE_UNJOINED = 904, - - /** - * This error indicates that a dependency would have been created which - * crosses the capture sequence boundary. Only implicit in-stream ordering - * dependencies are allowed to cross the boundary. - */ - HIP_ERROR_STREAM_CAPTURE_ISOLATION = 905, - - /** - * This error indicates a disallowed implicit dependency on a current capture - * sequence from HIPStreamLegacy. - */ - HIP_ERROR_STREAM_CAPTURE_IMPLICIT = 906, - - /** - * This error indicates that the operation is not permitted on an event which - * was last recorded in a capturing stream. - */ - HIP_ERROR_CAPTURED_EVENT = 907, - - /** - * A stream capture sequence not initiated with the ::HIP_STREAM_CAPTURE_MODE_RELAXED - * argument to ::HIPStreamBeginCapture was passed to ::hipStreamEndCapture in a - * different thread. - */ - HIP_ERROR_STREAM_CAPTURE_WRONG_THREAD = 908, - - /** - * This indicates that an unknown internal error has occurred. - */ - HIP_ERROR_UNKNOWN = 999 -} HIPresult; - -/** - * Online compiler and linker options - */ -typedef enum HIPjit_option_enum -{ - /** - * Max number of registers that a thread may use.\n - * Option type: unsigned int\n - * Applies to: compiler only - */ - HIP_JIT_MAX_REGISTERS = 0, - - /** - * IN: Specifies minimum number of threads per block to target compilation - * for\n - * OUT: Returns the number of threads the compiler actually targeted. - * This restricts the resource utilization fo the compiler (e.g. max - * registers) such that a block with the given number of threads should be - * able to launch based on register limitations. Note, this option does not - * currently take into account any other resource limitations, such as - * shared memory utilization.\n - * Cannot be combined with ::HIP_JIT_TARGET.\n - * Option type: unsigned int\n - * Applies to: compiler only - */ - HIP_JIT_THREADS_PER_BLOCK, - - /** - * Overwrites the option value with the total wall clock time, in - * milliseconds, spent in the compiler and linker\n - * Option type: float\n - * Applies to: compiler and linker - */ - HIP_JIT_WALL_TIME, - - /** - * Pointer to a buffer in which to print any log messages - * that are informational in nature (the buffer size is specified via - * option ::HIP_JIT_INFO_LOG_BUFFER_SIZE_BYTES)\n - * Option type: char *\n - * Applies to: compiler and linker - */ - HIP_JIT_INFO_LOG_BUFFER, - - /** - * IN: Log buffer size in bytes. Log messages will be capped at this size - * (including null terminator)\n - * OUT: Amount of log buffer filled with messages\n - * Option type: unsigned int\n - * Applies to: compiler and linker - */ - HIP_JIT_INFO_LOG_BUFFER_SIZE_BYTES, - - /** - * Pointer to a buffer in which to print any log messages that - * reflect errors (the buffer size is specified via option - * ::HIP_JIT_ERROR_LOG_BUFFER_SIZE_BYTES)\n - * Option type: char *\n - * Applies to: compiler and linker - */ - HIP_JIT_ERROR_LOG_BUFFER, - - /** - * IN: Log buffer size in bytes. Log messages will be capped at this size - * (including null terminator)\n - * OUT: Amount of log buffer filled with messages\n - * Option type: unsigned int\n - * Applies to: compiler and linker - */ - HIP_JIT_ERROR_LOG_BUFFER_SIZE_BYTES, - - /** - * Level of optimizations to apply to generated code (0 - 4), with 4 - * being the default and highest level of optimizations.\n - * Option type: unsigned int\n - * Applies to: compiler only - */ - HIP_JIT_OPTIMIZATION_LEVEL, - - /** - * No option value required. Determines the target based on the current - * attached context (default)\n - * Option type: No option value needed\n - * Applies to: compiler and linker - */ - HIP_JIT_TARGET_FROM_HIPCONTEXT, - - /** - * Target is chosen based on supplied ::HIPjit_target. Cannot be - * combined with ::HIP_JIT_THREADS_PER_BLOCK.\n - * Option type: unsigned int for enumerated type ::HIPjit_target\n - * Applies to: compiler and linker - */ - HIP_JIT_TARGET, - - /** - * Specifies choice of fallback strategy if matching HIPbin is not found. - * Choice is based on supplied ::HIPjit_fallback. This option cannot be - * used with HIPLink* APIs as the linker requires exact matches.\n - * Option type: unsigned int for enumerated type ::HIPjit_fallback\n - * Applies to: compiler only - */ - HIP_JIT_FALLBACK_STRATEGY, - - /** - * Specifies whether to create debug information in output (-g) - * (0: false, default)\n - * Option type: int\n - * Applies to: compiler and linker - */ - HIP_JIT_GENERATE_DEBUG_INFO, - - /** - * Generate verbose log messages (0: false, default)\n - * Option type: int\n - * Applies to: compiler and linker - */ - HIP_JIT_LOG_VERBOSE, - - /** - * Generate line number information (-lineinfo) (0: false, default)\n - * Option type: int\n - * Applies to: compiler only - */ - HIP_JIT_GENERATE_LINE_INFO, - - /** - * Specifies whether to enable caching explicitly (-dlcm) \n - * Choice is based on supplied ::HIPjit_cacheMode_enum.\n - * Option type: unsigned int for enumerated type ::HIPjit_cacheMode_enum\n - * Applies to: compiler only - */ - HIP_JIT_CACHE_MODE, - - /** - * The below jit options are used for internal purposes only, in this version of HIP - */ - HIP_JIT_NEW_SM3X_OPT, - HIP_JIT_FAST_COMPILE, - - /** - * Array of device symbol names that will be relocated to the corresponing - * host addresses stored in ::HIP_JIT_GLOBAL_SYMBOL_ADDRESSES.\n - * Must contain ::HIP_JIT_GLOBAL_SYMBOL_COUNT entries.\n - * When loding a device module, driver will relocate all encountered - * unresolved symbols to the host addresses.\n - * It is only allowed to register symbols that correspond to unresolved - * global variables.\n - * It is illegal to register the same device symbol at multiple addresses.\n - * Option type: const char **\n - * Applies to: dynamic linker only - */ - HIP_JIT_GLOBAL_SYMBOL_NAMES, - - /** - * Array of host addresses that will be used to relocate corresponding - * device symbols stored in ::HIP_JIT_GLOBAL_SYMBOL_NAMES.\n - * Must contain ::HIP_JIT_GLOBAL_SYMBOL_COUNT entries.\n - * Option type: void **\n - * Applies to: dynamic linker only - */ - HIP_JIT_GLOBAL_SYMBOL_ADDRESSES, - - /** - * Number of entries in ::HIP_JIT_GLOBAL_SYMBOL_NAMES and - * ::HIP_JIT_GLOBAL_SYMBOL_ADDRESSES arrays.\n - * Option type: unsigned int\n - * Applies to: dynamic linker only - */ - HIP_JIT_GLOBAL_SYMBOL_COUNT, - - HIP_JIT_NUM_OPTIONS - -} HIPjit_option; - - -/** - * Device properties - */ -typedef enum HIPdevice_attribute_enum { - - HIP_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK = 0, /**< Maximum number of threads per block */ - HIP_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X = 1, /**< Maximum block dimension X */ - HIP_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y = 2, /**< Maximum block dimension Y */ - HIP_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z = 3, /**< Maximum block dimension Z */ - HIP_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X = 4, /**< Maximum grid dimension X */ - HIP_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y = 5, /**< Maximum grid dimension Y */ - HIP_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z = 6, /**< Maximum grid dimension Z */ - HIP_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK = 7, /**< Maximum shared memory available per block in bytes */ - HIP_DEVICE_ATTRIBUTE_SHARED_MEMORY_PER_BLOCK = 7, /**< Deprecated, use HIP_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK */ - HIP_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK_OPTIN = 7, /**< Maximum optin shared memory per block */ - HIP_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY = 8, /**< Memory available on device for __constant__ variables in a HIP C kernel in bytes */ - HIP_DEVICE_ATTRIBUTE_WARP_SIZE = 9, /**< Warp size in threads */ - HIP_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK = 10, /**< Maximum number of 32-bit registers available per block */ - HIP_DEVICE_ATTRIBUTE_REGISTERS_PER_BLOCK = 10, /**< Deprecated, use HIP_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK */ - HIP_DEVICE_ATTRIBUTE_CLOCK_RATE = 11, /**< Typical clock frequency in kilohertz */ - HIP_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE = 12, /**< Peak memory clock frequency in kilohertz */ - HIP_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH = 13, /**< Global memory bus width in bits */ - HIP_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT = 14, /**< Number of multiprocessors on device */ - HIP_DEVICE_ATTRIBUTE_COMPUTE_MODE = 15, /**< Compute mode (See ::HIPcomputemode for details) */ - HIP_DEVICE_ATTRIBUTE_L2_CACHE_SIZE = 16, /**< Size of L2 cache in bytes */ - HIP_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR = 17, /**< Maximum resident threads per multiprocessor */ - HIP_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR = 18, /**< Major compute capability version number */ - HIP_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR = 19, /**< Minor compute capability version number */ - HIP_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS = 20, /**< Device can possibly execute multiple kernels concurrently */ - HIP_DEVICE_ATTRIBUTE_PCI_BUS_ID = 21, /**< PCI bus ID of the device */ - HIP_DEVICE_ATTRIBUTE_PCI_DEVICE_ID = 22, /**< PCI device ID of the device */ - HIP_DEVICE_ATTRIBUTE_PCI_DOMAIN_ID = 22, /**< PCI domain ID of the device */ - HIP_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_MULTIPROCESSOR = 23, /**< Maximum shared memory available per multiprocessor in bytes */ - HIP_DEVICE_ATTRIBUTE_MULTI_GPU_BOARD = 24, /**< Device is on a multi-GPU board */ - HIP_DEVICE_ATTRIBUTE_INTEGRATED = 25, /**< Device is integrated with host memory */ - HIP_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH = 26, /**< Device supports launching cooperative kernels via ::hipLaunchCooperativeKernel */ - HIP_DEVICE_ATTRIBUTE_COOPERATIVE_MULTI_DEVICE_LAUNCH = 27, /**< Device can participate in cooperative kernels launched via ::hipLaunchCooperativeKernelMultiDevice */ - HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH = 28, /**< Maximum 1D texture width */ - HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH = 29, /**< Maximum 2D texture width */ - HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT = 30, /**< Maximum 2D texture height */ - HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH = 31, /**< Maximum 3D texture width */ - HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT = 32, /**< Maximum 3D texture height */ - HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH = 33, /**< Maximum 3D texture depth */ - - HIP_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT = 37, /**< Alignment requirement for textures */ - HIP_DEVICE_ATTRIBUTE_TEXTURE_PITCH_ALIGNMENT = 38, /**< Pitch alignment requirement for textures */ - HIP_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT = 39, /**< Specifies whether there is a run time limit on kernels */ - HIP_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY = 40, /**< Device can map host memory into HIP address space */ - HIP_DEVICE_ATTRIBUTE_ECC_ENABLED = 41, /**< Device has ECC support enabled */ - - HIP_DEVICE_ATTRIBUTE_MANAGED_MEMORY = 47, /**< Device can allocate managed memory on this system */ - HIP_DEVICE_ATTRIBUTE_DIRECT_MANAGED_MEM_ACCESS_FROM_HOST = 48, /**< The host can directly access managed memory on the device without migration. */ - HIP_DEVICE_ATTRIBUTE_CONCURRENT_MANAGED_ACCESS = 49, /**< Device can coherently access managed memory concurrently with the CPU */ - HIP_DEVICE_ATTRIBUTE_PAGEABLE_MEMORY_ACCESS = 50, /**< Device supports coherently accessing pageable memory without calling HIPHostRegister on it */ - HIP_DEVICE_ATTRIBUTE_PAGEABLE_MEMORY_ACCESS_USES_HOST_PAGE_TABLES = 51, /**< Device accesses pageable memory via the host's page tables. */ - HIP_DEVICE_ATTRIBUTE_CAN_USE_STREAM_WAIT_VALUE_NOR = 52, /**< ::HIP_STREAM_WAIT_VALUE_NOR is supported. */ - - - // HIP_DEVICE_ATTRIBUTE_MAX_PITCH = , /**< Maximum pitch in bytes allowed by memory copies */ - // HIP_DEVICE_ATTRIBUTE_GPU_OVERLAP = , /**< Device can possibly copy memory and execute a kernel concurrently. Deprecated. Use instead HIP_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT. */ - // - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH = , /**< Maximum 2D layered texture width */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_HEIGHT = , /**< Maximum 2D layered texture height */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_LAYERS = , /**< Maximum layers in a 2D layered texture */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_ARRAY_WIDTH = , /**< Deprecated, use HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_ARRAY_HEIGHT = , /**< Deprecated, use HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_HEIGHT */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_ARRAY_NUMSLICES = , /**< Deprecated, use HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_LAYERS */ - // HIP_DEVICE_ATTRIBUTE_SURFACE_ALIGNMENT =, /**< Alignment requirement for surfaces */ - // HIP_DEVICE_ATTRIBUTE_TCC_DRIVER = , /**< Device is using TCC driver model */ - // HIP_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT = , /**< Number of asynchronous engines */ - // HIP_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING = , /**< Device shares a unified address space with the host */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LAYERED_WIDTH = , /**< Maximum 1D layered texture width */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LAYERED_LAYERS = , /**< Maximum layers in a 1D layered texture */ - // HIP_DEVICE_ATTRIBUTE_CAN_TEX2D_GATHER = , /**< Deprecated, do not use. */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_GATHER_WIDTH = , /**< Maximum 2D texture width if HIP_ARRAY3D_TEXTURE_GATHER is set */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_GATHER_HEIGHT = , /**< Maximum 2D texture height if HIP_ARRAY3D_TEXTURE_GATHER is set */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH_ALTERNATE = , /**< Alternate maximum 3D texture width */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT_ALTERNATE = ,/**< Alternate maximum 3D texture height */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH_ALTERNATE = , /**< Alternate maximum 3D texture depth */ - // - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURECUBEMAP_WIDTH = , /**< Maximum cubemap texture width/height */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURECUBEMAP_LAYERED_WIDTH = , /**< Maximum cubemap layered texture width/height */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURECUBEMAP_LAYERED_LAYERS = , /**< Maximum layers in a cubemap layered texture */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_WIDTH = , /**< Maximum 1D surface width */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_WIDTH = , /**< Maximum 2D surface width */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_HEIGHT = , /**< Maximum 2D surface height */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_WIDTH = , /**< Maximum 3D surface width */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_HEIGHT = , /**< Maximum 3D surface height */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE3D_DEPTH = , /**< Maximum 3D surface depth */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_LAYERED_WIDTH = , /**< Maximum 1D layered surface width */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_LAYERED_LAYERS = , /**< Maximum layers in a 1D layered surface */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_WIDTH = , /**< Maximum 2D layered surface width */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_HEIGHT = , /**< Maximum 2D layered surface height */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_LAYERS = , /**< Maximum layers in a 2D layered surface */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACECUBEMAP_WIDTH = , /**< Maximum cubemap surface width */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACECUBEMAP_LAYERED_WIDTH = , /**< Maximum cubemap layered surface width */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_SURFACECUBEMAP_LAYERED_LAYERS = , /**< Maximum layers in a cubemap layered surface */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LINEAR_WIDTH = , /**< Maximum 1D linear texture width */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_WIDTH = , /**< Maximum 2D linear texture width */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_HEIGHT = , /**< Maximum 2D linear texture height */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LINEAR_PITCH = , /**< Maximum 2D linear texture pitch in bytes */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_MIPMAPPED_WIDTH = , /**< Maximum mipmapped 2D texture width */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_MIPMAPPED_HEIGHT = ,/**< Maximum mipmapped 2D texture height */ - // HIP_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_MIPMAPPED_WIDTH = , /**< Maximum mipmapped 1D texture width */ - // HIP_DEVICE_ATTRIBUTE_STREAM_PRIORITIES_SUPPORTED = , /**< Device supports stream priorities */ - // HIP_DEVICE_ATTRIBUTE_GLOBAL_L1_CACHE_SUPPORTED = , /**< Device supports caching globals in L1 */ - // HIP_DEVICE_ATTRIBUTE_LOCAL_L1_CACHE_SUPPORTED = , /**< Device supports caching locals in L1 */ - // HIP_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_MULTIPROCESSOR = , /**< Maximum number of 32-bit registers available per multiprocessor */ - // HIP_DEVICE_ATTRIBUTE_MULTI_GPU_BOARD_GROUP_ID = , /**< Unique id for a group of devices on the same multi-GPU board */ - // HIP_DEVICE_ATTRIBUTE_HOST_NATIVE_ATOMIC_SUPPORTED = , /**< Link between the device and the host supports native atomic operations (this is a placeholder attribute, and is not supported on any current hardware)*/ - // HIP_DEVICE_ATTRIBUTE_SINGLE_TO_DOUBLE_PRECISION_PERF_RATIO = , /**< Ratio of single precision performance (in floating-point operations per second) to double precision performance */ - // HIP_DEVICE_ATTRIBUTE_COMPUTE_PREEMPTION_SUPPORTED = , /**< Device supports compute preemption. */ - // HIP_DEVICE_ATTRIBUTE_CAN_USE_HOST_POINTER_FOR_REGISTERED_MEM = , /**< Device can access host registered memory at the same virtual address as the CPU */ - // HIP_DEVICE_ATTRIBUTE_CAN_USE_STREAM_MEM_OPS = , /**< ::hipStreamBatchMemOp and related APIs are supported. */ - // HIP_DEVICE_ATTRIBUTE_CAN_USE_64_BIT_STREAM_MEM_OPS = , /**< 64-bit operations are supported in ::hipStreamBatchMemOp and related APIs. */ - // HIP_DEVICE_ATTRIBUTE_CAN_FLUSH_REMOTE_WRITES = , /**< Both the ::HIP_STREAM_WAIT_VALUE_FLUSH flag and the ::HIP_STREAM_MEM_OP_FLUSH_REMOTE_WRITES MemOp are supported on the device. See \ref HIP_MEMOP for additional details. */ - // HIP_DEVICE_ATTRIBUTE_HOST_REGISTER_SUPPORTED = , /**< Device supports host memory registration via ::HIPHostRegister. */ - // HIP_DEVICE_ATTRIBUTE_MAX -} HIPdevice_attribute; - -/** - * Function cache configurations - */ -typedef enum HIPfunc_cache_enum { - HIP_FUNC_CACHE_PREFER_NONE = 0x00, /**< no preference for shared memory or L1 (default) */ - HIP_FUNC_CACHE_PREFER_SHARED = 0x01, /**< prefer larger shared memory and smaller L1 cache */ - HIP_FUNC_CACHE_PREFER_L1 = 0x02, /**< prefer larger L1 cache and smaller shared memory */ - HIP_FUNC_CACHE_PREFER_EQUAL = 0x03 /**< prefer equal sized L1 cache and shared memory */ -} HIPfunc_cache; - -/** - * Shared memory configurations - */ -typedef enum HIPsharedconfig_enum { - HIP_SHARED_MEM_CONFIG_DEFAULT_BANK_SIZE = 0x00, /**< set default shared memory bank size */ - HIP_SHARED_MEM_CONFIG_FOUR_BYTE_BANK_SIZE = 0x01, /**< set shared memory bank width to four bytes */ - HIP_SHARED_MEM_CONFIG_EIGHT_BYTE_BANK_SIZE = 0x02 /**< set shared memory bank width to eight bytes */ -} HIPsharedconfig; - -/** - * Function properties - */ -typedef enum HIPfunction_attribute_enum { - /** - * The maximum number of threads per block, beyond which a launch of the - * function would fail. This number depends on both the function and the - * device on which the function is currently loaded. - */ - HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK = 0, - - /** - * The size in bytes of statically-allocated shared memory required by - * this function. This does not include dynamically-allocated shared - * memory requested by the user at runtime. - */ - HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES = 1, - - /** - * The size in bytes of user-allocated constant memory required by this - * function. - */ - HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES = 2, - - /** - * The size in bytes of local memory used by each thread of this function. - */ - HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES = 3, - - /** - * The number of registers used by each thread of this function. - */ - HIP_FUNC_ATTRIBUTE_NUM_REGS = 4, - - /** - * The PTX virtual architecture version for which the function was - * compiled. This value is the major PTX version * 10 + the minor PTX - * version, so a PTX version 1.3 function would return the value 13. - * Note that this may return the undefined value of 0 for cubins - * compiled prior to HIP 3.0. - */ - HIP_FUNC_ATTRIBUTE_PTX_VERSION = 5, - - /** - * The binary architecture version for which the function was compiled. - * This value is the major binary version * 10 + the minor binary version, - * so a binary version 1.3 function would return the value 13. Note that - * this will return a value of 10 for legacy cubins that do not have a - * properly-encoded binary architecture version. - */ - HIP_FUNC_ATTRIBUTE_BINARY_VERSION = 6, - - /** - * The attribute to indicate whether the function has been compiled with - * user specified option "-Xptxas --dlcm=ca" set . - */ - HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA = 7, - - /** - * The maximum size in bytes of dynamically-allocated shared memory that can be used by - * this function. If the user-specified dynamic shared memory size is larger than this - * value, the launch will fail. - * See ::hipFuncSetAttribute - */ - HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES = 8, - - /** - * On devices where the L1 cache and shared memory use the same hardware resources, - * this sets the shared memory carveout preference, in percent of the total shared memory. - * Refer to ::HIP_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_MULTIPROCESSOR. - * This is only a hint, and the driver can choose a different ratio if required to execute the function. - * See ::hipFuncSetAttribute - */ - HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT = 9, - - HIP_FUNC_ATTRIBUTE_MAX -} HIPfunction_attribute; - -/** - * Context creation flags - */ -typedef enum HIPctx_flags_enum { - HIP_CTX_SCHED_AUTO = 0x00, /**< Automatic scheduling */ - HIP_CTX_SCHED_SPIN = 0x01, /**< Set spin as default scheduling */ - HIP_CTX_SCHED_YIELD = 0x02, /**< Set yield as default scheduling */ - HIP_CTX_SCHED_BLOCKING_SYNC = 0x04, /**< Set blocking synchronization as default scheduling */ - HIP_CTX_BLOCKING_SYNC = 0x04, /**< Set blocking synchronization as default scheduling - * \deprecated This flag was deprecated as of HIP 4.0 - * and was replaced with ::HIP_CTX_SCHED_BLOCKING_SYNC. */ - HIP_CTX_SCHED_MASK = 0x07, - HIP_CTX_MAP_HOST = 0x08, /**< Support mapped pinned allocations */ - HIP_CTX_LMEM_RESIZE_TO_MAX = 0x10, /**< Keep local memory allocation after launch */ - HIP_CTX_FLAGS_MASK = 0x1f -} HIPctx_flags; - -/** - * Stream creation flags - */ -typedef enum HIPstream_flags_enum { - HIP_STREAM_DEFAULT = 0x0, /**< Default stream flag */ - HIP_STREAM_NON_BLOCKING = 0x1 /**< Stream does not synchronize with stream 0 (the NULL stream) */ -} HIPstream_flags; - -/** - * Event creation flags - */ -typedef enum HIPevent_flags_enum { - HIP_EVENT_DEFAULT = 0x0, /**< Default event flag */ - HIP_EVENT_BLOCKING_SYNC = 0x1, /**< Event uses blocking synchronization */ - HIP_EVENT_DISABLE_TIMING = 0x2, /**< Event will not record timing data */ - HIP_EVENT_INTERPROCESS = 0x4 /**< Event is suitable for interprocess use. HIP_EVENT_DISABLE_TIMING must be set */ -} HIPevent_flags; - -typedef enum HIPjitInputType_enum -{ - /** - * Compiled device-class-specific device code\n - * Applicable options: none - */ - HIP_JIT_INPUT_HIPBIN = 0, - - /** - * PTX source code\n - * Applicable options: PTX compiler options - */ - HIP_JIT_INPUT_PTX, - - /** - * Bundle of multiple cubins and/or PTX of some device code\n - * Applicable options: PTX compiler options, ::HIP_JIT_FALLBACK_STRATEGY - */ - HIP_JIT_INPUT_FATBINARY, - - /** - * Host object with embedded device code\n - * Applicable options: PTX compiler options, ::HIP_JIT_FALLBACK_STRATEGY - */ - HIP_JIT_INPUT_OBJECT, - - /** - * Archive of host objects with embedded device code\n - * Applicable options: PTX compiler options, ::HIP_JIT_FALLBACK_STRATEGY - */ - HIP_JIT_INPUT_LIBRARY, - - HIP_JIT_NUM_INPUT_TYPES -} HIPjitInputType; #ifdef _WIN32 #define HIPAPI __stdcall @@ -1004,66 +353,44 @@ typedef enum HIPjitInputType_enum #define HIP_API_CALL HIPAPI -typedef HIPresult (HIP_API_CALL *HIP_HIPCTXCREATE) (HIPcontext *, unsigned int, HIPdevice); -typedef HIPresult (HIP_API_CALL *HIP_HIPCTXDESTROY) (HIPcontext); -typedef HIPresult (HIP_API_CALL *HIP_HIPCTXGETCACHECONFIG) (HIPfunc_cache *); -typedef HIPresult (HIP_API_CALL *HIP_HIPCTXGETCURRENT) (HIPcontext *); -typedef HIPresult (HIP_API_CALL *HIP_HIPCTXGETSHAREDMEMCONFIG) (HIPsharedconfig *); -typedef HIPresult (HIP_API_CALL *HIP_HIPCTXPOPCURRENT) (HIPcontext *); -typedef HIPresult (HIP_API_CALL *HIP_HIPCTXPUSHCURRENT) (HIPcontext); -typedef HIPresult (HIP_API_CALL *HIP_HIPCTXSETCACHECONFIG) (HIPfunc_cache); -typedef HIPresult (HIP_API_CALL *HIP_HIPCTXSETCURRENT) (HIPcontext); -typedef HIPresult (HIP_API_CALL *HIP_HIPCTXSETSHAREDMEMCONFIG) (HIPsharedconfig); -typedef HIPresult (HIP_API_CALL *HIP_HIPCTXSYNCHRONIZE) (); -typedef HIPresult (HIP_API_CALL *HIP_HIPDEVICEGETATTRIBUTE) (int *, HIPdevice_attribute, HIPdevice); -typedef HIPresult (HIP_API_CALL *HIP_HIPDEVICEGETCOUNT) (int *); -typedef HIPresult (HIP_API_CALL *HIP_HIPDEVICEGET) (HIPdevice *, int); -typedef HIPresult (HIP_API_CALL *HIP_HIPDEVICEGETNAME) (char *, int, HIPdevice); -typedef HIPresult (HIP_API_CALL *HIP_HIPDEVICETOTALMEM) (size_t *, HIPdevice); -typedef HIPresult (HIP_API_CALL *HIP_HIPDRIVERGETVERSION) (int *); -typedef HIPresult (HIP_API_CALL *HIP_HIPEVENTCREATE) (HIPevent *, unsigned int); -typedef HIPresult (HIP_API_CALL *HIP_HIPEVENTDESTROY) (HIPevent); -typedef HIPresult (HIP_API_CALL *HIP_HIPEVENTELAPSEDTIME) (float *, HIPevent, HIPevent); -typedef HIPresult (HIP_API_CALL *HIP_HIPEVENTQUERY) (HIPevent); -typedef HIPresult (HIP_API_CALL *HIP_HIPEVENTRECORD) (HIPevent, HIPstream); -typedef HIPresult (HIP_API_CALL *HIP_HIPEVENTSYNCHRONIZE) (HIPevent); -typedef HIPresult (HIP_API_CALL *HIP_HIPFUNCGETATTRIBUTE) (int *, HIPfunction_attribute, HIPfunction); -typedef HIPresult (HIP_API_CALL *HIP_HIPFUNCSETATTRIBUTE) (HIPfunction, HIPfunction_attribute, int); -typedef HIPresult (HIP_API_CALL *HIP_HIPFUNCSETCACHECONFIG) (HIPfunction, HIPfunc_cache); -typedef HIPresult (HIP_API_CALL *HIP_HIPFUNCSETSHAREDMEMCONFIG) (HIPfunction, HIPsharedconfig); -typedef HIPresult (HIP_API_CALL *HIP_HIPGETERRORNAME) (HIPresult, const char **); -typedef HIPresult (HIP_API_CALL *HIP_HIPGETERRORSTRING) (HIPresult, const char **); -typedef HIPresult (HIP_API_CALL *HIP_HIPINIT) (unsigned int); -typedef HIPresult (HIP_API_CALL *HIP_HIPLAUNCHKERNEL) (HIPfunction, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, HIPstream, void **, void **); -typedef HIPresult (HIP_API_CALL *HIP_HIPMEMALLOC) (HIPdeviceptr *, size_t); -typedef HIPresult (HIP_API_CALL *HIP_HIPMEMALLOCHOST) (void **, size_t); -typedef HIPresult (HIP_API_CALL *HIP_HIPMEMCPYDTOD) (HIPdeviceptr, HIPdeviceptr, size_t); -typedef HIPresult (HIP_API_CALL *HIP_HIPMEMCPYDTODASYNC) (HIPdeviceptr, HIPdeviceptr, size_t, HIPstream); -typedef HIPresult (HIP_API_CALL *HIP_HIPMEMCPYDTOH) (void *, HIPdeviceptr, size_t); -typedef HIPresult (HIP_API_CALL *HIP_HIPMEMCPYDTOHASYNC) (void *, HIPdeviceptr, size_t, HIPstream); -typedef HIPresult (HIP_API_CALL *HIP_HIPMEMCPYHTOD) (HIPdeviceptr, const void *, size_t); -typedef HIPresult (HIP_API_CALL *HIP_HIPMEMCPYHTODASYNC) (HIPdeviceptr, const void *, size_t, HIPstream); -typedef HIPresult (HIP_API_CALL *HIP_HIPMEMFREE) (HIPdeviceptr); -typedef HIPresult (HIP_API_CALL *HIP_HIPMEMFREEHOST) (void *); -typedef HIPresult (HIP_API_CALL *HIP_HIPMEMGETINFO) (size_t *, size_t *); -typedef HIPresult (HIP_API_CALL *HIP_HIPMEMSETD32) (HIPdeviceptr, unsigned int, size_t); -typedef HIPresult (HIP_API_CALL *HIP_HIPMEMSETD8) (HIPdeviceptr, unsigned char, size_t); -typedef HIPresult (HIP_API_CALL *HIP_HIPMODULEGETFUNCTION) (HIPfunction *, HIPmodule, const char *); -typedef HIPresult (HIP_API_CALL *HIP_HIPMODULEGETGLOBAL) (HIPdeviceptr *, size_t *, HIPmodule, const char *); -typedef HIPresult (HIP_API_CALL *HIP_HIPMODULELOAD) (HIPmodule *, const char *); -typedef HIPresult (HIP_API_CALL *HIP_HIPMODULELOADDATA) (HIPmodule *, const void *); -typedef HIPresult (HIP_API_CALL *HIP_HIPMODULELOADDATAEX) (HIPmodule *, const void *, unsigned int, HIPjit_option *, void **); -typedef HIPresult (HIP_API_CALL *HIP_HIPMODULEUNLOAD) (HIPmodule); -typedef HIPresult (HIP_API_CALL *HIP_HIPPROFILERSTART) (); -typedef HIPresult (HIP_API_CALL *HIP_HIPPROFILERSTOP) (); -typedef HIPresult (HIP_API_CALL *HIP_HIPSTREAMCREATE) (HIPstream *, unsigned int); -typedef HIPresult (HIP_API_CALL *HIP_HIPSTREAMDESTROY) (HIPstream); -typedef HIPresult (HIP_API_CALL *HIP_HIPSTREAMSYNCHRONIZE) (HIPstream); -typedef HIPresult (HIP_API_CALL *HIP_HIPSTREAMWAITEVENT) (HIPstream, HIPevent, unsigned int); -typedef HIPresult (HIP_API_CALL *HIP_HIPLINKCREATE) (unsigned int, HIPjit_option *, void **, HIPlinkState *); -typedef HIPresult (HIP_API_CALL *HIP_HIPLINKADDDATA) (HIPlinkState, HIPjitInputType, void *, size_t, const char *, unsigned int, HIPjit_option *, void **); -typedef HIPresult (HIP_API_CALL *HIP_HIPLINKDESTROY) (HIPlinkState); -typedef HIPresult (HIP_API_CALL *HIP_HIPLINKCOMPLETE) (HIPlinkState, void **, size_t *); +typedef hipError_t (HIP_API_CALL *HIP_HIPCTXCREATE) (hipCtx_t *, unsigned int, hipDevice_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPCTXDESTROY) (hipCtx_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPCTXPOPCURRENT) (hipCtx_t *); +typedef hipError_t (HIP_API_CALL *HIP_HIPCTXPUSHCURRENT) (hipCtx_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPCTXSETCURRENT) (hipCtx_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPCTXSYNCHRONIZE) (); +typedef hipError_t (HIP_API_CALL *HIP_HIPDEVICEGETATTRIBUTE) (int *, hipDeviceAttribute_t, hipDevice_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPDEVICEGETCOUNT) (int *); +typedef hipError_t (HIP_API_CALL *HIP_HIPDEVICEGET) (hipDevice_t *, int); +typedef hipError_t (HIP_API_CALL *HIP_HIPDEVICEGETNAME) (char *, int, hipDevice_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPDEVICETOTALMEM) (size_t *, hipDevice_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPDRIVERGETVERSION) (int *); +typedef hipError_t (HIP_API_CALL *HIP_HIPEVENTCREATE) (hipEvent_t *, unsigned int); +typedef hipError_t (HIP_API_CALL *HIP_HIPEVENTDESTROY) (hipEvent_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPEVENTELAPSEDTIME) (float *, hipEvent_t, hipEvent_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPEVENTRECORD) (hipEvent_t, hipStream_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPEVENTSYNCHRONIZE) (hipEvent_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPFUNCGETATTRIBUTE) (int *, hipFunction_attribute, hipFunction_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPGETERRORNAME) (hipError_t, const char **); +typedef hipError_t (HIP_API_CALL *HIP_HIPGETERRORSTRING) (hipError_t, const char **); +typedef hipError_t (HIP_API_CALL *HIP_HIPINIT) (unsigned int); +typedef hipError_t (HIP_API_CALL *HIP_HIPLAUNCHKERNEL) (hipFunction_t, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, hipStream_t, void **, void **); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMALLOC) (hipDeviceptr_t *, size_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMFREE) (hipDeviceptr_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMGETINFO) (size_t *, size_t *); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOD) (hipDeviceptr_t, hipDeviceptr_t, size_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTODASYNC) (hipDeviceptr_t, hipDeviceptr_t, size_t, hipStream_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOH) (void *, hipDeviceptr_t, size_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOHASYNC) (void *, hipDeviceptr_t, size_t, hipStream_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYHTOD) (hipDeviceptr_t, const void *, size_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYHTODASYNC) (hipDeviceptr_t, const void *, size_t, hipStream_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPMODULEGETFUNCTION) (hipFunction_t *, hipModule_t, const char *); +typedef hipError_t (HIP_API_CALL *HIP_HIPMODULEGETGLOBAL) (hipDeviceptr_t *, size_t *, hipModule_t, const char *); +typedef hipError_t (HIP_API_CALL *HIP_HIPMODULELOADDATAEX) (hipModule_t *, const void *, unsigned int, hipJitOption *, void **); +typedef hipError_t (HIP_API_CALL *HIP_HIPMODULEUNLOAD) (hipModule_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPSTREAMCREATE) (hipStream_t *, unsigned int); +typedef hipError_t (HIP_API_CALL *HIP_HIPSTREAMDESTROY) (hipStream_t); +typedef hipError_t (HIP_API_CALL *HIP_HIPSTREAMSYNCHRONIZE) (hipStream_t); typedef struct hc_hip_lib { @@ -1071,14 +398,9 @@ typedef struct hc_hip_lib HIP_HIPCTXCREATE hipCtxCreate; HIP_HIPCTXDESTROY hipCtxDestroy; - HIP_HIPCTXGETCACHECONFIG hipCtxGetCacheConfig; - HIP_HIPCTXGETCURRENT hipCtxGetCurrent; - HIP_HIPCTXGETSHAREDMEMCONFIG hipCtxGetSharedMemConfig; HIP_HIPCTXPOPCURRENT hipCtxPopCurrent; HIP_HIPCTXPUSHCURRENT hipCtxPushCurrent; - HIP_HIPCTXSETCACHECONFIG hipCtxSetCacheConfig; HIP_HIPCTXSETCURRENT hipCtxSetCurrent; - HIP_HIPCTXSETSHAREDMEMCONFIG hipCtxSetSharedMemConfig; HIP_HIPCTXSYNCHRONIZE hipCtxSynchronize; HIP_HIPDEVICEGETATTRIBUTE hipDeviceGetAttribute; HIP_HIPDEVICEGETCOUNT hipDeviceGetCount; @@ -1089,46 +411,29 @@ typedef struct hc_hip_lib HIP_HIPEVENTCREATE hipEventCreate; HIP_HIPEVENTDESTROY hipEventDestroy; HIP_HIPEVENTELAPSEDTIME hipEventElapsedTime; - HIP_HIPEVENTQUERY hipEventQuery; HIP_HIPEVENTRECORD hipEventRecord; HIP_HIPEVENTSYNCHRONIZE hipEventSynchronize; HIP_HIPFUNCGETATTRIBUTE hipFuncGetAttribute; - HIP_HIPFUNCSETATTRIBUTE hipFuncSetAttribute; - HIP_HIPFUNCSETCACHECONFIG hipFuncSetCacheConfig; - HIP_HIPFUNCSETSHAREDMEMCONFIG hipFuncSetSharedMemConfig; HIP_HIPGETERRORNAME hipGetErrorName; HIP_HIPGETERRORSTRING hipGetErrorString; HIP_HIPINIT hipInit; HIP_HIPLAUNCHKERNEL hipLaunchKernel; HIP_HIPMEMALLOC hipMemAlloc; - HIP_HIPMEMALLOCHOST hipMemAllocHost; + HIP_HIPMEMFREE hipMemFree; + HIP_HIPMEMGETINFO hipMemGetInfo; HIP_HIPMEMCPYDTOD hipMemcpyDtoD; HIP_HIPMEMCPYDTODASYNC hipMemcpyDtoDAsync; HIP_HIPMEMCPYDTOH hipMemcpyDtoH; HIP_HIPMEMCPYDTOHASYNC hipMemcpyDtoHAsync; HIP_HIPMEMCPYHTOD hipMemcpyHtoD; HIP_HIPMEMCPYHTODASYNC hipMemcpyHtoDAsync; - HIP_HIPMEMFREE hipMemFree; - HIP_HIPMEMFREEHOST hipMemFreeHost; - HIP_HIPMEMGETINFO hipMemGetInfo; - HIP_HIPMEMSETD32 hipMemsetD32; - HIP_HIPMEMSETD8 hipMemsetD8; HIP_HIPMODULEGETFUNCTION hipModuleGetFunction; HIP_HIPMODULEGETGLOBAL hipModuleGetGlobal; - HIP_HIPMODULELOAD hipModuleLoad; - HIP_HIPMODULELOADDATA hipModuleLoadData; HIP_HIPMODULELOADDATAEX hipModuleLoadDataEx; HIP_HIPMODULEUNLOAD hipModuleUnload; - HIP_HIPPROFILERSTART hipProfilerStart; - HIP_HIPPROFILERSTOP hipProfilerStop; HIP_HIPSTREAMCREATE hipStreamCreate; HIP_HIPSTREAMDESTROY hipStreamDestroy; HIP_HIPSTREAMSYNCHRONIZE hipStreamSynchronize; - HIP_HIPSTREAMWAITEVENT hipStreamWaitEvent; - HIP_HIPLINKCREATE hipLinkCreate; - HIP_HIPLINKADDDATA hipLinkAddData; - HIP_HIPLINKDESTROY hipLinkDestroy; - HIP_HIPLINKCOMPLETE hipLinkComplete; } hc_hip_lib_t; diff --git a/include/ext_hiprtc.h b/include/ext_hiprtc.h index cd1be6c4b..347239c38 100644 --- a/include/ext_hiprtc.h +++ b/include/ext_hiprtc.h @@ -6,41 +6,26 @@ #ifndef _EXT_HIPRTC_H #define _EXT_HIPRTC_H -/** - * from hip_runtime.h (/opt/rocm/hip/include/hip/amd_detail/hiprtc.h) - */ +// start: amd_detail/hiprtc.h -/** - * \ingroup error - * \brief The enumerated type hiprtcResult defines API call result codes. - * HIPRTC API functions return hiprtcResult to indicate the call - * result. - */ -typedef enum { - HIPRTC_SUCCESS = 0, - HIPRTC_ERROR_OUT_OF_MEMORY = 1, - HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2, - HIPRTC_ERROR_INVALID_INPUT = 3, - HIPRTC_ERROR_INVALID_PROGRAM = 4, - HIPRTC_ERROR_INVALID_OPTION = 5, - HIPRTC_ERROR_COMPILATION = 6, - HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7, - HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8, - HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9, - HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10, - HIPRTC_ERROR_INTERNAL_ERROR = 11 +typedef enum hiprtcResult { + HIPRTC_SUCCESS = 0, + HIPRTC_ERROR_OUT_OF_MEMORY = 1, + HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2, + HIPRTC_ERROR_INVALID_INPUT = 3, + HIPRTC_ERROR_INVALID_PROGRAM = 4, + HIPRTC_ERROR_INVALID_OPTION = 5, + HIPRTC_ERROR_COMPILATION = 6, + HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7, + HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8, + HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9, + HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10, + HIPRTC_ERROR_INTERNAL_ERROR = 11 } hiprtcResult; -/** - * \ingroup compilation - * \brief hiprtcProgram is the unit of compilation, and an opaque handle for - * a program. - * - * To compile a CUDA program string, an instance of hiprtcProgram must be - * created first with ::hiprtcCreateProgram, then compiled with - * ::hiprtcCompileProgram. - */ -typedef struct _hiprtcProgram *hiprtcProgram; +typedef struct _hiprtcProgram* hiprtcProgram; + +// stop: amd_detail/hiprtc.h #ifdef _WIN32 #define HIPRTCAPI __stdcall @@ -54,13 +39,12 @@ typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCADDNAMEEXPRESSION) (hiprtc typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCCOMPILEPROGRAM) (hiprtcProgram, int, const char * const *); typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCCREATEPROGRAM) (hiprtcProgram *, const char *, const char *, int, const char * const *, const char * const *); typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCDESTROYPROGRAM) (hiprtcProgram *); +typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETCODE) (hiprtcProgram, char *); +typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETCODESIZE) (hiprtcProgram, size_t *); typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETLOWEREDNAME) (hiprtcProgram, const char * const, const char **); -typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETPTX) (hiprtcProgram, char *); -typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETPTXSIZE) (hiprtcProgram, size_t *); typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETPROGRAMLOG) (hiprtcProgram, char *); typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCGETPROGRAMLOGSIZE) (hiprtcProgram, size_t *); -typedef const char * (HIPRTC_API_CALL *HIPRTC_HIPRTCGETERRORSTRING) (hiprtcResult); -typedef hiprtcResult (HIPRTC_API_CALL *HIPRTC_HIPRTCVERSION) (int *, int *); +typedef const char * (HIPRTC_API_CALL *HIPRTC_HIPRTCGETERRORSTRING) (hiprtcResult); typedef struct hc_hiprtc_lib { @@ -70,13 +54,12 @@ typedef struct hc_hiprtc_lib HIPRTC_HIPRTCCOMPILEPROGRAM hiprtcCompileProgram; HIPRTC_HIPRTCCREATEPROGRAM hiprtcCreateProgram; HIPRTC_HIPRTCDESTROYPROGRAM hiprtcDestroyProgram; + HIPRTC_HIPRTCGETCODE hiprtcGetCode; + HIPRTC_HIPRTCGETCODESIZE hiprtcGetCodeSize; HIPRTC_HIPRTCGETLOWEREDNAME hiprtcGetLoweredName; - HIPRTC_HIPRTCGETPTX hiprtcGetCode; - HIPRTC_HIPRTCGETPTXSIZE hiprtcGetCodeSize; HIPRTC_HIPRTCGETPROGRAMLOG hiprtcGetProgramLog; HIPRTC_HIPRTCGETPROGRAMLOGSIZE hiprtcGetProgramLogSize; HIPRTC_HIPRTCGETERRORSTRING hiprtcGetErrorString; - HIPRTC_HIPRTCVERSION hiprtcVersion; } hc_hiprtc_lib_t; diff --git a/include/types.h b/include/types.h index ed22a95ee..efc56439b 100644 --- a/include/types.h +++ b/include/types.h @@ -1502,80 +1502,80 @@ typedef struct hc_device_param int hip_warp_size; - HIPdevice hip_device; - HIPcontext hip_context; - HIPstream hip_stream; + hipDevice_t hip_device; + hipCtx_t hip_context; + hipStream_t hip_stream; - HIPevent hip_event1; - HIPevent hip_event2; + hipEvent_t hip_event1; + hipEvent_t hip_event2; - HIPmodule hip_module; - HIPmodule hip_module_shared; - HIPmodule hip_module_mp; - HIPmodule hip_module_amp; + hipModule_t hip_module; + hipModule_t hip_module_shared; + hipModule_t hip_module_mp; + hipModule_t hip_module_amp; - HIPfunction hip_function1; - HIPfunction hip_function12; - HIPfunction hip_function2p; - HIPfunction hip_function2; - HIPfunction hip_function2e; - HIPfunction hip_function23; - HIPfunction hip_function3; - HIPfunction hip_function4; - HIPfunction hip_function_init2; - HIPfunction hip_function_loop2p; - HIPfunction hip_function_loop2; - HIPfunction hip_function_mp; - HIPfunction hip_function_mp_l; - HIPfunction hip_function_mp_r; - HIPfunction hip_function_amp; - HIPfunction hip_function_tm; - HIPfunction hip_function_memset; - HIPfunction hip_function_bzero; - HIPfunction hip_function_atinit; - HIPfunction hip_function_utf8toutf16le; - HIPfunction hip_function_decompress; - HIPfunction hip_function_aux1; - HIPfunction hip_function_aux2; - HIPfunction hip_function_aux3; - HIPfunction hip_function_aux4; + hipFunction_t hip_function1; + hipFunction_t hip_function12; + hipFunction_t hip_function2p; + hipFunction_t hip_function2; + hipFunction_t hip_function2e; + hipFunction_t hip_function23; + hipFunction_t hip_function3; + hipFunction_t hip_function4; + hipFunction_t hip_function_init2; + hipFunction_t hip_function_loop2p; + hipFunction_t hip_function_loop2; + hipFunction_t hip_function_mp; + hipFunction_t hip_function_mp_l; + hipFunction_t hip_function_mp_r; + hipFunction_t hip_function_amp; + hipFunction_t hip_function_tm; + hipFunction_t hip_function_memset; + hipFunction_t hip_function_bzero; + hipFunction_t hip_function_atinit; + hipFunction_t hip_function_utf8toutf16le; + hipFunction_t hip_function_decompress; + hipFunction_t hip_function_aux1; + hipFunction_t hip_function_aux2; + hipFunction_t hip_function_aux3; + hipFunction_t hip_function_aux4; - HIPdeviceptr hip_d_pws_buf; - HIPdeviceptr hip_d_pws_amp_buf; - HIPdeviceptr hip_d_pws_comp_buf; - HIPdeviceptr hip_d_pws_idx; - HIPdeviceptr hip_d_rules; - HIPdeviceptr hip_d_rules_c; - HIPdeviceptr hip_d_combs; - HIPdeviceptr hip_d_combs_c; - HIPdeviceptr hip_d_bfs; - HIPdeviceptr hip_d_bfs_c; - HIPdeviceptr hip_d_tm_c; - HIPdeviceptr hip_d_bitmap_s1_a; - HIPdeviceptr hip_d_bitmap_s1_b; - HIPdeviceptr hip_d_bitmap_s1_c; - HIPdeviceptr hip_d_bitmap_s1_d; - HIPdeviceptr hip_d_bitmap_s2_a; - HIPdeviceptr hip_d_bitmap_s2_b; - HIPdeviceptr hip_d_bitmap_s2_c; - HIPdeviceptr hip_d_bitmap_s2_d; - HIPdeviceptr hip_d_plain_bufs; - HIPdeviceptr hip_d_digests_buf; - HIPdeviceptr hip_d_digests_shown; - HIPdeviceptr hip_d_salt_bufs; - HIPdeviceptr hip_d_esalt_bufs; - HIPdeviceptr hip_d_tmps; - HIPdeviceptr hip_d_hooks; - HIPdeviceptr hip_d_result; - HIPdeviceptr hip_d_extra0_buf; - HIPdeviceptr hip_d_extra1_buf; - HIPdeviceptr hip_d_extra2_buf; - HIPdeviceptr hip_d_extra3_buf; - HIPdeviceptr hip_d_root_css_buf; - HIPdeviceptr hip_d_markov_css_buf; - HIPdeviceptr hip_d_st_digests_buf; - HIPdeviceptr hip_d_st_salts_buf; - HIPdeviceptr hip_d_st_esalts_buf; + hipDeviceptr_t hip_d_pws_buf; + hipDeviceptr_t hip_d_pws_amp_buf; + hipDeviceptr_t hip_d_pws_comp_buf; + hipDeviceptr_t hip_d_pws_idx; + hipDeviceptr_t hip_d_rules; + hipDeviceptr_t hip_d_rules_c; + hipDeviceptr_t hip_d_combs; + hipDeviceptr_t hip_d_combs_c; + hipDeviceptr_t hip_d_bfs; + hipDeviceptr_t hip_d_bfs_c; + hipDeviceptr_t hip_d_tm_c; + hipDeviceptr_t hip_d_bitmap_s1_a; + hipDeviceptr_t hip_d_bitmap_s1_b; + hipDeviceptr_t hip_d_bitmap_s1_c; + hipDeviceptr_t hip_d_bitmap_s1_d; + hipDeviceptr_t hip_d_bitmap_s2_a; + hipDeviceptr_t hip_d_bitmap_s2_b; + hipDeviceptr_t hip_d_bitmap_s2_c; + hipDeviceptr_t hip_d_bitmap_s2_d; + hipDeviceptr_t hip_d_plain_bufs; + hipDeviceptr_t hip_d_digests_buf; + hipDeviceptr_t hip_d_digests_shown; + hipDeviceptr_t hip_d_salt_bufs; + hipDeviceptr_t hip_d_esalt_bufs; + hipDeviceptr_t hip_d_tmps; + hipDeviceptr_t hip_d_hooks; + hipDeviceptr_t hip_d_result; + hipDeviceptr_t hip_d_extra0_buf; + hipDeviceptr_t hip_d_extra1_buf; + hipDeviceptr_t hip_d_extra2_buf; + hipDeviceptr_t hip_d_extra3_buf; + hipDeviceptr_t hip_d_root_css_buf; + hipDeviceptr_t hip_d_markov_css_buf; + hipDeviceptr_t hip_d_st_digests_buf; + hipDeviceptr_t hip_d_st_salts_buf; + hipDeviceptr_t hip_d_st_esalts_buf; // API: opencl @@ -1726,8 +1726,7 @@ typedef struct backend_ctx int rc_hip_init; int rc_hiprtc_init; - int hiprtc_driver_version; - int hip_driver_version; + int hip_driverVersion; // opencl diff --git a/src/backend.c b/src/backend.c index 31d6bade1..2982b7820 100644 --- a/src/backend.c +++ b/src/backend.c @@ -980,11 +980,11 @@ int hiprtc_init (hashcat_ctx_t *hashcat_ctx) memset (hiprtc, 0, sizeof (HIPRTC_PTR)); #if defined (_WIN) - hiprtc->lib = hc_dlopen ("fixme.dll"); + hiprtc->lib = hc_dlopen ("amdhip64.dll"); #elif defined (__APPLE__) hiprtc->lib = hc_dlopen ("fixme.dylib"); #elif defined (__CYGWIN__) - hiprtc->lib = hc_dlopen ("fixme.dll"); + hiprtc->lib = hc_dlopen ("amdhip64.dll"); #else hiprtc->lib = hc_dlopen ("libamdhip64.so"); @@ -998,12 +998,11 @@ int hiprtc_init (hashcat_ctx_t *hashcat_ctx) HC_LOAD_FUNC (hiprtc, hiprtcCreateProgram, HIPRTC_HIPRTCCREATEPROGRAM, HIPRTC, 1); HC_LOAD_FUNC (hiprtc, hiprtcDestroyProgram, HIPRTC_HIPRTCDESTROYPROGRAM, HIPRTC, 1); HC_LOAD_FUNC (hiprtc, hiprtcGetLoweredName, HIPRTC_HIPRTCGETLOWEREDNAME, HIPRTC, 1); - HC_LOAD_FUNC (hiprtc, hiprtcGetCode, HIPRTC_HIPRTCGETPTX, HIPRTC, 1); - HC_LOAD_FUNC (hiprtc, hiprtcGetCodeSize, HIPRTC_HIPRTCGETPTXSIZE, HIPRTC, 1); + HC_LOAD_FUNC (hiprtc, hiprtcGetCode, HIPRTC_HIPRTCGETCODE, HIPRTC, 1); + HC_LOAD_FUNC (hiprtc, hiprtcGetCodeSize, HIPRTC_HIPRTCGETCODESIZE, HIPRTC, 1); HC_LOAD_FUNC (hiprtc, hiprtcGetProgramLog, HIPRTC_HIPRTCGETPROGRAMLOG, HIPRTC, 1); HC_LOAD_FUNC (hiprtc, hiprtcGetProgramLogSize, HIPRTC_HIPRTCGETPROGRAMLOGSIZE, HIPRTC, 1); HC_LOAD_FUNC (hiprtc, hiprtcGetErrorString, HIPRTC_HIPRTCGETERRORSTRING, HIPRTC, 1); - HC_LOAD_FUNC (hiprtc, hiprtcVersion, HIPRTC_HIPRTCVERSION, HIPRTC, 1); return 0; } @@ -1069,11 +1068,6 @@ int hc_hiprtcCompileProgram (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, int HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; - #if 0 - for(int i =0; i< numOptions; i++) - printf("Option_%d = %s\n", i, options[i]); - #endif - const hiprtcResult HIPRTC_err = hiprtc->hiprtcCompileProgram (prog, numOptions, options); if (HIPRTC_err != HIPRTC_SUCCESS) @@ -1122,13 +1116,13 @@ int hc_hiprtcGetProgramLog (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char return 0; } -int hc_hiprtcGetCodeSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *ptxSizeRet) +int hc_hiprtcGetCodeSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t *codeSizeRet) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; - const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetCodeSize (prog, ptxSizeRet); + const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetCodeSize (prog, codeSizeRet); if (HIPRTC_err != HIPRTC_SUCCESS) { @@ -1140,13 +1134,13 @@ int hc_hiprtcGetCodeSize (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, size_t return 0; } -int hc_hiprtcGetCode (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *ptx) +int hc_hiprtcGetCode (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *code) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; - const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetCode (prog, ptx); + const hiprtcResult HIPRTC_err = hiprtc->hiprtcGetCode (prog, code); if (HIPRTC_err != HIPRTC_SUCCESS) { @@ -1158,24 +1152,6 @@ int hc_hiprtcGetCode (hashcat_ctx_t *hashcat_ctx, hiprtcProgram prog, char *ptx) return 0; } -int hc_hiprtcVersion (hashcat_ctx_t *hashcat_ctx, int *major, int *minor) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIPRTC_PTR *hiprtc = (HIPRTC_PTR *) backend_ctx->hiprtc; - - const hiprtcResult HIPRTC_err = hiprtc->hiprtcVersion (major, minor); - - if (HIPRTC_err != HIPRTC_SUCCESS) - { - event_log_error (hashcat_ctx, "hiprtcVersion(): %s", hiprtc->hiprtcGetErrorString (HIPRTC_err)); - - return -1; - } - - return 0; -} - // CUDA int cuda_init (hashcat_ctx_t *hashcat_ctx) @@ -2478,11 +2454,11 @@ int hip_init (hashcat_ctx_t *hashcat_ctx) memset (hip, 0, sizeof (HIP_PTR)); #if defined (_WIN) - hip->lib = hc_dlopen ("fixme.dll"); + hip->lib = hc_dlopen ("amdhip64.dll"); #elif defined (__APPLE__) hip->lib = hc_dlopen ("fixme.dylib"); #elif defined (__CYGWIN__) - hip->lib = hc_dlopen ("fixme.dll"); + hip->lib = hc_dlopen ("amdhip64.dll"); #else hip->lib = hc_dlopen ("libamdhip64.so"); @@ -2516,67 +2492,42 @@ int hip_init (hashcat_ctx_t *hashcat_ctx) HC_LOAD_FUNC_HIP (hip, hipCtxCreate, hipCtxCreate, HIP_HIPCTXCREATE, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipCtxDestroy, hipCtxDestroy, HIP_HIPCTXDESTROY, HIP, 1); - //HC_LOAD_FUNC_HIP (hip, hipCtxGetCacheConfig, hipCtxGetCacheConfig, HIP_HIPCTXGETCACHECONFIG, HIP, 1); - //HC_LOAD_FUNC_HIP (hip, hipCtxGetCurrent, hipCtxGetCurrent, HIP_HIPCTXGETCURRENT, HIP, 1); - //HC_LOAD_FUNC_HIP (hip, hipCtxGetSharedMemConfig, hipCtxGetSharedMemConfig, HIP_HIPCTXGETSHAREDMEMCONFIG, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipCtxPopCurrent, hipCtxPopCurrent, HIP_HIPCTXPOPCURRENT, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipCtxPushCurrent, hipCtxPushCurrent, HIP_HIPCTXPUSHCURRENT, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipCtxSetCacheConfig, hipCtxSetCacheConfig, HIP_HIPCTXSETCACHECONFIG, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipCtxSetCurrent, hipCtxSetCurrent, HIP_HIPCTXSETCURRENT, HIP, 1); - //HC_LOAD_FUNC_HIP (hip, hipCtxSetSharedMemConfig, hipCtxSetSharedMemConfig, HIP_HIPCTXSETSHAREDMEMCONFIG, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipCtxSynchronize, hipCtxSynchronize, HIP_HIPCTXSYNCHRONIZE, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipDeviceGet, hipDeviceGet, HIP_HIPDEVICEGET, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipDeviceGetAttribute, hipDeviceGetAttribute, HIP_HIPDEVICEGETATTRIBUTE, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipDeviceGetCount, hipGetDeviceCount, HIP_HIPDEVICEGETCOUNT, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipDeviceGet, hipDeviceGet, HIP_HIPDEVICEGET, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipDeviceGetName, hipDeviceGetName, HIP_HIPDEVICEGETNAME, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipDeviceTotalMem, hipDeviceTotalMem, HIP_HIPDEVICETOTALMEM, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipDriverGetVersion, hipDriverGetVersion, HIP_HIPDRIVERGETVERSION, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipEventCreate, hipEventCreateWithFlags, HIP_HIPEVENTCREATE, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipEventDestroy, hipEventDestroy, HIP_HIPEVENTDESTROY, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipEventElapsedTime, hipEventElapsedTime, HIP_HIPEVENTELAPSEDTIME, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipEventQuery, hipEventQuery, HIP_HIPEVENTQUERY, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipEventRecord, hipEventRecord, HIP_HIPEVENTRECORD, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipEventSynchronize, hipEventSynchronize, HIP_HIPEVENTSYNCHRONIZE, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipFuncGetAttribute, hipFuncGetAttribute, HIP_HIPFUNCGETATTRIBUTE, HIP, 1); - //HC_LOAD_FUNC_HIP (hip, hipFuncSetAttribute, hipFuncSetAttribute, HIP_HIPFUNCSETATTRIBUTE, HIP, 1); - //HC_LOAD_FUNC_HIP (hip, hipFuncSetCacheConfig, hipFuncSetCacheConfig, HIP_HIPFUNCSETCACHECONFIG, HIP, 1); - //HC_LOAD_FUNC_HIP (hip, hipFuncSetSharedMemConfig, hipFuncSetSharedMemConfig, HIP_HIPFUNCSETSHAREDMEMCONFIG, HIP, 1); - //HC_LOAD_FUNC_HIP (hip, hipGetErrorName, hipGetErrorName, HIP_HIPGETERRORNAME, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipGetErrorName, hipGetErrorName, HIP_HIPGETERRORNAME, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipGetErrorString, hipGetErrorString, HIP_HIPGETERRORSTRING, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipInit, hipInit, HIP_HIPINIT, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipLaunchKernel, hipModuleLaunchKernel, HIP_HIPLAUNCHKERNEL, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemAlloc, hipMalloc, HIP_HIPMEMALLOC, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipMemAllocHost, hipMemAllocHost, HIP_HIPMEMALLOCHOST, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemFree, hipFree, HIP_HIPMEMFREE, HIP, 1); + HC_LOAD_FUNC_HIP (hip, hipMemGetInfo, hipMemGetInfo, HIP_HIPMEMGETINFO, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoD, hipMemcpyDtoD, HIP_HIPMEMCPYDTOD, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoDAsync, hipMemcpyDtoDAsync, HIP_HIPMEMCPYDTODASYNC, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoH, hipMemcpyDtoH, HIP_HIPMEMCPYDTOH, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemcpyDtoHAsync, hipMemcpyDtoHAsync, HIP_HIPMEMCPYDTOHASYNC, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemcpyHtoD, hipMemcpyHtoD, HIP_HIPMEMCPYHTOD, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipMemcpyHtoDAsync, hipMemcpyHtoDAsync, HIP_HIPMEMCPYHTODASYNC, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipMemFree, hipFree, HIP_HIPMEMFREE, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipMemFreeHost, hipFreeHost, HIP_HIPMEMFREEHOST, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipMemGetInfo, hipMemGetInfo, HIP_HIPMEMGETINFO, HIP, 1); - //HC_LOAD_FUNC_HIP (hip, hipMemsetD32, hipMemsetD32, HIP_HIPMEMSETD32, HIP, 1); - //HC_LOAD_FUNC_HIP (hip, hipMemsetD8, hipMemsetD8, HIP_HIPMEMSETD8, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipModuleGetFunction, hipModuleGetFunction, HIP_HIPMODULEGETFUNCTION, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipModuleGetGlobal, hipModuleGetGlobal, HIP_HIPMODULEGETGLOBAL, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipModuleLoad, hipModuleLoad, HIP_HIPMODULELOAD, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipModuleLoadData, hipModuleLoadData, HIP_HIPMODULELOADDATA, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipModuleLoadDataEx, hipModuleLoadDataEx, HIP_HIPMODULELOADDATAEX, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipModuleUnload, hipModuleUnload, HIP_HIPMODULEUNLOAD, HIP, 1); - //HC_LOAD_FUNC_HIP (hip, hipProfilerStart, hipProfilerStart, HIP_HIPPROFILERSTART, HIP, 1); - //HC_LOAD_FUNC_HIP (hip, hipProfilerStop, hipProfilerStop, HIP_HIPPROFILERSTOP, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipStreamCreate, hipStreamCreate, HIP_HIPSTREAMCREATE, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipStreamDestroy, hipStreamDestroy, HIP_HIPSTREAMDESTROY, HIP, 1); HC_LOAD_FUNC_HIP (hip, hipStreamSynchronize, hipStreamSynchronize, HIP_HIPSTREAMSYNCHRONIZE, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipStreamWaitEvent, hipStreamWaitEvent, HIP_HIPSTREAMWAITEVENT, HIP, 1); - //TODO HIP? - #if defined (WITH_CUBINX) - HC_LOAD_FUNC_HIP (hip, hipLinkCreate, hipLinkCreate, HIP_HIPLINKCREATE, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipLinkAddData, hipLinkAddData, HIP_HIPLINKADDDATA, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipLinkDestroy, hipLinkDestroy, HIP_HIPLINKDESTROY, HIP, 1); - HC_LOAD_FUNC_HIP (hip, hipLinkComplete, hipLinkComplete, HIP_HIPLINKCOMPLETE, HIP, 1); - #endif return 0; } @@ -2600,25 +2551,25 @@ void hip_close (hashcat_ctx_t *hashcat_ctx) } } -int hc_hipInit (hashcat_ctx_t *hashcat_ctx, unsigned int Flags) +int hc_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, hipCtx_t *pctx, unsigned int flags, hipDevice_t dev) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipInit (Flags); + const hipError_t HIP_err = hip->hipCtxCreate (pctx, flags, dev); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { - event_log_error (hashcat_ctx, "hipInit(): %s", pStr); + event_log_error (hashcat_ctx, "hipCtxCreate(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipInit(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipCtxCreate(): %d", HIP_err); } return -1; @@ -2627,19 +2578,181 @@ int hc_hipInit (hashcat_ctx_t *hashcat_ctx, unsigned int Flags) return 0; } -int hc_hipDeviceGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPdevice_attribute attrib, HIPdevice dev) +int hc_hipCtxDestroy (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipDeviceGetAttribute (pi, attrib, dev); + const hipError_t HIP_err = hip->hipCtxDestroy (ctx); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipCtxDestroy(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipCtxDestroy(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipCtxPopCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t *pctx) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipCtxPopCurrent (pctx); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipCtxPopCurrent(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipCtxPopCurrent(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipCtxPushCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipCtxPushCurrent (ctx); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipCtxPushCurrent(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipCtxPushCurrent(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipCtxSetCurrent (hashcat_ctx_t *hashcat_ctx, hipCtx_t ctx) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipCtxSetCurrent (ctx); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipCtxSetCurrent(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipCtxSetCurrent(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipCtxSynchronize (hashcat_ctx_t *hashcat_ctx) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipCtxSynchronize (); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipCtxSynchronize(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipCtxSynchronize(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipDeviceGet (hashcat_ctx_t *hashcat_ctx, hipDevice_t* device, int ordinal) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipDeviceGet (device, ordinal); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipDeviceGet(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipDeviceGet(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipDeviceGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, hipDeviceAttribute_t attrib, hipDevice_t dev) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipDeviceGetAttribute (pi, attrib, dev); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { event_log_error (hashcat_ctx, "hipDeviceGetAttribute(): %s", pStr); } @@ -2660,13 +2773,13 @@ int hc_hipDeviceGetCount (hashcat_ctx_t *hashcat_ctx, int *count) HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipDeviceGetCount (count); + const hipError_t HIP_err = hip->hipDeviceGetCount (count); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { event_log_error (hashcat_ctx, "hipDeviceGetCount(): %s", pStr); } @@ -2681,46 +2794,19 @@ int hc_hipDeviceGetCount (hashcat_ctx_t *hashcat_ctx, int *count) return 0; } -int hc_hipDeviceGet (hashcat_ctx_t *hashcat_ctx, HIPdevice* device, int ordinal) +int hc_hipDeviceGetName (hashcat_ctx_t *hashcat_ctx, char *name, int len, hipDevice_t dev) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipDeviceGet (device, ordinal); + const hipError_t HIP_err = hip->hipDeviceGetName (name, len, dev); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipDeviceGet(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipDeviceGet(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipDeviceGetName (hashcat_ctx_t *hashcat_ctx, char *name, int len, HIPdevice dev) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipDeviceGetName (name, len, dev); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { event_log_error (hashcat_ctx, "hipDeviceGetName(): %s", pStr); } @@ -2735,19 +2821,19 @@ int hc_hipDeviceGetName (hashcat_ctx_t *hashcat_ctx, char *name, int len, HIPdev return 0; } -int hc_hipDeviceTotalMem (hashcat_ctx_t *hashcat_ctx, size_t *bytes, HIPdevice dev) +int hc_hipDeviceTotalMem (hashcat_ctx_t *hashcat_ctx, size_t *bytes, hipDevice_t dev) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipDeviceTotalMem (bytes, dev); + const hipError_t HIP_err = hip->hipDeviceTotalMem (bytes, dev); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { event_log_error (hashcat_ctx, "hipDeviceTotalMem(): %s", pStr); } @@ -2768,13 +2854,13 @@ int hc_hipDriverGetVersion (hashcat_ctx_t *hashcat_ctx, int *driverVersion) HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipDriverGetVersion (driverVersion); + const hipError_t HIP_err = hip->hipDriverGetVersion (driverVersion); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { event_log_error (hashcat_ctx, "hipDriverGetVersion(): %s", pStr); } @@ -2789,645 +2875,19 @@ int hc_hipDriverGetVersion (hashcat_ctx_t *hashcat_ctx, int *driverVersion) return 0; } -int hc_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx, unsigned int flags, HIPdevice dev) +int hc_hipEventCreate (hashcat_ctx_t *hashcat_ctx, hipEvent_t *phEvent, unsigned int Flags) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipCtxCreate (pctx, flags, dev); + const hipError_t HIP_err = hip->hipEventCreate (phEvent, Flags); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipCtxCreate(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipCtxCreate(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipCtxDestroy (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipCtxDestroy (ctx); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipCtxDestroy(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipCtxDestroy(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, HIPmodule *module, const void *image, unsigned int numOptions, HIPjit_option *options, void **optionValues) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipModuleLoadDataEx (module, image, numOptions, options, optionValues); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipModuleLoadDataEx(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipModuleLoadDataEx(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipModuleUnload (hashcat_ctx_t *hashcat_ctx, HIPmodule hmod) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipModuleUnload (hmod); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipModuleUnload(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipModuleUnload(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipCtxSetCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipCtxSetCurrent (ctx); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipCtxSetCurrent(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipCtxSetCurrent(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr *dptr, size_t bytesize) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipMemAlloc (dptr, bytesize); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipMemAlloc(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipMemAlloc(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipMemFree (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dptr) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipMemFree (dptr); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipMemFree(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipMemFree(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, HIPdeviceptr srcDevice, size_t ByteCount) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipMemcpyDtoH (dstHost, srcDevice, ByteCount); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipMemcpyDtoH(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipMemcpyDtoH(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipMemcpyDtoHAsync (hashcat_ctx_t *hashcat_ctx, void *dstHost, HIPdeviceptr srcDevice, size_t ByteCount, HIPstream hStream) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipMemcpyDtoHAsync (dstHost, srcDevice, ByteCount, hStream); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipMemcpyDtoHAsync(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipMemcpyDtoHAsync(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, HIPdeviceptr srcDevice, size_t ByteCount) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipMemcpyDtoD (dstDevice, srcDevice, ByteCount); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipMemcpyDtoDAsync (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, HIPdeviceptr srcDevice, size_t ByteCount, HIPstream hStream) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipMemcpyDtoDAsync (dstDevice, srcDevice, ByteCount, hStream); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipMemcpyDtoDAsync(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipMemcpyDtoDAsync(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, const void *srcHost, size_t ByteCount) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipMemcpyHtoD (dstDevice, srcHost, ByteCount); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipMemcpyHtoD(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipMemcpyHtoD(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipMemcpyHtoDAsync (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, const void *srcHost, size_t ByteCount, HIPstream hStream) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipMemcpyHtoDAsync (dstDevice, srcHost, ByteCount, hStream); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipMemcpyHtoDAsync(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipMemcpyHtoDAsync(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, HIPfunction *hfunc, HIPmodule hmod, const char *name) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipModuleGetFunction (hfunc, hmod, name); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipModuleGetFunction(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipModuleGetFunction(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipModuleGetGlobal (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr *dptr, size_t *bytes, HIPmodule hmod, const char *name) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipModuleGetGlobal (dptr, bytes, hmod, name); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipModuleGetGlobal(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipModuleGetGlobal(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipMemGetInfo (hashcat_ctx_t *hashcat_ctx, size_t *free, size_t *total) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipMemGetInfo (free, total); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipMemGetInfo(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipMemGetInfo(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPfunction_attribute attrib, HIPfunction hfunc) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipFuncGetAttribute (pi, attrib, hfunc); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipFuncGetAttribute(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipFuncGetAttribute(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -/* - -// ATTENTION, this one maps to cudaFuncSetAttribute not cuFuncSetAttribute !!! - -int hc_hipFuncSetAttribute (hashcat_ctx_t *hashcat_ctx, HIPfunction hfunc, HIPfunction_attribute attrib, int value) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipFuncSetAttribute (hfunc, attrib, value); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipFuncSetAttribute(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipFuncSetAttribute(): %d", HIP_err); - } - - return -1; - } - - return 0; -} -*/ - -int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, HIPstream *phStream, unsigned int Flags) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipStreamCreate (phStream, Flags); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipStreamCreate(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipStreamCreate(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipStreamDestroy (hashcat_ctx_t *hashcat_ctx, HIPstream hStream) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipStreamDestroy (hStream); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipStreamDestroy(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipStreamDestroy(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipStreamSynchronize (hashcat_ctx_t *hashcat_ctx, HIPstream hStream) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipStreamSynchronize (hStream); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipStreamSynchronize(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipStreamSynchronize(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipLaunchKernel (hashcat_ctx_t *hashcat_ctx, HIPfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, HIPstream hStream, void **kernelParams, void **extra) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipLaunchKernel (f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipLaunchKernel(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipLaunchKernel(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipCtxSynchronize (hashcat_ctx_t *hashcat_ctx) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipCtxSynchronize (); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipCtxSynchronize(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipCtxSynchronize(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipEventCreate (hashcat_ctx_t *hashcat_ctx, HIPevent *phEvent, unsigned int Flags) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipEventCreate (phEvent, Flags); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { event_log_error (hashcat_ctx, "hipEventCreate(): %s", pStr); } @@ -3442,19 +2902,19 @@ int hc_hipEventCreate (hashcat_ctx_t *hashcat_ctx, HIPevent *phEvent, unsigned i return 0; } -int hc_hipEventDestroy (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent) +int hc_hipEventDestroy (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipEventDestroy (hEvent); + const hipError_t HIP_err = hip->hipEventDestroy (hEvent); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { event_log_error (hashcat_ctx, "hipEventDestroy(): %s", pStr); } @@ -3469,19 +2929,19 @@ int hc_hipEventDestroy (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent) return 0; } -int hc_hipEventElapsedTime (hashcat_ctx_t *hashcat_ctx, float *pMilliseconds, HIPevent hStart, HIPevent hEnd) +int hc_hipEventElapsedTime (hashcat_ctx_t *hashcat_ctx, float *pMilliseconds, hipEvent_t hStart, hipEvent_t hEnd) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipEventElapsedTime (pMilliseconds, hStart, hEnd); + const hipError_t HIP_err = hip->hipEventElapsedTime (pMilliseconds, hStart, hEnd); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { event_log_error (hashcat_ctx, "hipEventElapsedTime(): %s", pStr); } @@ -3496,46 +2956,19 @@ int hc_hipEventElapsedTime (hashcat_ctx_t *hashcat_ctx, float *pMilliseconds, HI return 0; } -int hc_hipEventQuery (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent) +int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent, hipStream_t hStream) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipEventQuery (hEvent); + const hipError_t HIP_err = hip->hipEventRecord (hEvent, hStream); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipEventQuery(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipEventQuery(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent, HIPstream hStream) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipEventRecord (hEvent, hStream); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { event_log_error (hashcat_ctx, "hipEventRecord(): %s", pStr); } @@ -3550,19 +2983,19 @@ int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent, HIPstream hS return 0; } -int hc_hipEventSynchronize (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent) +int hc_hipEventSynchronize (hashcat_ctx_t *hashcat_ctx, hipEvent_t hEvent) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipEventSynchronize (hEvent); + const hipError_t HIP_err = hip->hipEventSynchronize (hEvent); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { event_log_error (hashcat_ctx, "hipEventSynchronize(): %s", pStr); } @@ -3577,25 +3010,25 @@ int hc_hipEventSynchronize (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent) return 0; } -int hc_hipCtxSetCacheConfig (hashcat_ctx_t *hashcat_ctx, HIPfunc_cache config) +int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, hipFunction_attribute attrib, hipFunction_t hfunc) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipCtxSetCacheConfig (config); + const hipError_t HIP_err = hip->hipFuncGetAttribute (pi, attrib, hfunc); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { - event_log_error (hashcat_ctx, "hipCtxSetCacheConfig(): %s", pStr); + event_log_error (hashcat_ctx, "hipFuncGetAttribute(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipCtxSetCacheConfig(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipFuncGetAttribute(): %d", HIP_err); } return -1; @@ -3604,25 +3037,25 @@ int hc_hipCtxSetCacheConfig (hashcat_ctx_t *hashcat_ctx, HIPfunc_cache config) return 0; } -int hc_hipCtxPushCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx) +int hc_hipLaunchKernel (hashcat_ctx_t *hashcat_ctx, hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipCtxPushCurrent (ctx); + const hipError_t HIP_err = hip->hipLaunchKernel (f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { - event_log_error (hashcat_ctx, "hipCtxPushCurrent(): %s", pStr); + event_log_error (hashcat_ctx, "hipLaunchKernel(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipCtxPushCurrent(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipLaunchKernel(): %d", HIP_err); } return -1; @@ -3631,25 +3064,25 @@ int hc_hipCtxPushCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx) return 0; } -int hc_hipCtxPopCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx) +int hc_hipInit (hashcat_ctx_t *hashcat_ctx, unsigned int Flags) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipCtxPopCurrent (pctx); + const hipError_t HIP_err = hip->hipInit (Flags); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { - event_log_error (hashcat_ctx, "hipCtxPopCurrent(): %s", pStr); + event_log_error (hashcat_ctx, "hipInit(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipCtxPopCurrent(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipInit(): %d", HIP_err); } return -1; @@ -3658,25 +3091,25 @@ int hc_hipCtxPopCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx) return 0; } -int hc_hipLinkCreate (hashcat_ctx_t *hashcat_ctx, unsigned int numOptions, HIPjit_option *options, void **optionValues, HIPlinkState *stateOut) +int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipLinkCreate (numOptions, options, optionValues, stateOut); + const hipError_t HIP_err = hip->hipMemAlloc (dptr, bytesize); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { - event_log_error (hashcat_ctx, "hipLinkCreate(): %s", pStr); + event_log_error (hashcat_ctx, "hipMemAlloc(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipLinkCreate(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipMemAlloc(): %d", HIP_err); } return -1; @@ -3685,25 +3118,25 @@ int hc_hipLinkCreate (hashcat_ctx_t *hashcat_ctx, unsigned int numOptions, HIPji return 0; } -int hc_hipLinkAddData (hashcat_ctx_t *hashcat_ctx, HIPlinkState state, HIPjitInputType type, void *data, size_t size, const char *name, unsigned int numOptions, HIPjit_option *options, void **optionValues) +int hc_hipMemFree (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dptr) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipLinkAddData (state, type, data, size, name, numOptions, options, optionValues); + const hipError_t HIP_err = hip->hipMemFree (dptr); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { - event_log_error (hashcat_ctx, "hipLinkAddData(): %s", pStr); + event_log_error (hashcat_ctx, "hipMemFree(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipLinkAddData(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipMemFree(): %d", HIP_err); } return -1; @@ -3712,25 +3145,25 @@ int hc_hipLinkAddData (hashcat_ctx_t *hashcat_ctx, HIPlinkState state, HIPjitInp return 0; } -int hc_hipLinkDestroy (hashcat_ctx_t *hashcat_ctx, HIPlinkState state) +int hc_hipMemGetInfo (hashcat_ctx_t *hashcat_ctx, size_t *free, size_t *total) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipLinkDestroy (state); + const hipError_t HIP_err = hip->hipMemGetInfo (free, total); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { - event_log_error (hashcat_ctx, "hipLinkDestroy(): %s", pStr); + event_log_error (hashcat_ctx, "hipMemGetInfo(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipLinkDestroy(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipMemGetInfo(): %d", HIP_err); } return -1; @@ -3739,25 +3172,349 @@ int hc_hipLinkDestroy (hashcat_ctx_t *hashcat_ctx, HIPlinkState state) return 0; } -int hc_hipLinkComplete (hashcat_ctx_t *hashcat_ctx, HIPlinkState state, void **hipbinOut, size_t *sizeOut) +int hc_hipMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount) { backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - const HIPresult HIP_err = hip->hipLinkComplete (state, hipbinOut, sizeOut); + const hipError_t HIP_err = hip->hipMemcpyDtoH (dstHost, srcDevice, ByteCount); - if (HIP_err != HIP_SUCCESS) + if (HIP_err != hipSuccess) { const char *pStr = NULL; - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { - event_log_error (hashcat_ctx, "hipLinkComplete(): %s", pStr); + event_log_error (hashcat_ctx, "hipMemcpyDtoH(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipLinkComplete(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipMemcpyDtoH(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemcpyDtoHAsync (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipMemcpyDtoHAsync (dstHost, srcDevice, ByteCount, hStream); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemcpyDtoHAsync(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyDtoHAsync(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipMemcpyDtoD (dstDevice, srcDevice, ByteCount); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemcpyDtoDAsync (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipMemcpyDtoDAsync (dstDevice, srcDevice, ByteCount, hStream); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemcpyDtoDAsync(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyDtoDAsync(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipMemcpyHtoD (dstDevice, srcHost, ByteCount); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemcpyHtoD(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyHtoD(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipMemcpyHtoDAsync (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount, hipStream_t hStream) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipMemcpyHtoDAsync (dstDevice, srcHost, ByteCount, hStream); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipMemcpyHtoDAsync(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipMemcpyHtoDAsync(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, hipFunction_t *hfunc, hipModule_t hmod, const char *name) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipModuleGetFunction (hfunc, hmod, name); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipModuleGetFunction(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipModuleGetFunction(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipModuleGetGlobal (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipModuleGetGlobal (dptr, bytes, hmod, name); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipModuleGetGlobal(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipModuleGetGlobal(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, hipModule_t *module, const void *image, unsigned int numOptions, hipJitOption *options, void **optionValues) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipModuleLoadDataEx (module, image, numOptions, options, optionValues); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipModuleLoadDataEx(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipModuleLoadDataEx(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipModuleUnload (hashcat_ctx_t *hashcat_ctx, hipModule_t hmod) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipModuleUnload (hmod); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipModuleUnload(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipModuleUnload(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, hipStream_t *phStream, unsigned int Flags) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipStreamCreate (phStream, Flags); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipStreamCreate(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipStreamCreate(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipStreamDestroy (hashcat_ctx_t *hashcat_ctx, hipStream_t hStream) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipStreamDestroy (hStream); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipStreamDestroy(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipStreamDestroy(): %d", HIP_err); + } + + return -1; + } + + return 0; +} + +int hc_hipStreamSynchronize (hashcat_ctx_t *hashcat_ctx, hipStream_t hStream) +{ + backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; + + HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; + + const hipError_t HIP_err = hip->hipStreamSynchronize (hStream); + + if (HIP_err != hipSuccess) + { + const char *pStr = NULL; + + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) + { + event_log_error (hashcat_ctx, "hipStreamSynchronize(): %s", pStr); + } + else + { + event_log_error (hashcat_ctx, "hipStreamSynchronize(): %d", HIP_err); } return -1; @@ -5245,7 +5002,7 @@ int run_cuda_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device return 0; } -int run_hip_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 num) +int run_hip_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 num) { u64 num_elements = num; @@ -5256,7 +5013,7 @@ int run_hip_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device num_elements = CEILDIV (num_elements, kernel_threads); - HIPfunction function = device_param->hip_function_atinit; + hipFunction_t function = device_param->hip_function_atinit; if (hc_hipLaunchKernel (hashcat_ctx, function, num_elements, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, device_param->kernel_params_atinit, NULL) == -1) return -1; @@ -5265,7 +5022,7 @@ int run_hip_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device return 0; } -int run_hip_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 num) +int run_hip_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 num) { u64 num_elements = num; @@ -5276,7 +5033,7 @@ int run_hip_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t num_elements = CEILDIV (num_elements, kernel_threads); - HIPfunction function = device_param->hip_function_utf8toutf16le; + hipFunction_t function = device_param->hip_function_utf8toutf16le; if (hc_hipLaunchKernel (hashcat_ctx, function, num_elements, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, device_param->kernel_params_utf8toutf16le, NULL) == -1) return -1; @@ -5285,7 +5042,7 @@ int run_hip_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t return 0; } -int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u32 value, const u64 size) +int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u32 value, const u64 size) { const u64 num16d = size / 16; const u64 num16m = size % 16; @@ -5301,7 +5058,7 @@ int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device u64 num_elements = CEILDIV (num16d, kernel_threads); - HIPfunction function = device_param->hip_function_memset; + hipFunction_t function = device_param->hip_function_memset; if (hc_hipLaunchKernel (hashcat_ctx, function, num_elements, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, device_param->kernel_params_memset, NULL) == -1) return -1; } @@ -5321,7 +5078,7 @@ int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device return 0; } -int run_hip_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, HIPdeviceptr buf, const u64 size) +int run_hip_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 size) { const u64 num16d = size / 16; const u64 num16m = size % 16; @@ -5336,7 +5093,7 @@ int run_hip_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_ u64 num_elements = CEILDIV(num16d, kernel_threads); - HIPfunction function = device_param->hip_function_bzero; + hipFunction_t function = device_param->hip_function_bzero; if (hc_hipLaunchKernel (hashcat_ctx, function, num_elements, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, device_param->kernel_params_bzero, NULL) == -1) return -1; } @@ -5686,7 +5443,7 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con if (device_param->is_hip == true) { - HIPfunction hip_function = NULL; + hipFunction_t hip_function = NULL; if (device_param->is_hip == true) { @@ -6014,7 +5771,7 @@ int run_kernel_mp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (device_param->is_hip == true) { - HIPfunction hip_function = NULL; + hipFunction_t hip_function = NULL; void **hip_args = NULL; @@ -6107,7 +5864,7 @@ int run_kernel_tm (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) if (device_param->is_hip == true) { - HIPfunction hip_function = device_param->hip_function_tm; + hipFunction_t hip_function = device_param->hip_function_tm; if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements / kernel_threads, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, device_param->kernel_params_tm, NULL) == -1) return -1; @@ -6154,7 +5911,7 @@ int run_kernel_amp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, { num_elements = CEILDIV (num_elements, kernel_threads); - HIPfunction hip_function = device_param->hip_function_amp; + hipFunction_t hip_function = device_param->hip_function_amp; if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, device_param->kernel_params_amp, NULL) == -1) return -1; @@ -6205,7 +5962,7 @@ int run_kernel_decompress (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device { num_elements = CEILDIV (num_elements, kernel_threads); - HIPfunction hip_function = device_param->hip_function_decompress; + hipFunction_t hip_function = device_param->hip_function_decompress; if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements, 1, 1, kernel_threads, 1, 1, 0, device_param->hip_stream, device_param->kernel_params_decompress, NULL) == -1) return -1; @@ -7407,38 +7164,17 @@ int backend_ctx_init (hashcat_ctx_t *hashcat_ctx) if ((rc_hip_init == 0) && (rc_hiprtc_init == 0)) { - // hiprtc version - - int hiprtc_major = 0; - int hiprtc_minor = 0; - - if (hc_hiprtcVersion (hashcat_ctx, &hiprtc_major, &hiprtc_minor) == -1) return -1; - - int hiprtc_driver_version = (hiprtc_major * 1000) + (hiprtc_minor * 10); - - backend_ctx->hiprtc_driver_version = hiprtc_driver_version; - - if (hiprtc_driver_version < 9000) - { - event_log_error (hashcat_ctx, "Outdated AMD HIPRTC driver version '%d' detected!", hiprtc_driver_version); - - event_log_warning (hashcat_ctx, "See hashcat.net for officially supported AMD HIP versions."); - event_log_warning (hashcat_ctx, NULL); - - return -1; - } - // hip version - int hip_driver_version = 10000; + int hip_driverVersion; - //if (hc_hipDriverGetVersion (hashcat_ctx, &hip_driver_version) == -1) return -1; + if (hc_hipDriverGetVersion (hashcat_ctx, &hip_driverVersion) == -1) return -1; - backend_ctx->hip_driver_version = hip_driver_version; + backend_ctx->hip_driverVersion = hip_driverVersion; - if (hip_driver_version < 9000) + if (hip_driverVersion < 404) { - event_log_error (hashcat_ctx, "Outdated AMD HIP driver version '%d' detected!", hip_driver_version); + event_log_error (hashcat_ctx, "Outdated AMD HIP driver version '%d' detected!", hip_driverVersion); event_log_warning (hashcat_ctx, "See hashcat.net for officially supported AMD HIP versions."); event_log_warning (hashcat_ctx, NULL); @@ -8289,7 +8025,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) backend_ctx->backend_device_from_hip[hip_devices_idx] = backend_devices_idx; - HIPdevice hip_device; + hipDevice_t hip_device; if (hc_hipDeviceGet (hashcat_ctx, &hip_device, hip_devices_idx) == -1) { @@ -8328,7 +8064,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) int device_processors = 0; - if (hc_hipDeviceGetAttribute (hashcat_ctx, &device_processors, HIP_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, hip_device) == -1) + if (hc_hipDeviceGetAttribute (hashcat_ctx, &device_processors, hipDeviceAttributeMultiprocessorCount, hip_device) == -1) { device_param->skipped = true; continue; @@ -8356,7 +8092,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) int hip_warp_size = 0; - if (hc_hipDeviceGetAttribute (hashcat_ctx, &hip_warp_size, HIP_DEVICE_ATTRIBUTE_WARP_SIZE, hip_device) == -1) + if (hc_hipDeviceGetAttribute (hashcat_ctx, &hip_warp_size, hipDeviceAttributeWarpSize, hip_device) == -1) { device_param->skipped = true; continue; @@ -8369,13 +8105,13 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) int sm_major = 0; int sm_minor = 0; - if (hc_hipDeviceGetAttribute (hashcat_ctx, &sm_major, HIP_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, hip_device) == -1) + if (hc_hipDeviceGetAttribute (hashcat_ctx, &sm_major, hipDeviceAttributeComputeCapabilityMajor, hip_device) == -1) { device_param->skipped = true; continue; } - if (hc_hipDeviceGetAttribute (hashcat_ctx, &sm_minor, HIP_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, hip_device) == -1) + if (hc_hipDeviceGetAttribute (hashcat_ctx, &sm_minor, hipDeviceAttributeComputeCapabilityMinor, hip_device) == -1) { device_param->skipped = true; continue; @@ -8388,7 +8124,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) int device_maxworkgroup_size = 0; - if (hc_hipDeviceGetAttribute (hashcat_ctx, &device_maxworkgroup_size, HIP_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, hip_device) == -1) + if (hc_hipDeviceGetAttribute (hashcat_ctx, &device_maxworkgroup_size, hipDeviceAttributeMaxThreadsPerBlock, hip_device) == -1) { device_param->skipped = true; continue; @@ -8400,7 +8136,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) int device_maxclock_frequency = 0; - if (hc_hipDeviceGetAttribute (hashcat_ctx, &device_maxclock_frequency, HIP_DEVICE_ATTRIBUTE_CLOCK_RATE, hip_device) == -1) + if (hc_hipDeviceGetAttribute (hashcat_ctx, &device_maxclock_frequency, hipDeviceAttributeClockRate, hip_device) == -1) { device_param->skipped = true; continue; @@ -8414,19 +8150,20 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) int pci_bus_id_nv = 0; int pci_slot_id_nv = 0; - if (hc_hipDeviceGetAttribute (hashcat_ctx, &pci_domain_id_nv, HIP_DEVICE_ATTRIBUTE_PCI_DOMAIN_ID, hip_device) == -1) + // Not supported by HIP + //if (hc_hipDeviceGetAttribute (hashcat_ctx, &pci_domain_id_nv, hipDeviceAttributePciDomainID, hip_device) == -1) + //{ + // device_param->skipped = true; + // continue; + //} + + if (hc_hipDeviceGetAttribute (hashcat_ctx, &pci_bus_id_nv, hipDeviceAttributePciBusId, hip_device) == -1) { device_param->skipped = true; continue; } - if (hc_hipDeviceGetAttribute (hashcat_ctx, &pci_bus_id_nv, HIP_DEVICE_ATTRIBUTE_PCI_BUS_ID, hip_device) == -1) - { - device_param->skipped = true; - continue; - } - - if (hc_hipDeviceGetAttribute (hashcat_ctx, &pci_slot_id_nv, HIP_DEVICE_ATTRIBUTE_PCI_DEVICE_ID, hip_device) == -1) + if (hc_hipDeviceGetAttribute (hashcat_ctx, &pci_slot_id_nv, hipDeviceAttributePciDeviceId, hip_device) == -1) { device_param->skipped = true; continue; @@ -8434,6 +8171,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) device_param->pcie_domain = (u8) (pci_domain_id_nv); device_param->pcie_bus = (u8) (pci_bus_id_nv); + device_param->pcie_device = (u8) (pci_slot_id_nv >> 3); device_param->pcie_function = (u8) (pci_slot_id_nv & 7); @@ -8441,7 +8179,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) int kernel_exec_timeout = 0; - if (hc_hipDeviceGetAttribute (hashcat_ctx, &kernel_exec_timeout, HIP_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT, hip_device) == -1) + if (hc_hipDeviceGetAttribute (hashcat_ctx, &kernel_exec_timeout, hipDeviceAttributeKernelExecTimeout, hip_device) == -1) { device_param->skipped = true; continue; @@ -8453,7 +8191,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) int warp_size = 0; - if (hc_hipDeviceGetAttribute (hashcat_ctx, &warp_size, HIP_DEVICE_ATTRIBUTE_WARP_SIZE, hip_device) == -1) + if (hc_hipDeviceGetAttribute (hashcat_ctx, &warp_size, hipDeviceAttributeWarpSize, hip_device) == -1) { device_param->skipped = true; continue; @@ -8465,7 +8203,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) int max_shared_memory_per_block = 0; - if (hc_hipDeviceGetAttribute (hashcat_ctx, &max_shared_memory_per_block, HIP_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK_OPTIN, hip_device) == -1) + if (hc_hipDeviceGetAttribute (hashcat_ctx, &max_shared_memory_per_block, hipDeviceAttributeMaxSharedMemoryPerBlock, hip_device) == -1) { device_param->skipped = true; continue; @@ -8484,7 +8222,7 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) int device_max_constant_buffer_size = 0; - if (hc_hipDeviceGetAttribute (hashcat_ctx, &device_max_constant_buffer_size, HIP_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY, hip_device) == -1) + if (hc_hipDeviceGetAttribute (hashcat_ctx, &device_max_constant_buffer_size, hipDeviceAttributeTotalConstantMemory, hip_device) == -1) { device_param->skipped = true; continue; @@ -8577,13 +8315,6 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) // instruction set - // bcrypt optimization? - //const int rc_cuCtxSetCacheConfig = hc_hipCtxSetCacheConfig (hashcat_ctx, HIP_FUNC_CACHE_PREFER_SHARED); - // - //if (rc_cuCtxSetCacheConfig == -1) return -1; - - // const int sm = (device_param->sm_major * 10) + device_param->sm_minor; - device_param->has_add = false; device_param->has_addc = false; device_param->has_sub = false; @@ -8595,9 +8326,9 @@ int backend_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) // device_available_mem - HIPcontext hip_context; + hipCtx_t hip_context; - if (hc_hipCtxCreate (hashcat_ctx, &hip_context, HIP_CTX_SCHED_BLOCKING_SYNC, device_param->hip_device) == -1) + if (hc_hipCtxCreate (hashcat_ctx, &hip_context, hipDeviceScheduleBlockingSync, device_param->hip_device) == -1) { device_param->skipped = true; continue; @@ -10280,7 +10011,7 @@ static int get_cuda_kernel_local_mem_size (hashcat_ctx_t *hashcat_ctx, CUfunctio return 0; } -static int get_hip_kernel_wgs (hashcat_ctx_t *hashcat_ctx, HIPfunction function, u32 *result) +static int get_hip_kernel_wgs (hashcat_ctx_t *hashcat_ctx, hipFunction_t function, u32 *result) { int max_threads_per_block; @@ -10291,7 +10022,7 @@ static int get_hip_kernel_wgs (hashcat_ctx_t *hashcat_ctx, HIPfunction function, return 0; } -static int get_hip_kernel_local_mem_size (hashcat_ctx_t *hashcat_ctx, HIPfunction function, u64 *result) +static int get_hip_kernel_local_mem_size (hashcat_ctx_t *hashcat_ctx, hipFunction_t function, u64 *result) { int shared_size_bytes; @@ -10410,7 +10141,7 @@ static u32 get_kernel_threads (const hc_device_param_t *device_param) return kernel_threads; } -static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const char *kernel_name, char *source_file, char *cached_file, const char *build_options_buf, const bool cache_disable, cl_program *opencl_program, CUmodule *cuda_module, HIPmodule *hip_module) +static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const char *kernel_name, char *source_file, char *cached_file, const char *build_options_buf, const bool cache_disable, cl_program *opencl_program, CUmodule *cuda_module, hipModule_t *hip_module) { const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; const folder_config_t *folder_config = hashcat_ctx->folder_config; @@ -10698,18 +10429,12 @@ static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_p char **hiprtc_options = (char **) hccalloc (7 + strlen (build_options_buf) + 1, sizeof (char *)); // ... - //hiprtc_options[0] = "--restrict"; - //hiprtc_options[1] = "--device-as-default-execution-space"; - //hiprtc_options[2] = "--gpu-architecture"; - hc_asprintf (&hiprtc_options[0], "--gpu-max-threads-per-block=%d", (user_options->kernel_threads_chgd == true) ? user_options->kernel_threads : device_param->kernel_preferred_wgs_multiple); - //hiprtc_options[0] = "--gpu-max-threads-per-block=64"; hiprtc_options[1] = "-nocudainc"; hiprtc_options[2] = "-nocudalib"; hiprtc_options[3] = ""; hiprtc_options[4] = ""; - hiprtc_options[5] = "-I"; hiprtc_options[6] = folder_config->cpath_real; @@ -10772,137 +10497,27 @@ static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_p int mod_cnt = 6; - HIPjit_option mod_opts[7]; - void *mod_vals[7]; + hipJitOption mod_opts[6]; + void *mod_vals[6]; - mod_opts[0] = HIP_JIT_TARGET_FROM_HIPCONTEXT; + mod_opts[0] = hipJitOptionTargetFromContext; mod_vals[0] = (void *) 0; - mod_opts[1] = HIP_JIT_LOG_VERBOSE; + mod_opts[1] = hipJitOptionLogVerbose; mod_vals[1] = (void *) 1; - mod_opts[2] = HIP_JIT_INFO_LOG_BUFFER; + mod_opts[2] = hipJitOptionInfoLogBuffer; mod_vals[2] = (void *) mod_info_log; - mod_opts[3] = HIP_JIT_INFO_LOG_BUFFER_SIZE_BYTES; + mod_opts[3] = hipJitOptionInfoLogBufferSizeBytes; mod_vals[3] = (void *) LOG_SIZE; - mod_opts[4] = HIP_JIT_ERROR_LOG_BUFFER; + mod_opts[4] = hipJitOptionErrorLogBuffer; mod_vals[4] = (void *) mod_error_log; - mod_opts[5] = HIP_JIT_ERROR_LOG_BUFFER_SIZE_BYTES; + mod_opts[5] = hipJitOptionErrorLogBufferSizeBytes; mod_vals[5] = (void *) LOG_SIZE; - if (hashconfig->opti_type & OPTI_TYPE_REGISTER_LIMIT) - { - mod_opts[6] = HIP_JIT_MAX_REGISTERS; - mod_vals[6] = (void *) 128; - - mod_cnt++; - } - - #if defined (WITH_HIPBIN) - - char *jit_info_log = (char *) hcmalloc (LOG_SIZE + 1); - char *jit_error_log = (char *) hcmalloc (LOG_SIZE + 1); - - int jit_cnt = 6; - - HIPjit_option jit_opts[7]; - void *jit_vals[7]; - - jit_opts[0] = HIP_JIT_TARGET_FROM_HIPCONTEXT; - jit_vals[0] = (void *) 0; - - jit_opts[1] = HIP_JIT_LOG_VERBOSE; - jit_vals[1] = (void *) 1; - - jit_opts[2] = HIP_JIT_INFO_LOG_BUFFER; - jit_vals[2] = (void *) jit_info_log; - - jit_opts[3] = HIP_JIT_INFO_LOG_BUFFER_SIZE_BYTES; - jit_vals[3] = (void *) LOG_SIZE; - - jit_opts[4] = HIP_JIT_ERROR_LOG_BUFFER; - jit_vals[4] = (void *) jit_error_log; - - jit_opts[5] = HIP_JIT_ERROR_LOG_BUFFER_SIZE_BYTES; - jit_vals[5] = (void *) LOG_SIZE; - - if (hashconfig->opti_type & OPTI_TYPE_REGISTER_LIMIT) - { - jit_opts[6] = HIP_JIT_MAX_REGISTERS; - jit_vals[6] = (void *) 128; - - jit_cnt++; - } - - HIPlinkState state; - - if (hc_cuLinkCreate (hashcat_ctx, jit_cnt, jit_opts, jit_vals, &state) == -1) - { - event_log_error (hashcat_ctx, "* Device #%u: Kernel %s link failed. Error Log:", device_param->device_id + 1, source_file); - event_log_error (hashcat_ctx, "%s", jit_error_log); - event_log_error (hashcat_ctx, NULL); - - return false; - } - - if (hc_cuLinkAddData (hashcat_ctx, state, HIP_JIT_INPUT_CODE, binary, binary_size, kernel_name, 0, NULL, NULL) == -1) - { - event_log_error (hashcat_ctx, "* Device #%u: Kernel %s link failed. Error Log:", device_param->device_id + 1, source_file); - event_log_error (hashcat_ctx, "%s", jit_error_log); - event_log_error (hashcat_ctx, NULL); - - return false; - } - - void *cubin = NULL; - - size_t cubin_size = 0; - - if (hc_cuLinkComplete (hashcat_ctx, state, &cubin, &cubin_size) == -1) - { - event_log_error (hashcat_ctx, "* Device #%u: Kernel %s link failed. Error Log:", device_param->device_id + 1, source_file); - event_log_error (hashcat_ctx, "%s", jit_error_log); - event_log_error (hashcat_ctx, NULL); - - return false; - } - - #if defined (DEBUG) - event_log_info (hashcat_ctx, "* Device #%u: Kernel %s link successful. Info Log:", device_param->device_id + 1, source_file); - event_log_info (hashcat_ctx, "%s", jit_info_log); - event_log_info (hashcat_ctx, NULL); - #endif - - if (hc_cuModuleLoadDataEx (hashcat_ctx, hip_module, cubin, mod_cnt, mod_opts, mod_vals) == -1) - { - event_log_error (hashcat_ctx, "* Device #%u: Kernel %s load failed. Error Log:", device_param->device_id + 1, source_file); - event_log_error (hashcat_ctx, "%s", mod_error_log); - event_log_error (hashcat_ctx, NULL); - - return false; - } - - #if defined (DEBUG) - event_log_info (hashcat_ctx, "* Device #%u: Kernel %s load successful. Info Log:", device_param->device_id + 1, source_file); - event_log_info (hashcat_ctx, "%s", mod_info_log); - event_log_info (hashcat_ctx, NULL); - #endif - - if (cache_disable == false) - { - if (write_kernel_binary (hashcat_ctx, cached_file, cubin, cubin_size) == false) return false; - } - - if (hc_hipLinkDestroy (hashcat_ctx, state) == -1) return false; - - hcfree (jit_info_log); - hcfree (jit_error_log); - - #else - if (hc_hipModuleLoadDataEx (hashcat_ctx, hip_module, binary, mod_cnt, mod_opts, mod_vals) == -1) { event_log_error (hashcat_ctx, "* Device #%u: Kernel %s load failed. Error Log:", device_param->device_id + 1, source_file); @@ -10923,8 +10538,6 @@ static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_p if (write_kernel_binary (hashcat_ctx, cached_file, binary, binary_size) == false) return false; } - #endif - hcfree (mod_info_log); hcfree (mod_error_log); @@ -11072,35 +10685,27 @@ static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_p int mod_cnt = 6; - HIPjit_option mod_opts[7]; - void *mod_vals[7]; + hipJitOption mod_opts[6]; + void *mod_vals[6]; - mod_opts[0] = HIP_JIT_TARGET_FROM_HIPCONTEXT; + mod_opts[0] = hipJitOptionTargetFromContext; mod_vals[0] = (void *) 0; - mod_opts[1] = HIP_JIT_LOG_VERBOSE; + mod_opts[1] = hipJitOptionLogVerbose; mod_vals[1] = (void *) 1; - mod_opts[2] = HIP_JIT_INFO_LOG_BUFFER; + mod_opts[2] = hipJitOptionInfoLogBuffer; mod_vals[2] = (void *) mod_info_log; - mod_opts[3] = HIP_JIT_INFO_LOG_BUFFER_SIZE_BYTES; + mod_opts[3] = hipJitOptionInfoLogBufferSizeBytes; mod_vals[3] = (void *) LOG_SIZE; - mod_opts[4] = HIP_JIT_ERROR_LOG_BUFFER; + mod_opts[4] = hipJitOptionErrorLogBuffer; mod_vals[4] = (void *) mod_error_log; - mod_opts[5] = HIP_JIT_ERROR_LOG_BUFFER_SIZE_BYTES; + mod_opts[5] = hipJitOptionErrorLogBufferSizeBytes; mod_vals[5] = (void *) LOG_SIZE; - if (hashconfig->opti_type & OPTI_TYPE_REGISTER_LIMIT) - { - mod_opts[6] = HIP_JIT_MAX_REGISTERS; - mod_vals[6] = (void *) 128; - - mod_cnt++; - } - if (hc_hipModuleLoadDataEx (hashcat_ctx, hip_module, kernel_sources[0], mod_cnt, mod_opts, mod_vals) == -1) { event_log_error (hashcat_ctx, "* Device #%u: Kernel %s load failed. Error Log:", device_param->device_id + 1, source_file); @@ -11491,7 +11096,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (device_param->is_hip == true) { - if (hc_hipCtxCreate (hashcat_ctx, &device_param->hip_context, HIP_CTX_SCHED_BLOCKING_SYNC, device_param->hip_device) == -1) + if (hc_hipCtxCreate (hashcat_ctx, &device_param->hip_context, hipDeviceScheduleBlockingSync, device_param->hip_device) == -1) { device_param->skipped = true; continue; @@ -11555,7 +11160,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (device_param->is_hip == true) { - if (hc_hipStreamCreate (hashcat_ctx, &device_param->hip_stream, HIP_STREAM_DEFAULT) == -1) + if (hc_hipStreamCreate (hashcat_ctx, &device_param->hip_stream, hipStreamDefault) == -1) { device_param->skipped = true; continue; @@ -11587,13 +11192,13 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) if (device_param->is_hip == true) { - if (hc_hipEventCreate (hashcat_ctx, &device_param->hip_event1, HIP_EVENT_BLOCKING_SYNC) == -1) + if (hc_hipEventCreate (hashcat_ctx, &device_param->hip_event1, hipEventBlockingSync) == -1) { device_param->skipped = true; continue; } - if (hc_hipEventCreate (hashcat_ctx, &device_param->hip_event2, HIP_EVENT_BLOCKING_SYNC) == -1) + if (hc_hipEventCreate (hashcat_ctx, &device_param->hip_event2, hipEventBlockingSync) == -1) { device_param->skipped = true; continue; @@ -11781,7 +11386,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) const size_t dnclen_amp_mp = snprintf (device_name_chksum_amp_mp, HCBUFSIZ_TINY, "%d-%d-%d-%d-%u-%s-%s-%s-%u", backend_ctx->comptime, backend_ctx->cuda_driver_version, - backend_ctx->hip_driver_version, + backend_ctx->hip_driverVersion, device_param->is_opencl, device_param->opencl_platform_vendor_id, device_param->device_name, @@ -12109,7 +11714,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) const size_t dnclen = snprintf (device_name_chksum, HCBUFSIZ_TINY, "%d-%d-%d-%d-%u-%s-%s-%s-%d-%u-%u-%u-%s", backend_ctx->comptime, backend_ctx->cuda_driver_version, - backend_ctx->hip_driver_version, + backend_ctx->hip_driverVersion, device_param->is_opencl, device_param->opencl_platform_vendor_id, device_param->device_name, diff --git a/src/terminal.c b/src/terminal.c index 167b8b821..397eace9f 100644 --- a/src/terminal.c +++ b/src/terminal.c @@ -818,9 +818,9 @@ void backend_info (hashcat_ctx_t *hashcat_ctx) event_log_info (hashcat_ctx, NULL); int hip_devices_cnt = backend_ctx->hip_devices_cnt; - int hip_driver_version = backend_ctx->hip_driver_version; + int hip_driverVersion = backend_ctx->hip_driverVersion; - event_log_info (hashcat_ctx, "HIP.Version.: %d.%d", hip_driver_version / 1000, (hip_driver_version % 100) / 10); + event_log_info (hashcat_ctx, "HIP.Version.: %d.%d", hip_driverVersion / 100, hip_driverVersion % 10); event_log_info (hashcat_ctx, NULL); for (int hip_devices_idx = 0; hip_devices_idx < hip_devices_cnt; hip_devices_idx++) @@ -1014,10 +1014,10 @@ void backend_info_compact (hashcat_ctx_t *hashcat_ctx) if (backend_ctx->hip) { - int hip_devices_cnt = backend_ctx->hip_devices_cnt; - int hip_driver_version = backend_ctx->hip_driver_version; + int hip_devices_cnt = backend_ctx->hip_devices_cnt; + int hip_driverVersion = backend_ctx->hip_driverVersion; - const size_t len = event_log_info (hashcat_ctx, "HIP API (HIP %d.%d)", hip_driver_version / 1000, (hip_driver_version % 100) / 10); + const size_t len = event_log_info (hashcat_ctx, "HIP API (HIP %d.%d)", hip_driverVersion / 100, hip_driverVersion % 10); char line[HCBUFSIZ_TINY] = { 0 }; From 3f6c5a0042c3c8cd7e1c95708761b1998ad7279f Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Fri, 23 Jul 2021 21:09:55 +0200 Subject: [PATCH 2/4] Update module_unstable_warning() for -m 172xx on HIP --- OpenCL/inc_platform.cl | 13 +++++++------ OpenCL/inc_types.h | 2 +- OpenCL/inc_vendor.h | 4 ---- OpenCL/inc_zip_inflate.cl | 4 +--- src/modules/module_17200.c | 5 ----- src/modules/module_17220.c | 5 ----- src/modules/module_17225.c | 5 ----- 7 files changed, 9 insertions(+), 29 deletions(-) diff --git a/OpenCL/inc_platform.cl b/OpenCL/inc_platform.cl index 47cfc9a84..8ccb034aa 100644 --- a/OpenCL/inc_platform.cl +++ b/OpenCL/inc_platform.cl @@ -208,19 +208,20 @@ DECLSPEC u32 hc_atomic_or (GLOBAL_AS u32 *p, volatile const u32 val) return atomicOr (p, val); } -DECLSPEC size_t get_global_id (const u32 dimindx) +DECLSPEC size_t get_global_id (const u32 dimindx __attribute__((unused))) { - return (__ockl_get_group_id (dimindx) * __ockl_get_local_size (dimindx)) + __ockl_get_local_id (dimindx); + return (blockIdx.x * blockDim.x) + threadIdx.x; } -DECLSPEC size_t get_local_id (const u32 dimindx) +DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused))) { - return __ockl_get_local_id (dimindx); + return threadIdx.x; } -DECLSPEC size_t get_local_size (const u32 dimindx) +DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused))) { - return __ockl_get_local_size (dimindx); + // verify + return blockDim.x; } DECLSPEC u32x rotl32 (const u32x a, const int n) diff --git a/OpenCL/inc_types.h b/OpenCL/inc_types.h index 4caf595de..295caabc0 100644 --- a/OpenCL/inc_types.h +++ b/OpenCL/inc_types.h @@ -833,7 +833,7 @@ typedef __device_builtin__ struct u64x u64x; typedef VTYPE(uchar, VECT_SIZE) u8x; typedef VTYPE(ushort, VECT_SIZE) u16x; typedef VTYPE(uint, VECT_SIZE) u32x; -typedef VTYPE(ulong, VECT_SIZE) u64x; +typedef VTYPE(ullong, VECT_SIZE) u64x; #define make_u8x (u8x) #define make_u16x (u16x) diff --git a/OpenCL/inc_vendor.h b/OpenCL/inc_vendor.h index d44ac87c9..bbd7e23d8 100644 --- a/OpenCL/inc_vendor.h +++ b/OpenCL/inc_vendor.h @@ -32,10 +32,6 @@ #define LOCAL_AS #define KERNEL_FQ extern "C" __global__ #elif defined IS_HIP -#define __device__ __attribute__((device)) -#define __constant__ __attribute__((constant)) -#define __shared__ __attribute__((shared)) -#define __global__ __attribute__((global)) #define CONSTANT_VK __constant__ #define CONSTANT_AS #define GLOBAL_AS diff --git a/OpenCL/inc_zip_inflate.cl b/OpenCL/inc_zip_inflate.cl index d05f6a792..fcc31e76b 100644 --- a/OpenCL/inc_zip_inflate.cl +++ b/OpenCL/inc_zip_inflate.cl @@ -73,18 +73,16 @@ enum{ MZ_VERSION_ERROR = -6, MZ_PARAM_ERROR = -10000 }; -typedef unsigned long mz_ulong; +typedef ullong mz_ulong; #ifndef MINIZ_NO_ZLIB_COMPATIBLE_NAMES typedef unsigned char Byte; typedef unsigned int uInt; -typedef mz_ulong uLong; typedef Byte Bytef; typedef uInt uIntf; typedef char charf; typedef int intf; typedef void *voidpf; -typedef uLong uLongf; typedef void *voidp; typedef void *const voidpc; #define Z_NULL 0 diff --git a/src/modules/module_17200.c b/src/modules/module_17200.c index 44798110d..fa510ba27 100644 --- a/src/modules/module_17200.c +++ b/src/modules/module_17200.c @@ -170,11 +170,6 @@ bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE // it leads to CL_KERNEL_WORK_GROUP_SIZE to return 0 and later we will divide with 0 // workaround would be to rewrite kernel to use global memory - if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP) - { - return true; - } - if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD) { return true; diff --git a/src/modules/module_17220.c b/src/modules/module_17220.c index 9028040d9..8f1beaf1a 100644 --- a/src/modules/module_17220.c +++ b/src/modules/module_17220.c @@ -170,11 +170,6 @@ bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE // it leads to CL_KERNEL_WORK_GROUP_SIZE to return 0 and later we will divide with 0 // workaround would be to rewrite kernel to use global memory - if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP) - { - return true; - } - if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD) { return true; diff --git a/src/modules/module_17225.c b/src/modules/module_17225.c index 75c376c9e..3b3291d5f 100644 --- a/src/modules/module_17225.c +++ b/src/modules/module_17225.c @@ -170,11 +170,6 @@ bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE // it leads to CL_KERNEL_WORK_GROUP_SIZE to return 0 and later we will divide with 0 // workaround would be to rewrite kernel to use global memory - if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP) - { - return true; - } - if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD) { return true; From c9d79374a59bc4e2eed0620c59d11f9b07ab7ca4 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Tue, 27 Jul 2021 09:54:15 +0200 Subject: [PATCH 3/4] Add missing u64 vector datatypes mapping for OpenCL Improve u64 usage in OpenCL/inc_zip_inflate.cl --- OpenCL/inc_types.h | 6 +++++- OpenCL/inc_zip_inflate.cl | 28 +++++++++++++++------------- 2 files changed, 20 insertions(+), 14 deletions(-) diff --git a/OpenCL/inc_types.h b/OpenCL/inc_types.h index 295caabc0..9611ec05d 100644 --- a/OpenCL/inc_types.h +++ b/OpenCL/inc_types.h @@ -26,7 +26,11 @@ typedef unsigned long long ullong; #endif #ifdef IS_OPENCL -typedef ulong ullong; +typedef ulong ullong; +typedef ulong2 ullong2; +typedef ulong4 ullong4; +typedef ulong8 ullong8; +typedef ulong16 ullong16; #endif #ifdef KERNEL_STATIC diff --git a/OpenCL/inc_zip_inflate.cl b/OpenCL/inc_zip_inflate.cl index fcc31e76b..00f762d81 100644 --- a/OpenCL/inc_zip_inflate.cl +++ b/OpenCL/inc_zip_inflate.cl @@ -202,10 +202,6 @@ DECLSPEC void *memset(u8 *s, int c, u32 len){ #define MZ_MIN(a, b) (((a) < (b)) ? (a) : (b)) #define MZ_DEFAULT_WINDOW_BITS 15 #define TINFL_LZ_DICT_SIZE 32768 -#define TINFL_MEMCPY(d, s, l) memcpy(d, s, l) -#define TINFL_MEMCPY_G(d, s, l, p) memcpy_g(d, s, l, p) -#define TINFL_MEMSET(p, c, l) memset(p, c, (u32)l) -#define MZ_CLEAR_OBJ(obj) memset(&(obj), 0, sizeof(obj)) // hashcat-patched/hashcat-specific: #ifdef CRC32_IN_INFLATE @@ -581,7 +577,7 @@ DECLSPEC tinfl_status tinfl_decompress(tinfl_decompressor *r, MAYBE_GLOBAL const TINFL_CR_RETURN(38, (decomp_flags & TINFL_FLAG_HAS_MORE_INPUT) ? TINFL_STATUS_NEEDS_MORE_INPUT : TINFL_STATUS_FAILED_CANNOT_MAKE_PROGRESS); } n = MZ_MIN(MZ_MIN((size_t)(pOut_buf_end - pOut_buf_cur), (size_t)(pIn_buf_end - pIn_buf_cur)), counter); - TINFL_MEMCPY_G(pOut_buf_cur, pIn_buf_cur, n, pStream); + memcpy_g(pOut_buf_cur, pIn_buf_cur, n, pStream); pIn_buf_cur += n; pOut_buf_cur += n; counter -= (mz_uint)n; @@ -599,7 +595,7 @@ DECLSPEC tinfl_status tinfl_decompress(tinfl_decompressor *r, MAYBE_GLOBAL const mz_uint i; r->m_table_sizes[0] = 288; r->m_table_sizes[1] = 32; - TINFL_MEMSET(r->m_tables[1].m_code_size, 5, 32); + memset(r->m_tables[1].m_code_size, 5, 32); for (i = 0; i <= 143; ++i) *p++ = 8; for (; i <= 255; ++i) @@ -616,7 +612,8 @@ DECLSPEC tinfl_status tinfl_decompress(tinfl_decompressor *r, MAYBE_GLOBAL const TINFL_GET_BITS(11, r->m_table_sizes[counter], "\05\05\04"[counter]); r->m_table_sizes[counter] += s_min_table_sizes[counter]; } - MZ_CLEAR_OBJ(r->m_tables[2].m_code_size); + memset(r->m_tables[2].m_code_size, 0, TINFL_MAX_HUFF_SYMBOLS_0); + for (counter = 0; counter < r->m_table_sizes[2]; counter++) { mz_uint s; @@ -631,9 +628,11 @@ DECLSPEC tinfl_status tinfl_decompress(tinfl_decompressor *r, MAYBE_GLOBAL const tinfl_huff_table *pTable; mz_uint i, j, used_syms, total, sym_index, next_code[17], total_syms[16]; pTable = &r->m_tables[r->m_type]; - MZ_CLEAR_OBJ(total_syms); - MZ_CLEAR_OBJ(pTable->m_look_up); - MZ_CLEAR_OBJ(pTable->m_tree); + + memset((u8 *) total_syms, 0, 64); + memset((u8 *) pTable->m_look_up, 0, TINFL_FAST_LOOKUP_SIZE * 2); + memset((u8 *) pTable->m_tree, 0, TINFL_MAX_HUFF_SYMBOLS_0 * 2 * 2); + for (i = 0; i < r->m_table_sizes[r->m_type]; ++i) total_syms[pTable->m_code_size[i]]++; used_syms = 0, total = 0; @@ -705,15 +704,18 @@ DECLSPEC tinfl_status tinfl_decompress(tinfl_decompressor *r, MAYBE_GLOBAL const num_extra = "\02\03\07"[dist - 16]; TINFL_GET_BITS(18, s, num_extra); s += "\03\03\013"[dist - 16]; - TINFL_MEMSET(r->m_len_codes + counter, (dist == 16) ? r->m_len_codes[counter - 1] : 0, s); + + memset(r->m_len_codes + counter, (dist == 16) ? r->m_len_codes[counter - 1] : 0, s); + + counter += s; } if ((r->m_table_sizes[0] + r->m_table_sizes[1]) != counter) { TINFL_CR_RETURN_FOREVER(21, TINFL_STATUS_FAILED); } - TINFL_MEMCPY(r->m_tables[0].m_code_size, r->m_len_codes, r->m_table_sizes[0]); - TINFL_MEMCPY(r->m_tables[1].m_code_size, r->m_len_codes + r->m_table_sizes[0], r->m_table_sizes[1]); + memcpy(r->m_tables[0].m_code_size, r->m_len_codes, r->m_table_sizes[0]); + memcpy(r->m_tables[1].m_code_size, r->m_len_codes + r->m_table_sizes[0], r->m_table_sizes[1]); } } for (;;) From 9c0a37accf9709dfb4613d1b7d29fb62743fa253 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Wed, 28 Jul 2021 07:56:34 +0200 Subject: [PATCH 4/4] Update driver requirement --- docs/readme.txt | 4 ++-- src/backend.c | 8 ++++---- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/docs/readme.txt b/docs/readme.txt index ca10a4644..d70edfcb7 100644 --- a/docs/readme.txt +++ b/docs/readme.txt @@ -10,8 +10,8 @@ hashcat v6.2.3 ============== -AMD GPUs on Linux require "RadeonOpenCompute (ROCm)" Software Platform (3.1 or later) -AMD GPUs on Windows require "AMD Radeon Adrenalin 2020 Edition" (20.2.2 or later) +AMD GPUs on Linux require "AMD ROCm" (4.4 or later) +AMD GPUs on Windows require "AMD Radeon Adrenalin 2020 Edition" (21.2.1 or later) Intel CPUs require "OpenCL Runtime for Intel Core and Intel Xeon Processors" (16.1.1 or later) NVIDIA GPUs require "NVIDIA Driver" (440.64 or later) and "CUDA Toolkit" (9.0 or later) diff --git a/src/backend.c b/src/backend.c index cddd8cb5d..eef57d412 100644 --- a/src/backend.c +++ b/src/backend.c @@ -7224,10 +7224,10 @@ int backend_ctx_init (hashcat_ctx_t *hashcat_ctx) #if defined (__linux__) event_log_warning (hashcat_ctx, "* AMD GPUs on Linux require this driver:"); - event_log_warning (hashcat_ctx, " \"RadeonOpenCompute (ROCm)\" Software Platform (3.1 or later)"); + event_log_warning (hashcat_ctx, " \"AMD ROCm\" (4.4 or later)"); #elif defined (_WIN) event_log_warning (hashcat_ctx, "* AMD GPUs on Windows require this driver:"); - event_log_warning (hashcat_ctx, " \"AMD Radeon Adrenalin 2020 Edition\" (20.2.2 or later)"); + event_log_warning (hashcat_ctx, " \"AMD Radeon Adrenalin 2020 Edition\" (21.2.1 or later)"); #endif event_log_warning (hashcat_ctx, "* Intel CPUs require this runtime:"); @@ -7549,10 +7549,10 @@ int backend_ctx_init (hashcat_ctx_t *hashcat_ctx) #if defined (__linux__) event_log_warning (hashcat_ctx, "* AMD GPUs on Linux require this driver:"); - event_log_warning (hashcat_ctx, " \"RadeonOpenCompute (ROCm)\" Software Platform (3.1 or later)"); + event_log_warning (hashcat_ctx, " \"AMD ROCm\" (4.4 or later)"); #elif defined (_WIN) event_log_warning (hashcat_ctx, "* AMD GPUs on Windows require this driver:"); - event_log_warning (hashcat_ctx, " \"AMD Radeon Adrenalin 2020 Edition\" (20.2.2 or later)"); + event_log_warning (hashcat_ctx, " \"AMD Radeon Adrenalin 2020 Edition\" (21.2.1 or later)"); #endif event_log_warning (hashcat_ctx, "* Intel CPUs require this runtime:");