Get rid of hip/hip_runtime.h dependancy

OpenCL/inc_platform.cl

@@ -193,35 +193,47 @@ DECLSPEC u32 hc_atomic_dec (GLOBAL_AS u32 *p)
 {
   volatile const u32 val = 1;
 
-  return atomicSub (p, val);
+  return __atomic_fetch_sub (p, val, __ATOMIC_RELAXED);
 }
 
 DECLSPEC u32 hc_atomic_inc (GLOBAL_AS u32 *p)
 {
   volatile const u32 val = 1;
 
-  return atomicAdd (p, val);
+  return __atomic_fetch_add (p, val, __ATOMIC_RELAXED);
 }
 
 DECLSPEC u32 hc_atomic_or (GLOBAL_AS u32 *p, volatile const u32 val)
 {
-  return atomicOr (p, 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);
+}
+
+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 __attribute__((unused)))
 {
-  return (blockIdx.x * blockDim.x) + threadIdx.x;
+  return (__ockl_get_group_id (0) * __ockl_get_local_size (0)) + __ockl_get_local_id (0);
 }
 
 DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused)))
 {
-  return threadIdx.x;
+  return __ockl_get_local_id (0);
 }
 
 DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
 {
   // verify
-  return blockDim.x;
+  return __ockl_get_local_size (0);
 }
 
 DECLSPEC u32x rotl32 (const u32x a, const int n)

OpenCL/inc_types.h

@@ -24,6 +24,17 @@ typedef unsigned int       uint;
 typedef unsigned long long ulong;
 #endif
 
+#ifdef IS_HIP
+typedef unsigned char      uchar;
+typedef unsigned short     ushort;
+typedef unsigned int       uint;
+typedef unsigned long long ulong;
+
+typedef __SIZE_TYPE__ size_t;
+
+typedef uint uint4 __attribute__((ext_vector_type(4)));
+#endif
+
 #ifdef KERNEL_STATIC
 typedef uchar  u8;
 typedef ushort u16;

OpenCL/inc_vendor.h

@@ -16,9 +16,6 @@
 #define IS_OPENCL
 #endif
 
-#ifdef IS_HIP
-#include <hip/hip_runtime.h>
-#endif
 
 #if defined IS_NATIVE
 #define CONSTANT_VK
@@ -35,6 +32,10 @@
 #define LOCAL_AS
 #define KERNEL_FQ   extern "C" __global__
 #elif defined IS_HIP
+#define __device__   __attribute__((device))
+#define __constant__ __attribute__((constant))
+#define __shared__   __attribute__((shared))
+#define __global__   __attribute__((global))
 #define CONSTANT_VK __constant__
 #define CONSTANT_AS
 #define GLOBAL_AS

src/backend.c

@@ -10495,7 +10495,7 @@ static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_p
 
       if (hc_hiprtcCreateProgram (hashcat_ctx, &program, kernel_sources[0], kernel_name, 0, NULL, NULL) == -1) return false;
 
-      char **hiprtc_options = (char **) hccalloc (6 + strlen (build_options_buf) + 1, sizeof (char *)); // ...
+      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";
@@ -10504,16 +10504,17 @@ static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_p
       //hc_asprintf (&hiprtc_options[3], "compute_%d%d", device_param->sm_major, device_param->sm_minor);
 
       hiprtc_options[0] = "--gpu-max-threads-per-block=64";
-      hiprtc_options[1] = "";
-      hiprtc_options[2] = "";
+      hiprtc_options[1] = "-nocudainc";
+      hiprtc_options[2] = "-nocudalib";
       hiprtc_options[3] = "";
+      hiprtc_options[4] = "";
 
-      hiprtc_options[4] = "-I";
-      hiprtc_options[5] = folder_config->cpath_real;
+      hiprtc_options[5] = "-I";
+      hiprtc_options[6] = folder_config->cpath_real;
 
       char *hiprtc_options_string = hcstrdup (build_options_buf);
 
-      const int num_options = 6 + hiprtc_make_options_array_from_string (hiprtc_options_string, hiprtc_options + 6);
+      const int num_options = 7 + hiprtc_make_options_array_from_string (hiprtc_options_string, hiprtc_options + 7);
 
       const int rc_hiprtcCompileProgram = hc_hiprtcCompileProgram (hashcat_ctx, program, num_options, (const char * const *) hiprtc_options);