From 70d1343d5745c1ec93e1d638a6c269c7ace930d1 Mon Sep 17 00:00:00 2001 From: jsteube Date: Sat, 16 Mar 2019 16:51:54 +0100 Subject: [PATCH] Fix variables s_Ch and s_Cl in whirlpool hashes in non REAL_SHM mode --- OpenCL/m06100_a0-optimized.cl | 8 ++++---- OpenCL/m06100_a0-pure.cl | 8 ++++---- OpenCL/m06100_a1-optimized.cl | 8 ++++---- OpenCL/m06100_a1-pure.cl | 8 ++++---- OpenCL/m06100_a3-optimized.cl | 16 ++++++++-------- OpenCL/m06100_a3-pure.cl | 8 ++++---- OpenCL/m06231-pure.cl | 12 ++++++------ OpenCL/m06232-pure.cl | 12 ++++++------ OpenCL/m06233-pure.cl | 12 ++++++------ OpenCL/m13731-pure.cl | 16 ++++++++-------- OpenCL/m13732-pure.cl | 16 ++++++++-------- OpenCL/m13733-pure.cl | 16 ++++++++-------- 12 files changed, 70 insertions(+), 70 deletions(-) diff --git a/OpenCL/m06100_a0-optimized.cl b/OpenCL/m06100_a0-optimized.cl index e53523854..bccb54b28 100644 --- a/OpenCL/m06100_a0-optimized.cl +++ b/OpenCL/m06100_a0-optimized.cl @@ -64,8 +64,8 @@ __kernel void m06100_m04 (KERN_ATTR_RULES ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -204,8 +204,8 @@ __kernel void m06100_s04 (KERN_ATTR_RULES ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif diff --git a/OpenCL/m06100_a0-pure.cl b/OpenCL/m06100_a0-pure.cl index b32fe287a..020170aa3 100644 --- a/OpenCL/m06100_a0-pure.cl +++ b/OpenCL/m06100_a0-pure.cl @@ -59,8 +59,8 @@ __kernel void m06100_mxx (KERN_ATTR_RULES ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -143,8 +143,8 @@ __kernel void m06100_sxx (KERN_ATTR_RULES ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif diff --git a/OpenCL/m06100_a1-optimized.cl b/OpenCL/m06100_a1-optimized.cl index 764cb85a9..c43528a76 100644 --- a/OpenCL/m06100_a1-optimized.cl +++ b/OpenCL/m06100_a1-optimized.cl @@ -62,8 +62,8 @@ __kernel void m06100_m04 (KERN_ATTR_BASIC ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -260,8 +260,8 @@ __kernel void m06100_s04 (KERN_ATTR_BASIC ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif diff --git a/OpenCL/m06100_a1-pure.cl b/OpenCL/m06100_a1-pure.cl index 4772f68b8..ecc2f199d 100644 --- a/OpenCL/m06100_a1-pure.cl +++ b/OpenCL/m06100_a1-pure.cl @@ -57,8 +57,8 @@ __kernel void m06100_mxx (KERN_ATTR_BASIC ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -139,8 +139,8 @@ __kernel void m06100_sxx (KERN_ATTR_BASIC ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif diff --git a/OpenCL/m06100_a3-optimized.cl b/OpenCL/m06100_a3-optimized.cl index f6dd778f9..fc8cd5504 100644 --- a/OpenCL/m06100_a3-optimized.cl +++ b/OpenCL/m06100_a3-optimized.cl @@ -212,8 +212,8 @@ __kernel void m06100_m04 (KERN_ATTR_BASIC ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -304,8 +304,8 @@ __kernel void m06100_m08 (KERN_ATTR_BASIC ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -400,8 +400,8 @@ __kernel void m06100_s04 (KERN_ATTR_BASIC ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -492,8 +492,8 @@ __kernel void m06100_s08 (KERN_ATTR_BASIC ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif diff --git a/OpenCL/m06100_a3-pure.cl b/OpenCL/m06100_a3-pure.cl index 7f9e3de87..ba002f18a 100644 --- a/OpenCL/m06100_a3-pure.cl +++ b/OpenCL/m06100_a3-pure.cl @@ -57,8 +57,8 @@ __kernel void m06100_mxx (KERN_ATTR_VECTOR ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -152,8 +152,8 @@ __kernel void m06100_sxx (KERN_ATTR_VECTOR ()) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif diff --git a/OpenCL/m06231-pure.cl b/OpenCL/m06231-pure.cl index 1660acc93..1092562b4 100644 --- a/OpenCL/m06231-pure.cl +++ b/OpenCL/m06231-pure.cl @@ -194,8 +194,8 @@ __kernel void m06231_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -406,8 +406,8 @@ __kernel void m06231_loop (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -660,8 +660,8 @@ __kernel void m06231_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif diff --git a/OpenCL/m06232-pure.cl b/OpenCL/m06232-pure.cl index d4515b4a4..2c4cf0bd0 100644 --- a/OpenCL/m06232-pure.cl +++ b/OpenCL/m06232-pure.cl @@ -194,8 +194,8 @@ __kernel void m06232_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -406,8 +406,8 @@ __kernel void m06232_loop (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -660,8 +660,8 @@ __kernel void m06232_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif diff --git a/OpenCL/m06233-pure.cl b/OpenCL/m06233-pure.cl index 3b844392f..30315c8c4 100644 --- a/OpenCL/m06233-pure.cl +++ b/OpenCL/m06233-pure.cl @@ -194,8 +194,8 @@ __kernel void m06233_init (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -406,8 +406,8 @@ __kernel void m06233_loop (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -660,8 +660,8 @@ __kernel void m06233_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif diff --git a/OpenCL/m13731-pure.cl b/OpenCL/m13731-pure.cl index 1b6ac9431..84affaf5e 100644 --- a/OpenCL/m13731-pure.cl +++ b/OpenCL/m13731-pure.cl @@ -5,8 +5,8 @@ //#define NEW_SIMD_CODE -#undef LOCAL_MEM_TYPE -#define LOCAL_MEM_TYPE LOCAL_MEM_TYPE_GLOBAL +//#undef LOCAL_MEM_TYPE +//#define LOCAL_MEM_TYPE LOCAL_MEM_TYPE_GLOBAL #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -238,8 +238,8 @@ __kernel void m13731_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -501,8 +501,8 @@ __kernel void m13731_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -812,8 +812,8 @@ __kernel void m13731_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif diff --git a/OpenCL/m13732-pure.cl b/OpenCL/m13732-pure.cl index 522ab2c1d..105f062f4 100644 --- a/OpenCL/m13732-pure.cl +++ b/OpenCL/m13732-pure.cl @@ -5,8 +5,8 @@ //#define NEW_SIMD_CODE -#undef LOCAL_MEM_TYPE -#define LOCAL_MEM_TYPE LOCAL_MEM_TYPE_GLOBAL +//#undef LOCAL_MEM_TYPE +//#define LOCAL_MEM_TYPE LOCAL_MEM_TYPE_GLOBAL #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -289,8 +289,8 @@ __kernel void m13732_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -552,8 +552,8 @@ __kernel void m13732_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -864,8 +864,8 @@ __kernel void m13732_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif diff --git a/OpenCL/m13733-pure.cl b/OpenCL/m13733-pure.cl index 9a8bbaa3c..affc31f51 100644 --- a/OpenCL/m13733-pure.cl +++ b/OpenCL/m13733-pure.cl @@ -5,8 +5,8 @@ //#define NEW_SIMD_CODE -#undef LOCAL_MEM_TYPE -#define LOCAL_MEM_TYPE LOCAL_MEM_TYPE_GLOBAL +//#undef LOCAL_MEM_TYPE +//#define LOCAL_MEM_TYPE LOCAL_MEM_TYPE_GLOBAL #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -354,8 +354,8 @@ __kernel void m13733_init (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -617,8 +617,8 @@ __kernel void m13733_loop (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif @@ -930,8 +930,8 @@ __kernel void m13733_comp (KERN_ATTR_TMPS_ESALT (vc_tmp_t, vc_t)) #else - __constant u32a *s_Ch = Ch; - __constant u32a *s_Cl = Cl; + __constant u32a (*s_Ch)[256] = Ch; + __constant u32a (*s_Cl)[256] = Cl; #endif