Fix -m 6100 in optimized mode for use with REAL_SHM

OpenCL/m06100_a0-optimized.cl

@@ -15,7 +15,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_whirlpool.cl"
 
-DECLSPEC void whirlpool_transform_transport_vector (const u32x *w, u32x *digest, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
+DECLSPEC void whirlpool_transform_transport_vector (const u32x *w, u32x *digest, SHM_TYPE u32 (*s_Ch)[256], SHM_TYPE u32 (*s_Cl)[256])
 {
   whirlpool_transform_vector (w + 0, w + 4, w + 8, w + 12, digest, s_Ch, s_Cl);
 }
@@ -31,9 +31,11 @@ __kernel void m06100_m04 (KERN_ATTR_RULES ())
   const u64 lsz = get_local_size (0);
 
   /**
-   * shared
+   * Whirlpool shared
    */
 
+  #ifdef REAL_SHM
+
   __local u32 s_Ch[8][256];
   __local u32 s_Cl[8][256];
 
@@ -60,6 +62,13 @@ __kernel void m06100_m04 (KERN_ATTR_RULES ())
 
   barrier (CLK_LOCAL_MEM_FENCE);
 
+  #else
+
+  __constant u32a *s_Ch = Ch;
+  __constant u32a *s_Cl = Cl;
+
+  #endif
+
   if (gid >= gid_max) return;
 
   /**
@@ -162,9 +171,11 @@ __kernel void m06100_s04 (KERN_ATTR_RULES ())
   const u64 lsz = get_local_size (0);
 
   /**
-   * shared
+   * Whirlpool shared
    */
 
+  #ifdef REAL_SHM
+
   __local u32 s_Ch[8][256];
   __local u32 s_Cl[8][256];
 
@@ -191,6 +202,13 @@ __kernel void m06100_s04 (KERN_ATTR_RULES ())
 
   barrier (CLK_LOCAL_MEM_FENCE);
 
+  #else
+
+  __constant u32a *s_Ch = Ch;
+  __constant u32a *s_Cl = Cl;
+
+  #endif
+
   if (gid >= gid_max) return;
 
   /**

OpenCL/m06100_a1-optimized.cl

@@ -13,7 +13,7 @@
 #include "inc_simd.cl"
 #include "inc_hash_whirlpool.cl"
 
-DECLSPEC void whirlpool_transform_transport_vector (const u32x *w, u32x *digest, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
+DECLSPEC void whirlpool_transform_transport_vector (const u32x *w, u32x *digest, SHM_TYPE u32 (*s_Ch)[256], SHM_TYPE u32 (*s_Cl)[256])
 {
   whirlpool_transform_vector (w + 0, w + 4, w + 8, w + 12, digest, s_Ch, s_Cl);
 }
@@ -29,9 +29,11 @@ __kernel void m06100_m04 (KERN_ATTR_BASIC ())
   const u64 lsz = get_local_size (0);
 
   /**
-   * shared
+   * Whirlpool shared
    */
 
+  #ifdef REAL_SHM
+
   __local u32 s_Ch[8][256];
   __local u32 s_Cl[8][256];
 
@@ -58,6 +60,13 @@ __kernel void m06100_m04 (KERN_ATTR_BASIC ())
 
   barrier (CLK_LOCAL_MEM_FENCE);
 
+  #else
+
+  __constant u32a *s_Ch = Ch;
+  __constant u32a *s_Cl = Cl;
+
+  #endif
+
   if (gid >= gid_max) return;
 
   /**
@@ -218,9 +227,11 @@ __kernel void m06100_s04 (KERN_ATTR_BASIC ())
   const u64 lsz = get_local_size (0);
 
   /**
-   * shared
+   * Whirlpool shared
    */
 
+  #ifdef REAL_SHM
+
   __local u32 s_Ch[8][256];
   __local u32 s_Cl[8][256];
 
@@ -247,6 +258,13 @@ __kernel void m06100_s04 (KERN_ATTR_BASIC ())
 
   barrier (CLK_LOCAL_MEM_FENCE);
 
+  #else
+
+  __constant u32a *s_Ch = Ch;
+  __constant u32a *s_Cl = Cl;
+
+  #endif
+
   if (gid >= gid_max) return;
 
   /**

OpenCL/m06100_a3-optimized.cl

@@ -13,12 +13,12 @@
 #include "inc_simd.cl"
 #include "inc_hash_whirlpool.cl"
 
-DECLSPEC void whirlpool_transform_transport_vector (const u32x *w, u32x *digest, __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256])
+DECLSPEC void whirlpool_transform_transport_vector (const u32x *w, u32x *digest, SHM_TYPE u32 (*s_Ch)[256], SHM_TYPE u32 (*s_Cl)[256])
 {
   whirlpool_transform_vector (w + 0, w + 4, w + 8, w + 12, digest, s_Ch, s_Cl);
 }
 
-DECLSPEC void m06100m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_BASIC (), __local u32 (*s_Cl)[256], __local u32 (*s_Ch)[256])
+DECLSPEC void m06100m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_BASIC (), SHM_TYPE u32 (*s_Cl)[256], SHM_TYPE u32 (*s_Ch)[256])
 {
   /**
    * modifier
@@ -87,7 +87,7 @@ DECLSPEC void m06100m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
   }
 }
 
-DECLSPEC void m06100s (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_BASIC (), __local u32 (*s_Cl)[256], __local u32 (*s_Ch)[256])
+DECLSPEC void m06100s (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KERN_ATTR_BASIC (), SHM_TYPE u32 (*s_Cl)[256], SHM_TYPE u32 (*s_Ch)[256])
 {
   /**
    * modifier
@@ -179,9 +179,11 @@ __kernel void m06100_m04 (KERN_ATTR_BASIC ())
   const u64 lsz = get_local_size (0);
 
   /**
-   * shared
+   * Whirlpool shared
    */
 
+  #ifdef REAL_SHM
+
   __local u32 s_Ch[8][256];
   __local u32 s_Cl[8][256];
 
@@ -208,6 +210,13 @@ __kernel void m06100_m04 (KERN_ATTR_BASIC ())
 
   barrier (CLK_LOCAL_MEM_FENCE);
 
+  #else
+
+  __constant u32a *s_Ch = Ch;
+  __constant u32a *s_Cl = Cl;
+
+  #endif
+
   if (gid >= gid_max) return;
 
   /**
@@ -262,9 +271,11 @@ __kernel void m06100_m08 (KERN_ATTR_BASIC ())
   const u64 lsz = get_local_size (0);
 
   /**
-   * shared
+   * Whirlpool shared
    */
 
+  #ifdef REAL_SHM
+
   __local u32 s_Ch[8][256];
   __local u32 s_Cl[8][256];
 
@@ -291,6 +302,13 @@ __kernel void m06100_m08 (KERN_ATTR_BASIC ())
 
   barrier (CLK_LOCAL_MEM_FENCE);
 
+  #else
+
+  __constant u32a *s_Ch = Ch;
+  __constant u32a *s_Cl = Cl;
+
+  #endif
+
   if (gid >= gid_max) return;
 
   /**
@@ -349,9 +367,11 @@ __kernel void m06100_s04 (KERN_ATTR_BASIC ())
   const u64 lsz = get_local_size (0);
 
   /**
-   * shared
+   * Whirlpool shared
    */
 
+  #ifdef REAL_SHM
+
   __local u32 s_Ch[8][256];
   __local u32 s_Cl[8][256];
 
@@ -378,6 +398,13 @@ __kernel void m06100_s04 (KERN_ATTR_BASIC ())
 
   barrier (CLK_LOCAL_MEM_FENCE);
 
+  #else
+
+  __constant u32a *s_Ch = Ch;
+  __constant u32a *s_Cl = Cl;
+
+  #endif
+
   if (gid >= gid_max) return;
 
   /**
@@ -432,9 +459,11 @@ __kernel void m06100_s08 (KERN_ATTR_BASIC ())
   const u64 lsz = get_local_size (0);
 
   /**
-   * shared
+   * Whirlpool shared
    */
 
+  #ifdef REAL_SHM
+
   __local u32 s_Ch[8][256];
   __local u32 s_Cl[8][256];
 
@@ -461,6 +490,13 @@ __kernel void m06100_s08 (KERN_ATTR_BASIC ())
 
   barrier (CLK_LOCAL_MEM_FENCE);
 
+  #else
+
+  __constant u32a *s_Ch = Ch;
+  __constant u32a *s_Cl = Cl;
+
+  #endif
+
   if (gid >= gid_max) return;
 
   /**