diff --git a/OpenCL/inc_vendor.h b/OpenCL/inc_vendor.h
index dc6a41d4a..d98a85053 100644
--- a/OpenCL/inc_vendor.h
+++ b/OpenCL/inc_vendor.h
@@ -131,7 +131,7 @@
 #if defined IS_AMD && defined IS_GPU
 #define DECLSPEC inline static
 #elif defined IS_HIP
-#define DECLSPEC __device__
+#define DECLSPEC inline static __device__
 #else
 #define DECLSPEC
 #endif
diff --git a/src/modules/module_01460.c b/src/modules/module_01460.c
index 6cb814ed2..f2952aa36 100644
--- a/src/modules/module_01460.c
+++ b/src/modules/module_01460.c
@@ -59,6 +59,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_01500.c b/src/modules/module_01500.c
index 274d9541a..ea01dab96 100644
--- a/src/modules/module_01500.c
+++ b/src/modules/module_01500.c
@@ -179,6 +179,14 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
       hc_asprintf (&jit_build_options, "-DDESCRYPT_SALT=%u -D _unroll", hashes->salts_buf[0].salt_buf[0] & 0xfff);
     }
   }
+  // ROCM
+  else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    if ((user_options->attack_mode == ATTACK_MODE_BF) && (hashes->salts_cnt == 1) && (user_options->slow_candidates == false))
+    {
+      hc_asprintf (&jit_build_options, "-DDESCRYPT_SALT=%u -D _unroll", hashes->salts_buf[0].salt_buf[0] & 0xfff);
+    }
+  }
   else
   {
     if ((user_options->attack_mode == ATTACK_MODE_BF) && (hashes->salts_cnt == 1) && (user_options->slow_candidates == false))
diff --git a/src/modules/module_01700.c b/src/modules/module_01700.c
index 9a7f1d34f..04f2762c1 100644
--- a/src/modules/module_01700.c
+++ b/src/modules/module_01700.c
@@ -58,6 +58,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_01720.c b/src/modules/module_01720.c
index 6833f4405..3fdc77653 100644
--- a/src/modules/module_01720.c
+++ b/src/modules/module_01720.c
@@ -58,6 +58,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_01722.c b/src/modules/module_01722.c
index 4585e2dbb..3264c5f46 100644
--- a/src/modules/module_01722.c
+++ b/src/modules/module_01722.c
@@ -59,6 +59,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_01800.c b/src/modules/module_01800.c
index 1cc1781c2..aefab6e3b 100644
--- a/src/modules/module_01800.c
+++ b/src/modules/module_01800.c
@@ -438,6 +438,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-fno-unroll-loops");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_03000.c b/src/modules/module_03000.c
index a3373e09b..c9b616ab5 100644
--- a/src/modules/module_03000.c
+++ b/src/modules/module_03000.c
@@ -81,6 +81,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_05200.c b/src/modules/module_05200.c
index 470411378..6fb3f08d8 100644
--- a/src/modules/module_05200.c
+++ b/src/modules/module_05200.c
@@ -81,6 +81,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_06211.c b/src/modules/module_06211.c
index b7aa35874..5cb417d26 100644
--- a/src/modules/module_06211.c
+++ b/src/modules/module_06211.c
@@ -83,25 +83,6 @@ bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE
   return false;
 }
 
