From 970e5f35181296b344b3664b90104750b346825b Mon Sep 17 00:00:00 2001 From: jsteube Date: Fri, 15 Mar 2019 23:27:44 +0100 Subject: [PATCH] Fix -m 6100 in optimized mode for use with REAL_SHM --- OpenCL/m06100_a0-optimized.cl | 24 ++++++++++++++--- OpenCL/m06100_a1-optimized.cl | 24 ++++++++++++++--- OpenCL/m06100_a3-optimized.cl | 50 ++++++++++++++++++++++++++++++----- 3 files changed, 85 insertions(+), 13 deletions(-) diff --git a/OpenCL/m06100_a0-optimized.cl b/OpenCL/m06100_a0-optimized.cl index e2f466263..e53523854 100644 --- a/OpenCL/m06100_a0-optimized.cl +++ b/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; /** diff --git a/OpenCL/m06100_a1-optimized.cl b/OpenCL/m06100_a1-optimized.cl index 02d090baa..764cb85a9 100644 --- a/OpenCL/m06100_a1-optimized.cl +++ b/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; /** diff --git a/OpenCL/m06100_a3-optimized.cl b/OpenCL/m06100_a3-optimized.cl index 6f68d1e77..f6dd778f9 100644 --- a/OpenCL/m06100_a3-optimized.cl +++ b/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; /**