From 3df7d7cf150904ab9781c56eba6a148489a7f3c7 Mon Sep 17 00:00:00 2001 From: djm34 Date: Tue, 11 Aug 2015 04:08:45 +0200 Subject: [PATCH] algo cleaning --- kernel/Lyra2REv2.cl | 32 ++++++++++++++++---------------- kernel/bmw256.cl | 39 ++++++++++++++++++++++++++++++++++++--- 2 files changed, 52 insertions(+), 19 deletions(-) diff --git a/kernel/Lyra2REv2.cl b/kernel/Lyra2REv2.cl index fb2a261b..fd4b1125 100644 --- a/kernel/Lyra2REv2.cl +++ b/kernel/Lyra2REv2.cl @@ -45,8 +45,8 @@ typedef unsigned int sph_u32; typedef int sph_s32; #ifndef __OPENCL_VERSION__ -typedef unsigned long long sph_u64; -typedef long long sph_s64; +typedef unsigned long sph_u64; +typedef long sph_s64; #else typedef unsigned long sph_u64; typedef long sph_s64; @@ -484,22 +484,22 @@ __kernel void search6(__global uchar* hashes, __global uint* output, const ulong __global hash_t *hash = (__global hash_t *)(hashes + (4 * sizeof(ulong)* (get_global_id(0) % MAX_GLOBAL_THREADS))); uint dh[16] = { - (0x40414243), (0x44454647), - (0x48494A4B), (0x4C4D4E4F), - (0x50515253), (0x54555657), - (0x58595A5B), (0x5C5D5E5F), - (0x60616263), (0x64656667), - (0x68696A6B), (0x6C6D6E6F), - (0x70717273), (0x74757677), - (0x78797A7B), (0x7C7D7E7F) + 0x40414243, 0x44454647, + 0x48494A4B, 0x4C4D4E4F, + 0x50515253, 0x54555657, + 0x58595A5B, 0x5C5D5E5F, + 0x60616263, 0x64656667, + 0x68696A6B, 0x6C6D6E6F, + 0x70717273, 0x74757677, + 0x78797A7B, 0x7C7D7E7F }; uint final_s[16] = { - (0xaaaaaaa0), (0xaaaaaaa1), (0xaaaaaaa2), - (0xaaaaaaa3), (0xaaaaaaa4), (0xaaaaaaa5), - (0xaaaaaaa6), (0xaaaaaaa7), (0xaaaaaaa8), - (0xaaaaaaa9), (0xaaaaaaaa), (0xaaaaaaab), - (0xaaaaaaac), (0xaaaaaaad), (0xaaaaaaae), - (0xaaaaaaaf) + 0xaaaaaaa0, 0xaaaaaaa1, 0xaaaaaaa2, + 0xaaaaaaa3, 0xaaaaaaa4, 0xaaaaaaa5, + 0xaaaaaaa6, 0xaaaaaaa7, 0xaaaaaaa8, + 0xaaaaaaa9, 0xaaaaaaaa, 0xaaaaaaab, + 0xaaaaaaac, 0xaaaaaaad, 0xaaaaaaae, + 0xaaaaaaaf }; uint message[16]; diff --git a/kernel/bmw256.cl b/kernel/bmw256.cl index 9d625340..313752b9 100644 --- a/kernel/bmw256.cl +++ b/kernel/bmw256.cl @@ -1,3 +1,36 @@ +/* +* bmw256 kernel implementation. +* +* ==========================(LICENSE BEGIN)============================ +* Copyright (c) 2015 djm34 +* +* +* Permission is hereby granted, free of charge, to any person obtaining +* a copy of this software and associated documentation files (the +* "Software"), to deal in the Software without restriction, including +* without limitation the rights to use, copy, modify, merge, publish, +* distribute, sublicense, and/or sell copies of the Software, and to +* permit persons to whom the Software is furnished to do so, subject to +* the following conditions: +* +* The above copyright notice and this permission notice shall be +* included in all copies or substantial portions of the Software. +* +* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, +* EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF +* MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. +* IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY +* CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, +* TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE +* SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. +* +* ===========================(LICENSE END)============================= +* +* @author djm34 +*/ + + + #define shl(x, n) ((x) << (n)) #define shr(x, n) ((x) >> (n)) //#define SHR(x, n) SHR2(x, n) @@ -20,7 +53,7 @@ #define rs7(x) SPH_ROTL32((x), 27) /* Message expansion function 1 */ -static uint expand32_1(int i, uint *M32, uint *H, uint *Q) +uint expand32_1(int i, uint *M32, uint *H, uint *Q) { return (ss1(Q[i - 16]) + ss2(Q[i - 15]) + ss3(Q[i - 14]) + ss0(Q[i - 13]) @@ -32,7 +65,7 @@ static uint expand32_1(int i, uint *M32, uint *H, uint *Q) } /* Message expansion function 2 */ -static uint expand32_2(int i, uint *M32, uint *H, uint *Q) +uint expand32_2(int i, uint *M32, uint *H, uint *Q) { return (Q[i - 16] + rs1(Q[i - 15]) + Q[i - 14] + rs2(Q[i - 13]) @@ -43,7 +76,7 @@ static uint expand32_2(int i, uint *M32, uint *H, uint *Q) } -static void Compression256(uint *M32, uint *H) +void Compression256(uint *M32, uint *H) { int i;