-char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
-{
-  char *jit_build_options = NULL;
-
-  // Extra treatment for Apple systems
-  if (device_param->opencl_platform_vendor_id == VENDOR_ID_APPLE)
-  {
-    return jit_build_options;
-  }
-
-  // ROCM
-  if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
-  {
-    hc_asprintf (&jit_build_options, "-D _unroll");
-  }
-
-  return jit_build_options;
-}
-
 bool module_potfile_disable (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
 {
   const bool potfile_disable = true;
@@ -303,7 +284,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = module_jit_build_options;
+  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;
diff --git a/src/modules/module_06212.c b/src/modules/module_06212.c
index 9ac3487c7..ceb18f192 100644
--- a/src/modules/module_06212.c
+++ b/src/modules/module_06212.c
@@ -83,25 +83,6 @@ bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE
   return false;
 }
 
-char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
-{
-  char *jit_build_options = NULL;
-
-  // Extra treatment for Apple systems
-  if (device_param->opencl_platform_vendor_id == VENDOR_ID_APPLE)
-  {
-    return jit_build_options;
-  }
-
-  // ROCM
-  if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
-  {
-    hc_asprintf (&jit_build_options, "-D _unroll");
-  }
-
-  return jit_build_options;
-}
-
 bool module_potfile_disable (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
 {
   const bool potfile_disable = true;
@@ -303,7 +284,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = module_jit_build_options;
+  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;
diff --git a/src/modules/module_06213.c b/src/modules/module_06213.c
index 04430ec31..20323fc62 100644
--- a/src/modules/module_06213.c
+++ b/src/modules/module_06213.c
@@ -83,25 +83,6 @@ bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE
   return false;
 }
 
-char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
-{
-  char *jit_build_options = NULL;
-
-  // Extra treatment for Apple systems
-  if (device_param->opencl_platform_vendor_id == VENDOR_ID_APPLE)
-  {
-    return jit_build_options;
-  }
-
-  // ROCM
-  if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
-  {
-    hc_asprintf (&jit_build_options, "-D _unroll");
-  }
-
-  return jit_build_options;
-}
-
 bool module_potfile_disable (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
 {
   const bool potfile_disable = true;
@@ -301,7 +282,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = module_jit_build_options;
+  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;
diff --git a/src/modules/module_06800.c b/src/modules/module_06800.c
index 5a79ca8e7..0f25fa29d 100644
--- a/src/modules/module_06800.c
+++ b/src/modules/module_06800.c
@@ -78,6 +78,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_07400.c b/src/modules/module_07400.c
index f8ebca33d..3be47f898 100644
--- a/src/modules/module_07400.c
+++ b/src/modules/module_07400.c
@@ -244,6 +244,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_07401.c b/src/modules/module_07401.c
index b96318d8e..ba71bf179 100644
--- a/src/modules/module_07401.c
+++ b/src/modules/module_07401.c
@@ -245,6 +245,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_07500.c b/src/modules/module_07500.c
index 931cc6b47..1681fb4a8 100644
--- a/src/modules/module_07500.c
+++ b/src/modules/module_07500.c
@@ -80,6 +80,17 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
         native_threads = 64;
       }
     }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      if (device_param->device_local_mem_size < 49152)
+      {
+        native_threads = 32;
+      }
+      else
+      {
+        native_threads = 64;
+      }
+    }
     else
     {
       native_threads = 32;
diff --git a/src/modules/module_07900.c b/src/modules/module_07900.c
index 67944eb5a..a51efecd9 100644
--- a/src/modules/module_07900.c
+++ b/src/modules/module_07900.c
@@ -79,6 +79,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_08700.c b/src/modules/module_08700.c
index 6f75c6e02..be902d527 100644
--- a/src/modules/module_08700.c
+++ b/src/modules/module_08700.c
@@ -60,6 +60,19 @@ u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED con
   return pw_max;
 }
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-fno-unroll-loops");
+  }
+
+  return jit_build_options;
+}
+
 int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
 {
   u32 *digest = (u32 *) digest_buf;
@@ -179,7 +192,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;
diff --git a/src/modules/module_09600.c b/src/modules/module_09600.c
index abfe9fdee..ecf6dc6fb 100644
--- a/src/modules/module_09600.c
+++ b/src/modules/module_09600.c
@@ -84,6 +84,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_09700.c b/src/modules/module_09700.c
index f5e2f1138..cad911186 100644
--- a/src/modules/module_09700.c
+++ b/src/modules/module_09700.c
@@ -77,6 +77,10 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     {
       native_threads = 64;
     }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      native_threads = 64;
+    }
     else
     {
       native_threads = 32;
diff --git a/src/modules/module_09710.c b/src/modules/module_09710.c
index 250bc3863..033f77ea0 100644
--- a/src/modules/module_09710.c
+++ b/src/modules/module_09710.c
@@ -77,6 +77,10 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     {
       native_threads = 64;
     }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      native_threads = 64;
+    }
     else
     {
       native_threads = 32;
diff --git a/src/modules/module_09720.c b/src/modules/module_09720.c
index 7db204dc6..04e99201f 100644
--- a/src/modules/module_09720.c
+++ b/src/modules/module_09720.c
@@ -58,6 +58,41 @@ static const char *SIGNATURE_OLDOFFICE  = "$oldoffice$";
 static const char *SIGNATURE_OLDOFFICE0 = "$oldoffice$0";
 static const char *SIGNATURE_OLDOFFICE1 = "$oldoffice$1";
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  u32 native_threads = 0;
+
+  if (device_param->opencl_device_type & CL_DEVICE_TYPE_CPU)
+  {
+    native_threads = 1;
+  }
+  else if (device_param->opencl_device_type & CL_DEVICE_TYPE_GPU)
+  {
+    if (device_param->opencl_device_vendor_id == VENDOR_ID_INTEL_SDK)
+    {
+      native_threads = 8;
+    }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD)
+    {
+      native_threads = 64;
+    }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      native_threads = 64;
+    }
+    else
+    {
+      native_threads = 32;
+    }
+  }
+
+  hc_asprintf (&jit_build_options, "-D FIXED_LOCAL_SIZE=%u -D _unroll", native_threads);
+
+  return jit_build_options;
+}
+
 u64 module_esalt_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
 {
   const u64 esalt_size = (const u64) sizeof (oldoffice01_t);
@@ -273,7 +308,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;
diff --git a/src/modules/module_09800.c b/src/modules/module_09800.c
index 4508fcd5b..2eb7fab05 100644
--- a/src/modules/module_09800.c
+++ b/src/modules/module_09800.c
@@ -79,6 +79,10 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     {
       native_threads = 64;
     }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      native_threads = 64;
+    }
     else
     {
       native_threads = 32;
diff --git a/src/modules/module_09810.c b/src/modules/module_09810.c
index e1a434cf3..2a1074b2c 100644
--- a/src/modules/module_09810.c
+++ b/src/modules/module_09810.c
@@ -78,6 +78,10 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     {
       native_threads = 64;
     }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      native_threads = 64;
+    }
     else
     {
       native_threads = 32;
diff --git a/src/modules/module_09820.c b/src/modules/module_09820.c
index f3f7ddee2..ea3dfe22b 100644
--- a/src/modules/module_09820.c
+++ b/src/modules/module_09820.c
@@ -80,6 +80,10 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     {
       native_threads = 64;
     }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      native_threads = 64;
+    }
     else
     {
       native_threads = 32;
diff --git a/src/modules/module_10400.c b/src/modules/module_10400.c
index c782d9c04..77416f5ce 100644
--- a/src/modules/module_10400.c
+++ b/src/modules/module_10400.c
@@ -84,6 +84,10 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     {
       native_threads = 64;
     }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      native_threads = 64;
+    }
     else
     {
       native_threads = 32;
diff --git a/src/modules/module_10410.c b/src/modules/module_10410.c
index df1d40d2d..b2c98363f 100644
--- a/src/modules/module_10410.c
+++ b/src/modules/module_10410.c
@@ -85,6 +85,10 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     {
       native_threads = 64;
     }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      native_threads = 64;
+    }
     else
     {
       native_threads = 32;
diff --git a/src/modules/module_10420.c b/src/modules/module_10420.c
index 6f182a436..23e537bf5 100644
--- a/src/modules/module_10420.c
+++ b/src/modules/module_10420.c
@@ -64,6 +64,41 @@ typedef struct pdf
 
 static const char *SIGNATURE_PDF = "$pdf$";
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  u32 native_threads = 0;
+
+  if (device_param->opencl_device_type & CL_DEVICE_TYPE_CPU)
+  {
+    native_threads = 1;
+  }
+  else if (device_param->opencl_device_type & CL_DEVICE_TYPE_GPU)
+  {
+    if (device_param->opencl_device_vendor_id == VENDOR_ID_INTEL_SDK)
+    {
+      native_threads = 8;
+    }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD)
+    {
+      native_threads = 64;
+    }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      native_threads = 64;
+    }
+    else
+    {
+      native_threads = 32;
+    }
+  }
+
+  hc_asprintf (&jit_build_options, "-D FIXED_LOCAL_SIZE=%u -D _unroll", native_threads);
+
+  return jit_build_options;
+}
+
 u64 module_esalt_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
 {
   const u64 esalt_size = (const u64) sizeof (pdf_t);
@@ -369,7 +404,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;
diff --git a/src/modules/module_10500.c b/src/modules/module_10500.c
index fbb1af6a3..80a8478ef 100644
--- a/src/modules/module_10500.c
+++ b/src/modules/module_10500.c
@@ -108,6 +108,10 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     {
       native_threads = 64;
     }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      native_threads = 64;
+    }
     else
     {
       native_threads = 32;
diff --git a/src/modules/module_10800.c b/src/modules/module_10800.c
index 65cff2b7e..1765bddac 100644
--- a/src/modules/module_10800.c
+++ b/src/modules/module_10800.c
@@ -72,6 +72,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_10810.c b/src/modules/module_10810.c
index 797c2cf17..10d1443f4 100644
--- a/src/modules/module_10810.c
+++ b/src/modules/module_10810.c
@@ -72,6 +72,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_10820.c b/src/modules/module_10820.c
index d9b76cc5c..82987fe39 100644
--- a/src/modules/module_10820.c
+++ b/src/modules/module_10820.c
@@ -72,6 +72,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_10830.c b/src/modules/module_10830.c
index 91a70b707..f431762f8 100644
--- a/src/modules/module_10830.c
+++ b/src/modules/module_10830.c
@@ -73,6 +73,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_10840.c b/src/modules/module_10840.c
index 4cbb7db28..f60d3ea13 100644
--- a/src/modules/module_10840.c
+++ b/src/modules/module_10840.c
@@ -73,6 +73,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_10870.c b/src/modules/module_10870.c
index 52a70afac..047c67242 100644
--- a/src/modules/module_10870.c
+++ b/src/modules/module_10870.c
@@ -73,6 +73,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_10900.c b/src/modules/module_10900.c
index b0634ec4e..efde01301 100644
--- a/src/modules/module_10900.c
+++ b/src/modules/module_10900.c
@@ -77,6 +77,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_11300.c b/src/modules/module_11300.c
index 9cb3ae217..981a0b471 100644
--- a/src/modules/module_11300.c
+++ b/src/modules/module_11300.c
@@ -81,6 +81,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_11600.c b/src/modules/module_11600.c
index 7694b71f3..25fe732a3 100644
--- a/src/modules/module_11600.c
+++ b/src/modules/module_11600.c
@@ -111,6 +111,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_12200.c b/src/modules/module_12200.c
index 15b6c0c85..3f6b57821 100644
--- a/src/modules/module_12200.c
+++ b/src/modules/module_12200.c
@@ -81,6 +81,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_12800.c b/src/modules/module_12800.c
index 22658f2c6..cd2099cdc 100644
--- a/src/modules/module_12800.c
+++ b/src/modules/module_12800.c
@@ -69,6 +69,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_12900.c b/src/modules/module_12900.c
index cabff4977..4c5a9892b 100644
--- a/src/modules/module_12900.c
+++ b/src/modules/module_12900.c
@@ -69,6 +69,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_13000.c b/src/modules/module_13000.c
index 2d441994a..ab389431d 100644
--- a/src/modules/module_13000.c
+++ b/src/modules/module_13000.c
@@ -75,6 +75,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_13100.c b/src/modules/module_13100.c
index bab0dbf26..fac5cb24c 100644
--- a/src/modules/module_13100.c
+++ b/src/modules/module_13100.c
@@ -79,6 +79,17 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
         native_threads = 64;
       }
     }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      if (device_param->device_local_mem_size < 49152)
