From 3f6c5a0042c3c8cd7e1c95708761b1998ad7279f Mon Sep 17 00:00:00 2001
From: Jens Steube <jens.steube@gmail.com>
Date: Fri, 23 Jul 2021 21:09:55 +0200
Subject: [PATCH] Update module_unstable_warning() for -m 172xx on HIP

---
 OpenCL/inc_platform.cl     | 13 +++++++------
 OpenCL/inc_types.h         |  2 +-
 OpenCL/inc_vendor.h        |  4 ----
 OpenCL/inc_zip_inflate.cl  |  4 +---
 src/modules/module_17200.c |  5 -----
 src/modules/module_17220.c |  5 -----
 src/modules/module_17225.c |  5 -----
 7 files changed, 9 insertions(+), 29 deletions(-)

diff --git a/OpenCL/inc_platform.cl b/OpenCL/inc_platform.cl
index 47cfc9a84..8ccb034aa 100644
--- a/OpenCL/inc_platform.cl
+++ b/OpenCL/inc_platform.cl
@@ -208,19 +208,20 @@ DECLSPEC u32 hc_atomic_or (GLOBAL_AS u32 *p, volatile const u32 val)
   return atomicOr (p, val);
 }
 
-DECLSPEC size_t get_global_id  (const u32 dimindx)
+DECLSPEC size_t get_global_id  (const u32 dimindx __attribute__((unused)))
 {
-  return (__ockl_get_group_id (dimindx) * __ockl_get_local_size (dimindx)) + __ockl_get_local_id (dimindx);
+  return (blockIdx.x * blockDim.x) + threadIdx.x;
 }
 
-DECLSPEC size_t get_local_id (const u32 dimindx)
+DECLSPEC size_t get_local_id (const u32 dimindx __attribute__((unused)))
 {
-  return __ockl_get_local_id (dimindx);
+  return threadIdx.x;
 }
 
-DECLSPEC size_t get_local_size (const u32 dimindx)
+DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
 {
-  return __ockl_get_local_size (dimindx);
+  // verify
+  return blockDim.x;
 }
 
 DECLSPEC u32x rotl32 (const u32x a, const int n)
diff --git a/OpenCL/inc_types.h b/OpenCL/inc_types.h
index 4caf595de..295caabc0 100644
--- a/OpenCL/inc_types.h
+++ b/OpenCL/inc_types.h
@@ -833,7 +833,7 @@ typedef __device_builtin__ struct u64x u64x;
 typedef VTYPE(uchar,  VECT_SIZE) u8x;
 typedef VTYPE(ushort, VECT_SIZE) u16x;
 typedef VTYPE(uint,   VECT_SIZE) u32x;
-typedef VTYPE(ulong,  VECT_SIZE) u64x;
+typedef VTYPE(ullong, VECT_SIZE) u64x;
 
 #define make_u8x  (u8x)
 #define make_u16x (u16x)
diff --git a/OpenCL/inc_vendor.h b/OpenCL/inc_vendor.h
index d44ac87c9..bbd7e23d8 100644
--- a/OpenCL/inc_vendor.h
+++ b/OpenCL/inc_vendor.h
@@ -32,10 +32,6 @@
 #define LOCAL_AS
 #define KERNEL_FQ   extern "C" __global__
 #elif defined IS_HIP
-#define __device__   __attribute__((device))
-#define __constant__ __attribute__((constant))
-#define __shared__   __attribute__((shared))
-#define __global__   __attribute__((global))
 #define CONSTANT_VK __constant__
 #define CONSTANT_AS
 #define GLOBAL_AS
diff --git a/OpenCL/inc_zip_inflate.cl b/OpenCL/inc_zip_inflate.cl
index d05f6a792..fcc31e76b 100644
--- a/OpenCL/inc_zip_inflate.cl
+++ b/OpenCL/inc_zip_inflate.cl
@@ -73,18 +73,16 @@ enum{
     MZ_VERSION_ERROR = -6,
     MZ_PARAM_ERROR = -10000
 };
-typedef unsigned long mz_ulong;
+typedef ullong mz_ulong;
 
 #ifndef MINIZ_NO_ZLIB_COMPATIBLE_NAMES
 typedef unsigned char Byte;
 typedef unsigned int uInt;
-typedef mz_ulong uLong;
 typedef Byte Bytef;
 typedef uInt uIntf;
 typedef char charf;
 typedef int intf;
 typedef void *voidpf;
-typedef uLong uLongf;
 typedef void *voidp;
 typedef void *const voidpc;
 #define Z_NULL 0
diff --git a/src/modules/module_17200.c b/src/modules/module_17200.c
index 44798110d..fa510ba27 100644
--- a/src/modules/module_17200.c
+++ b/src/modules/module_17200.c
@@ -170,11 +170,6 @@ bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE
   // it leads to CL_KERNEL_WORK_GROUP_SIZE to return 0 and later we will divide with 0
   // workaround would be to rewrite kernel to use global memory
 
-  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
-  {
-    return true;
-  }
-
   if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD)
   {
     return true;
diff --git a/src/modules/module_17220.c b/src/modules/module_17220.c
index 9028040d9..8f1beaf1a 100644
--- a/src/modules/module_17220.c
+++ b/src/modules/module_17220.c
@@ -170,11 +170,6 @@ bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE
   // it leads to CL_KERNEL_WORK_GROUP_SIZE to return 0 and later we will divide with 0
   // workaround would be to rewrite kernel to use global memory
 
-  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
-  {
-    return true;
-  }
-
   if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD)
   {
     return true;
diff --git a/src/modules/module_17225.c b/src/modules/module_17225.c
index 75c376c9e..3b3291d5f 100644
--- a/src/modules/module_17225.c
+++ b/src/modules/module_17225.c
@@ -170,11 +170,6 @@ bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE
   // it leads to CL_KERNEL_WORK_GROUP_SIZE to return 0 and later we will divide with 0
   // workaround would be to rewrite kernel to use global memory
 
-  if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD_USE_HIP)
-  {
-    return true;
-  }
-
   if (device_param->opencl_device_vendor_id == VENDOR_ID_AMD)
   {
     return true;