Update module_unstable_warning() for -m 172xx on HIP

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)

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)

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

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

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;

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;

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;