+      {
+        native_threads = 32;
+      }
+      else
+      {
+        native_threads = 64;
+      }
+    }
     else
     {
       native_threads = 32;
diff --git a/src/modules/module_14000.c b/src/modules/module_14000.c
index c6854d6c9..013888bcc 100644
--- a/src/modules/module_14000.c
+++ b/src/modules/module_14000.c
@@ -76,6 +76,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_14500.c b/src/modules/module_14500.c
index 286117f86..ba4edab17 100644
--- a/src/modules/module_14500.c
+++ b/src/modules/module_14500.c
@@ -113,6 +113,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_15000.c b/src/modules/module_15000.c
index 481f88cb3..9fae0e769 100644
--- a/src/modules/module_15000.c
+++ b/src/modules/module_15000.c
@@ -69,6 +69,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_15600.c b/src/modules/module_15600.c
index 168609688..c7acdb8d3 100644
--- a/src/modules/module_15600.c
+++ b/src/modules/module_15600.c
@@ -78,6 +78,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_16200.c b/src/modules/module_16200.c
index 7f1b34959..19cbbac8d 100644
--- a/src/modules/module_16200.c
+++ b/src/modules/module_16200.c
@@ -78,6 +78,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_16300.c b/src/modules/module_16300.c
index 8ce5e668d..33997b1ed 100644
--- a/src/modules/module_16300.c
+++ b/src/modules/module_16300.c
@@ -79,6 +79,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_16900.c b/src/modules/module_16900.c
index 5cfe5aeb7..93915b592 100644
--- a/src/modules/module_16900.c
+++ b/src/modules/module_16900.c
@@ -79,6 +79,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_18200.c b/src/modules/module_18200.c
index b95ddab6f..e6596306b 100644
--- a/src/modules/module_18200.c
+++ b/src/modules/module_18200.c
@@ -79,6 +79,17 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
         native_threads = 64;
       }
     }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      if (device_param->device_local_mem_size < 49152)
