From 96a368715fa176eb128a0df53969f75488c9ed0c Mon Sep 17 00:00:00 2001 From: jsteube Date: Sat, 19 Dec 2015 18:30:34 +0100 Subject: [PATCH] Fix -m 11700 and -m 11800 --- OpenCL/m11700_a0.cl | 36 ++++++++++++++++++------------------ OpenCL/m11700_a1.cl | 36 ++++++++++++++++++------------------ OpenCL/m11700_a3.cl | 36 ++++++++++++++++++------------------ OpenCL/m11800_a0.cl | 36 ++++++++++++++++++------------------ OpenCL/m11800_a1.cl | 36 ++++++++++++++++++------------------ OpenCL/m11800_a3.cl | 36 ++++++++++++++++++------------------ 6 files changed, 108 insertions(+), 108 deletions(-) diff --git a/OpenCL/m11700_a0.cl b/OpenCL/m11700_a0.cl index 305424de8..93162548d 100644 --- a/OpenCL/m11700_a0.cl +++ b/OpenCL/m11700_a0.cl @@ -2403,14 +2403,14 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m04 (__glo m[6] = hl32_to_64 (w[ 3], w[ 2]); m[7] = hl32_to_64 (w[ 1], w[ 0]); - m[0] = swap32 (m[0]); - m[1] = swap32 (m[1]); - m[2] = swap32 (m[2]); - m[3] = swap32 (m[3]); - m[4] = swap32 (m[4]); - m[5] = swap32 (m[5]); - m[6] = swap32 (m[6]); - m[7] = swap32 (m[7]); + m[0] = swap64 (m[0]); + m[1] = swap64 (m[1]); + m[2] = swap64 (m[2]); + m[3] = swap64 (m[3]); + m[4] = swap64 (m[4]); + m[5] = swap64 (m[5]); + m[6] = swap64 (m[6]); + m[7] = swap64 (m[7]); // state buffer (hash) @@ -2436,7 +2436,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m04 (__glo z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap32 ((u64) (out_len * 8)); + z[7] = swap64 ((u64) (out_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64); @@ -2590,14 +2590,14 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s04 (__glo m[6] = hl32_to_64 (w[ 3], w[ 2]); m[7] = hl32_to_64 (w[ 1], w[ 0]); - m[0] = swap32 (m[0]); - m[1] = swap32 (m[1]); - m[2] = swap32 (m[2]); - m[3] = swap32 (m[3]); - m[4] = swap32 (m[4]); - m[5] = swap32 (m[5]); - m[6] = swap32 (m[6]); - m[7] = swap32 (m[7]); + m[0] = swap64 (m[0]); + m[1] = swap64 (m[1]); + m[2] = swap64 (m[2]); + m[3] = swap64 (m[3]); + m[4] = swap64 (m[4]); + m[5] = swap64 (m[5]); + m[6] = swap64 (m[6]); + m[7] = swap64 (m[7]); // state buffer (hash) @@ -2623,7 +2623,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s04 (__glo z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap32 ((u64) (out_len * 8)); + z[7] = swap64 ((u64) (out_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64); diff --git a/OpenCL/m11700_a1.cl b/OpenCL/m11700_a1.cl index 9dc6cfa2c..1aec6a56d 100644 --- a/OpenCL/m11700_a1.cl +++ b/OpenCL/m11700_a1.cl @@ -2457,14 +2457,14 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m04 (__glo m[6] = hl32_to_64 (w[ 3], w[ 2]); m[7] = hl32_to_64 (w[ 1], w[ 0]); - m[0] = swap32 (m[0]); - m[1] = swap32 (m[1]); - m[2] = swap32 (m[2]); - m[3] = swap32 (m[3]); - m[4] = swap32 (m[4]); - m[5] = swap32 (m[5]); - m[6] = swap32 (m[6]); - m[7] = swap32 (m[7]); + m[0] = swap64 (m[0]); + m[1] = swap64 (m[1]); + m[2] = swap64 (m[2]); + m[3] = swap64 (m[3]); + m[4] = swap64 (m[4]); + m[5] = swap64 (m[5]); + m[6] = swap64 (m[6]); + m[7] = swap64 (m[7]); // state buffer (hash) @@ -2490,7 +2490,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_m04 (__glo z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap32 ((u64) (pw_len * 8)); + z[7] = swap64 ((u64) (pw_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64); @@ -2700,14 +2700,14 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s04 (__glo m[6] = hl32_to_64 (w[ 3], w[ 2]); m[7] = hl32_to_64 (w[ 1], w[ 0]); - m[0] = swap32 (m[0]); - m[1] = swap32 (m[1]); - m[2] = swap32 (m[2]); - m[3] = swap32 (m[3]); - m[4] = swap32 (m[4]); - m[5] = swap32 (m[5]); - m[6] = swap32 (m[6]); - m[7] = swap32 (m[7]); + m[0] = swap64 (m[0]); + m[1] = swap64 (m[1]); + m[2] = swap64 (m[2]); + m[3] = swap64 (m[3]); + m[4] = swap64 (m[4]); + m[5] = swap64 (m[5]); + m[6] = swap64 (m[6]); + m[7] = swap64 (m[7]); // state buffer (hash) @@ -2733,7 +2733,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11700_s04 (__glo z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap32 ((u64) (pw_len * 8)); + z[7] = swap64 ((u64) (pw_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64); diff --git a/OpenCL/m11700_a3.cl b/OpenCL/m11700_a3.cl index 8b5fa9305..f6fa9d776 100644 --- a/OpenCL/m11700_a3.cl +++ b/OpenCL/m11700_a3.cl @@ -2318,14 +2318,14 @@ static void m11700m (__local u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_le m[6] = hl32_to_64 (w[ 3], w[ 2]); m[7] = hl32_to_64 (w[ 1], w[ 0]); - m[0] = swap32 (m[0]); - m[1] = swap32 (m[1]); - m[2] = swap32 (m[2]); - m[3] = swap32 (m[3]); - m[4] = swap32 (m[4]); - m[5] = swap32 (m[5]); - m[6] = swap32 (m[6]); - m[7] = swap32 (m[7]); + m[0] = swap64 (m[0]); + m[1] = swap64 (m[1]); + m[2] = swap64 (m[2]); + m[3] = swap64 (m[3]); + m[4] = swap64 (m[4]); + m[5] = swap64 (m[5]); + m[6] = swap64 (m[6]); + m[7] = swap64 (m[7]); // state buffer (hash) @@ -2351,7 +2351,7 @@ static void m11700m (__local u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_le z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap32 ((u64) (pw_len * 8)); + z[7] = swap64 ((u64) (pw_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64); @@ -2413,14 +2413,14 @@ static void m11700s (__local u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_le m[6] = hl32_to_64 (w[ 3], w[ 2]); m[7] = hl32_to_64 (w[ 1], w[ 0]); - m[0] = swap32 (m[0]); - m[1] = swap32 (m[1]); - m[2] = swap32 (m[2]); - m[3] = swap32 (m[3]); - m[4] = swap32 (m[4]); - m[5] = swap32 (m[5]); - m[6] = swap32 (m[6]); - m[7] = swap32 (m[7]); + m[0] = swap64 (m[0]); + m[1] = swap64 (m[1]); + m[2] = swap64 (m[2]); + m[3] = swap64 (m[3]); + m[4] = swap64 (m[4]); + m[5] = swap64 (m[5]); + m[6] = swap64 (m[6]); + m[7] = swap64 (m[7]); // state buffer (hash) @@ -2446,7 +2446,7 @@ static void m11700s (__local u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_le z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap32 ((u64) (pw_len * 8)); + z[7] = swap64 ((u64) (pw_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64); diff --git a/OpenCL/m11800_a0.cl b/OpenCL/m11800_a0.cl index ed31f1819..e01764ff5 100644 --- a/OpenCL/m11800_a0.cl +++ b/OpenCL/m11800_a0.cl @@ -2403,14 +2403,14 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11800_m04 (__glo m[6] = hl32_to_64 (w[ 3], w[ 2]); m[7] = hl32_to_64 (w[ 1], w[ 0]); - m[0] = swap32 (m[0]); - m[1] = swap32 (m[1]); - m[2] = swap32 (m[2]); - m[3] = swap32 (m[3]); - m[4] = swap32 (m[4]); - m[5] = swap32 (m[5]); - m[6] = swap32 (m[6]); - m[7] = swap32 (m[7]); + m[0] = swap64 (m[0]); + m[1] = swap64 (m[1]); + m[2] = swap64 (m[2]); + m[3] = swap64 (m[3]); + m[4] = swap64 (m[4]); + m[5] = swap64 (m[5]); + m[6] = swap64 (m[6]); + m[7] = swap64 (m[7]); // state buffer (hash) @@ -2436,7 +2436,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11800_m04 (__glo z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap32 ((u64) (out_len * 8)); + z[7] = swap64 ((u64) (out_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64); @@ -2590,14 +2590,14 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11800_s04 (__glo m[6] = hl32_to_64 (w[ 3], w[ 2]); m[7] = hl32_to_64 (w[ 1], w[ 0]); - m[0] = swap32 (m[0]); - m[1] = swap32 (m[1]); - m[2] = swap32 (m[2]); - m[3] = swap32 (m[3]); - m[4] = swap32 (m[4]); - m[5] = swap32 (m[5]); - m[6] = swap32 (m[6]); - m[7] = swap32 (m[7]); + m[0] = swap64 (m[0]); + m[1] = swap64 (m[1]); + m[2] = swap64 (m[2]); + m[3] = swap64 (m[3]); + m[4] = swap64 (m[4]); + m[5] = swap64 (m[5]); + m[6] = swap64 (m[6]); + m[7] = swap64 (m[7]); // state buffer (hash) @@ -2623,7 +2623,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11800_s04 (__glo z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap32 ((u64) (out_len * 8)); + z[7] = swap64 ((u64) (out_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64); diff --git a/OpenCL/m11800_a1.cl b/OpenCL/m11800_a1.cl index 64232d3b7..db967473d 100644 --- a/OpenCL/m11800_a1.cl +++ b/OpenCL/m11800_a1.cl @@ -2458,14 +2458,14 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11800_m04 (__glo m[6] = hl32_to_64 (w[ 3], w[ 2]); m[7] = hl32_to_64 (w[ 1], w[ 0]); - m[0] = swap32 (m[0]); - m[1] = swap32 (m[1]); - m[2] = swap32 (m[2]); - m[3] = swap32 (m[3]); - m[4] = swap32 (m[4]); - m[5] = swap32 (m[5]); - m[6] = swap32 (m[6]); - m[7] = swap32 (m[7]); + m[0] = swap64 (m[0]); + m[1] = swap64 (m[1]); + m[2] = swap64 (m[2]); + m[3] = swap64 (m[3]); + m[4] = swap64 (m[4]); + m[5] = swap64 (m[5]); + m[6] = swap64 (m[6]); + m[7] = swap64 (m[7]); // state buffer (hash) @@ -2491,7 +2491,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11800_m04 (__glo z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap32 ((u64) (pw_len * 8)); + z[7] = swap64 ((u64) (pw_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64); @@ -2701,14 +2701,14 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11800_s04 (__glo m[6] = hl32_to_64 (w[ 3], w[ 2]); m[7] = hl32_to_64 (w[ 1], w[ 0]); - m[0] = swap32 (m[0]); - m[1] = swap32 (m[1]); - m[2] = swap32 (m[2]); - m[3] = swap32 (m[3]); - m[4] = swap32 (m[4]); - m[5] = swap32 (m[5]); - m[6] = swap32 (m[6]); - m[7] = swap32 (m[7]); + m[0] = swap64 (m[0]); + m[1] = swap64 (m[1]); + m[2] = swap64 (m[2]); + m[3] = swap64 (m[3]); + m[4] = swap64 (m[4]); + m[5] = swap64 (m[5]); + m[6] = swap64 (m[6]); + m[7] = swap64 (m[7]); // state buffer (hash) @@ -2734,7 +2734,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11800_s04 (__glo z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap32 ((u64) (pw_len * 8)); + z[7] = swap64 ((u64) (pw_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64); diff --git a/OpenCL/m11800_a3.cl b/OpenCL/m11800_a3.cl index 6a702fc48..faa4e96e7 100644 --- a/OpenCL/m11800_a3.cl +++ b/OpenCL/m11800_a3.cl @@ -2318,14 +2318,14 @@ static void m11800m (__local u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_le m[6] = hl32_to_64 (w[ 3], w[ 2]); m[7] = hl32_to_64 (w[ 1], w[ 0]); - m[0] = swap32 (m[0]); - m[1] = swap32 (m[1]); - m[2] = swap32 (m[2]); - m[3] = swap32 (m[3]); - m[4] = swap32 (m[4]); - m[5] = swap32 (m[5]); - m[6] = swap32 (m[6]); - m[7] = swap32 (m[7]); + m[0] = swap64 (m[0]); + m[1] = swap64 (m[1]); + m[2] = swap64 (m[2]); + m[3] = swap64 (m[3]); + m[4] = swap64 (m[4]); + m[5] = swap64 (m[5]); + m[6] = swap64 (m[6]); + m[7] = swap64 (m[7]); // state buffer (hash) @@ -2351,7 +2351,7 @@ static void m11800m (__local u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_le z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap32 ((u64) (pw_len * 8)); + z[7] = swap64 ((u64) (pw_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64); @@ -2413,14 +2413,14 @@ static void m11800s (__local u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_le m[6] = hl32_to_64 (w[ 3], w[ 2]); m[7] = hl32_to_64 (w[ 1], w[ 0]); - m[0] = swap32 (m[0]); - m[1] = swap32 (m[1]); - m[2] = swap32 (m[2]); - m[3] = swap32 (m[3]); - m[4] = swap32 (m[4]); - m[5] = swap32 (m[5]); - m[6] = swap32 (m[6]); - m[7] = swap32 (m[7]); + m[0] = swap64 (m[0]); + m[1] = swap64 (m[1]); + m[2] = swap64 (m[2]); + m[3] = swap64 (m[3]); + m[4] = swap64 (m[4]); + m[5] = swap64 (m[5]); + m[6] = swap64 (m[6]); + m[7] = swap64 (m[7]); // state buffer (hash) @@ -2446,7 +2446,7 @@ static void m11800s (__local u64 s_sbob_sl64[8][256], u32 w[16], const u32 pw_le z[4] = 0; z[5] = 0; z[6] = 0; - z[7] = swap32 ((u64) (pw_len * 8)); + z[7] = swap64 ((u64) (pw_len * 8)); streebog_g (h, z, s_sbob_sl64); streebog_g (h, m, s_sbob_sl64);