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); -} - -extern "C" __device__ __attribute__((pure)) double __ocml_log2_f64(double); - -DECLSPEC double log2 (double x) -{ - return __ocml_log2_f64 (x); + return atomicOr (p, val); } -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_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_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, 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 - */ - -#define __HIP_API_VERSION 4221131 - -/** - * 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 - -#if defined(_WIN64) || defined(__LP64__) -typedef unsigned long long HIPdeviceptr; -#else -typedef unsigned int HIPdeviceptr; -#endif - -#endif /* __HIP_API_VERSION >= 3020 */ - -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, - +// 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 + +// start: amd_driver_types.h + +typedef void* hipDeviceptr_t; + +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; - -/** - * 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; +}hipFunction_attribute; -typedef enum HIPjitInputType_enum -{ - /** - * Compiled device-class-specific device code\n - * Applicable options: none - */ - HIP_JIT_INPUT_HIPBIN = 0, +// stop: amd_driver_types.h - /** - * PTX source code\n - * Applicable options: PTX compiler options - */ - HIP_JIT_INPUT_PTX, +// start: hip_runtime_api.h - /** - * 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, +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; - /** - * Host object with embedded device code\n - * Applicable options: PTX compiler options, ::HIP_JIT_FALLBACK_STRATEGY - */ - HIP_JIT_INPUT_OBJECT, +// 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 +#define __HIP_NODISCARD +#endif - /** - * Archive of host objects with embedded device code\n - * Applicable options: PTX compiler options, ::HIP_JIT_FALLBACK_STRATEGY - */ - HIP_JIT_INPUT_LIBRARY, +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 - 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; - - HIPevent hip_event1; - HIPevent hip_event2; - - HIPmodule hip_module; - HIPmodule hip_module_shared; - HIPmodule hip_module_mp; - HIPmodule 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; - - 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; + hipDevice_t hip_device; + hipCtx_t hip_context; + hipStream_t hip_stream; + + hipEvent_t hip_event1; + hipEvent_t hip_event2; + + hipModule_t hip_module; + hipModule_t hip_module_shared; + hipModule_t hip_module_mp; + hipModule_t hip_module_amp; + + 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_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,208 +2551,19 @@ 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) - { - event_log_error (hashcat_ctx, "hipInit(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipInit(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipDeviceGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPdevice_attribute attrib, HIPdevice dev) -{ - 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); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipDeviceGetAttribute(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipDeviceGetAttribute(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipDeviceGetCount (hashcat_ctx_t *hashcat_ctx, int *count) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipDeviceGetCount (count); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipDeviceGetCount(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipDeviceGetCount(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipDeviceGet (hashcat_ctx_t *hashcat_ctx, HIPdevice* device, int ordinal) -{ - 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); - - if (HIP_err != HIP_SUCCESS) - { - 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) - { - event_log_error (hashcat_ctx, "hipDeviceGetName(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipDeviceGetName(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipDeviceTotalMem (hashcat_ctx_t *hashcat_ctx, size_t *bytes, HIPdevice 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); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipDeviceTotalMem(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipDeviceTotalMem(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipDriverGetVersion (hashcat_ctx_t *hashcat_ctx, int *driverVersion) -{ - backend_ctx_t *backend_ctx = hashcat_ctx->backend_ctx; - - HIP_PTR *hip = (HIP_PTR *) backend_ctx->hip; - - const HIPresult HIP_err = hip->hipDriverGetVersion (driverVersion); - - if (HIP_err != HIP_SUCCESS) - { - const char *pStr = NULL; - - if (hip->hipGetErrorString (HIP_err, &pStr) == HIP_SUCCESS) - { - event_log_error (hashcat_ctx, "hipDriverGetVersion(): %s", pStr); - } - else - { - event_log_error (hashcat_ctx, "hipDriverGetVersion(): %d", HIP_err); - } - - return -1; - } - - return 0; -} - -int hc_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx, unsigned int flags, HIPdevice dev) -{ - 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); - - 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, "hipCtxCreate(): %s", pStr); } @@ -2816,19 +2578,19 @@ int hc_hipCtxCreate (hashcat_ctx_t *hashcat_ctx, HIPcontext *pctx, unsigned int return 0; } -int hc_hipCtxDestroy (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx) +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->hipCtxDestroy (ctx); + 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); } @@ -2843,25 +2605,25 @@ int hc_hipCtxDestroy (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx) return 0; } -int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, HIPmodule *module, const void *image, unsigned int numOptions, HIPjit_option *options, void **optionValues) +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 HIPresult HIP_err = hip->hipModuleLoadDataEx (module, image, numOptions, options, optionValues); + const hipError_t HIP_err = hip->hipCtxPopCurrent (pctx); - 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, "hipModuleLoadDataEx(): %s", pStr); + event_log_error (hashcat_ctx, "hipCtxPopCurrent(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipModuleLoadDataEx(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipCtxPopCurrent(): %d", HIP_err); } return -1; @@ -2870,25 +2632,25 @@ int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, HIPmodule *module, const return 0; } -int hc_hipModuleUnload (hashcat_ctx_t *hashcat_ctx, HIPmodule hmod) +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 HIPresult HIP_err = hip->hipModuleUnload (hmod); + const hipError_t HIP_err = hip->hipCtxPushCurrent (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, "hipModuleUnload(): %s", pStr); + event_log_error (hashcat_ctx, "hipCtxPushCurrent(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipModuleUnload(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipCtxPushCurrent(): %d", HIP_err); } return -1; @@ -2897,19 +2659,19 @@ int hc_hipModuleUnload (hashcat_ctx_t *hashcat_ctx, HIPmodule hmod) return 0; } -int hc_hipCtxSetCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx) +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 HIPresult HIP_err = hip->hipCtxSetCurrent (ctx); + const hipError_t HIP_err = hip->hipCtxSetCurrent (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, "hipCtxSetCurrent(): %s", pStr); } @@ -2924,25 +2686,25 @@ int hc_hipCtxSetCurrent (hashcat_ctx_t *hashcat_ctx, HIPcontext ctx) return 0; } -int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr *dptr, size_t bytesize) +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->hipMemAlloc (dptr, bytesize); + const hipError_t HIP_err = hip->hipCtxSynchronize (); - 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, "hipMemAlloc(): %s", pStr); + event_log_error (hashcat_ctx, "hipCtxSynchronize(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipMemAlloc(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipCtxSynchronize(): %d", HIP_err); } return -1; @@ -2951,25 +2713,25 @@ int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr *dptr, size_t bytes return 0; } -int hc_hipMemFree (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dptr) +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 HIPresult HIP_err = hip->hipMemFree (dptr); + const hipError_t HIP_err = hip->hipDeviceGet (device, ordinal); - 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, "hipMemFree(): %s", pStr); + event_log_error (hashcat_ctx, "hipDeviceGet(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipMemFree(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipDeviceGet(): %d", HIP_err); } return -1; @@ -2978,25 +2740,25 @@ int hc_hipMemFree (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dptr) return 0; } -int hc_hipMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, HIPdeviceptr srcDevice, size_t ByteCount) +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 HIPresult HIP_err = hip->hipMemcpyDtoH (dstHost, srcDevice, ByteCount); + const hipError_t HIP_err = hip->hipDeviceGetAttribute (pi, attrib, 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, "hipMemcpyDtoH(): %s", pStr); + event_log_error (hashcat_ctx, "hipDeviceGetAttribute(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipMemcpyDtoH(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipDeviceGetAttribute(): %d", HIP_err); } return -1; @@ -3005,25 +2767,25 @@ int hc_hipMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, HIPdeviceptr sr return 0; } -int hc_hipMemcpyDtoHAsync (hashcat_ctx_t *hashcat_ctx, void *dstHost, HIPdeviceptr srcDevice, size_t ByteCount, HIPstream hStream) +int hc_hipDeviceGetCount (hashcat_ctx_t *hashcat_ctx, int *count) { 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); + 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, "hipMemcpyDtoHAsync(): %s", pStr); + event_log_error (hashcat_ctx, "hipDeviceGetCount(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipMemcpyDtoHAsync(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipDeviceGetCount(): %d", HIP_err); } return -1; @@ -3032,25 +2794,25 @@ int hc_hipMemcpyDtoHAsync (hashcat_ctx_t *hashcat_ctx, void *dstHost, HIPdevicep return 0; } -int hc_hipMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, HIPdeviceptr srcDevice, size_t ByteCount) +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->hipMemcpyDtoD (dstDevice, srcDevice, ByteCount); + 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) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { - event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %s", pStr); + event_log_error (hashcat_ctx, "hipDeviceGetName(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipDeviceGetName(): %d", HIP_err); } return -1; @@ -3059,25 +2821,25 @@ int hc_hipMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, HIPdev return 0; } -int hc_hipMemcpyDtoDAsync (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, HIPdeviceptr srcDevice, size_t ByteCount, HIPstream hStream) +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->hipMemcpyDtoDAsync (dstDevice, srcDevice, ByteCount, hStream); + 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, "hipMemcpyDtoDAsync(): %s", pStr); + event_log_error (hashcat_ctx, "hipDeviceTotalMem(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipMemcpyDtoDAsync(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipDeviceTotalMem(): %d", HIP_err); } return -1; @@ -3086,25 +2848,25 @@ int hc_hipMemcpyDtoDAsync (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, H return 0; } -int hc_hipMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, const void *srcHost, size_t ByteCount) +int hc_hipDriverGetVersion (hashcat_ctx_t *hashcat_ctx, int *driverVersion) { 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); + 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, "hipMemcpyHtoD(): %s", pStr); + event_log_error (hashcat_ctx, "hipDriverGetVersion(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipMemcpyHtoD(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipDriverGetVersion(): %d", HIP_err); } return -1; @@ -3113,25 +2875,25 @@ int hc_hipMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, const return 0; } -int hc_hipMemcpyHtoDAsync (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, const void *srcHost, size_t ByteCount, HIPstream hStream) +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->hipMemcpyHtoDAsync (dstDevice, srcHost, ByteCount, hStream); + 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) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { - event_log_error (hashcat_ctx, "hipMemcpyHtoDAsync(): %s", pStr); + event_log_error (hashcat_ctx, "hipEventCreate(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipMemcpyHtoDAsync(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipEventCreate(): %d", HIP_err); } return -1; @@ -3140,25 +2902,25 @@ int hc_hipMemcpyHtoDAsync (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr dstDevice, c return 0; } -int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, HIPfunction *hfunc, HIPmodule hmod, const char *name) +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->hipModuleGetFunction (hfunc, hmod, name); + 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, "hipModuleGetFunction(): %s", pStr); + event_log_error (hashcat_ctx, "hipEventDestroy(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipModuleGetFunction(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipEventDestroy(): %d", HIP_err); } return -1; @@ -3167,25 +2929,25 @@ int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, HIPfunction *hfunc, HIP return 0; } -int hc_hipModuleGetGlobal (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr *dptr, size_t *bytes, HIPmodule hmod, const char *name) +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->hipModuleGetGlobal (dptr, bytes, hmod, name); + 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, "hipModuleGetGlobal(): %s", pStr); + event_log_error (hashcat_ctx, "hipEventElapsedTime(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipModuleGetGlobal(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipEventElapsedTime(): %d", HIP_err); } return -1; @@ -3194,25 +2956,25 @@ int hc_hipModuleGetGlobal (hashcat_ctx_t *hashcat_ctx, HIPdeviceptr *dptr, size_ return 0; } -int hc_hipMemGetInfo (hashcat_ctx_t *hashcat_ctx, size_t *free, size_t *total) +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->hipMemGetInfo (free, total); + 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) + if (hip->hipGetErrorString (HIP_err, &pStr) == hipSuccess) { - event_log_error (hashcat_ctx, "hipMemGetInfo(): %s", pStr); + event_log_error (hashcat_ctx, "hipEventRecord(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipMemGetInfo(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipEventRecord(): %d", HIP_err); } return -1; @@ -3221,25 +2983,25 @@ int hc_hipMemGetInfo (hashcat_ctx_t *hashcat_ctx, size_t *free, size_t *total) return 0; } -int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPfunction_attribute attrib, HIPfunction hfunc) +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->hipFuncGetAttribute (pi, attrib, hfunc); + 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, "hipFuncGetAttribute(): %s", pStr); + event_log_error (hashcat_ctx, "hipEventSynchronize(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipFuncGetAttribute(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipEventSynchronize(): %d", HIP_err); } return -1; @@ -3248,29 +3010,25 @@ int hc_hipFuncGetAttribute (hashcat_ctx_t *hashcat_ctx, int *pi, HIPfunction_att 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) +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->hipFuncSetAttribute (hfunc, attrib, value); + 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, "hipFuncSetAttribute(): %s", pStr); + event_log_error (hashcat_ctx, "hipFuncGetAttribute(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipFuncSetAttribute(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipFuncGetAttribute(): %d", HIP_err); } return -1; @@ -3278,27 +3036,26 @@ int hc_hipFuncSetAttribute (hashcat_ctx_t *hashcat_ctx, HIPfunction hfunc, HIPfu return 0; } -*/ -int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, HIPstream *phStream, unsigned int Flags) +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->hipStreamCreate (phStream, Flags); + 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, "hipStreamCreate(): %s", pStr); + event_log_error (hashcat_ctx, "hipLaunchKernel(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipStreamCreate(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipLaunchKernel(): %d", HIP_err); } return -1; @@ -3307,25 +3064,25 @@ int hc_hipStreamCreate (hashcat_ctx_t *hashcat_ctx, HIPstream *phStream, unsigne return 0; } -int hc_hipStreamDestroy (hashcat_ctx_t *hashcat_ctx, HIPstream hStream) +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->hipStreamDestroy (hStream); + 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, "hipStreamDestroy(): %s", pStr); + event_log_error (hashcat_ctx, "hipInit(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipStreamDestroy(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipInit(): %d", HIP_err); } return -1; @@ -3334,25 +3091,25 @@ int hc_hipStreamDestroy (hashcat_ctx_t *hashcat_ctx, HIPstream hStream) return 0; } -int hc_hipStreamSynchronize (hashcat_ctx_t *hashcat_ctx, HIPstream hStream) +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->hipStreamSynchronize (hStream); + 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, "hipStreamSynchronize(): %s", pStr); + event_log_error (hashcat_ctx, "hipMemAlloc(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipStreamSynchronize(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipMemAlloc(): %d", HIP_err); } return -1; @@ -3361,25 +3118,25 @@ int hc_hipStreamSynchronize (hashcat_ctx_t *hashcat_ctx, HIPstream hStream) 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) +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->hipLaunchKernel (f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra); + 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, "hipLaunchKernel(): %s", pStr); + event_log_error (hashcat_ctx, "hipMemFree(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipLaunchKernel(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipMemFree(): %d", HIP_err); } return -1; @@ -3388,25 +3145,25 @@ int hc_hipLaunchKernel (hashcat_ctx_t *hashcat_ctx, HIPfunction f, unsigned int return 0; } -int hc_hipCtxSynchronize (hashcat_ctx_t *hashcat_ctx) +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->hipCtxSynchronize (); + 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, "hipCtxSynchronize(): %s", pStr); + event_log_error (hashcat_ctx, "hipMemGetInfo(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipCtxSynchronize(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipMemGetInfo(): %d", HIP_err); } return -1; @@ -3415,25 +3172,25 @@ int hc_hipCtxSynchronize (hashcat_ctx_t *hashcat_ctx) return 0; } -int hc_hipEventCreate (hashcat_ctx_t *hashcat_ctx, HIPevent *phEvent, unsigned int Flags) +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->hipEventCreate (phEvent, Flags); + 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, "hipEventCreate(): %s", pStr); + event_log_error (hashcat_ctx, "hipMemcpyDtoH(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipEventCreate(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipMemcpyDtoH(): %d", HIP_err); } return -1; @@ -3442,25 +3199,25 @@ 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_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 HIPresult HIP_err = hip->hipEventDestroy (hEvent); + const hipError_t HIP_err = hip->hipMemcpyDtoHAsync (dstHost, srcDevice, ByteCount, hStream); - 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); + event_log_error (hashcat_ctx, "hipMemcpyDtoHAsync(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipEventDestroy(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipMemcpyDtoHAsync(): %d", HIP_err); } return -1; @@ -3469,25 +3226,25 @@ 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_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 HIPresult HIP_err = hip->hipEventElapsedTime (pMilliseconds, hStart, hEnd); + const hipError_t HIP_err = hip->hipMemcpyDtoD (dstDevice, 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, "hipEventElapsedTime(): %s", pStr); + event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipEventElapsedTime(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipMemcpyDtoD(): %d", HIP_err); } return -1; @@ -3496,25 +3253,25 @@ 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_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 HIPresult HIP_err = hip->hipEventQuery (hEvent); + const hipError_t HIP_err = hip->hipMemcpyDtoDAsync (dstDevice, srcDevice, ByteCount, hStream); - 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, "hipEventQuery(): %s", pStr); + event_log_error (hashcat_ctx, "hipMemcpyDtoDAsync(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipEventQuery(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipMemcpyDtoDAsync(): %d", HIP_err); } return -1; @@ -3523,25 +3280,25 @@ int hc_hipEventQuery (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent) return 0; } -int hc_hipEventRecord (hashcat_ctx_t *hashcat_ctx, HIPevent hEvent, HIPstream hStream) +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 HIPresult HIP_err = hip->hipEventRecord (hEvent, hStream); + const hipError_t HIP_err = hip->hipMemcpyHtoD (dstDevice, srcHost, 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, "hipEventRecord(): %s", pStr); + event_log_error (hashcat_ctx, "hipMemcpyHtoD(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipEventRecord(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipMemcpyHtoD(): %d", HIP_err); } return -1; @@ -3550,25 +3307,25 @@ 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_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 HIPresult HIP_err = hip->hipEventSynchronize (hEvent); + const hipError_t HIP_err = hip->hipMemcpyHtoDAsync (dstDevice, srcHost, ByteCount, hStream); - 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); + event_log_error (hashcat_ctx, "hipMemcpyHtoDAsync(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipEventSynchronize(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipMemcpyHtoDAsync(): %d", HIP_err); } return -1; @@ -3577,25 +3334,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_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 HIPresult HIP_err = hip->hipCtxSetCacheConfig (config); + const hipError_t HIP_err = hip->hipModuleGetFunction (hfunc, hmod, name); - 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, "hipModuleGetFunction(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipCtxSetCacheConfig(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipModuleGetFunction(): %d", HIP_err); } return -1; @@ -3604,25 +3361,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_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 HIPresult HIP_err = hip->hipCtxPushCurrent (ctx); + const hipError_t HIP_err = hip->hipModuleGetGlobal (dptr, bytes, hmod, name); - 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, "hipModuleGetGlobal(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipCtxPushCurrent(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipModuleGetGlobal(): %d", HIP_err); } return -1; @@ -3631,25 +3388,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_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 HIPresult HIP_err = hip->hipCtxPopCurrent (pctx); + const hipError_t HIP_err = hip->hipModuleLoadDataEx (module, image, numOptions, options, optionValues); - 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, "hipModuleLoadDataEx(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipCtxPopCurrent(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipModuleLoadDataEx(): %d", HIP_err); } return -1; @@ -3658,25 +3415,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_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 HIPresult HIP_err = hip->hipLinkCreate (numOptions, options, optionValues, stateOut); + const hipError_t HIP_err = hip->hipModuleUnload (hmod); - 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, "hipModuleUnload(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipLinkCreate(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipModuleUnload(): %d", HIP_err); } return -1; @@ -3685,25 +3442,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_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 HIPresult HIP_err = hip->hipLinkAddData (state, type, data, size, name, numOptions, options, optionValues); + const hipError_t HIP_err = hip->hipStreamCreate (phStream, 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, "hipLinkAddData(): %s", pStr); + event_log_error (hashcat_ctx, "hipStreamCreate(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipLinkAddData(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipStreamCreate(): %d", HIP_err); } return -1; @@ -3712,25 +3469,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_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 HIPresult HIP_err = hip->hipLinkDestroy (state); + const hipError_t HIP_err = hip->hipStreamDestroy (hStream); - 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, "hipStreamDestroy(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipLinkDestroy(): %d", HIP_err); + event_log_error (hashcat_ctx, "hipStreamDestroy(): %d", HIP_err); } return -1; @@ -3739,25 +3496,25 @@ 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_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 HIPresult HIP_err = hip->hipLinkComplete (state, hipbinOut, sizeOut); + const hipError_t HIP_err = hip->hipStreamSynchronize (hStream); - 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, "hipStreamSynchronize(): %s", pStr); } else { - event_log_error (hashcat_ctx, "hipLinkComplete(): %d", HIP_err); + 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) - { - device_param->skipped = true; - continue; - } + // 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, HIP_DEVICE_ATTRIBUTE_PCI_BUS_ID, hip_device) == -1) + if (hc_hipDeviceGetAttribute (hashcat_ctx, &pci_bus_id_nv, hipDeviceAttributePciBusId, 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 };