+      {
+        native_threads = 32;
+      }
+      else
+      {
+        native_threads = 64;
+      }
+    }
     else
     {
       native_threads = 32;
diff --git a/src/modules/module_18300.c b/src/modules/module_18300.c
index 592081296..b58ef35f5 100644
--- a/src/modules/module_18300.c
+++ b/src/modules/module_18300.c
@@ -78,6 +78,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_18800.c b/src/modules/module_18800.c
index 5bb6132a7..6847edde2 100644
--- a/src/modules/module_18800.c
+++ b/src/modules/module_18800.c
@@ -80,6 +80,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_20600.c b/src/modules/module_20600.c
index a7debbc58..e270fde50 100644
--- a/src/modules/module_20600.c
+++ b/src/modules/module_20600.c
@@ -71,6 +71,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_21000.c b/src/modules/module_21000.c
index e482be5e1..7df98beb4 100644
--- a/src/modules/module_21000.c
+++ b/src/modules/module_21000.c
@@ -72,6 +72,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_22100.c b/src/modules/module_22100.c
index cd79bc7f8..47b72d7df 100644
--- a/src/modules/module_22100.c
+++ b/src/modules/module_22100.c
@@ -86,6 +86,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_22200.c b/src/modules/module_22200.c
index 1f034ef50..5ea525fbe 100644
--- a/src/modules/module_22200.c
+++ b/src/modules/module_22200.c
@@ -72,6 +72,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_22400.c b/src/modules/module_22400.c
index 567dec821..a3ab81101 100644
--- a/src/modules/module_22400.c
+++ b/src/modules/module_22400.c
@@ -96,6 +96,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_23400.c b/src/modules/module_23400.c
index 4ba636bfa..5921ca9db 100644
--- a/src/modules/module_23400.c
+++ b/src/modules/module_23400.c
@@ -69,6 +69,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_24200.c b/src/modules/module_24200.c
index b1da70c9a..bd93b36e0 100644
--- a/src/modules/module_24200.c
+++ b/src/modules/module_24200.c
@@ -79,6 +79,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_25300.c b/src/modules/module_25300.c
index 6bed1017c..e21d55881 100644
--- a/src/modules/module_25300.c
+++ b/src/modules/module_25300.c
@@ -72,6 +72,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     hc_asprintf (&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_25400.c b/src/modules/module_25400.c
index 5dbbe8dc3..341837786 100644
--- a/src/modules/module_25400.c
+++ b/src/modules/module_25400.c
@@ -118,6 +118,17 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
         native_threads = 64;
       }
     }
+    else if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+    {
+      if (device_param->device_local_mem_size < 49152)
+      {
+        native_threads = 32;
+      }
+      else
+      {
+        native_threads = 64;
+      }
+    }
     else
     {
       native_threads = 32;
diff --git a/src/modules/module_25500.c b/src/modules/module_25500.c
index 10edd7203..e6853e951 100644
--- a/src/modules/module_25500.c
+++ b/src/modules/module_25500.c
@@ -74,6 +74,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_25900.c b/src/modules/module_25900.c
index efe7f39f2..c44a1642b 100644
--- a/src/modules/module_25900.c
+++ b/src/modules/module_25900.c
@@ -81,6 +81,12 @@ char* module_jit_build_options(MAYBE_UNUSED const hashconfig_t *hashconfig, MAYB
     hc_asprintf(&jit_build_options, "-D _unroll");
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_26200.c b/src/modules/module_26200.c
index 63989796e..8b082f2d5 100644
--- a/src/modules/module_26200.c
+++ b/src/modules/module_26200.c
@@ -74,6 +74,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {
diff --git a/src/modules/module_26600.c b/src/modules/module_26600.c
index 61ccac983..c9e04958c 100644
--- a/src/modules/module_26600.c
+++ b/src/modules/module_26600.c
@@ -74,6 +74,12 @@ char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAY
     return jit_build_options;
   }
 
+  // HIP
+  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
+  {
+    hc_asprintf (&jit_build_options, "-D _unroll");
+  }
+
   // ROCM
   if ((device_param->opencl_device_vendor_id == VENDOR_ID_AMD) && (device_param->has_vperm == true))
   {