GOSTSec
/
sgminer
mirror of https://github.com/GOSTSec/sgminer
1
0
Fork 0
Code Issues Releases Wiki Activity
Browse Source

Added support for DarkCoin kernel. Modified stats display to handle low difficulty values.

djm34
phm 11 years ago
parent
80e6307033
commit
9585d2a8e6
24 changed files with 8266 additions and 39 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 5
      Makefile.am
  2. 8
      api.c
  3. 2
      configure.ac
  4. 225
      darkcoin.c
  5. 10
      darkcoin.h
  6. 33
      driver-opencl.c
  7. 389
      kernel/aes_helper.cl
  8. 375
      kernel/blake.cl
  9. 380
      kernel/bmw.cl
  10. 328
      kernel/cubehash.cl
  11. 721
      kernel/darkcoin.cl
  12. 169
      kernel/echo.cl
  13. 1418
      kernel/groestl.cl
  14. 320
      kernel/jh.cl
  15. 683
      kernel/keccak.cl
  16. 488
      kernel/luffa.cl
  17. 607
      kernel/shavite.cl
  18. 1694
      kernel/simd.cl
  19. 350
      kernel/skein.cl
  20. 14
      miner.h
  21. 8
      ocl.c
  22. 2
      ocl.h
  23. 74
      sgminer.c
  24. 2
      util.c

5
Makefile.am
Unescape View File

@ -8,7 +8,7 @@ EXTRA_DIST = example.conf m4/gnulib-cache.m4 \ @@ -8,7 +8,7 @@ EXTRA_DIST = example.conf m4/gnulib-cache.m4 \
API.class API.java api-example.c hexdump.c \
doc/API doc/FAQ doc/GPU doc/SCRYPT doc/windows-build.txt
SUBDIRS = lib compat ccan
SUBDIRS = lib compat ccan sph
bin_PROGRAMS = sgminer
@ -16,7 +16,7 @@ sgminer_CPPFLAGS = $(PTHREAD_FLAGS) -fno-strict-aliasing $(JANSSON_CPPFLAGS) @@ -16,7 +16,7 @@ sgminer_CPPFLAGS = $(PTHREAD_FLAGS) -fno-strict-aliasing $(JANSSON_CPPFLAGS)
sgminer_LDFLAGS = $(PTHREAD_FLAGS)
sgminer_LDADD = $(DLOPEN_FLAGS) @LIBCURL_LIBS@ @JANSSON_LIBS@ @PTHREAD_LIBS@ \
@OPENCL_LIBS@ @NCURSES_LIBS@ @PDCURSES_LIBS@ @WS2_LIBS@ \
@MM_LIBS@ @RT_LIBS@ @MATH_LIBS@ lib/libgnu.a ccan/libccan.a
@MM_LIBS@ @RT_LIBS@ @MATH_LIBS@ lib/libgnu.a ccan/libccan.a sph/libsph.a
sgminer_CPPFLAGS += -I$(top_builddir)/lib -I$(top_srcdir)/lib @OPENCL_FLAGS@
@ -42,6 +42,7 @@ sgminer_SOURCES += ocl.c ocl.h @@ -42,6 +42,7 @@ sgminer_SOURCES += ocl.c ocl.h
sgminer_SOURCES += findnonce.c findnonce.h
sgminer_SOURCES += adl.c adl.h adl_functions.h
sgminer_SOURCES += scrypt.c scrypt.h
sgminer_SOURCES += darkcoin.c darkcoin.h
sgminer_SOURCES += kernel/*.cl
bin_SCRIPTS = $(top_srcdir)/kernel/*.cl

8
api.c
Unescape View File

@ -1695,7 +1695,7 @@ static void gpustatus(struct io_data *io_data, int gpu, bool isjson, bool precom @@ -1695,7 +1695,7 @@ static void gpustatus(struct io_data *io_data, int gpu, bool isjson, bool precom
root = api_add_int(root, "Last Share Pool", &last_share_pool, false);
root = api_add_time(root, "Last Share Time", &(cgpu->last_share_pool_time), false);
root = api_add_mhtotal(root, "Total MH", &(cgpu->total_mhashes), false);
root = api_add_int(root, "Diff1 Work", &(cgpu->diff1), false);
root = api_add_double(root, "Diff1 Work", &(cgpu->diff1), false);
root = api_add_diff(root, "Difficulty Accepted", &(cgpu->diff_accepted), false);
root = api_add_diff(root, "Difficulty Rejected", &(cgpu->diff_rejected), false);
root = api_add_diff(root, "Last Share Difficulty", &(cgpu->last_share_diff), false);
@ -1835,7 +1835,7 @@ static void poolstatus(struct io_data *io_data, __maybe_unused SOCKETTYPE c, __m @@ -1835,7 +1835,7 @@ static void poolstatus(struct io_data *io_data, __maybe_unused SOCKETTYPE c, __m
root = api_add_uint(root, "Remote Failures", &(pool->remotefail_occasions), false);
root = api_add_escape(root, "User", pool->rpc_user, false);
root = api_add_time(root, "Last Share Time", &(pool->last_share_time), false);
root = api_add_int(root, "Diff1 Shares", &(pool->diff1), false);
root = api_add_double(root, "Diff1 Shares", &(pool->diff1), false);
if (pool->rpc_proxy) {
root = api_add_const(root, "Proxy Type", proxytype(pool->rpc_proxytype), false);
root = api_add_escape(root, "Proxy", pool->rpc_proxy, false);
@ -1854,7 +1854,7 @@ static void poolstatus(struct io_data *io_data, __maybe_unused SOCKETTYPE c, __m @@ -1854,7 +1854,7 @@ static void poolstatus(struct io_data *io_data, __maybe_unused SOCKETTYPE c, __m
else
root = api_add_const(root, "Stratum URL", BLANK, false);
root = api_add_bool(root, "Has GBT", &(pool->has_gbt), false);
root = api_add_uint64(root, "Best Share", &(pool->best_diff), true);
root = api_add_double(root, "Best Share", &(pool->best_diff), true);
double rejp = (pool->diff_accepted + pool->diff_rejected + pool->diff_stale) ?
(double)(pool->diff_rejected) / (double)(pool->diff_accepted + pool->diff_rejected + pool->diff_stale) : 0;
root = api_add_percent(root, "Pool Rejected%", &rejp, false);
@ -1915,7 +1915,7 @@ static void summary(struct io_data *io_data, __maybe_unused SOCKETTYPE c, __mayb @@ -1915,7 +1915,7 @@ static void summary(struct io_data *io_data, __maybe_unused SOCKETTYPE c, __mayb
root = api_add_diff(root, "Difficulty Accepted", &(total_diff_accepted), true);
root = api_add_diff(root, "Difficulty Rejected", &(total_diff_rejected), true);
root = api_add_diff(root, "Difficulty Stale", &(total_diff_stale), true);
root = api_add_uint64(root, "Best Share", &(best_diff), true);
root = api_add_double(root, "Best Share", &(best_diff), true);
double hwp = (hw_errors + total_diff1) ?
(double)(hw_errors) / (double)(hw_errors + total_diff1) : 0;
root = api_add_percent(root, "Device Hardware%", &hwp, false);

2
configure.ac
Unescape View File

@ -345,6 +345,7 @@ AC_DEFINE_UNQUOTED([ALEXKAROLD_KERNNAME], ["alexkarold"], [Filename for Alexey K @@ -345,6 +345,7 @@ AC_DEFINE_UNQUOTED([ALEXKAROLD_KERNNAME], ["alexkarold"], [Filename for Alexey K
AC_DEFINE_UNQUOTED([CKOLIVAS_KERNNAME], ["ckolivas"], [Filename for original scrypt kernel])
AC_DEFINE_UNQUOTED([ZUIKKIS_KERNNAME], ["zuikkis"], [Filename for Zuikkis' optimised kernel])
AC_DEFINE_UNQUOTED([PSW_KERNNAME], ["psw"], [Filename for psw's experimental kernel])
AC_DEFINE_UNQUOTED([DARKCOIN_KERNNAME], ["darkcoin"], [Filename for DarkCoin optimised kernel])
AC_SUBST(OPENCL_LIBS)
AC_SUBST(OPENCL_FLAGS)
@ -365,6 +366,7 @@ AC_CONFIG_FILES([ @@ -365,6 +366,7 @@ AC_CONFIG_FILES([
compat/Makefile
ccan/Makefile
lib/Makefile
sph/Makefile
])
AC_OUTPUT

225
darkcoin.c
Unescape View File

@ -0,0 +1,225 @@ @@ -0,0 +1,225 @@
/*-
* Copyright 2009 Colin Percival, 2011 ArtForz
* All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions
* are met:
* 1. Redistributions of source code must retain the above copyright
* notice, this list of conditions and the following disclaimer.
* 2. Redistributions in binary form must reproduce the above copyright
* notice, this list of conditions and the following disclaimer in the
* documentation and/or other materials provided with the distribution.
*
* THIS SOFTWARE IS PROVIDED BY THE AUTHOR AND CONTRIBUTORS ``AS IS'' AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
* ARE DISCLAIMED. IN NO EVENT SHALL THE AUTHOR OR CONTRIBUTORS BE LIABLE
* FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
* DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS
* OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION)
* HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
* LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY
* OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF
* SUCH DAMAGE.
*
* This file was originally written by Colin Percival as part of the Tarsnap
* online backup system.
*/
#include "config.h"
#include "miner.h"
#include <stdlib.h>
#include <stdint.h>
#include <string.h>
#include "sph/sph_blake.h"
#include "sph/sph_bmw.h"
#include "sph/sph_groestl.h"
#include "sph/sph_jh.h"
#include "sph/sph_keccak.h"
#include "sph/sph_skein.h"
#include "sph/sph_luffa.h"
#include "sph/sph_cubehash.h"
#include "sph/sph_shavite.h"
#include "sph/sph_simd.h"
#include "sph/sph_echo.h"
/* Move init out of loop, so init once externally, and then use one single memcpy with that bigger memory block */
typedef struct {
sph_blake512_context blake1;
sph_bmw512_context bmw1;
sph_groestl512_context groestl1;
sph_skein512_context skein1;
sph_jh512_context jh1;
sph_keccak512_context keccak1;
sph_luffa512_context luffa1;
sph_cubehash512_context cubehash1;
sph_shavite512_context shavite1;
sph_simd512_context simd1;
sph_echo512_context echo1;
} Xhash_context_holder;
Xhash_context_holder base_contexts;
void init_Xhash_contexts()
{
sph_blake512_init(&base_contexts.blake1);
sph_bmw512_init(&base_contexts.bmw1);
sph_groestl512_init(&base_contexts.groestl1);
sph_skein512_init(&base_contexts.skein1);
sph_jh512_init(&base_contexts.jh1);
sph_keccak512_init(&base_contexts.keccak1);
sph_luffa512_init(&base_contexts.luffa1);
sph_cubehash512_init(&base_contexts.cubehash1);
sph_shavite512_init(&base_contexts.shavite1);
sph_simd512_init(&base_contexts.simd1);
sph_echo512_init(&base_contexts.echo1);
}
/*
* Encode a length len/4 vector of (uint32_t) into a length len vector of
* (unsigned char) in big-endian form. Assumes len is a multiple of 4.
*/
static inline void
be32enc_vect(uint32_t *dst, const uint32_t *src, uint32_t len)
{
uint32_t i;
for (i = 0; i < len; i++)
dst[i] = htobe32(src[i]);
}
inline void xhash(void *state, const void *input)
{
init_Xhash_contexts();
Xhash_context_holder ctx;
uint32_t hashA[16], hashB[16];
//blake-bmw-groestl-sken-jh-meccak-luffa-cubehash-shivite-simd-echo
memcpy(&ctx, &base_contexts, sizeof(base_contexts));
sph_blake512 (&ctx.blake1, input, 80);
sph_blake512_close (&ctx.blake1, hashA);
sph_bmw512 (&ctx.bmw1, hashA, 64);
sph_bmw512_close(&ctx.bmw1, hashB);
sph_groestl512 (&ctx.groestl1, hashB, 64);
sph_groestl512_close(&ctx.groestl1, hashA);
sph_skein512 (&ctx.skein1, hashA, 64);
sph_skein512_close(&ctx.skein1, hashB);
sph_jh512 (&ctx.jh1, hashB, 64);
sph_jh512_close(&ctx.jh1, hashA);
sph_keccak512 (&ctx.keccak1, hashA, 64);
sph_keccak512_close(&ctx.keccak1, hashB);
sph_luffa512 (&ctx.luffa1, hashB, 64);
sph_luffa512_close (&ctx.luffa1, hashA);
sph_cubehash512 (&ctx.cubehash1, hashA, 64);
sph_cubehash512_close(&ctx.cubehash1, hashB);
sph_shavite512 (&ctx.shavite1, hashB, 64);
sph_shavite512_close(&ctx.shavite1, hashA);
sph_simd512 (&ctx.simd1, hashA, 64);
sph_simd512_close(&ctx.simd1, hashB);
sph_echo512 (&ctx.echo1, hashB, 64);
sph_echo512_close(&ctx.echo1, hashA);
memcpy(state, hashA, 32);
}
static const uint32_t diff1targ = 0x0000ffff;
/* Used externally as confirmation of correct OCL code */
int darkcoin_test(unsigned char *pdata, const unsigned char *ptarget, uint32_t nonce)
{
uint32_t tmp_hash7, Htarg = le32toh(((const uint32_t *)ptarget)[7]);
uint32_t data[20], ohash[8];
//char *scratchbuf;
be32enc_vect(data, (const uint32_t *)pdata, 19);
data[19] = htobe32(nonce);
//scratchbuf = alloca(SCRATCHBUF_SIZE);
xhash(ohash, data);
tmp_hash7 = be32toh(ohash[7]);
applog(LOG_DEBUG, "htarget %08lx diff1 %08lx hash %08lx",
(long unsigned int)Htarg,
(long unsigned int)diff1targ,
(long unsigned int)tmp_hash7);
if (tmp_hash7 > diff1targ)
return -1;
if (tmp_hash7 > Htarg)
return 0;
return 1;
}
void darkcoin_regenhash(struct work *work)
{
uint32_t data[20];
char *scratchbuf;
uint32_t *nonce = (uint32_t *)(work->data + 76);
uint32_t *ohash = (uint32_t *)(work->hash);
be32enc_vect(data, (const uint32_t *)work->data, 19);
data[19] = htobe32(*nonce);
xhash(ohash, data);
}
bool scanhash_darkcoin(struct thr_info *thr, const unsigned char __maybe_unused *pmidstate,
unsigned char *pdata, unsigned char __maybe_unused *phash1,
unsigned char __maybe_unused *phash, const unsigned char *ptarget,
uint32_t max_nonce, uint32_t *last_nonce, uint32_t n)
{
uint32_t *nonce = (uint32_t *)(pdata + 76);
char *scratchbuf;
uint32_t data[20];
uint32_t tmp_hash7;
uint32_t Htarg = le32toh(((const uint32_t *)ptarget)[7]);
bool ret = false;
be32enc_vect(data, (const uint32_t *)pdata, 19);
while(1) {
uint32_t ostate[8];
*nonce = ++n;
data[19] = (n);
xhash(ostate, data);
tmp_hash7 = (ostate[7]);
applog(LOG_INFO, "data7 %08lx",
(long unsigned int)data[7]);
if (unlikely(tmp_hash7 <= Htarg)) {
((uint32_t *)pdata)[19] = htobe32(n);
*last_nonce = n;
ret = true;
break;
}
if (unlikely((n >= max_nonce) || thr->work_restart)) {
*last_nonce = n;
break;
}
}
return ret;
}

10
darkcoin.h
Unescape View File

@ -0,0 +1,10 @@ @@ -0,0 +1,10 @@
#ifndef DARKCOIN_H
#define DARKCOIN_H
#include "miner.h"
extern int darkcoin_test(unsigned char *pdata, const unsigned char *ptarget,
uint32_t nonce);
extern void darkcoin_regenhash(struct work *work);
#endif /* DARKCOIN_H */

33
driver-opencl.c
Unescape View File

@ -48,6 +48,7 @@ extern bool opt_loginput; @@ -48,6 +48,7 @@ extern bool opt_loginput;
extern char *opt_kernel_path;
extern int gpur_thr_id;
extern bool opt_noadl;
extern bool is_scrypt;
extern void *miner_thread(void *userdata);
extern int dev_from_id(int thr_id);
@ -206,6 +207,8 @@ static enum cl_kernels select_kernel(char *arg) @@ -206,6 +207,8 @@ static enum cl_kernels select_kernel(char *arg)
return KL_ZUIKKIS;
if (!strcmp(arg, PSW_KERNNAME))
return KL_PSW;
if (!strcmp(arg, DARKCOIN_KERNNAME))
return KL_DARKCOIN;
return KL_NONE;
}
@ -223,6 +226,8 @@ char *set_kernel(char *arg) @@ -223,6 +226,8 @@ char *set_kernel(char *arg)
if (kern == KL_NONE)
return "Invalid parameter to set_kernel";
gpus[device++].kernel = kern;
if (kern >= KL_DARKCOIN)
is_scrypt = false;
while ((nextptr = strtok(NULL, ",")) != NULL) {
kern = select_kernel(nextptr);
@ -1021,7 +1026,7 @@ static cl_int queue_scrypt_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_u @@ -1021,7 +1026,7 @@ static cl_int queue_scrypt_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_u
cl_int status = 0;
le_target = *(cl_uint *)(blk->work->device_target + 28);
clState->cldata = blk->work->data;
memcpy(clState->cldata, blk->work->data, 80);
status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 80, clState->cldata, 0, NULL,NULL);
CL_SET_ARG(clState->CLbuffer0);
@ -1034,6 +1039,26 @@ static cl_int queue_scrypt_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_u @@ -1034,6 +1039,26 @@ static cl_int queue_scrypt_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_u
return status;
}
static cl_int queue_sph_kernel(_clState *clState, dev_blk_ctx *blk, __maybe_unused cl_uint threads)
{
unsigned char *midstate = blk->work->midstate;
cl_kernel *kernel = &clState->kernel;
unsigned int num = 0;
cl_ulong le_target;
cl_int status = 0;
le_target = *(cl_ulong *)(blk->work->device_target + 24);
flip80(clState->cldata, blk->work->data);
status = clEnqueueWriteBuffer(clState->commandQueue, clState->CLbuffer0, true, 0, 80, clState->cldata, 0, NULL,NULL);
CL_SET_ARG(clState->CLbuffer0);
CL_SET_ARG(clState->outputBuffer);
CL_SET_ARG(le_target);
return status;
}
static void set_threads_hashes(unsigned int vectors, unsigned int compute_shaders, int64_t *hashes, size_t *globalThreads,
unsigned int minthreads, __maybe_unused int *intensity, __maybe_unused int *xintensity, __maybe_unused int *rawintensity)
{
@ -1320,6 +1345,9 @@ static bool opencl_thread_prepare(struct thr_info *thr) @@ -1320,6 +1345,9 @@ static bool opencl_thread_prepare(struct thr_info *thr)
case KL_PSW:
cgpu->kname = PSW_KERNNAME;
break;
case KL_DARKCOIN:
cgpu->kname = DARKCOIN_KERNNAME;
break;
default:
break;
}
@ -1355,6 +1383,9 @@ static bool opencl_thread_init(struct thr_info *thr) @@ -1355,6 +1383,9 @@ static bool opencl_thread_init(struct thr_info *thr)
case KL_ZUIKKIS:
thrdata->queue_kernel_parameters = &queue_scrypt_kernel;
break;
case KL_DARKCOIN:
thrdata->queue_kernel_parameters = &queue_sph_kernel;
break;
default:
applog(LOG_ERR, "Failed to choose kernel in opencl_thread_init");
break;

389
kernel/aes_helper.cl
Unescape View File

@ -0,0 +1,389 @@ @@ -0,0 +1,389 @@
#ifndef AES_HELPER_H
#define AES_HELPER_H
/* $Id: aes_helper.c 220 2010-06-09 09:21:50Z tp $ */
/*
* AES tables. This file is not meant to be compiled by itself; it
* is included by some hash function implementations. It contains
* the precomputed tables and helper macros for evaluating an AES
* round, optionally with a final XOR with a subkey.
*
* By default, this file defines the tables and macros for little-endian
* processing (i.e. it is assumed that the input bytes have been read
* from memory and assembled with the little-endian convention). If
* the 'AES_BIG_ENDIAN' macro is defined (to a non-zero integer value)
* when this file is included, then the tables and macros for big-endian
* processing are defined instead. The big-endian tables and macros have
* names distinct from the little-endian tables and macros, hence it is
* possible to have both simultaneously, by including this file twice
* (with and without the AES_BIG_ENDIAN macro).
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2007-2010 Projet RNRT SAPHIR
*
* 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 Thomas Pornin <thomas.pornin@cryptolog.com>
*/
#if AES_BIG_ENDIAN
#define AESx(x) ( ((SPH_C32(x) >> 24) & SPH_C32(0x000000FF)) \
| ((SPH_C32(x) >> 8) & SPH_C32(0x0000FF00)) \
| ((SPH_C32(x) << 8) & SPH_C32(0x00FF0000)) \
| ((SPH_C32(x) << 24) & SPH_C32(0xFF000000)))
#define AES0 AES0_BE
#define AES1 AES1_BE
#define AES2 AES2_BE
#define AES3 AES3_BE
#define AES_ROUND_BE(X0, X1, X2, X3, K0, K1, K2, K3, Y0, Y1, Y2, Y3) do { \
(Y0) = AES0[((X0) >> 24) & 0xFF] \
^ AES1[((X1) >> 16) & 0xFF] \
^ AES2[((X2) >> 8) & 0xFF] \
^ AES3[(X3) & 0xFF] ^ (K0); \
(Y1) = AES0[((X1) >> 24) & 0xFF] \
^ AES1[((X2) >> 16) & 0xFF] \
^ AES2[((X3) >> 8) & 0xFF] \
^ AES3[(X0) & 0xFF] ^ (K1); \
(Y2) = AES0[((X2) >> 24) & 0xFF] \
^ AES1[((X3) >> 16) & 0xFF] \
^ AES2[((X0) >> 8) & 0xFF] \
^ AES3[(X1) & 0xFF] ^ (K2); \
(Y3) = AES0[((X3) >> 24) & 0xFF] \
^ AES1[((X0) >> 16) & 0xFF] \
^ AES2[((X1) >> 8) & 0xFF] \
^ AES3[(X2) & 0xFF] ^ (K3); \
} while (0)
#define AES_ROUND_NOKEY_BE(X0, X1, X2, X3, Y0, Y1, Y2, Y3) \
AES_ROUND_BE(X0, X1, X2, X3, 0, 0, 0, 0, Y0, Y1, Y2, Y3)
#else
#define AESx(x) SPH_C32(x)
#define AES0 AES0_LE
#define AES1 AES1_LE
#define AES2 AES2_LE
#define AES3 AES3_LE
#define AES_ROUND_LE(X0, X1, X2, X3, K0, K1, K2, K3, Y0, Y1, Y2, Y3) do { \
(Y0) = AES0[(X0) & 0xFF] \
^ AES1[((X1) >> 8) & 0xFF] \
^ AES2[((X2) >> 16) & 0xFF] \
^ AES3[((X3) >> 24) & 0xFF] ^ (K0); \
(Y1) = AES0[(X1) & 0xFF] \
^ AES1[((X2) >> 8) & 0xFF] \
^ AES2[((X3) >> 16) & 0xFF] \
^ AES3[((X0) >> 24) & 0xFF] ^ (K1); \
(Y2) = AES0[(X2) & 0xFF] \
^ AES1[((X3) >> 8) & 0xFF] \
^ AES2[((X0) >> 16) & 0xFF] \
^ AES3[((X1) >> 24) & 0xFF] ^ (K2); \
(Y3) = AES0[(X3) & 0xFF] \
^ AES1[((X0) >> 8) & 0xFF] \
^ AES2[((X1) >> 16) & 0xFF] \
^ AES3[((X2) >> 24) & 0xFF] ^ (K3); \
} while (0)
#define AES_ROUND_NOKEY_LE(X0, X1, X2, X3, Y0, Y1, Y2, Y3) \
AES_ROUND_LE(X0, X1, X2, X3, 0, 0, 0, 0, Y0, Y1, Y2, Y3)
#endif
/*
* The AES*[] tables allow us to perform a fast evaluation of an AES
* round; table AESi[] combines SubBytes for a byte at row i, and
* MixColumns for the column where that byte goes after ShiftRows.
*/
__constant static const sph_u32 AES0_C[256] = {
AESx(0xA56363C6), AESx(0x847C7CF8), AESx(0x997777EE), AESx(0x8D7B7BF6),
AESx(0x0DF2F2FF), AESx(0xBD6B6BD6), AESx(0xB16F6FDE), AESx(0x54C5C591),
AESx(0x50303060), AESx(0x03010102), AESx(0xA96767CE), AESx(0x7D2B2B56),
AESx(0x19FEFEE7), AESx(0x62D7D7B5), AESx(0xE6ABAB4D), AESx(0x9A7676EC),
AESx(0x45CACA8F), AESx(0x9D82821F), AESx(0x40C9C989), AESx(0x877D7DFA),
AESx(0x15FAFAEF), AESx(0xEB5959B2), AESx(0xC947478E), AESx(0x0BF0F0FB),
AESx(0xECADAD41), AESx(0x67D4D4B3), AESx(0xFDA2A25F), AESx(0xEAAFAF45),
AESx(0xBF9C9C23), AESx(0xF7A4A453), AESx(0x967272E4), AESx(0x5BC0C09B),
AESx(0xC2B7B775), AESx(0x1CFDFDE1), AESx(0xAE93933D), AESx(0x6A26264C),
AESx(0x5A36366C), AESx(0x413F3F7E), AESx(0x02F7F7F5), AESx(0x4FCCCC83),
AESx(0x5C343468), AESx(0xF4A5A551), AESx(0x34E5E5D1), AESx(0x08F1F1F9),
AESx(0x937171E2), AESx(0x73D8D8AB), AESx(0x53313162), AESx(0x3F15152A),
AESx(0x0C040408), AESx(0x52C7C795), AESx(0x65232346), AESx(0x5EC3C39D),
AESx(0x28181830), AESx(0xA1969637), AESx(0x0F05050A), AESx(0xB59A9A2F),
AESx(0x0907070E), AESx(0x36121224), AESx(0x9B80801B), AESx(0x3DE2E2DF),
AESx(0x26EBEBCD), AESx(0x6927274E), AESx(0xCDB2B27F), AESx(0x9F7575EA),
AESx(0x1B090912), AESx(0x9E83831D), AESx(0x742C2C58), AESx(0x2E1A1A34),
AESx(0x2D1B1B36), AESx(0xB26E6EDC), AESx(0xEE5A5AB4), AESx(0xFBA0A05B),
AESx(0xF65252A4), AESx(0x4D3B3B76), AESx(0x61D6D6B7), AESx(0xCEB3B37D),
AESx(0x7B292952), AESx(0x3EE3E3DD), AESx(0x712F2F5E), AESx(0x97848413),
AESx(0xF55353A6), AESx(0x68D1D1B9), AESx(0x00000000), AESx(0x2CEDEDC1),
AESx(0x60202040), AESx(0x1FFCFCE3), AESx(0xC8B1B179), AESx(0xED5B5BB6),
AESx(0xBE6A6AD4), AESx(0x46CBCB8D), AESx(0xD9BEBE67), AESx(0x4B393972),
AESx(0xDE4A4A94), AESx(0xD44C4C98), AESx(0xE85858B0), AESx(0x4ACFCF85),
AESx(0x6BD0D0BB), AESx(0x2AEFEFC5), AESx(0xE5AAAA4F), AESx(0x16FBFBED),
AESx(0xC5434386), AESx(0xD74D4D9A), AESx(0x55333366), AESx(0x94858511),
AESx(0xCF45458A), AESx(0x10F9F9E9), AESx(0x06020204), AESx(0x817F7FFE),
AESx(0xF05050A0), AESx(0x443C3C78), AESx(0xBA9F9F25), AESx(0xE3A8A84B),
AESx(0xF35151A2), AESx(0xFEA3A35D), AESx(0xC0404080), AESx(0x8A8F8F05),
AESx(0xAD92923F), AESx(0xBC9D9D21), AESx(0x48383870), AESx(0x04F5F5F1),
AESx(0xDFBCBC63), AESx(0xC1B6B677), AESx(0x75DADAAF), AESx(0x63212142),
AESx(0x30101020), AESx(0x1AFFFFE5), AESx(0x0EF3F3FD), AESx(0x6DD2D2BF),
AESx(0x4CCDCD81), AESx(0x140C0C18), AESx(0x35131326), AESx(0x2FECECC3),
AESx(0xE15F5FBE), AESx(0xA2979735), AESx(0xCC444488), AESx(0x3917172E),
AESx(0x57C4C493), AESx(0xF2A7A755), AESx(0x827E7EFC), AESx(0x473D3D7A),
AESx(0xAC6464C8), AESx(0xE75D5DBA), AESx(0x2B191932), AESx(0x957373E6),
AESx(0xA06060C0), AESx(0x98818119), AESx(0xD14F4F9E), AESx(0x7FDCDCA3),
AESx(0x66222244), AESx(0x7E2A2A54), AESx(0xAB90903B), AESx(0x8388880B),
AESx(0xCA46468C), AESx(0x29EEEEC7), AESx(0xD3B8B86B), AESx(0x3C141428),
AESx(0x79DEDEA7), AESx(0xE25E5EBC), AESx(0x1D0B0B16), AESx(0x76DBDBAD),
AESx(0x3BE0E0DB), AESx(0x56323264), AESx(0x4E3A3A74), AESx(0x1E0A0A14),
AESx(0xDB494992), AESx(0x0A06060C), AESx(0x6C242448), AESx(0xE45C5CB8),
AESx(0x5DC2C29F), AESx(0x6ED3D3BD), AESx(0xEFACAC43), AESx(0xA66262C4),
AESx(0xA8919139), AESx(0xA4959531), AESx(0x37E4E4D3), AESx(0x8B7979F2),
AESx(0x32E7E7D5), AESx(0x43C8C88B), AESx(0x5937376E), AESx(0xB76D6DDA),
AESx(0x8C8D8D01), AESx(0x64D5D5B1), AESx(0xD24E4E9C), AESx(0xE0A9A949),
AESx(0xB46C6CD8), AESx(0xFA5656AC), AESx(0x07F4F4F3), AESx(0x25EAEACF),
AESx(0xAF6565CA), AESx(0x8E7A7AF4), AESx(0xE9AEAE47), AESx(0x18080810),
AESx(0xD5BABA6F), AESx(0x887878F0), AESx(0x6F25254A), AESx(0x722E2E5C),
AESx(0x241C1C38), AESx(0xF1A6A657), AESx(0xC7B4B473), AESx(0x51C6C697),
AESx(0x23E8E8CB), AESx(0x7CDDDDA1), AESx(0x9C7474E8), AESx(0x211F1F3E),
AESx(0xDD4B4B96), AESx(0xDCBDBD61), AESx(0x868B8B0D), AESx(0x858A8A0F),
AESx(0x907070E0), AESx(0x423E3E7C), AESx(0xC4B5B571), AESx(0xAA6666CC),
AESx(0xD8484890), AESx(0x05030306), AESx(0x01F6F6F7), AESx(0x120E0E1C),
AESx(0xA36161C2), AESx(0x5F35356A), AESx(0xF95757AE), AESx(0xD0B9B969),
AESx(0x91868617), AESx(0x58C1C199), AESx(0x271D1D3A), AESx(0xB99E9E27),
AESx(0x38E1E1D9), AESx(0x13F8F8EB), AESx(0xB398982B), AESx(0x33111122),
AESx(0xBB6969D2), AESx(0x70D9D9A9), AESx(0x898E8E07), AESx(0xA7949433),
AESx(0xB69B9B2D), AESx(0x221E1E3C), AESx(0x92878715), AESx(0x20E9E9C9),
AESx(0x49CECE87), AESx(0xFF5555AA), AESx(0x78282850), AESx(0x7ADFDFA5),
AESx(0x8F8C8C03), AESx(0xF8A1A159), AESx(0x80898909), AESx(0x170D0D1A),
AESx(0xDABFBF65), AESx(0x31E6E6D7), AESx(0xC6424284), AESx(0xB86868D0),
AESx(0xC3414182), AESx(0xB0999929), AESx(0x772D2D5A), AESx(0x110F0F1E),
AESx(0xCBB0B07B), AESx(0xFC5454A8), AESx(0xD6BBBB6D), AESx(0x3A16162C)
};
__constant static const sph_u32 AES1_C[256] = {
AESx(0x6363C6A5), AESx(0x7C7CF884), AESx(0x7777EE99), AESx(0x7B7BF68D),
AESx(0xF2F2FF0D), AESx(0x6B6BD6BD), AESx(0x6F6FDEB1), AESx(0xC5C59154),
AESx(0x30306050), AESx(0x01010203), AESx(0x6767CEA9), AESx(0x2B2B567D),
AESx(0xFEFEE719), AESx(0xD7D7B562), AESx(0xABAB4DE6), AESx(0x7676EC9A),
AESx(0xCACA8F45), AESx(0x82821F9D), AESx(0xC9C98940), AESx(0x7D7DFA87),
AESx(0xFAFAEF15), AESx(0x5959B2EB), AESx(0x47478EC9), AESx(0xF0F0FB0B),
AESx(0xADAD41EC), AESx(0xD4D4B367), AESx(0xA2A25FFD), AESx(0xAFAF45EA),
AESx(0x9C9C23BF), AESx(0xA4A453F7), AESx(0x7272E496), AESx(0xC0C09B5B),
AESx(0xB7B775C2), AESx(0xFDFDE11C), AESx(0x93933DAE), AESx(0x26264C6A),
AESx(0x36366C5A), AESx(0x3F3F7E41), AESx(0xF7F7F502), AESx(0xCCCC834F),
AESx(0x3434685C), AESx(0xA5A551F4), AESx(0xE5E5D134), AESx(0xF1F1F908),
AESx(0x7171E293), AESx(0xD8D8AB73), AESx(0x31316253), AESx(0x15152A3F),
AESx(0x0404080C), AESx(0xC7C79552), AESx(0x23234665), AESx(0xC3C39D5E),
AESx(0x18183028), AESx(0x969637A1), AESx(0x05050A0F), AESx(0x9A9A2FB5),
AESx(0x07070E09), AESx(0x12122436), AESx(0x80801B9B), AESx(0xE2E2DF3D),
AESx(0xEBEBCD26), AESx(0x27274E69), AESx(0xB2B27FCD), AESx(0x7575EA9F),
AESx(0x0909121B), AESx(0x83831D9E), AESx(0x2C2C5874), AESx(0x1A1A342E),
AESx(0x1B1B362D), AESx(0x6E6EDCB2), AESx(0x5A5AB4EE), AESx(0xA0A05BFB),
AESx(0x5252A4F6), AESx(0x3B3B764D), AESx(0xD6D6B761), AESx(0xB3B37DCE),
AESx(0x2929527B), AESx(0xE3E3DD3E), AESx(0x2F2F5E71), AESx(0x84841397),
AESx(0x5353A6F5), AESx(0xD1D1B968), AESx(0x00000000), AESx(0xEDEDC12C),
AESx(0x20204060), AESx(0xFCFCE31F), AESx(0xB1B179C8), AESx(0x5B5BB6ED),
AESx(0x6A6AD4BE), AESx(0xCBCB8D46), AESx(0xBEBE67D9), AESx(0x3939724B),
AESx(0x4A4A94DE), AESx(0x4C4C98D4), AESx(0x5858B0E8), AESx(0xCFCF854A),
AESx(0xD0D0BB6B), AESx(0xEFEFC52A), AESx(0xAAAA4FE5), AESx(0xFBFBED16),
AESx(0x434386C5), AESx(0x4D4D9AD7), AESx(0x33336655), AESx(0x85851194),
AESx(0x45458ACF), AESx(0xF9F9E910), AESx(0x02020406), AESx(0x7F7FFE81),
AESx(0x5050A0F0), AESx(0x3C3C7844), AESx(0x9F9F25BA), AESx(0xA8A84BE3),
AESx(0x5151A2F3), AESx(0xA3A35DFE), AESx(0x404080C0), AESx(0x8F8F058A),
AESx(0x92923FAD), AESx(0x9D9D21BC), AESx(0x38387048), AESx(0xF5F5F104),
AESx(0xBCBC63DF), AESx(0xB6B677C1), AESx(0xDADAAF75), AESx(0x21214263),
AESx(0x10102030), AESx(0xFFFFE51A), AESx(0xF3F3FD0E), AESx(0xD2D2BF6D),
AESx(0xCDCD814C), AESx(0x0C0C1814), AESx(0x13132635), AESx(0xECECC32F),
AESx(0x5F5FBEE1), AESx(0x979735A2), AESx(0x444488CC), AESx(0x17172E39),
AESx(0xC4C49357), AESx(0xA7A755F2), AESx(0x7E7EFC82), AESx(0x3D3D7A47),
AESx(0x6464C8AC), AESx(0x5D5DBAE7), AESx(0x1919322B), AESx(0x7373E695),
AESx(0x6060C0A0), AESx(0x81811998), AESx(0x4F4F9ED1), AESx(0xDCDCA37F),
AESx(0x22224466), AESx(0x2A2A547E), AESx(0x90903BAB), AESx(0x88880B83),
AESx(0x46468CCA), AESx(0xEEEEC729), AESx(0xB8B86BD3), AESx(0x1414283C),
AESx(0xDEDEA779), AESx(0x5E5EBCE2), AESx(0x0B0B161D), AESx(0xDBDBAD76),
AESx(0xE0E0DB3B), AESx(0x32326456), AESx(0x3A3A744E), AESx(0x0A0A141E),
AESx(0x494992DB), AESx(0x06060C0A), AESx(0x2424486C), AESx(0x5C5CB8E4),
AESx(0xC2C29F5D), AESx(0xD3D3BD6E), AESx(0xACAC43EF), AESx(0x6262C4A6),
AESx(0x919139A8), AESx(0x959531A4), AESx(0xE4E4D337), AESx(0x7979F28B),
AESx(0xE7E7D532), AESx(0xC8C88B43), AESx(0x37376E59), AESx(0x6D6DDAB7),
AESx(0x8D8D018C), AESx(0xD5D5B164), AESx(0x4E4E9CD2), AESx(0xA9A949E0),
AESx(0x6C6CD8B4), AESx(0x5656ACFA), AESx(0xF4F4F307), AESx(0xEAEACF25),
AESx(0x6565CAAF), AESx(0x7A7AF48E), AESx(0xAEAE47E9), AESx(0x08081018),
AESx(0xBABA6FD5), AESx(0x7878F088), AESx(0x25254A6F), AESx(0x2E2E5C72),
AESx(0x1C1C3824), AESx(0xA6A657F1), AESx(0xB4B473C7), AESx(0xC6C69751),
AESx(0xE8E8CB23), AESx(0xDDDDA17C), AESx(0x7474E89C), AESx(0x1F1F3E21),
AESx(0x4B4B96DD), AESx(0xBDBD61DC), AESx(0x8B8B0D86), AESx(0x8A8A0F85),
AESx(0x7070E090), AESx(0x3E3E7C42), AESx(0xB5B571C4), AESx(0x6666CCAA),
AESx(0x484890D8), AESx(0x03030605), AESx(0xF6F6F701), AESx(0x0E0E1C12),
AESx(0x6161C2A3), AESx(0x35356A5F), AESx(0x5757AEF9), AESx(0xB9B969D0),
AESx(0x86861791), AESx(0xC1C19958), AESx(0x1D1D3A27), AESx(0x9E9E27B9),
AESx(0xE1E1D938), AESx(0xF8F8EB13), AESx(0x98982BB3), AESx(0x11112233),
AESx(0x6969D2BB), AESx(0xD9D9A970), AESx(0x8E8E0789), AESx(0x949433A7),
AESx(0x9B9B2DB6), AESx(0x1E1E3C22), AESx(0x87871592), AESx(0xE9E9C920),
AESx(0xCECE8749), AESx(0x5555AAFF), AESx(0x28285078), AESx(0xDFDFA57A),
AESx(0x8C8C038F), AESx(0xA1A159F8), AESx(0x89890980), AESx(0x0D0D1A17),
AESx(0xBFBF65DA), AESx(0xE6E6D731), AESx(0x424284C6), AESx(0x6868D0B8),
AESx(0x414182C3), AESx(0x999929B0), AESx(0x2D2D5A77), AESx(0x0F0F1E11),
AESx(0xB0B07BCB), AESx(0x5454A8FC), AESx(0xBBBB6DD6), AESx(0x16162C3A)
};
__constant static const sph_u32 AES2_C[256] = {
AESx(0x63C6A563), AESx(0x7CF8847C), AESx(0x77EE9977), AESx(0x7BF68D7B),
AESx(0xF2FF0DF2), AESx(0x6BD6BD6B), AESx(0x6FDEB16F), AESx(0xC59154C5),
AESx(0x30605030), AESx(0x01020301), AESx(0x67CEA967), AESx(0x2B567D2B),
AESx(0xFEE719FE), AESx(0xD7B562D7), AESx(0xAB4DE6AB), AESx(0x76EC9A76),
AESx(0xCA8F45CA), AESx(0x821F9D82), AESx(0xC98940C9), AESx(0x7DFA877D),
AESx(0xFAEF15FA), AESx(0x59B2EB59), AESx(0x478EC947), AESx(0xF0FB0BF0),
AESx(0xAD41ECAD), AESx(0xD4B367D4), AESx(0xA25FFDA2), AESx(0xAF45EAAF),
AESx(0x9C23BF9C), AESx(0xA453F7A4), AESx(0x72E49672), AESx(0xC09B5BC0),
AESx(0xB775C2B7), AESx(0xFDE11CFD), AESx(0x933DAE93), AESx(0x264C6A26),
AESx(0x366C5A36), AESx(0x3F7E413F), AESx(0xF7F502F7), AESx(0xCC834FCC),
AESx(0x34685C34), AESx(0xA551F4A5), AESx(0xE5D134E5), AESx(0xF1F908F1),
AESx(0x71E29371), AESx(0xD8AB73D8), AESx(0x31625331), AESx(0x152A3F15),
AESx(0x04080C04), AESx(0xC79552C7), AESx(0x23466523), AESx(0xC39D5EC3),
AESx(0x18302818), AESx(0x9637A196), AESx(0x050A0F05), AESx(0x9A2FB59A),
AESx(0x070E0907), AESx(0x12243612), AESx(0x801B9B80), AESx(0xE2DF3DE2),
AESx(0xEBCD26EB), AESx(0x274E6927), AESx(0xB27FCDB2), AESx(0x75EA9F75),
AESx(0x09121B09), AESx(0x831D9E83), AESx(0x2C58742C), AESx(0x1A342E1A),
AESx(0x1B362D1B), AESx(0x6EDCB26E), AESx(0x5AB4EE5A), AESx(0xA05BFBA0),
AESx(0x52A4F652), AESx(0x3B764D3B), AESx(0xD6B761D6), AESx(0xB37DCEB3),
AESx(0x29527B29), AESx(0xE3DD3EE3), AESx(0x2F5E712F), AESx(0x84139784),
AESx(0x53A6F553), AESx(0xD1B968D1), AESx(0x00000000), AESx(0xEDC12CED),
AESx(0x20406020), AESx(0xFCE31FFC), AESx(0xB179C8B1), AESx(0x5BB6ED5B),
AESx(0x6AD4BE6A), AESx(0xCB8D46CB), AESx(0xBE67D9BE), AESx(0x39724B39),
AESx(0x4A94DE4A), AESx(0x4C98D44C), AESx(0x58B0E858), AESx(0xCF854ACF),
AESx(0xD0BB6BD0), AESx(0xEFC52AEF), AESx(0xAA4FE5AA), AESx(0xFBED16FB),
AESx(0x4386C543), AESx(0x4D9AD74D), AESx(0x33665533), AESx(0x85119485),
AESx(0x458ACF45), AESx(0xF9E910F9), AESx(0x02040602), AESx(0x7FFE817F),
AESx(0x50A0F050), AESx(0x3C78443C), AESx(0x9F25BA9F), AESx(0xA84BE3A8),
AESx(0x51A2F351), AESx(0xA35DFEA3), AESx(0x4080C040), AESx(0x8F058A8F),
AESx(0x923FAD92), AESx(0x9D21BC9D), AESx(0x38704838), AESx(0xF5F104F5),
AESx(0xBC63DFBC), AESx(0xB677C1B6), AESx(0xDAAF75DA), AESx(0x21426321),
AESx(0x10203010), AESx(0xFFE51AFF), AESx(0xF3FD0EF3), AESx(0xD2BF6DD2),
AESx(0xCD814CCD), AESx(0x0C18140C), AESx(0x13263513), AESx(0xECC32FEC),
AESx(0x5FBEE15F), AESx(0x9735A297), AESx(0x4488CC44), AESx(0x172E3917),
AESx(0xC49357C4), AESx(0xA755F2A7), AESx(0x7EFC827E), AESx(0x3D7A473D),
AESx(0x64C8AC64), AESx(0x5DBAE75D), AESx(0x19322B19), AESx(0x73E69573),
AESx(0x60C0A060), AESx(0x81199881), AESx(0x4F9ED14F), AESx(0xDCA37FDC),
AESx(0x22446622), AESx(0x2A547E2A), AESx(0x903BAB90), AESx(0x880B8388),
AESx(0x468CCA46), AESx(0xEEC729EE), AESx(0xB86BD3B8), AESx(0x14283C14),
AESx(0xDEA779DE), AESx(0x5EBCE25E), AESx(0x0B161D0B), AESx(0xDBAD76DB),
AESx(0xE0DB3BE0), AESx(0x32645632), AESx(0x3A744E3A), AESx(0x0A141E0A),
AESx(0x4992DB49), AESx(0x060C0A06), AESx(0x24486C24), AESx(0x5CB8E45C),
AESx(0xC29F5DC2), AESx(0xD3BD6ED3), AESx(0xAC43EFAC), AESx(0x62C4A662),
AESx(0x9139A891), AESx(0x9531A495), AESx(0xE4D337E4), AESx(0x79F28B79),
AESx(0xE7D532E7), AESx(0xC88B43C8), AESx(0x376E5937), AESx(0x6DDAB76D),
AESx(0x8D018C8D), AESx(0xD5B164D5), AESx(0x4E9CD24E), AESx(0xA949E0A9),
AESx(0x6CD8B46C), AESx(0x56ACFA56), AESx(0xF4F307F4), AESx(0xEACF25EA),
AESx(0x65CAAF65), AESx(0x7AF48E7A), AESx(0xAE47E9AE), AESx(0x08101808),
AESx(0xBA6FD5BA), AESx(0x78F08878), AESx(0x254A6F25), AESx(0x2E5C722E),
AESx(0x1C38241C), AESx(0xA657F1A6), AESx(0xB473C7B4), AESx(0xC69751C6),
AESx(0xE8CB23E8), AESx(0xDDA17CDD), AESx(0x74E89C74), AESx(0x1F3E211F),
AESx(0x4B96DD4B), AESx(0xBD61DCBD), AESx(0x8B0D868B), AESx(0x8A0F858A),
AESx(0x70E09070), AESx(0x3E7C423E), AESx(0xB571C4B5), AESx(0x66CCAA66),
AESx(0x4890D848), AESx(0x03060503), AESx(0xF6F701F6), AESx(0x0E1C120E),
AESx(0x61C2A361), AESx(0x356A5F35), AESx(0x57AEF957), AESx(0xB969D0B9),
AESx(0x86179186), AESx(0xC19958C1), AESx(0x1D3A271D), AESx(0x9E27B99E),
AESx(0xE1D938E1), AESx(0xF8EB13F8), AESx(0x982BB398), AESx(0x11223311),
AESx(0x69D2BB69), AESx(0xD9A970D9), AESx(0x8E07898E), AESx(0x9433A794),
AESx(0x9B2DB69B), AESx(0x1E3C221E), AESx(0x87159287), AESx(0xE9C920E9),
AESx(0xCE8749CE), AESx(0x55AAFF55), AESx(0x28507828), AESx(0xDFA57ADF),
AESx(0x8C038F8C), AESx(0xA159F8A1), AESx(0x89098089), AESx(0x0D1A170D),
AESx(0xBF65DABF), AESx(0xE6D731E6), AESx(0x4284C642), AESx(0x68D0B868),
AESx(0x4182C341), AESx(0x9929B099), AESx(0x2D5A772D), AESx(0x0F1E110F),
AESx(0xB07BCBB0), AESx(0x54A8FC54), AESx(0xBB6DD6BB), AESx(0x162C3A16)
};
__constant static const sph_u32 AES3_C[256] = {
AESx(0xC6A56363), AESx(0xF8847C7C), AESx(0xEE997777), AESx(0xF68D7B7B),
AESx(0xFF0DF2F2), AESx(0xD6BD6B6B), AESx(0xDEB16F6F), AESx(0x9154C5C5),
AESx(0x60503030), AESx(0x02030101), AESx(0xCEA96767), AESx(0x567D2B2B),
AESx(0xE719FEFE), AESx(0xB562D7D7), AESx(0x4DE6ABAB), AESx(0xEC9A7676),
AESx(0x8F45CACA), AESx(0x1F9D8282), AESx(0x8940C9C9), AESx(0xFA877D7D),
AESx(0xEF15FAFA), AESx(0xB2EB5959), AESx(0x8EC94747), AESx(0xFB0BF0F0),
AESx(0x41ECADAD), AESx(0xB367D4D4), AESx(0x5FFDA2A2), AESx(0x45EAAFAF),
AESx(0x23BF9C9C), AESx(0x53F7A4A4), AESx(0xE4967272), AESx(0x9B5BC0C0),
AESx(0x75C2B7B7), AESx(0xE11CFDFD), AESx(0x3DAE9393), AESx(0x4C6A2626),
AESx(0x6C5A3636), AESx(0x7E413F3F), AESx(0xF502F7F7), AESx(0x834FCCCC),
AESx(0x685C3434), AESx(0x51F4A5A5), AESx(0xD134E5E5), AESx(0xF908F1F1),
AESx(0xE2937171), AESx(0xAB73D8D8), AESx(0x62533131), AESx(0x2A3F1515),
AESx(0x080C0404), AESx(0x9552C7C7), AESx(0x46652323), AESx(0x9D5EC3C3),
AESx(0x30281818), AESx(0x37A19696), AESx(0x0A0F0505), AESx(0x2FB59A9A),
AESx(0x0E090707), AESx(0x24361212), AESx(0x1B9B8080), AESx(0xDF3DE2E2),
AESx(0xCD26EBEB), AESx(0x4E692727), AESx(0x7FCDB2B2), AESx(0xEA9F7575),
AESx(0x121B0909), AESx(0x1D9E8383), AESx(0x58742C2C), AESx(0x342E1A1A),
AESx(0x362D1B1B), AESx(0xDCB26E6E), AESx(0xB4EE5A5A), AESx(0x5BFBA0A0),
AESx(0xA4F65252), AESx(0x764D3B3B), AESx(0xB761D6D6), AESx(0x7DCEB3B3),
AESx(0x527B2929), AESx(0xDD3EE3E3), AESx(0x5E712F2F), AESx(0x13978484),
AESx(0xA6F55353), AESx(0xB968D1D1), AESx(0x00000000), AESx(0xC12CEDED),
AESx(0x40602020), AESx(0xE31FFCFC), AESx(0x79C8B1B1), AESx(0xB6ED5B5B),
AESx(0xD4BE6A6A), AESx(0x8D46CBCB), AESx(0x67D9BEBE), AESx(0x724B3939),
AESx(0x94DE4A4A), AESx(0x98D44C4C), AESx(0xB0E85858), AESx(0x854ACFCF),
AESx(0xBB6BD0D0), AESx(0xC52AEFEF), AESx(0x4FE5AAAA), AESx(0xED16FBFB),
AESx(0x86C54343), AESx(0x9AD74D4D), AESx(0x66553333), AESx(0x11948585),
AESx(0x8ACF4545), AESx(0xE910F9F9), AESx(0x04060202), AESx(0xFE817F7F),
AESx(0xA0F05050), AESx(0x78443C3C), AESx(0x25BA9F9F), AESx(0x4BE3A8A8),
AESx(0xA2F35151), AESx(0x5DFEA3A3), AESx(0x80C04040), AESx(0x058A8F8F),
AESx(0x3FAD9292), AESx(0x21BC9D9D), AESx(0x70483838), AESx(0xF104F5F5),
AESx(0x63DFBCBC), AESx(0x77C1B6B6), AESx(0xAF75DADA), AESx(0x42632121),
AESx(0x20301010), AESx(0xE51AFFFF), AESx(0xFD0EF3F3), AESx(0xBF6DD2D2),
AESx(0x814CCDCD), AESx(0x18140C0C), AESx(0x26351313), AESx(0xC32FECEC),
AESx(0xBEE15F5F), AESx(0x35A29797), AESx(0x88CC4444), AESx(0x2E391717),
AESx(0x9357C4C4), AESx(0x55F2A7A7), AESx(0xFC827E7E), AESx(0x7A473D3D),
AESx(0xC8AC6464), AESx(0xBAE75D5D), AESx(0x322B1919), AESx(0xE6957373),
AESx(0xC0A06060), AESx(0x19988181), AESx(0x9ED14F4F), AESx(0xA37FDCDC),
AESx(0x44662222), AESx(0x547E2A2A), AESx(0x3BAB9090), AESx(0x0B838888),
AESx(0x8CCA4646), AESx(0xC729EEEE), AESx(0x6BD3B8B8), AESx(0x283C1414),
AESx(0xA779DEDE), AESx(0xBCE25E5E), AESx(0x161D0B0B), AESx(0xAD76DBDB),
AESx(0xDB3BE0E0), AESx(0x64563232), AESx(0x744E3A3A), AESx(0x141E0A0A),
AESx(0x92DB4949), AESx(0x0C0A0606), AESx(0x486C2424), AESx(0xB8E45C5C),
AESx(0x9F5DC2C2), AESx(0xBD6ED3D3), AESx(0x43EFACAC), AESx(0xC4A66262),
AESx(0x39A89191), AESx(0x31A49595), AESx(0xD337E4E4), AESx(0xF28B7979),
AESx(0xD532E7E7), AESx(0x8B43C8C8), AESx(0x6E593737), AESx(0xDAB76D6D),
AESx(0x018C8D8D), AESx(0xB164D5D5), AESx(0x9CD24E4E), AESx(0x49E0A9A9),
AESx(0xD8B46C6C), AESx(0xACFA5656), AESx(0xF307F4F4), AESx(0xCF25EAEA),
AESx(0xCAAF6565), AESx(0xF48E7A7A), AESx(0x47E9AEAE), AESx(0x10180808),
AESx(0x6FD5BABA), AESx(0xF0887878), AESx(0x4A6F2525), AESx(0x5C722E2E),
AESx(0x38241C1C), AESx(0x57F1A6A6), AESx(0x73C7B4B4), AESx(0x9751C6C6),
AESx(0xCB23E8E8), AESx(0xA17CDDDD), AESx(0xE89C7474), AESx(0x3E211F1F),
AESx(0x96DD4B4B), AESx(0x61DCBDBD), AESx(0x0D868B8B), AESx(0x0F858A8A),
AESx(0xE0907070), AESx(0x7C423E3E), AESx(0x71C4B5B5), AESx(0xCCAA6666),
AESx(0x90D84848), AESx(0x06050303), AESx(0xF701F6F6), AESx(0x1C120E0E),
AESx(0xC2A36161), AESx(0x6A5F3535), AESx(0xAEF95757), AESx(0x69D0B9B9),
AESx(0x17918686), AESx(0x9958C1C1), AESx(0x3A271D1D), AESx(0x27B99E9E),
AESx(0xD938E1E1), AESx(0xEB13F8F8), AESx(0x2BB39898), AESx(0x22331111),
AESx(0xD2BB6969), AESx(0xA970D9D9), AESx(0x07898E8E), AESx(0x33A79494),
AESx(0x2DB69B9B), AESx(0x3C221E1E), AESx(0x15928787), AESx(0xC920E9E9),
AESx(0x8749CECE), AESx(0xAAFF5555), AESx(0x50782828), AESx(0xA57ADFDF),
AESx(0x038F8C8C), AESx(0x59F8A1A1), AESx(0x09808989), AESx(0x1A170D0D),
AESx(0x65DABFBF), AESx(0xD731E6E6), AESx(0x84C64242), AESx(0xD0B86868),
AESx(0x82C34141), AESx(0x29B09999), AESx(0x5A772D2D), AESx(0x1E110F0F),
AESx(0x7BCBB0B0), AESx(0xA8FC5454), AESx(0x6DD6BBBB), AESx(0x2C3A1616)
};
#endif

375
kernel/blake.cl
Unescape View File

@ -0,0 +1,375 @@ @@ -0,0 +1,375 @@
/* $Id: blake.c 252 2011-06-07 17:55:14Z tp $ */
/*
* BLAKE implementation.
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2007-2010 Projet RNRT SAPHIR
*
* 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 Thomas Pornin <thomas.pornin@cryptolog.com>
*/
__constant static const sph_u64 BLAKE_IV512[8] = {
SPH_C64(0x6A09E667F3BCC908), SPH_C64(0xBB67AE8584CAA73B),
SPH_C64(0x3C6EF372FE94F82B), SPH_C64(0xA54FF53A5F1D36F1),
SPH_C64(0x510E527FADE682D1), SPH_C64(0x9B05688C2B3E6C1F),
SPH_C64(0x1F83D9ABFB41BD6B), SPH_C64(0x5BE0CD19137E2179)
};
#define Z00 0
#define Z01 1
#define Z02 2
#define Z03 3
#define Z04 4
#define Z05 5
#define Z06 6
#define Z07 7
#define Z08 8
#define Z09 9
#define Z0A A
#define Z0B B
#define Z0C C
#define Z0D D
#define Z0E E
#define Z0F F
#define Z10 E
#define Z11 A
#define Z12 4
#define Z13 8
#define Z14 9
#define Z15 F
#define Z16 D
#define Z17 6
#define Z18 1
#define Z19 C
#define Z1A 0
#define Z1B 2
#define Z1C B
#define Z1D 7
#define Z1E 5
#define Z1F 3
#define Z20 B
#define Z21 8
#define Z22 C
#define Z23 0
#define Z24 5
#define Z25 2
#define Z26 F
#define Z27 D
#define Z28 A
#define Z29 E
#define Z2A 3
#define Z2B 6
#define Z2C 7
#define Z2D 1
#define Z2E 9
#define Z2F 4
#define Z30 7
#define Z31 9
#define Z32 3
#define Z33 1
#define Z34 D
#define Z35 C
#define Z36 B
#define Z37 E
#define Z38 2
#define Z39 6
#define Z3A 5
#define Z3B A
#define Z3C 4
#define Z3D 0
#define Z3E F
#define Z3F 8
#define Z40 9
#define Z41 0
#define Z42 5
#define Z43 7
#define Z44 2
#define Z45 4
#define Z46 A
#define Z47 F
#define Z48 E
#define Z49 1
#define Z4A B
#define Z4B C
#define Z4C 6
#define Z4D 8
#define Z4E 3
#define Z4F D
#define Z50 2
#define Z51 C
#define Z52 6
#define Z53 A
#define Z54 0
#define Z55 B
#define Z56 8
#define Z57 3
#define Z58 4
#define Z59 D
#define Z5A 7
#define Z5B 5
#define Z5C F
#define Z5D E
#define Z5E 1
#define Z5F 9
#define Z60 C
#define Z61 5
#define Z62 1
#define Z63 F
#define Z64 E
#define Z65 D
#define Z66 4
#define Z67 A
#define Z68 0
#define Z69 7
#define Z6A 6
#define Z6B 3
#define Z6C 9
#define Z6D 2
#define Z6E 8
#define Z6F B
#define Z70 D
#define Z71 B
#define Z72 7
#define Z73 E
#define Z74 C
#define Z75 1
#define Z76 3
#define Z77 9
#define Z78 5
#define Z79 0
#define Z7A F
#define Z7B 4
#define Z7C 8
#define Z7D 6
#define Z7E 2
#define Z7F A
#define Z80 6
#define Z81 F
#define Z82 E
#define Z83 9
#define Z84 B
#define Z85 3
#define Z86 0
#define Z87 8
#define Z88 C
#define Z89 2
#define Z8A D
#define Z8B 7
#define Z8C 1
#define Z8D 4
#define Z8E A
#define Z8F 5
#define Z90 A
#define Z91 2
#define Z92 8
#define Z93 4
#define Z94 7
#define Z95 6
#define Z96 1
#define Z97 5
#define Z98 F
#define Z99 B
#define Z9A 9
#define Z9B E
#define Z9C 3
#define Z9D C
#define Z9E D
#define Z9F 0
#define Mx(r, i) Mx_(Z ## r ## i)
#define Mx_(n) Mx__(n)
#define Mx__(n) M ## n
#define CSx(r, i) CSx_(Z ## r ## i)
#define CSx_(n) CSx__(n)
#define CSx__(n) CS ## n
#define CS0 SPH_C32(0x243F6A88)
#define CS1 SPH_C32(0x85A308D3)
#define CS2 SPH_C32(0x13198A2E)
#define CS3 SPH_C32(0x03707344)
#define CS4 SPH_C32(0xA4093822)
#define CS5 SPH_C32(0x299F31D0)
#define CS6 SPH_C32(0x082EFA98)
#define CS7 SPH_C32(0xEC4E6C89)
#define CS8 SPH_C32(0x452821E6)
#define CS9 SPH_C32(0x38D01377)
#define CSA SPH_C32(0xBE5466CF)
#define CSB SPH_C32(0x34E90C6C)
#define CSC SPH_C32(0xC0AC29B7)
#define CSD SPH_C32(0xC97C50DD)
#define CSE SPH_C32(0x3F84D5B5)
#define CSF SPH_C32(0xB5470917)
#if SPH_64
#define CBx(r, i) CBx_(Z ## r ## i)
#define CBx_(n) CBx__(n)
#define CBx__(n) CB ## n
#define CB0 SPH_C64(0x243F6A8885A308D3)
#define CB1 SPH_C64(0x13198A2E03707344)
#define CB2 SPH_C64(0xA4093822299F31D0)
#define CB3 SPH_C64(0x082EFA98EC4E6C89)
#define CB4 SPH_C64(0x452821E638D01377)
#define CB5 SPH_C64(0xBE5466CF34E90C6C)
#define CB6 SPH_C64(0xC0AC29B7C97C50DD)
#define CB7 SPH_C64(0x3F84D5B5B5470917)
#define CB8 SPH_C64(0x9216D5D98979FB1B)
#define CB9 SPH_C64(0xD1310BA698DFB5AC)
#define CBA SPH_C64(0x2FFD72DBD01ADFB7)
#define CBB SPH_C64(0xB8E1AFED6A267E96)
#define CBC SPH_C64(0xBA7C9045F12C7F99)
#define CBD SPH_C64(0x24A19947B3916CF7)
#define CBE SPH_C64(0x0801F2E2858EFC16)
#define CBF SPH_C64(0x636920D871574E69)
#endif
#if SPH_64
#define GB(m0, m1, c0, c1, a, b, c, d) do { \
a = SPH_T64(a + b + (m0 ^ c1)); \
d = SPH_ROTR64(d ^ a, 32); \
c = SPH_T64(c + d); \
b = SPH_ROTR64(b ^ c, 25); \
a = SPH_T64(a + b + (m1 ^ c0)); \
d = SPH_ROTR64(d ^ a, 16); \
c = SPH_T64(c + d); \
b = SPH_ROTR64(b ^ c, 11); \
} while (0)
#define ROUND_B(r) do { \
GB(Mx(r, 0), Mx(r, 1), CBx(r, 0), CBx(r, 1), V0, V4, V8, VC); \
GB(Mx(r, 2), Mx(r, 3), CBx(r, 2), CBx(r, 3), V1, V5, V9, VD); \
GB(Mx(r, 4), Mx(r, 5), CBx(r, 4), CBx(r, 5), V2, V6, VA, VE); \
GB(Mx(r, 6), Mx(r, 7), CBx(r, 6), CBx(r, 7), V3, V7, VB, VF); \
GB(Mx(r, 8), Mx(r, 9), CBx(r, 8), CBx(r, 9), V0, V5, VA, VF); \
GB(Mx(r, A), Mx(r, B), CBx(r, A), CBx(r, B), V1, V6, VB, VC); \
GB(Mx(r, C), Mx(r, D), CBx(r, C), CBx(r, D), V2, V7, V8, VD); \
GB(Mx(r, E), Mx(r, F), CBx(r, E), CBx(r, F), V3, V4, V9, VE); \
} while (0)
#endif
#if SPH_64
#define BLAKE_DECL_STATE64 \
sph_u64 H0, H1, H2, H3, H4, H5, H6, H7; \
sph_u64 S0, S1, S2, S3, T0, T1;
#define BLAKE_READ_STATE64(state) do { \
H0 = (state)->H[0]; \
H1 = (state)->H[1]; \
H2 = (state)->H[2]; \
H3 = (state)->H[3]; \
H4 = (state)->H[4]; \
H5 = (state)->H[5]; \
H6 = (state)->H[6]; \
H7 = (state)->H[7]; \
S0 = (state)->S[0]; \
S1 = (state)->S[1]; \
S2 = (state)->S[2]; \
S3 = (state)->S[3]; \
T0 = (state)->T0; \
T1 = (state)->T1; \
} while (0)
#define BLAKE_WRITE_STATE64(state) do { \
(state)->H[0] = H0; \
(state)->H[1] = H1; \
(state)->H[2] = H2; \
(state)->H[3] = H3; \
(state)->H[4] = H4; \
(state)->H[5] = H5; \
(state)->H[6] = H6; \
(state)->H[7] = H7; \
(state)->S[0] = S0; \
(state)->S[1] = S1; \
(state)->S[2] = S2; \
(state)->S[3] = S3; \
(state)->T0 = T0; \
(state)->T1 = T1; \
} while (0)
#define COMPRESS64 do { \
V0 = H0; \
V1 = H1; \
V2 = H2; \
V3 = H3; \
V4 = H4; \
V5 = H5; \
V6 = H6; \
V7 = H7; \
V8 = S0 ^ CB0; \
V9 = S1 ^ CB1; \
VA = S2 ^ CB2; \
VB = S3 ^ CB3; \
VC = T0 ^ CB4; \
VD = T0 ^ CB5; \
VE = T1 ^ CB6; \
VF = T1 ^ CB7; \
ROUND_B(0); \
ROUND_B(1); \
ROUND_B(2); \
ROUND_B(3); \
ROUND_B(4); \
ROUND_B(5); \
ROUND_B(6); \
ROUND_B(7); \
ROUND_B(8); \
ROUND_B(9); \
ROUND_B(0); \
ROUND_B(1); \
ROUND_B(2); \
ROUND_B(3); \
ROUND_B(4); \
ROUND_B(5); \
H0 ^= S0 ^ V0 ^ V8; \
H1 ^= S1 ^ V1 ^ V9; \
H2 ^= S2 ^ V2 ^ VA; \
H3 ^= S3 ^ V3 ^ VB; \
H4 ^= S0 ^ V4 ^ VC; \
H5 ^= S1 ^ V5 ^ VD; \
H6 ^= S2 ^ V6 ^ VE; \
H7 ^= S3 ^ V7 ^ VF; \
} while (0)
#endif
__constant static const sph_u64 salt_zero_big[4] = { 0, 0, 0, 0 };
#ifdef __cplusplus
}
#endif

380
kernel/bmw.cl
Unescape View File

@ -0,0 +1,380 @@ @@ -0,0 +1,380 @@
/* $Id: bmw.c 227 2010-06-16 17:28:38Z tp $ */
/*
* BMW implementation.
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2007-2010 Projet RNRT SAPHIR
*
* 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 Thomas Pornin <thomas.pornin@cryptolog.com>
*/
__constant static const sph_u64 BMW_IV512[] = {
SPH_C64(0x8081828384858687), SPH_C64(0x88898A8B8C8D8E8F),
SPH_C64(0x9091929394959697), SPH_C64(0x98999A9B9C9D9E9F),
SPH_C64(0xA0A1A2A3A4A5A6A7), SPH_C64(0xA8A9AAABACADAEAF),
SPH_C64(0xB0B1B2B3B4B5B6B7), SPH_C64(0xB8B9BABBBCBDBEBF),
SPH_C64(0xC0C1C2C3C4C5C6C7), SPH_C64(0xC8C9CACBCCCDCECF),
SPH_C64(0xD0D1D2D3D4D5D6D7), SPH_C64(0xD8D9DADBDCDDDEDF),
SPH_C64(0xE0E1E2E3E4E5E6E7), SPH_C64(0xE8E9EAEBECEDEEEF),
SPH_C64(0xF0F1F2F3F4F5F6F7), SPH_C64(0xF8F9FAFBFCFDFEFF)
};
#define XCAT(x, y) XCAT_(x, y)
#define XCAT_(x, y) x ## y
#define LPAR (
#define I16_16 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15
#define I16_17 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16
#define I16_18 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17
#define I16_19 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18
#define I16_20 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19
#define I16_21 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20
#define I16_22 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21
#define I16_23 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22
#define I16_24 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23
#define I16_25 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24
#define I16_26 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25
#define I16_27 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26
#define I16_28 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27
#define I16_29 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28
#define I16_30 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29
#define I16_31 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30
#define M16_16 0, 1, 3, 4, 7, 10, 11
#define M16_17 1, 2, 4, 5, 8, 11, 12
#define M16_18 2, 3, 5, 6, 9, 12, 13
#define M16_19 3, 4, 6, 7, 10, 13, 14
#define M16_20 4, 5, 7, 8, 11, 14, 15
#define M16_21 5, 6, 8, 9, 12, 15, 16
#define M16_22 6, 7, 9, 10, 13, 0, 1
#define M16_23 7, 8, 10, 11, 14, 1, 2
#define M16_24 8, 9, 11, 12, 15, 2, 3
#define M16_25 9, 10, 12, 13, 0, 3, 4
#define M16_26 10, 11, 13, 14, 1, 4, 5
#define M16_27 11, 12, 14, 15, 2, 5, 6
#define M16_28 12, 13, 15, 16, 3, 6, 7
#define M16_29 13, 14, 0, 1, 4, 7, 8
#define M16_30 14, 15, 1, 2, 5, 8, 9
#define M16_31 15, 16, 2, 3, 6, 9, 10
#define ss0(x) (((x) >> 1) ^ SPH_T32((x) << 3) \
^ SPH_ROTL32(x, 4) ^ SPH_ROTL32(x, 19))
#define ss1(x) (((x) >> 1) ^ SPH_T32((x) << 2) \
^ SPH_ROTL32(x, 8) ^ SPH_ROTL32(x, 23))
#define ss2(x) (((x) >> 2) ^ SPH_T32((x) << 1) \
^ SPH_ROTL32(x, 12) ^ SPH_ROTL32(x, 25))
#define ss3(x) (((x) >> 2) ^ SPH_T32((x) << 2) \
^ SPH_ROTL32(x, 15) ^ SPH_ROTL32(x, 29))
#define ss4(x) (((x) >> 1) ^ (x))
#define ss5(x) (((x) >> 2) ^ (x))
#define rs1(x) SPH_ROTL32(x, 3)
#define rs2(x) SPH_ROTL32(x, 7)
#define rs3(x) SPH_ROTL32(x, 13)
#define rs4(x) SPH_ROTL32(x, 16)
#define rs5(x) SPH_ROTL32(x, 19)
#define rs6(x) SPH_ROTL32(x, 23)
#define rs7(x) SPH_ROTL32(x, 27)
#define Ks(j) SPH_T32((sph_u32)(j) * SPH_C32(0x05555555))
#define add_elt_s(mf, hf, j0m, j1m, j3m, j4m, j7m, j10m, j11m, j16) \
(SPH_T32(SPH_ROTL32(mf(j0m), j1m) + SPH_ROTL32(mf(j3m), j4m) \
- SPH_ROTL32(mf(j10m), j11m) + Ks(j16)) ^ hf(j7m))
#define expand1s_inner(qf, mf, hf, i16, \
i0, i1, i2, i3, i4, i5, i6, i7, i8, \
i9, i10, i11, i12, i13, i14, i15, \
i0m, i1m, i3m, i4m, i7m, i10m, i11m) \
SPH_T32(ss1(qf(i0)) + ss2(qf(i1)) + ss3(qf(i2)) + ss0(qf(i3)) \
+ ss1(qf(i4)) + ss2(qf(i5)) + ss3(qf(i6)) + ss0(qf(i7)) \
+ ss1(qf(i8)) + ss2(qf(i9)) + ss3(qf(i10)) + ss0(qf(i11)) \
+ ss1(qf(i12)) + ss2(qf(i13)) + ss3(qf(i14)) + ss0(qf(i15)) \
+ add_elt_s(mf, hf, i0m, i1m, i3m, i4m, i7m, i10m, i11m, i16))
#define expand1s(qf, mf, hf, i16) \
expand1s_(qf, mf, hf, i16, I16_ ## i16, M16_ ## i16)
#define expand1s_(qf, mf, hf, i16, ix, iy) \
expand1s_inner LPAR qf, mf, hf, i16, ix, iy)
#define expand2s_inner(qf, mf, hf, i16, \
i0, i1, i2, i3, i4, i5, i6, i7, i8, \
i9, i10, i11, i12, i13, i14, i15, \
i0m, i1m, i3m, i4m, i7m, i10m, i11m) \
SPH_T32(qf(i0) + rs1(qf(i1)) + qf(i2) + rs2(qf(i3)) \
+ qf(i4) + rs3(qf(i5)) + qf(i6) + rs4(qf(i7)) \
+ qf(i8) + rs5(qf(i9)) + qf(i10) + rs6(qf(i11)) \
+ qf(i12) + rs7(qf(i13)) + ss4(qf(i14)) + ss5(qf(i15)) \
+ add_elt_s(mf, hf, i0m, i1m, i3m, i4m, i7m, i10m, i11m, i16))
#define expand2s(qf, mf, hf, i16) \
expand2s_(qf, mf, hf, i16, I16_ ## i16, M16_ ## i16)
#define expand2s_(qf, mf, hf, i16, ix, iy) \
expand2s_inner LPAR qf, mf, hf, i16, ix, iy)
#if SPH_64
#define sb0(x) (((x) >> 1) ^ SPH_T64((x) << 3) \
^ SPH_ROTL64(x, 4) ^ SPH_ROTL64(x, 37))
#define sb1(x) (((x) >> 1) ^ SPH_T64((x) << 2) \
^ SPH_ROTL64(x, 13) ^ SPH_ROTL64(x, 43))
#define sb2(x) (((x) >> 2) ^ SPH_T64((x) << 1) \
^ SPH_ROTL64(x, 19) ^ SPH_ROTL64(x, 53))
#define sb3(x) (((x) >> 2) ^ SPH_T64((x) << 2) \
^ SPH_ROTL64(x, 28) ^ SPH_ROTL64(x, 59))
#define sb4(x) (((x) >> 1) ^ (x))
#define sb5(x) (((x) >> 2) ^ (x))
#define rb1(x) SPH_ROTL64(x, 5)
#define rb2(x) SPH_ROTL64(x, 11)
#define rb3(x) SPH_ROTL64(x, 27)
#define rb4(x) SPH_ROTL64(x, 32)
#define rb5(x) SPH_ROTL64(x, 37)
#define rb6(x) SPH_ROTL64(x, 43)
#define rb7(x) SPH_ROTL64(x, 53)
#define Kb(j) SPH_T64((sph_u64)(j) * SPH_C64(0x0555555555555555))
#define add_elt_b(mf, hf, j0m, j1m, j3m, j4m, j7m, j10m, j11m, j16) \
(SPH_T64(SPH_ROTL64(mf(j0m), j1m) + SPH_ROTL64(mf(j3m), j4m) \
- SPH_ROTL64(mf(j10m), j11m) + Kb(j16)) ^ hf(j7m))
#define expand1b_inner(qf, mf, hf, i16, \
i0, i1, i2, i3, i4, i5, i6, i7, i8, \
i9, i10, i11, i12, i13, i14, i15, \
i0m, i1m, i3m, i4m, i7m, i10m, i11m) \
SPH_T64(sb1(qf(i0)) + sb2(qf(i1)) + sb3(qf(i2)) + sb0(qf(i3)) \
+ sb1(qf(i4)) + sb2(qf(i5)) + sb3(qf(i6)) + sb0(qf(i7)) \
+ sb1(qf(i8)) + sb2(qf(i9)) + sb3(qf(i10)) + sb0(qf(i11)) \
+ sb1(qf(i12)) + sb2(qf(i13)) + sb3(qf(i14)) + sb0(qf(i15)) \
+ add_elt_b(mf, hf, i0m, i1m, i3m, i4m, i7m, i10m, i11m, i16))
#define expand1b(qf, mf, hf, i16) \
expand1b_(qf, mf, hf, i16, I16_ ## i16, M16_ ## i16)
#define expand1b_(qf, mf, hf, i16, ix, iy) \
expand1b_inner LPAR qf, mf, hf, i16, ix, iy)
#define expand2b_inner(qf, mf, hf, i16, \
i0, i1, i2, i3, i4, i5, i6, i7, i8, \
i9, i10, i11, i12, i13, i14, i15, \
i0m, i1m, i3m, i4m, i7m, i10m, i11m) \
SPH_T64(qf(i0) + rb1(qf(i1)) + qf(i2) + rb2(qf(i3)) \
+ qf(i4) + rb3(qf(i5)) + qf(i6) + rb4(qf(i7)) \
+ qf(i8) + rb5(qf(i9)) + qf(i10) + rb6(qf(i11)) \
+ qf(i12) + rb7(qf(i13)) + sb4(qf(i14)) + sb5(qf(i15)) \
+ add_elt_b(mf, hf, i0m, i1m, i3m, i4m, i7m, i10m, i11m, i16))
#define expand2b(qf, mf, hf, i16) \
expand2b_(qf, mf, hf, i16, I16_ ## i16, M16_ ## i16)
#define expand2b_(qf, mf, hf, i16, ix, iy) \
expand2b_inner LPAR qf, mf, hf, i16, ix, iy)
#endif
#define MAKE_W(tt, i0, op01, i1, op12, i2, op23, i3, op34, i4) \
tt((M(i0) ^ H(i0)) op01 (M(i1) ^ H(i1)) op12 (M(i2) ^ H(i2)) \
op23 (M(i3) ^ H(i3)) op34 (M(i4) ^ H(i4)))
#define Ws0 MAKE_W(SPH_T32, 5, -, 7, +, 10, +, 13, +, 14)
#define Ws1 MAKE_W(SPH_T32, 6, -, 8, +, 11, +, 14, -, 15)
#define Ws2 MAKE_W(SPH_T32, 0, +, 7, +, 9, -, 12, +, 15)
#define Ws3 MAKE_W(SPH_T32, 0, -, 1, +, 8, -, 10, +, 13)
#define Ws4 MAKE_W(SPH_T32, 1, +, 2, +, 9, -, 11, -, 14)
#define Ws5 MAKE_W(SPH_T32, 3, -, 2, +, 10, -, 12, +, 15)
#define Ws6 MAKE_W(SPH_T32, 4, -, 0, -, 3, -, 11, +, 13)
#define Ws7 MAKE_W(SPH_T32, 1, -, 4, -, 5, -, 12, -, 14)
#define Ws8 MAKE_W(SPH_T32, 2, -, 5, -, 6, +, 13, -, 15)
#define Ws9 MAKE_W(SPH_T32, 0, -, 3, +, 6, -, 7, +, 14)
#define Ws10 MAKE_W(SPH_T32, 8, -, 1, -, 4, -, 7, +, 15)
#define Ws11 MAKE_W(SPH_T32, 8, -, 0, -, 2, -, 5, +, 9)
#define Ws12 MAKE_W(SPH_T32, 1, +, 3, -, 6, -, 9, +, 10)
#define Ws13 MAKE_W(SPH_T32, 2, +, 4, +, 7, +, 10, +, 11)
#define Ws14 MAKE_W(SPH_T32, 3, -, 5, +, 8, -, 11, -, 12)
#define Ws15 MAKE_W(SPH_T32, 12, -, 4, -, 6, -, 9, +, 13)
#define MAKE_Qas do { \
qt[ 0] = SPH_T32(ss0(Ws0 ) + H( 1)); \
qt[ 1] = SPH_T32(ss1(Ws1 ) + H( 2)); \
qt[ 2] = SPH_T32(ss2(Ws2 ) + H( 3)); \
qt[ 3] = SPH_T32(ss3(Ws3 ) + H( 4)); \
qt[ 4] = SPH_T32(ss4(Ws4 ) + H( 5)); \
qt[ 5] = SPH_T32(ss0(Ws5 ) + H( 6)); \
qt[ 6] = SPH_T32(ss1(Ws6 ) + H( 7)); \
qt[ 7] = SPH_T32(ss2(Ws7 ) + H( 8)); \
qt[ 8] = SPH_T32(ss3(Ws8 ) + H( 9)); \
qt[ 9] = SPH_T32(ss4(Ws9 ) + H(10)); \
qt[10] = SPH_T32(ss0(Ws10) + H(11)); \
qt[11] = SPH_T32(ss1(Ws11) + H(12)); \
qt[12] = SPH_T32(ss2(Ws12) + H(13)); \
qt[13] = SPH_T32(ss3(Ws13) + H(14)); \
qt[14] = SPH_T32(ss4(Ws14) + H(15)); \
qt[15] = SPH_T32(ss0(Ws15) + H( 0)); \
} while (0)
#define MAKE_Qbs do { \
qt[16] = expand1s(Qs, M, H, 16); \
qt[17] = expand1s(Qs, M, H, 17); \
qt[18] = expand2s(Qs, M, H, 18); \
qt[19] = expand2s(Qs, M, H, 19); \
qt[20] = expand2s(Qs, M, H, 20); \
qt[21] = expand2s(Qs, M, H, 21); \
qt[22] = expand2s(Qs, M, H, 22); \
qt[23] = expand2s(Qs, M, H, 23); \
qt[24] = expand2s(Qs, M, H, 24); \
qt[25] = expand2s(Qs, M, H, 25); \
qt[26] = expand2s(Qs, M, H, 26); \
qt[27] = expand2s(Qs, M, H, 27); \
qt[28] = expand2s(Qs, M, H, 28); \
qt[29] = expand2s(Qs, M, H, 29); \
qt[30] = expand2s(Qs, M, H, 30); \
qt[31] = expand2s(Qs, M, H, 31); \
} while (0)
#define MAKE_Qs do { \
MAKE_Qas; \
MAKE_Qbs; \
} while (0)
#define Qs(j) (qt[j])
#if SPH_64
#define Wb0 MAKE_W(SPH_T64, 5, -, 7, +, 10, +, 13, +, 14)
#define Wb1 MAKE_W(SPH_T64, 6, -, 8, +, 11, +, 14, -, 15)
#define Wb2 MAKE_W(SPH_T64, 0, +, 7, +, 9, -, 12, +, 15)
#define Wb3 MAKE_W(SPH_T64, 0, -, 1, +, 8, -, 10, +, 13)
#define Wb4 MAKE_W(SPH_T64, 1, +, 2, +, 9, -, 11, -, 14)
#define Wb5 MAKE_W(SPH_T64, 3, -, 2, +, 10, -, 12, +, 15)
#define Wb6 MAKE_W(SPH_T64, 4, -, 0, -, 3, -, 11, +, 13)
#define Wb7 MAKE_W(SPH_T64, 1, -, 4, -, 5, -, 12, -, 14)
#define Wb8 MAKE_W(SPH_T64, 2, -, 5, -, 6, +, 13, -, 15)
#define Wb9 MAKE_W(SPH_T64, 0, -, 3, +, 6, -, 7, +, 14)
#define Wb10 MAKE_W(SPH_T64, 8, -, 1, -, 4, -, 7, +, 15)
#define Wb11 MAKE_W(SPH_T64, 8, -, 0, -, 2, -, 5, +, 9)
#define Wb12 MAKE_W(SPH_T64, 1, +, 3, -, 6, -, 9, +, 10)
#define Wb13 MAKE_W(SPH_T64, 2, +, 4, +, 7, +, 10, +, 11)
#define Wb14 MAKE_W(SPH_T64, 3, -, 5, +, 8, -, 11, -, 12)
#define Wb15 MAKE_W(SPH_T64, 12, -, 4, -, 6, -, 9, +, 13)
#define MAKE_Qab do { \
qt[ 0] = SPH_T64(sb0(Wb0 ) + H( 1)); \
qt[ 1] = SPH_T64(sb1(Wb1 ) + H( 2)); \
qt[ 2] = SPH_T64(sb2(Wb2 ) + H( 3)); \
qt[ 3] = SPH_T64(sb3(Wb3 ) + H( 4)); \
qt[ 4] = SPH_T64(sb4(Wb4 ) + H( 5)); \
qt[ 5] = SPH_T64(sb0(Wb5 ) + H( 6)); \
qt[ 6] = SPH_T64(sb1(Wb6 ) + H( 7)); \
qt[ 7] = SPH_T64(sb2(Wb7 ) + H( 8)); \
qt[ 8] = SPH_T64(sb3(Wb8 ) + H( 9)); \
qt[ 9] = SPH_T64(sb4(Wb9 ) + H(10)); \
qt[10] = SPH_T64(sb0(Wb10) + H(11)); \
qt[11] = SPH_T64(sb1(Wb11) + H(12)); \
qt[12] = SPH_T64(sb2(Wb12) + H(13)); \
qt[13] = SPH_T64(sb3(Wb13) + H(14)); \
qt[14] = SPH_T64(sb4(Wb14) + H(15)); \
qt[15] = SPH_T64(sb0(Wb15) + H( 0)); \
} while (0)
#define MAKE_Qbb do { \
qt[16] = expand1b(Qb, M, H, 16); \
qt[17] = expand1b(Qb, M, H, 17); \
qt[18] = expand2b(Qb, M, H, 18); \
qt[19] = expand2b(Qb, M, H, 19); \
qt[20] = expand2b(Qb, M, H, 20); \
qt[21] = expand2b(Qb, M, H, 21); \
qt[22] = expand2b(Qb, M, H, 22); \
qt[23] = expand2b(Qb, M, H, 23); \
qt[24] = expand2b(Qb, M, H, 24); \
qt[25] = expand2b(Qb, M, H, 25); \
qt[26] = expand2b(Qb, M, H, 26); \
qt[27] = expand2b(Qb, M, H, 27); \
qt[28] = expand2b(Qb, M, H, 28); \
qt[29] = expand2b(Qb, M, H, 29); \
qt[30] = expand2b(Qb, M, H, 30); \
qt[31] = expand2b(Qb, M, H, 31); \
} while (0)
#define MAKE_Qb do { \
MAKE_Qab; \
MAKE_Qbb; \
} while (0)
#define Qb(j) (qt[j])
#endif
#define FOLD(type, mkQ, tt, rol, mf, qf, dhf) do { \
type qt[32], xl, xh; \
mkQ; \
xl = qf(16) ^ qf(17) ^ qf(18) ^ qf(19) \
^ qf(20) ^ qf(21) ^ qf(22) ^ qf(23); \
xh = xl ^ qf(24) ^ qf(25) ^ qf(26) ^ qf(27) \
^ qf(28) ^ qf(29) ^ qf(30) ^ qf(31); \
dhf( 0) = tt(((xh << 5) ^ (qf(16) >> 5) ^ mf( 0)) \
+ (xl ^ qf(24) ^ qf( 0))); \
dhf( 1) = tt(((xh >> 7) ^ (qf(17) << 8) ^ mf( 1)) \
+ (xl ^ qf(25) ^ qf( 1))); \
dhf( 2) = tt(((xh >> 5) ^ (qf(18) << 5) ^ mf( 2)) \
+ (xl ^ qf(26) ^ qf( 2))); \
dhf( 3) = tt(((xh >> 1) ^ (qf(19) << 5) ^ mf( 3)) \
+ (xl ^ qf(27) ^ qf( 3))); \
dhf( 4) = tt(((xh >> 3) ^ (qf(20) << 0) ^ mf( 4)) \
+ (xl ^ qf(28) ^ qf( 4))); \
dhf( 5) = tt(((xh << 6) ^ (qf(21) >> 6) ^ mf( 5)) \
+ (xl ^ qf(29) ^ qf( 5))); \
dhf( 6) = tt(((xh >> 4) ^ (qf(22) << 6) ^ mf( 6)) \
+ (xl ^ qf(30) ^ qf( 6))); \
dhf( 7) = tt(((xh >> 11) ^ (qf(23) << 2) ^ mf( 7)) \
+ (xl ^ qf(31) ^ qf( 7))); \
dhf( 8) = tt(rol(dhf(4), 9) + (xh ^ qf(24) ^ mf( 8)) \
+ ((xl << 8) ^ qf(23) ^ qf( 8))); \
dhf( 9) = tt(rol(dhf(5), 10) + (xh ^ qf(25) ^ mf( 9)) \
+ ((xl >> 6) ^ qf(16) ^ qf( 9))); \
dhf(10) = tt(rol(dhf(6), 11) + (xh ^ qf(26) ^ mf(10)) \
+ ((xl << 6) ^ qf(17) ^ qf(10))); \
dhf(11) = tt(rol(dhf(7), 12) + (xh ^ qf(27) ^ mf(11)) \
+ ((xl << 4) ^ qf(18) ^ qf(11))); \
dhf(12) = tt(rol(dhf(0), 13) + (xh ^ qf(28) ^ mf(12)) \
+ ((xl >> 3) ^ qf(19) ^ qf(12))); \
dhf(13) = tt(rol(dhf(1), 14) + (xh ^ qf(29) ^ mf(13)) \
+ ((xl >> 4) ^ qf(20) ^ qf(13))); \
dhf(14) = tt(rol(dhf(2), 15) + (xh ^ qf(30) ^ mf(14)) \
+ ((xl >> 7) ^ qf(21) ^ qf(14))); \
dhf(15) = tt(rol(dhf(3), 16) + (xh ^ qf(31) ^ mf(15)) \
+ ((xl >> 2) ^ qf(22) ^ qf(15))); \
} while (0)
#define FOLDb FOLD(sph_u64, MAKE_Qb, SPH_T64, SPH_ROTL64, M, Qb, dH)
__constant static const sph_u64 final_b[16] = {
SPH_C64(0xaaaaaaaaaaaaaaa0), SPH_C64(0xaaaaaaaaaaaaaaa1),
SPH_C64(0xaaaaaaaaaaaaaaa2), SPH_C64(0xaaaaaaaaaaaaaaa3),
SPH_C64(0xaaaaaaaaaaaaaaa4), SPH_C64(0xaaaaaaaaaaaaaaa5),
SPH_C64(0xaaaaaaaaaaaaaaa6), SPH_C64(0xaaaaaaaaaaaaaaa7),
SPH_C64(0xaaaaaaaaaaaaaaa8), SPH_C64(0xaaaaaaaaaaaaaaa9),
SPH_C64(0xaaaaaaaaaaaaaaaa), SPH_C64(0xaaaaaaaaaaaaaaab),
SPH_C64(0xaaaaaaaaaaaaaaac), SPH_C64(0xaaaaaaaaaaaaaaad),
SPH_C64(0xaaaaaaaaaaaaaaae), SPH_C64(0xaaaaaaaaaaaaaaaf)
};

328
kernel/cubehash.cl
Unescape View File

@ -0,0 +1,328 @@ @@ -0,0 +1,328 @@
/* $Id: cubehash.c 227 2010-06-16 17:28:38Z tp $ */
/*
* CubeHash implementation.
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2007-2010 Projet RNRT SAPHIR
*
* 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 Thomas Pornin <thomas.pornin@cryptolog.com>
*/
/*
* Some tests were conducted on an Intel Core2 Q6600 (32-bit and 64-bit
* mode), a PowerPC G3, and a MIPS-compatible CPU (Broadcom BCM3302).
* It appears that the optimal settings are:
* -- full unroll, no state copy on the "big" systems (x86, PowerPC)
* -- unroll to 4 or 8, state copy on the "small" system (MIPS)
*/
#if !defined SPH_CUBEHASH_UNROLL
#define SPH_CUBEHASH_UNROLL 0
#endif
__constant static const sph_u32 CUBEHASH_IV512[] = {
SPH_C32(0x2AEA2A61), SPH_C32(0x50F494D4), SPH_C32(0x2D538B8B),
SPH_C32(0x4167D83E), SPH_C32(0x3FEE2313), SPH_C32(0xC701CF8C),
SPH_C32(0xCC39968E), SPH_C32(0x50AC5695), SPH_C32(0x4D42C787),
SPH_C32(0xA647A8B3), SPH_C32(0x97CF0BEF), SPH_C32(0x825B4537),
SPH_C32(0xEEF864D2), SPH_C32(0xF22090C4), SPH_C32(0xD0E5CD33),
SPH_C32(0xA23911AE), SPH_C32(0xFCD398D9), SPH_C32(0x148FE485),
SPH_C32(0x1B017BEF), SPH_C32(0xB6444532), SPH_C32(0x6A536159),
SPH_C32(0x2FF5781C), SPH_C32(0x91FA7934), SPH_C32(0x0DBADEA9),
SPH_C32(0xD65C8A2B), SPH_C32(0xA5A70E75), SPH_C32(0xB1C62456),
SPH_C32(0xBC796576), SPH_C32(0x1921C8F7), SPH_C32(0xE7989AF1),
SPH_C32(0x7795D246), SPH_C32(0xD43E3B44)
};
#define T32 SPH_T32
#define ROTL32 SPH_ROTL32
#define ROUND_EVEN do { \
xg = T32(x0 + xg); \
x0 = ROTL32(x0, 7); \
xh = T32(x1 + xh); \
x1 = ROTL32(x1, 7); \
xi = T32(x2 + xi); \
x2 = ROTL32(x2, 7); \
xj = T32(x3 + xj); \
x3 = ROTL32(x3, 7); \
xk = T32(x4 + xk); \
x4 = ROTL32(x4, 7); \
xl = T32(x5 + xl); \
x5 = ROTL32(x5, 7); \
xm = T32(x6 + xm); \
x6 = ROTL32(x6, 7); \
xn = T32(x7 + xn); \
x7 = ROTL32(x7, 7); \
xo = T32(x8 + xo); \
x8 = ROTL32(x8, 7); \
xp = T32(x9 + xp); \
x9 = ROTL32(x9, 7); \
xq = T32(xa + xq); \
xa = ROTL32(xa, 7); \
xr = T32(xb + xr); \
xb = ROTL32(xb, 7); \
xs = T32(xc + xs); \
xc = ROTL32(xc, 7); \
xt = T32(xd + xt); \
xd = ROTL32(xd, 7); \
xu = T32(xe + xu); \
xe = ROTL32(xe, 7); \
xv = T32(xf + xv); \
xf = ROTL32(xf, 7); \
x8 ^= xg; \
x9 ^= xh; \
xa ^= xi; \
xb ^= xj; \
xc ^= xk; \
xd ^= xl; \
xe ^= xm; \
xf ^= xn; \
x0 ^= xo; \
x1 ^= xp; \
x2 ^= xq; \
x3 ^= xr; \
x4 ^= xs; \
x5 ^= xt; \
x6 ^= xu; \
x7 ^= xv; \
xi = T32(x8 + xi); \
x8 = ROTL32(x8, 11); \
xj = T32(x9 + xj); \
x9 = ROTL32(x9, 11); \
xg = T32(xa + xg); \
xa = ROTL32(xa, 11); \
xh = T32(xb + xh); \
xb = ROTL32(xb, 11); \
xm = T32(xc + xm); \
xc = ROTL32(xc, 11); \
xn = T32(xd + xn); \
xd = ROTL32(xd, 11); \
xk = T32(xe + xk); \
xe = ROTL32(xe, 11); \
xl = T32(xf + xl); \
xf = ROTL32(xf, 11); \
xq = T32(x0 + xq); \
x0 = ROTL32(x0, 11); \
xr = T32(x1 + xr); \
x1 = ROTL32(x1, 11); \
xo = T32(x2 + xo); \
x2 = ROTL32(x2, 11); \
xp = T32(x3 + xp); \
x3 = ROTL32(x3, 11); \
xu = T32(x4 + xu); \
x4 = ROTL32(x4, 11); \
xv = T32(x5 + xv); \
x5 = ROTL32(x5, 11); \
xs = T32(x6 + xs); \
x6 = ROTL32(x6, 11); \
xt = T32(x7 + xt); \
x7 = ROTL32(x7, 11); \
xc ^= xi; \
xd ^= xj; \
xe ^= xg; \
xf ^= xh; \
x8 ^= xm; \
x9 ^= xn; \
xa ^= xk; \
xb ^= xl; \
x4 ^= xq; \
x5 ^= xr; \
x6 ^= xo; \
x7 ^= xp; \
x0 ^= xu; \
x1 ^= xv; \
x2 ^= xs; \
x3 ^= xt; \
} while (0)
#define ROUND_ODD do { \
xj = T32(xc + xj); \
xc = ROTL32(xc, 7); \
xi = T32(xd + xi); \
xd = ROTL32(xd, 7); \
xh = T32(xe + xh); \
xe = ROTL32(xe, 7); \
xg = T32(xf + xg); \
xf = ROTL32(xf, 7); \
xn = T32(x8 + xn); \
x8 = ROTL32(x8, 7); \
xm = T32(x9 + xm); \
x9 = ROTL32(x9, 7); \
xl = T32(xa + xl); \
xa = ROTL32(xa, 7); \
xk = T32(xb + xk); \
xb = ROTL32(xb, 7); \
xr = T32(x4 + xr); \
x4 = ROTL32(x4, 7); \
xq = T32(x5 + xq); \
x5 = ROTL32(x5, 7); \
xp = T32(x6 + xp); \
x6 = ROTL32(x6, 7); \
xo = T32(x7 + xo); \
x7 = ROTL32(x7, 7); \
xv = T32(x0 + xv); \
x0 = ROTL32(x0, 7); \
xu = T32(x1 + xu); \
x1 = ROTL32(x1, 7); \
xt = T32(x2 + xt); \
x2 = ROTL32(x2, 7); \
xs = T32(x3 + xs); \
x3 = ROTL32(x3, 7); \
x4 ^= xj; \
x5 ^= xi; \
x6 ^= xh; \
x7 ^= xg; \
x0 ^= xn; \
x1 ^= xm; \
x2 ^= xl; \
x3 ^= xk; \
xc ^= xr; \
xd ^= xq; \
xe ^= xp; \
xf ^= xo; \
x8 ^= xv; \
x9 ^= xu; \
xa ^= xt; \
xb ^= xs; \
xh = T32(x4 + xh); \
x4 = ROTL32(x4, 11); \
xg = T32(x5 + xg); \
x5 = ROTL32(x5, 11); \
xj = T32(x6 + xj); \
x6 = ROTL32(x6, 11); \
xi = T32(x7 + xi); \
x7 = ROTL32(x7, 11); \
xl = T32(x0 + xl); \
x0 = ROTL32(x0, 11); \
xk = T32(x1 + xk); \
x1 = ROTL32(x1, 11); \
xn = T32(x2 + xn); \
x2 = ROTL32(x2, 11); \
xm = T32(x3 + xm); \
x3 = ROTL32(x3, 11); \
xp = T32(xc + xp); \
xc = ROTL32(xc, 11); \
xo = T32(xd + xo); \
xd = ROTL32(xd, 11); \
xr = T32(xe + xr); \
xe = ROTL32(xe, 11); \
xq = T32(xf + xq); \
xf = ROTL32(xf, 11); \
xt = T32(x8 + xt); \
x8 = ROTL32(x8, 11); \
xs = T32(x9 + xs); \
x9 = ROTL32(x9, 11); \
xv = T32(xa + xv); \
xa = ROTL32(xa, 11); \
xu = T32(xb + xu); \
xb = ROTL32(xb, 11); \
x0 ^= xh; \
x1 ^= xg; \
x2 ^= xj; \
x3 ^= xi; \
x4 ^= xl; \
x5 ^= xk; \
x6 ^= xn; \
x7 ^= xm; \
x8 ^= xp; \
x9 ^= xo; \
xa ^= xr; \
xb ^= xq; \
xc ^= xt; \
xd ^= xs; \
xe ^= xv; \
xf ^= xu; \
} while (0)
/*
* There is no need to unroll all 16 rounds. The word-swapping permutation
* is an involution, so we need to unroll an even number of rounds. On
* "big" systems, unrolling 4 rounds yields about 97% of the speed
* achieved with full unrolling; and it keeps the code more compact
* for small architectures.
*/
#if SPH_CUBEHASH_UNROLL == 2
#define SIXTEEN_ROUNDS do { \
int j; \
for (j = 0; j < 8; j ++) { \
ROUND_EVEN; \
ROUND_ODD; \
} \
} while (0)
#elif SPH_CUBEHASH_UNROLL == 4
#define SIXTEEN_ROUNDS do { \
int j; \
for (j = 0; j < 4; j ++) { \
ROUND_EVEN; \
ROUND_ODD; \
ROUND_EVEN; \
ROUND_ODD; \
} \
} while (0)
#elif SPH_CUBEHASH_UNROLL == 8
#define SIXTEEN_ROUNDS do { \
int j; \
for (j = 0; j < 2; j ++) { \
ROUND_EVEN; \
ROUND_ODD; \
ROUND_EVEN; \
ROUND_ODD; \
ROUND_EVEN; \
ROUND_ODD; \
ROUND_EVEN; \
ROUND_ODD; \
} \
} while (0)
#else
#define SIXTEEN_ROUNDS do { \
ROUND_EVEN; \
ROUND_ODD; \
ROUND_EVEN; \
ROUND_ODD; \
ROUND_EVEN; \
ROUND_ODD; \
ROUND_EVEN; \
ROUND_ODD; \
ROUND_EVEN; \
ROUND_ODD; \
ROUND_EVEN; \
ROUND_ODD; \
ROUND_EVEN; \
ROUND_ODD; \
ROUND_EVEN; \
ROUND_ODD; \
} while (0)
#endif

721
kernel/darkcoin.cl
Unescape View File

@ -0,0 +1,721 @@ @@ -0,0 +1,721 @@
/*
* X11 kernel implementation.
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2014 phm
*
* 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 phm <phm@inbox.com>
*/
#ifndef DARKCOIN_CL
#define DARKCOIN_CL
#if __ENDIAN_LITTLE__
#define SPH_LITTLE_ENDIAN 1
#else
#define SPH_BIG_ENDIAN 1
#endif
#define SPH_UPTR sph_u64
typedef unsigned int sph_u32;
typedef int sph_s32;
#ifndef __OPENCL_VERSION__
typedef unsigned long long sph_u64;
typedef long long sph_s64;
#else
typedef unsigned long sph_u64;
typedef long sph_s64;
#endif
#define SPH_64 1
#define SPH_64_TRUE 1
#define SPH_C32(x) ((sph_u32)(x ## U))
#define SPH_T32(x) ((x) & SPH_C32(0xFFFFFFFF))
#define SPH_ROTL32(x, n) SPH_T32(((x) << (n)) | ((x) >> (32 - (n))))
#define SPH_ROTR32(x, n) SPH_ROTL32(x, (32 - (n)))
#define SPH_C64(x) ((sph_u64)(x ## UL))
#define SPH_T64(x) ((x) & SPH_C64(0xFFFFFFFFFFFFFFFF))
#define SPH_ROTL64(x, n) SPH_T64(((x) << (n)) | ((x) >> (64 - (n))))
#define SPH_ROTR64(x, n) SPH_ROTL64(x, (64 - (n)))
#define SPH_ECHO_64 1
#define SPH_KECCAK_64 1
#define SPH_JH_64 1
#define SPH_SIMD_NOCOPY 0
#define SPH_KECCAK_NOCOPY 0
#define SPH_COMPACT_BLAKE_64 0
#define SPH_LUFFA_PARALLEL 0
#define SPH_SMALL_FOOTPRINT_GROESTL 0
#define SPH_GROESTL_BIG_ENDIAN 0
#define SPH_CUBEHASH_UNROLL 0
#define SPH_KECCAK_UNROLL 0
#include "blake.cl"
#include "bmw.cl"
#include "groestl.cl"
#include "jh.cl"
#include "keccak.cl"
#include "skein.cl"
#include "luffa.cl"
#include "cubehash.cl"
#include "shavite.cl"
#include "simd.cl"
#include "echo.cl"
#define SWAP4(x) as_uint(as_uchar4(x).wzyx)
#define SWAP8(x) as_ulong(as_uchar8(x).s76543210)
#if SPH_BIG_ENDIAN
#define DEC64E(x) (x)
#define DEC64BE(x) (*(const __global sph_u64 *) (x));
#else
#define DEC64E(x) SWAP8(x)
#define DEC64BE(x) SWAP8(*(const __global sph_u64 *) (x));
#endif
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
__kernel void search(__global unsigned char* block, volatile __global uint* output, const ulong target)
{
uint gid = get_global_id(0);
union {
unsigned char h1[64];
uint h4[16];
ulong h8[8];
} hash;
__local sph_u32 AES0[256], AES1[256], AES2[256], AES3[256];
int init = get_local_id(0);
int step = get_local_size(0);
for (int i = init; i < 256; i += step)
{
AES0[i] = AES0_C[i];
AES1[i] = AES1_C[i];
AES2[i] = AES2_C[i];
AES3[i] = AES3_C[i];
}
barrier(CLK_LOCAL_MEM_FENCE);
// blake
{
sph_u64 H0 = SPH_C64(0x6A09E667F3BCC908), H1 = SPH_C64(0xBB67AE8584CAA73B);
sph_u64 H2 = SPH_C64(0x3C6EF372FE94F82B), H3 = SPH_C64(0xA54FF53A5F1D36F1);
sph_u64 H4 = SPH_C64(0x510E527FADE682D1), H5 = SPH_C64(0x9B05688C2B3E6C1F);
sph_u64 H6 = SPH_C64(0x1F83D9ABFB41BD6B), H7 = SPH_C64(0x5BE0CD19137E2179);
sph_u64 S0 = 0, S1 = 0, S2 = 0, S3 = 0;
sph_u64 T0 = SPH_C64(0xFFFFFFFFFFFFFC00) + (80 << 3), T1 = 0xFFFFFFFFFFFFFFFF;;
if ((T0 = SPH_T64(T0 + 1024)) < 1024)
{
T1 = SPH_T64(T1 + 1);
}
sph_u64 M0, M1, M2, M3, M4, M5, M6, M7;
sph_u64 M8, M9, MA, MB, MC, MD, ME, MF;
sph_u64 V0, V1, V2, V3, V4, V5, V6, V7;
sph_u64 V8, V9, VA, VB, VC, VD, VE, VF;
M0 = DEC64BE(block + 0);
M1 = DEC64BE(block + 8);
M2 = DEC64BE(block + 16);
M3 = DEC64BE(block + 24);
M4 = DEC64BE(block + 32);
M5 = DEC64BE(block + 40);
M6 = DEC64BE(block + 48);
M7 = DEC64BE(block + 56);
M8 = DEC64BE(block + 64);
M9 = DEC64BE(block + 72);
M9 &= 0xFFFFFFFF00000000;
M9 ^= SWAP4(gid);
MA = 0x8000000000000000;
MB = 0;
MC = 0;
MD = 1;
ME = 0;
MF = 0x280;
COMPRESS64;
hash.h8[0] = H0;
hash.h8[1] = H1;
hash.h8[2] = H2;
hash.h8[3] = H3;
hash.h8[4] = H4;
hash.h8[5] = H5;
hash.h8[6] = H6;
hash.h8[7] = H7;
}
// bmw
sph_u64 BMW_H[16];
for(unsigned u = 0; u < 16; u++)
BMW_H[u] = BMW_IV512[u];
sph_u64 BMW_h1[16], BMW_h2[16];
sph_u64 mv[16];
mv[ 0] = SWAP8(hash.h8[0]);
mv[ 1] = SWAP8(hash.h8[1]);
mv[ 2] = SWAP8(hash.h8[2]);
mv[ 3] = SWAP8(hash.h8[3]);
mv[ 4] = SWAP8(hash.h8[4]);
mv[ 5] = SWAP8(hash.h8[5]);
mv[ 6] = SWAP8(hash.h8[6]);
mv[ 7] = SWAP8(hash.h8[7]);
mv[ 8] = 0x80;
mv[ 9] = 0;
mv[10] = 0;
mv[11] = 0;
mv[12] = 0;
mv[13] = 0;
mv[14] = 0;
mv[15] = 0x200;
#define M(x) (mv[x])
#define H(x) (BMW_H[x])
#define dH(x) (BMW_h2[x])
FOLDb;
#undef M
#undef H
#undef dH
#define M(x) (BMW_h2[x])
#define H(x) (final_b[x])
#define dH(x) (BMW_h1[x])
FOLDb;
#undef M
#undef H
#undef dH
hash.h8[0] = SWAP8(BMW_h1[8]);
hash.h8[1] = SWAP8(BMW_h1[9]);
hash.h8[2] = SWAP8(BMW_h1[10]);
hash.h8[3] = SWAP8(BMW_h1[11]);
hash.h8[4] = SWAP8(BMW_h1[12]);
hash.h8[5] = SWAP8(BMW_h1[13]);
hash.h8[6] = SWAP8(BMW_h1[14]);
hash.h8[7] = SWAP8(BMW_h1[15]);
// groestl
sph_u64 H[16];
for (unsigned int u = 0; u < 15; u ++)
H[u] = 0;
#if USE_LE
H[15] = ((sph_u64)(512 & 0xFF) << 56) | ((sph_u64)(512 & 0xFF00) << 40);
#else
H[15] = (sph_u64)512;
#endif
sph_u64 g[16], m[16];
m[0] = DEC64E(hash.h8[0]);
m[1] = DEC64E(hash.h8[1]);
m[2] = DEC64E(hash.h8[2]);
m[3] = DEC64E(hash.h8[3]);
m[4] = DEC64E(hash.h8[4]);
m[5] = DEC64E(hash.h8[5]);
m[6] = DEC64E(hash.h8[6]);
m[7] = DEC64E(hash.h8[7]);
for (unsigned int u = 0; u < 16; u ++)
g[u] = m[u] ^ H[u];
m[8] = 0x80; g[8] = m[8] ^ H[8];
m[9] = 0; g[9] = m[9] ^ H[9];
m[10] = 0; g[10] = m[10] ^ H[10];
m[11] = 0; g[11] = m[11] ^ H[11];
m[12] = 0; g[12] = m[12] ^ H[12];
m[13] = 0; g[13] = m[13] ^ H[13];
m[14] = 0; g[14] = m[14] ^ H[14];
m[15] = 0x100000000000000; g[15] = m[15] ^ H[15];
PERM_BIG_P(g);
PERM_BIG_Q(m);
for (unsigned int u = 0; u < 16; u ++)
H[u] ^= g[u] ^ m[u];
sph_u64 xH[16];
for (unsigned int u = 0; u < 16; u ++)
xH[u] = H[u];
PERM_BIG_P(xH);
for (unsigned int u = 0; u < 16; u ++)
H[u] ^= xH[u];
for (unsigned int u = 0; u < 8; u ++)
hash.h8[u] = DEC64E(H[u + 8]);
// skein
sph_u64 h0 = SPH_C64(0x4903ADFF749C51CE), h1 = SPH_C64(0x0D95DE399746DF03), h2 = SPH_C64(0x8FD1934127C79BCE), h3 = SPH_C64(0x9A255629FF352CB1), h4 = SPH_C64(0x5DB62599DF6CA7B0), h5 = SPH_C64(0xEABE394CA9D5C3F4), h6 = SPH_C64(0x991112C71A75B523), h7 = SPH_C64(0xAE18A40B660FCC33);
sph_u64 m0, m1, m2, m3, m4, m5, m6, m7;
sph_u64 bcount = 0;
m0 = SWAP8(hash.h8[0]);
m1 = SWAP8(hash.h8[1]);
m2 = SWAP8(hash.h8[2]);
m3 = SWAP8(hash.h8[3]);
m4 = SWAP8(hash.h8[4]);
m5 = SWAP8(hash.h8[5]);
m6 = SWAP8(hash.h8[6]);
m7 = SWAP8(hash.h8[7]);
UBI_BIG(480, 64);
bcount = 0;
m0 = m1 = m2 = m3 = m4 = m5 = m6 = m7 = 0;
UBI_BIG(510, 8);
hash.h8[0] = SWAP8(h0);
hash.h8[1] = SWAP8(h1);
hash.h8[2] = SWAP8(h2);
hash.h8[3] = SWAP8(h3);
hash.h8[4] = SWAP8(h4);
hash.h8[5] = SWAP8(h5);
hash.h8[6] = SWAP8(h6);
hash.h8[7] = SWAP8(h7);
// jh
sph_u64 h0h = C64e(0x6fd14b963e00aa17), h0l = C64e(0x636a2e057a15d543), h1h = C64e(0x8a225e8d0c97ef0b), h1l = C64e(0xe9341259f2b3c361), h2h = C64e(0x891da0c1536f801e), h2l = C64e(0x2aa9056bea2b6d80), h3h = C64e(0x588eccdb2075baa6), h3l = C64e(0xa90f3a76baf83bf7);
sph_u64 h4h = C64e(0x0169e60541e34a69), h4l = C64e(0x46b58a8e2e6fe65a), h5h = C64e(0x1047a7d0c1843c24), h5l = C64e(0x3b6e71b12d5ac199), h6h = C64e(0xcf57f6ec9db1f856), h6l = C64e(0xa706887c5716b156), h7h = C64e(0xe3c2fcdfe68517fb), h7l = C64e(0x545a4678cc8cdd4b);
sph_u64 tmp;
for(int i = 0; i < 2; i++)
{
if (i == 0) {
h0h ^= DEC64E(hash.h8[0]);
h0l ^= DEC64E(hash.h8[1]);
h1h ^= DEC64E(hash.h8[2]);
h1l ^= DEC64E(hash.h8[3]);
h2h ^= DEC64E(hash.h8[4]);
h2l ^= DEC64E(hash.h8[5]);
h3h ^= DEC64E(hash.h8[6]);
h3l ^= DEC64E(hash.h8[7]);
} else if(i == 1) {
h4h ^= DEC64E(hash.h8[0]);
h4l ^= DEC64E(hash.h8[1]);
h5h ^= DEC64E(hash.h8[2]);
h5l ^= DEC64E(hash.h8[3]);
h6h ^= DEC64E(hash.h8[4]);
h6l ^= DEC64E(hash.h8[5]);
h7h ^= DEC64E(hash.h8[6]);
h7l ^= DEC64E(hash.h8[7]);
h0h ^= 0x80;
h3l ^= 0x2000000000000;
}
E8;
}
h4h ^= 0x80;
h7l ^= 0x2000000000000;
hash.h8[0] = DEC64E(h4h);
hash.h8[1] = DEC64E(h4l);
hash.h8[2] = DEC64E(h5h);
hash.h8[3] = DEC64E(h5l);
hash.h8[4] = DEC64E(h6h);
hash.h8[5] = DEC64E(h6l);
hash.h8[6] = DEC64E(h7h);
hash.h8[7] = DEC64E(h7l);
// keccak
sph_u64 a00 = 0, a01 = 0, a02 = 0, a03 = 0, a04 = 0;
sph_u64 a10 = 0, a11 = 0, a12 = 0, a13 = 0, a14 = 0;
sph_u64 a20 = 0, a21 = 0, a22 = 0, a23 = 0, a24 = 0;
sph_u64 a30 = 0, a31 = 0, a32 = 0, a33 = 0, a34 = 0;
sph_u64 a40 = 0, a41 = 0, a42 = 0, a43 = 0, a44 = 0;
a10 = SPH_C64(0xFFFFFFFFFFFFFFFF);
a20 = SPH_C64(0xFFFFFFFFFFFFFFFF);
a31 = SPH_C64(0xFFFFFFFFFFFFFFFF);
a22 = SPH_C64(0xFFFFFFFFFFFFFFFF);
a23 = SPH_C64(0xFFFFFFFFFFFFFFFF);
a04 = SPH_C64(0xFFFFFFFFFFFFFFFF);
a00 ^= SWAP8(hash.h8[0]);
a10 ^= SWAP8(hash.h8[1]);
a20 ^= SWAP8(hash.h8[2]);
a30 ^= SWAP8(hash.h8[3]);
a40 ^= SWAP8(hash.h8[4]);
a01 ^= SWAP8(hash.h8[5]);
a11 ^= SWAP8(hash.h8[6]);
a21 ^= SWAP8(hash.h8[7]);
a31 ^= 0x8000000000000001;
KECCAK_F_1600;
// Finalize the "lane complement"
a10 = ~a10;
a20 = ~a20;
hash.h8[0] = SWAP8(a00);
hash.h8[1] = SWAP8(a10);
hash.h8[2] = SWAP8(a20);
hash.h8[3] = SWAP8(a30);
hash.h8[4] = SWAP8(a40);
hash.h8[5] = SWAP8(a01);
hash.h8[6] = SWAP8(a11);
hash.h8[7] = SWAP8(a21);
// luffa
sph_u32 V00 = SPH_C32(0x6d251e69), V01 = SPH_C32(0x44b051e0), V02 = SPH_C32(0x4eaa6fb4), V03 = SPH_C32(0xdbf78465), V04 = SPH_C32(0x6e292011), V05 = SPH_C32(0x90152df4), V06 = SPH_C32(0xee058139), V07 = SPH_C32(0xdef610bb);
sph_u32 V10 = SPH_C32(0xc3b44b95), V11 = SPH_C32(0xd9d2f256), V12 = SPH_C32(0x70eee9a0), V13 = SPH_C32(0xde099fa3), V14 = SPH_C32(0x5d9b0557), V15 = SPH_C32(0x8fc944b3), V16 = SPH_C32(0xcf1ccf0e), V17 = SPH_C32(0x746cd581);
sph_u32 V20 = SPH_C32(0xf7efc89d), V21 = SPH_C32(0x5dba5781), V22 = SPH_C32(0x04016ce5), V23 = SPH_C32(0xad659c05), V24 = SPH_C32(0x0306194f), V25 = SPH_C32(0x666d1836), V26 = SPH_C32(0x24aa230a), V27 = SPH_C32(0x8b264ae7);
sph_u32 V30 = SPH_C32(0x858075d5), V31 = SPH_C32(0x36d79cce), V32 = SPH_C32(0xe571f7d7), V33 = SPH_C32(0x204b1f67), V34 = SPH_C32(0x35870c6a), V35 = SPH_C32(0x57e9e923), V36 = SPH_C32(0x14bcb808), V37 = SPH_C32(0x7cde72ce);
sph_u32 V40 = SPH_C32(0x6c68e9be), V41 = SPH_C32(0x5ec41e22), V42 = SPH_C32(0xc825b7c7), V43 = SPH_C32(0xaffb4363), V44 = SPH_C32(0xf5df3999), V45 = SPH_C32(0x0fc688f1), V46 = SPH_C32(0xb07224cc), V47 = SPH_C32(0x03e86cea);
DECL_TMP8(M);
M0 = hash.h4[1];
M1 = hash.h4[0];
M2 = hash.h4[3];
M3 = hash.h4[2];
M4 = hash.h4[5];
M5 = hash.h4[4];
M6 = hash.h4[7];
M7 = hash.h4[6];
for(uint i = 0; i < 5; i++)
{
MI5;
LUFFA_P5;
if(i == 0) {
M0 = hash.h4[9];
M1 = hash.h4[8];
M2 = hash.h4[11];
M3 = hash.h4[10];
M4 = hash.h4[13];
M5 = hash.h4[12];
M6 = hash.h4[15];
M7 = hash.h4[14];
} else if(i == 1) {
M0 = 0x80000000;
M1 = M2 = M3 = M4 = M5 = M6 = M7 = 0;
} else if(i == 2) {
M0 = M1 = M2 = M3 = M4 = M5 = M6 = M7 = 0;
} else if(i == 3) {
hash.h4[1] = V00 ^ V10 ^ V20 ^ V30 ^ V40;
hash.h4[0] = V01 ^ V11 ^ V21 ^ V31 ^ V41;
hash.h4[3] = V02 ^ V12 ^ V22 ^ V32 ^ V42;
hash.h4[2] = V03 ^ V13 ^ V23 ^ V33 ^ V43;
hash.h4[5] = V04 ^ V14 ^ V24 ^ V34 ^ V44;
hash.h4[4] = V05 ^ V15 ^ V25 ^ V35 ^ V45;
hash.h4[7] = V06 ^ V16 ^ V26 ^ V36 ^ V46;
hash.h4[6] = V07 ^ V17 ^ V27 ^ V37 ^ V47;
}
}
hash.h4[9] = V00 ^ V10 ^ V20 ^ V30 ^ V40;
hash.h4[8] = V01 ^ V11 ^ V21 ^ V31 ^ V41;
hash.h4[11] = V02 ^ V12 ^ V22 ^ V32 ^ V42;
hash.h4[10] = V03 ^ V13 ^ V23 ^ V33 ^ V43;
hash.h4[13] = V04 ^ V14 ^ V24 ^ V34 ^ V44;
hash.h4[12] = V05 ^ V15 ^ V25 ^ V35 ^ V45;
hash.h4[15] = V06 ^ V16 ^ V26 ^ V36 ^ V46;
hash.h4[14] = V07 ^ V17 ^ V27 ^ V37 ^ V47;
// cubehash.h1
sph_u32 x0 = SPH_C32(0x2AEA2A61), x1 = SPH_C32(0x50F494D4), x2 = SPH_C32(0x2D538B8B), x3 = SPH_C32(0x4167D83E);
sph_u32 x4 = SPH_C32(0x3FEE2313), x5 = SPH_C32(0xC701CF8C), x6 = SPH_C32(0xCC39968E), x7 = SPH_C32(0x50AC5695);
sph_u32 x8 = SPH_C32(0x4D42C787), x9 = SPH_C32(0xA647A8B3), xa = SPH_C32(0x97CF0BEF), xb = SPH_C32(0x825B4537);
sph_u32 xc = SPH_C32(0xEEF864D2), xd = SPH_C32(0xF22090C4), xe = SPH_C32(0xD0E5CD33), xf = SPH_C32(0xA23911AE);
sph_u32 xg = SPH_C32(0xFCD398D9), xh = SPH_C32(0x148FE485), xi = SPH_C32(0x1B017BEF), xj = SPH_C32(0xB6444532);
sph_u32 xk = SPH_C32(0x6A536159), xl = SPH_C32(0x2FF5781C), xm = SPH_C32(0x91FA7934), xn = SPH_C32(0x0DBADEA9);
sph_u32 xo = SPH_C32(0xD65C8A2B), xp = SPH_C32(0xA5A70E75), xq = SPH_C32(0xB1C62456), xr = SPH_C32(0xBC796576);
sph_u32 xs = SPH_C32(0x1921C8F7), xt = SPH_C32(0xE7989AF1), xu = SPH_C32(0x7795D246), xv = SPH_C32(0xD43E3B44);
x0 ^= SWAP4(hash.h4[1]);
x1 ^= SWAP4(hash.h4[0]);
x2 ^= SWAP4(hash.h4[3]);
x3 ^= SWAP4(hash.h4[2]);
x4 ^= SWAP4(hash.h4[5]);
x5 ^= SWAP4(hash.h4[4]);
x6 ^= SWAP4(hash.h4[7]);
x7 ^= SWAP4(hash.h4[6]);
for (int i = 0; i < 13; i ++) {
SIXTEEN_ROUNDS;
if (i == 0) {
x0 ^= SWAP4(hash.h4[9]);
x1 ^= SWAP4(hash.h4[8]);
x2 ^= SWAP4(hash.h4[11]);
x3 ^= SWAP4(hash.h4[10]);
x4 ^= SWAP4(hash.h4[13]);
x5 ^= SWAP4(hash.h4[12]);
x6 ^= SWAP4(hash.h4[15]);
x7 ^= SWAP4(hash.h4[14]);
} else if(i == 1) {
x0 ^= 0x80;
} else if (i == 2) {
xv ^= SPH_C32(1);
}
}
hash.h4[0] = x0;
hash.h4[1] = x1;
hash.h4[2] = x2;
hash.h4[3] = x3;
hash.h4[4] = x4;
hash.h4[5] = x5;
hash.h4[6] = x6;
hash.h4[7] = x7;
hash.h4[8] = x8;
hash.h4[9] = x9;
hash.h4[10] = xa;
hash.h4[11] = xb;
hash.h4[12] = xc;
hash.h4[13] = xd;
hash.h4[14] = xe;
hash.h4[15] = xf;
// shavite
{
// IV
sph_u32 h0 = SPH_C32(0x72FCCDD8), h1 = SPH_C32(0x79CA4727), h2 = SPH_C32(0x128A077B), h3 = SPH_C32(0x40D55AEC);
sph_u32 h4 = SPH_C32(0xD1901A06), h5 = SPH_C32(0x430AE307), h6 = SPH_C32(0xB29F5CD1), h7 = SPH_C32(0xDF07FBFC);
sph_u32 h8 = SPH_C32(0x8E45D73D), h9 = SPH_C32(0x681AB538), hA = SPH_C32(0xBDE86578), hB = SPH_C32(0xDD577E47);
sph_u32 hC = SPH_C32(0xE275EADE), hD = SPH_C32(0x502D9FCD), hE = SPH_C32(0xB9357178), hF = SPH_C32(0x022A4B9A);
// state
sph_u32 rk00, rk01, rk02, rk03, rk04, rk05, rk06, rk07;
sph_u32 rk08, rk09, rk0A, rk0B, rk0C, rk0D, rk0E, rk0F;
sph_u32 rk10, rk11, rk12, rk13, rk14, rk15, rk16, rk17;
sph_u32 rk18, rk19, rk1A, rk1B, rk1C, rk1D, rk1E, rk1F;
sph_u32 sc_count0 = (64 << 3), sc_count1 = 0, sc_count2 = 0, sc_count3 = 0;
rk00 = hash.h4[0];
rk01 = hash.h4[1];
rk02 = hash.h4[2];
rk03 = hash.h4[3];
rk04 = hash.h4[4];
rk05 = hash.h4[5];
rk06 = hash.h4[6];
rk07 = hash.h4[7];
rk08 = hash.h4[8];
rk09 = hash.h4[9];
rk0A = hash.h4[10];
rk0B = hash.h4[11];
rk0C = hash.h4[12];
rk0D = hash.h4[13];
rk0E = hash.h4[14];
rk0F = hash.h4[15];
rk10 = 0x80;
rk11 = rk12 = rk13 = rk14 = rk15 = rk16 = rk17 = rk18 = rk19 = rk1A = 0;
rk1B = 0x2000000;
rk1C = rk1D = rk1E = 0;
rk1F = 0x2000000;
c512(buf);
hash.h4[0] = h0;
hash.h4[1] = h1;
hash.h4[2] = h2;
hash.h4[3] = h3;
hash.h4[4] = h4;
hash.h4[5] = h5;
hash.h4[6] = h6;
hash.h4[7] = h7;
hash.h4[8] = h8;
hash.h4[9] = h9;
hash.h4[10] = hA;
hash.h4[11] = hB;
hash.h4[12] = hC;
hash.h4[13] = hD;
hash.h4[14] = hE;
hash.h4[15] = hF;
}
// simd
s32 q[256];
unsigned char x[128];
for(unsigned int i = 0; i < 64; i++)
x[i] = hash.h1[i];
for(unsigned int i = 64; i < 128; i++)
x[i] = 0;
u32 A0 = C32(0x0BA16B95), A1 = C32(0x72F999AD), A2 = C32(0x9FECC2AE), A3 = C32(0xBA3264FC), A4 = C32(0x5E894929), A5 = C32(0x8E9F30E5), A6 = C32(0x2F1DAA37), A7 = C32(0xF0F2C558);
u32 B0 = C32(0xAC506643), B1 = C32(0xA90635A5), B2 = C32(0xE25B878B), B3 = C32(0xAAB7878F), B4 = C32(0x88817F7A), B5 = C32(0x0A02892B), B6 = C32(0x559A7550), B7 = C32(0x598F657E);
u32 C0 = C32(0x7EEF60A1), C1 = C32(0x6B70E3E8), C2 = C32(0x9C1714D1), C3 = C32(0xB958E2A8), C4 = C32(0xAB02675E), C5 = C32(0xED1C014F), C6 = C32(0xCD8D65BB), C7 = C32(0xFDB7A257);
u32 D0 = C32(0x09254899), D1 = C32(0xD699C7BC), D2 = C32(0x9019B6DC), D3 = C32(0x2B9022E4), D4 = C32(0x8FA14956), D5 = C32(0x21BF9BD3), D6 = C32(0xB94D0943), D7 = C32(0x6FFDDC22);
FFT256(0, 1, 0, ll1);
for (int i = 0; i < 256; i ++) {
s32 tq;
tq = q[i] + yoff_b_n[i];
tq = REDS2(tq);
tq = REDS1(tq);
tq = REDS1(tq);
q[i] = (tq <= 128 ? tq : tq - 257);
}
A0 ^= hash.h4[0];
A1 ^= hash.h4[1];
A2 ^= hash.h4[2];
A3 ^= hash.h4[3];
A4 ^= hash.h4[4];
A5 ^= hash.h4[5];
A6 ^= hash.h4[6];
A7 ^= hash.h4[7];
B0 ^= hash.h4[8];
B1 ^= hash.h4[9];
B2 ^= hash.h4[10];
B3 ^= hash.h4[11];
B4 ^= hash.h4[12];
B5 ^= hash.h4[13];
B6 ^= hash.h4[14];
B7 ^= hash.h4[15];
ONE_ROUND_BIG(0_, 0, 3, 23, 17, 27);
ONE_ROUND_BIG(1_, 1, 28, 19, 22, 7);
ONE_ROUND_BIG(2_, 2, 29, 9, 15, 5);
ONE_ROUND_BIG(3_, 3, 4, 13, 10, 25);
STEP_BIG(
C32(0x0BA16B95), C32(0x72F999AD), C32(0x9FECC2AE), C32(0xBA3264FC),
C32(0x5E894929), C32(0x8E9F30E5), C32(0x2F1DAA37), C32(0xF0F2C558),
IF, 4, 13, PP8_4_);
STEP_BIG(
C32(0xAC506643), C32(0xA90635A5), C32(0xE25B878B), C32(0xAAB7878F),
C32(0x88817F7A), C32(0x0A02892B), C32(0x559A7550), C32(0x598F657E),
IF, 13, 10, PP8_5_);
STEP_BIG(
C32(0x7EEF60A1), C32(0x6B70E3E8), C32(0x9C1714D1), C32(0xB958E2A8),
C32(0xAB02675E), C32(0xED1C014F), C32(0xCD8D65BB), C32(0xFDB7A257),
IF, 10, 25, PP8_6_);
STEP_BIG(
C32(0x09254899), C32(0xD699C7BC), C32(0x9019B6DC), C32(0x2B9022E4),
C32(0x8FA14956), C32(0x21BF9BD3), C32(0xB94D0943), C32(0x6FFDDC22),
IF, 25, 4, PP8_0_);
u32 COPY_A0 = A0, COPY_A1 = A1, COPY_A2 = A2, COPY_A3 = A3, COPY_A4 = A4, COPY_A5 = A5, COPY_A6 = A6, COPY_A7 = A7;
u32 COPY_B0 = B0, COPY_B1 = B1, COPY_B2 = B2, COPY_B3 = B3, COPY_B4 = B4, COPY_B5 = B5, COPY_B6 = B6, COPY_B7 = B7;
u32 COPY_C0 = C0, COPY_C1 = C1, COPY_C2 = C2, COPY_C3 = C3, COPY_C4 = C4, COPY_C5 = C5, COPY_C6 = C6, COPY_C7 = C7;
u32 COPY_D0 = D0, COPY_D1 = D1, COPY_D2 = D2, COPY_D3 = D3, COPY_D4 = D4, COPY_D5 = D5, COPY_D6 = D6, COPY_D7 = D7;
#define q SIMD_Q
A0 ^= 0x200;
ONE_ROUND_BIG(0_, 0, 3, 23, 17, 27);
ONE_ROUND_BIG(1_, 1, 28, 19, 22, 7);
ONE_ROUND_BIG(2_, 2, 29, 9, 15, 5);
ONE_ROUND_BIG(3_, 3, 4, 13, 10, 25);
STEP_BIG(
COPY_A0, COPY_A1, COPY_A2, COPY_A3,
COPY_A4, COPY_A5, COPY_A6, COPY_A7,
IF, 4, 13, PP8_4_);
STEP_BIG(
COPY_B0, COPY_B1, COPY_B2, COPY_B3,
COPY_B4, COPY_B5, COPY_B6, COPY_B7,
IF, 13, 10, PP8_5_);
STEP_BIG(
COPY_C0, COPY_C1, COPY_C2, COPY_C3,
COPY_C4, COPY_C5, COPY_C6, COPY_C7,
IF, 10, 25, PP8_6_);
STEP_BIG(
COPY_D0, COPY_D1, COPY_D2, COPY_D3,
COPY_D4, COPY_D5, COPY_D6, COPY_D7,
IF, 25, 4, PP8_0_);
#undef q
hash.h4[0] = A0;
hash.h4[1] = A1;
hash.h4[2] = A2;
hash.h4[3] = A3;
hash.h4[4] = A4;
hash.h4[5] = A5;
hash.h4[6] = A6;
hash.h4[7] = A7;
hash.h4[8] = B0;
hash.h4[9] = B1;
hash.h4[10] = B2;
hash.h4[11] = B3;
hash.h4[12] = B4;
hash.h4[13] = B5;
hash.h4[14] = B6;
hash.h4[15] = B7;
// echo
sph_u64 W00, W01, W10, W11, W20, W21, W30, W31, W40, W41, W50, W51, W60, W61, W70, W71, W80, W81, W90, W91, WA0, WA1, WB0, WB1, WC0, WC1, WD0, WD1, WE0, WE1, WF0, WF1;
sph_u64 Vb00, Vb01, Vb10, Vb11, Vb20, Vb21, Vb30, Vb31, Vb40, Vb41, Vb50, Vb51, Vb60, Vb61, Vb70, Vb71;
Vb00 = Vb10 = Vb20 = Vb30 = Vb40 = Vb50 = Vb60 = Vb70 = 512UL;
Vb01 = Vb11 = Vb21 = Vb31 = Vb41 = Vb51 = Vb61 = Vb71 = 0;
sph_u32 K0 = 512;
sph_u32 K1 = 0;
sph_u32 K2 = 0;
sph_u32 K3 = 0;
W00 = Vb00;
W01 = Vb01;
W10 = Vb10;
W11 = Vb11;
W20 = Vb20;
W21 = Vb21;
W30 = Vb30;
W31 = Vb31;
W40 = Vb40;
W41 = Vb41;
W50 = Vb50;
W51 = Vb51;
W60 = Vb60;
W61 = Vb61;
W70 = Vb70;
W71 = Vb71;
W80 = hash.h8[0];
W81 = hash.h8[1];
W90 = hash.h8[2];
W91 = hash.h8[3];
WA0 = hash.h8[4];
WA1 = hash.h8[5];
WB0 = hash.h8[6];
WB1 = hash.h8[7];
WC0 = 0x80;
WC1 = 0;
WD0 = 0;
WD1 = 0;
WE0 = 0;
WE1 = 0x200000000000000;
WF0 = 0x200;
WF1 = 0;
for (unsigned u = 0; u < 10; u ++) {
BIG_ROUND;
}
Vb00 ^= hash.h8[0] ^ W00 ^ W80;
Vb01 ^= hash.h8[1] ^ W01 ^ W81;
Vb10 ^= hash.h8[2] ^ W10 ^ W90;
Vb11 ^= hash.h8[3] ^ W11 ^ W91;
Vb20 ^= hash.h8[4] ^ W20 ^ WA0;
Vb21 ^= hash.h8[5] ^ W21 ^ WA1;
Vb30 ^= hash.h8[6] ^ W30 ^ WB0;
Vb31 ^= hash.h8[7] ^ W31 ^ WB1;
bool result = (Vb11 <= target);
if (result)
output[output[0xFF]++] = SWAP4(gid);
}
#endif // DARKCOIN_CL

169
kernel/echo.cl
Unescape View File

@ -0,0 +1,169 @@ @@ -0,0 +1,169 @@
/* $Id: echo.c 227 2010-06-16 17:28:38Z tp $ */
/*
* ECHO implementation.
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2007-2010 Projet RNRT SAPHIR
*
* 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 Thomas Pornin <thomas.pornin@cryptolog.com>
*/
#define T32 SPH_T32
#define C32 SPH_C32
#if SPH_64
#define C64 SPH_C64
#endif
#define AES_BIG_ENDIAN 0
#include "aes_helper.cl"
#define ECHO_DECL_STATE_BIG \
sph_u64 W00, W01, W10, W11, W20, W21, W30, W31, W40, W41, W50, W51, W60, W61, W70, W71, W80, W81, W90, W91, WA0, WA1, WB0, WB1, WC0, WC1, WD0, WD1, WE0, WE1, WF0, WF1;
#define AES_2ROUNDS(XX, XY) do { \
sph_u32 X0 = (sph_u32)(XX); \
sph_u32 X1 = (sph_u32)(XX >> 32); \
sph_u32 X2 = (sph_u32)(XY); \
sph_u32 X3 = (sph_u32)(XY >> 32); \
sph_u32 Y0, Y1, Y2, Y3; \
AES_ROUND_LE(X0, X1, X2, X3, K0, K1, K2, K3, Y0, Y1, Y2, Y3); \
AES_ROUND_NOKEY_LE(Y0, Y1, Y2, Y3, X0, X1, X2, X3); \
XX = (sph_u64)X0 | ((sph_u64)X1 << 32); \
XY = (sph_u64)X2 | ((sph_u64)X3 << 32); \
if ((K0 = T32(K0 + 1)) == 0) { \
if ((K1 = T32(K1 + 1)) == 0) \
if ((K2 = T32(K2 + 1)) == 0) \
K3 = T32(K3 + 1); \
} \
} while (0)
#define BIG_SUB_WORDS do { \
AES_2ROUNDS(W00, W01); \
AES_2ROUNDS(W10, W11); \
AES_2ROUNDS(W20, W21); \
AES_2ROUNDS(W30, W31); \
AES_2ROUNDS(W40, W41); \
AES_2ROUNDS(W50, W51); \
AES_2ROUNDS(W60, W61); \
AES_2ROUNDS(W70, W71); \
AES_2ROUNDS(W80, W81); \
AES_2ROUNDS(W90, W91); \
AES_2ROUNDS(WA0, WA1); \
AES_2ROUNDS(WB0, WB1); \
AES_2ROUNDS(WC0, WC1); \
AES_2ROUNDS(WD0, WD1); \
AES_2ROUNDS(WE0, WE1); \
AES_2ROUNDS(WF0, WF1); \
} while (0)
#define SHIFT_ROW1(a, b, c, d) do { \
sph_u64 tmp; \
tmp = W ## a ## 0; \
W ## a ## 0 = W ## b ## 0; \
W ## b ## 0 = W ## c ## 0; \
W ## c ## 0 = W ## d ## 0; \
W ## d ## 0 = tmp; \
tmp = W ## a ## 1; \
W ## a ## 1 = W ## b ## 1; \
W ## b ## 1 = W ## c ## 1; \
W ## c ## 1 = W ## d ## 1; \
W ## d ## 1 = tmp; \
} while (0)
#define SHIFT_ROW2(a, b, c, d) do { \
sph_u64 tmp; \
tmp = W ## a ## 0; \
W ## a ## 0 = W ## c ## 0; \
W ## c ## 0 = tmp; \
tmp = W ## b ## 0; \
W ## b ## 0 = W ## d ## 0; \
W ## d ## 0 = tmp; \
tmp = W ## a ## 1; \
W ## a ## 1 = W ## c ## 1; \
W ## c ## 1 = tmp; \
tmp = W ## b ## 1; \
W ## b ## 1 = W ## d ## 1; \
W ## d ## 1 = tmp; \
} while (0)
#define SHIFT_ROW3(a, b, c, d) SHIFT_ROW1(d, c, b, a)
#define BIG_SHIFT_ROWS do { \
SHIFT_ROW1(1, 5, 9, D); \
SHIFT_ROW2(2, 6, A, E); \
SHIFT_ROW3(3, 7, B, F); \
} while (0)
#define MIX_COLUMN1(ia, ib, ic, id, n) do { \
sph_u64 a = W ## ia ## n; \
sph_u64 b = W ## ib ## n; \
sph_u64 c = W ## ic ## n; \
sph_u64 d = W ## id ## n; \
sph_u64 ab = a ^ b; \
sph_u64 bc = b ^ c; \
sph_u64 cd = c ^ d; \
sph_u64 abx = ((ab & C64(0x8080808080808080)) >> 7) * 27U \
^ ((ab & C64(0x7F7F7F7F7F7F7F7F)) << 1); \
sph_u64 bcx = ((bc & C64(0x8080808080808080)) >> 7) * 27U \
^ ((bc & C64(0x7F7F7F7F7F7F7F7F)) << 1); \
sph_u64 cdx = ((cd & C64(0x8080808080808080)) >> 7) * 27U \
^ ((cd & C64(0x7F7F7F7F7F7F7F7F)) << 1); \
W ## ia ## n = abx ^ bc ^ d; \
W ## ib ## n = bcx ^ a ^ cd; \
W ## ic ## n = cdx ^ ab ^ d; \
W ## id ## n = abx ^ bcx ^ cdx ^ ab ^ c; \
} while (0)
#define MIX_COLUMN(a, b, c, d) do { \
MIX_COLUMN1(a, b, c, d, 0); \
MIX_COLUMN1(a, b, c, d, 1); \
} while (0)
#define BIG_MIX_COLUMNS do { \
MIX_COLUMN(0, 1, 2, 3); \
MIX_COLUMN(4, 5, 6, 7); \
MIX_COLUMN(8, 9, A, B); \
MIX_COLUMN(C, D, E, F); \
} while (0)
#define BIG_ROUND do { \
BIG_SUB_WORDS; \
BIG_SHIFT_ROWS; \
BIG_MIX_COLUMNS; \
} while (0)
#define ECHO_COMPRESS_BIG(sc) do { \
sph_u32 K0 = sc->C0; \
sph_u32 K1 = sc->C1; \
sph_u32 K2 = sc->C2; \
sph_u32 K3 = sc->C3; \
unsigned u; \
INPUT_BLOCK_BIG(sc); \
for (u = 0; u < 10; u ++) { \
BIG_ROUND; \
} \
ECHO_FINAL_BIG; \
} while (0)

1418
kernel/groestl.cl
View File

File diff suppressed because it is too large Load Diff

320
kernel/jh.cl
Unescape View File

@ -0,0 +1,320 @@ @@ -0,0 +1,320 @@
/* $Id: jh.c 255 2011-06-07 19:50:20Z tp $ */
/*
* JH implementation.
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2007-2010 Projet RNRT SAPHIR
*
* 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 Thomas Pornin <thomas.pornin@cryptolog.com>
*/
#if !defined SPH_JH_64 && SPH_64_TRUE
#define SPH_JH_64 1
#endif
/*
* The internal bitslice representation may use either big-endian or
* little-endian (true bitslice operations do not care about the bit
* ordering, and the bit-swapping linear operations in JH happen to
* be invariant through endianness-swapping). The constants must be
* defined according to the chosen endianness; we use some
* byte-swapping macros for that.
*/
#if SPH_LITTLE_ENDIAN
#define C32e(x) ((SPH_C32(x) >> 24) \
| ((SPH_C32(x) >> 8) & SPH_C32(0x0000FF00)) \
| ((SPH_C32(x) << 8) & SPH_C32(0x00FF0000)) \
| ((SPH_C32(x) << 24) & SPH_C32(0xFF000000)))
#define dec32e_aligned sph_dec32le_aligned
#define enc32e sph_enc32le
#define C64e(x) ((SPH_C64(x) >> 56) \
| ((SPH_C64(x) >> 40) & SPH_C64(0x000000000000FF00)) \
| ((SPH_C64(x) >> 24) & SPH_C64(0x0000000000FF0000)) \
| ((SPH_C64(x) >> 8) & SPH_C64(0x00000000FF000000)) \
| ((SPH_C64(x) << 8) & SPH_C64(0x000000FF00000000)) \
| ((SPH_C64(x) << 24) & SPH_C64(0x0000FF0000000000)) \
| ((SPH_C64(x) << 40) & SPH_C64(0x00FF000000000000)) \
| ((SPH_C64(x) << 56) & SPH_C64(0xFF00000000000000)))
#define dec64e_aligned sph_dec64le_aligned
#define enc64e sph_enc64le
#else
#define C32e(x) SPH_C32(x)
#define dec32e_aligned sph_dec32be_aligned
#define enc32e sph_enc32be
#define C64e(x) SPH_C64(x)
#define dec64e_aligned sph_dec64be_aligned
#define enc64e sph_enc64be
#endif
#define Sb(x0, x1, x2, x3, c) do { \
x3 = ~x3; \
x0 ^= (c) & ~x2; \
tmp = (c) ^ (x0 & x1); \
x0 ^= x2 & x3; \
x3 ^= ~x1 & x2; \
x1 ^= x0 & x2; \
x2 ^= x0 & ~x3; \
x0 ^= x1 | x3; \
x3 ^= x1 & x2; \
x1 ^= tmp & x0; \
x2 ^= tmp; \
} while (0)
#define Lb(x0, x1, x2, x3, x4, x5, x6, x7) do { \
x4 ^= x1; \
x5 ^= x2; \
x6 ^= x3 ^ x0; \
x7 ^= x0; \
x0 ^= x5; \
x1 ^= x6; \
x2 ^= x7 ^ x4; \
x3 ^= x4; \
} while (0)
__constant static const sph_u64 C[] = {
C64e(0x72d5dea2df15f867), C64e(0x7b84150ab7231557),
C64e(0x81abd6904d5a87f6), C64e(0x4e9f4fc5c3d12b40),
C64e(0xea983ae05c45fa9c), C64e(0x03c5d29966b2999a),
C64e(0x660296b4f2bb538a), C64e(0xb556141a88dba231),
C64e(0x03a35a5c9a190edb), C64e(0x403fb20a87c14410),
C64e(0x1c051980849e951d), C64e(0x6f33ebad5ee7cddc),
C64e(0x10ba139202bf6b41), C64e(0xdc786515f7bb27d0),
C64e(0x0a2c813937aa7850), C64e(0x3f1abfd2410091d3),
C64e(0x422d5a0df6cc7e90), C64e(0xdd629f9c92c097ce),
C64e(0x185ca70bc72b44ac), C64e(0xd1df65d663c6fc23),
C64e(0x976e6c039ee0b81a), C64e(0x2105457e446ceca8),
C64e(0xeef103bb5d8e61fa), C64e(0xfd9697b294838197),
C64e(0x4a8e8537db03302f), C64e(0x2a678d2dfb9f6a95),
C64e(0x8afe7381f8b8696c), C64e(0x8ac77246c07f4214),
C64e(0xc5f4158fbdc75ec4), C64e(0x75446fa78f11bb80),
C64e(0x52de75b7aee488bc), C64e(0x82b8001e98a6a3f4),
C64e(0x8ef48f33a9a36315), C64e(0xaa5f5624d5b7f989),
C64e(0xb6f1ed207c5ae0fd), C64e(0x36cae95a06422c36),
C64e(0xce2935434efe983d), C64e(0x533af974739a4ba7),
C64e(0xd0f51f596f4e8186), C64e(0x0e9dad81afd85a9f),
C64e(0xa7050667ee34626a), C64e(0x8b0b28be6eb91727),
C64e(0x47740726c680103f), C64e(0xe0a07e6fc67e487b),
C64e(0x0d550aa54af8a4c0), C64e(0x91e3e79f978ef19e),
C64e(0x8676728150608dd4), C64e(0x7e9e5a41f3e5b062),
C64e(0xfc9f1fec4054207a), C64e(0xe3e41a00cef4c984),
C64e(0x4fd794f59dfa95d8), C64e(0x552e7e1124c354a5),
C64e(0x5bdf7228bdfe6e28), C64e(0x78f57fe20fa5c4b2),
C64e(0x05897cefee49d32e), C64e(0x447e9385eb28597f),
C64e(0x705f6937b324314a), C64e(0x5e8628f11dd6e465),
C64e(0xc71b770451b920e7), C64e(0x74fe43e823d4878a),
C64e(0x7d29e8a3927694f2), C64e(0xddcb7a099b30d9c1),
C64e(0x1d1b30fb5bdc1be0), C64e(0xda24494ff29c82bf),
C64e(0xa4e7ba31b470bfff), C64e(0x0d324405def8bc48),
C64e(0x3baefc3253bbd339), C64e(0x459fc3c1e0298ba0),
C64e(0xe5c905fdf7ae090f), C64e(0x947034124290f134),
C64e(0xa271b701e344ed95), C64e(0xe93b8e364f2f984a),
C64e(0x88401d63a06cf615), C64e(0x47c1444b8752afff),
C64e(0x7ebb4af1e20ac630), C64e(0x4670b6c5cc6e8ce6),
C64e(0xa4d5a456bd4fca00), C64e(0xda9d844bc83e18ae),
C64e(0x7357ce453064d1ad), C64e(0xe8a6ce68145c2567),
C64e(0xa3da8cf2cb0ee116), C64e(0x33e906589a94999a),
C64e(0x1f60b220c26f847b), C64e(0xd1ceac7fa0d18518),
C64e(0x32595ba18ddd19d3), C64e(0x509a1cc0aaa5b446),
C64e(0x9f3d6367e4046bba), C64e(0xf6ca19ab0b56ee7e),
C64e(0x1fb179eaa9282174), C64e(0xe9bdf7353b3651ee),
C64e(0x1d57ac5a7550d376), C64e(0x3a46c2fea37d7001),
C64e(0xf735c1af98a4d842), C64e(0x78edec209e6b6779),
C64e(0x41836315ea3adba8), C64e(0xfac33b4d32832c83),
C64e(0xa7403b1f1c2747f3), C64e(0x5940f034b72d769a),
C64e(0xe73e4e6cd2214ffd), C64e(0xb8fd8d39dc5759ef),
C64e(0x8d9b0c492b49ebda), C64e(0x5ba2d74968f3700d),
C64e(0x7d3baed07a8d5584), C64e(0xf5a5e9f0e4f88e65),
C64e(0xa0b8a2f436103b53), C64e(0x0ca8079e753eec5a),
C64e(0x9168949256e8884f), C64e(0x5bb05c55f8babc4c),
C64e(0xe3bb3b99f387947b), C64e(0x75daf4d6726b1c5d),
C64e(0x64aeac28dc34b36d), C64e(0x6c34a550b828db71),
C64e(0xf861e2f2108d512a), C64e(0xe3db643359dd75fc),
C64e(0x1cacbcf143ce3fa2), C64e(0x67bbd13c02e843b0),
C64e(0x330a5bca8829a175), C64e(0x7f34194db416535c),
C64e(0x923b94c30e794d1e), C64e(0x797475d7b6eeaf3f),
C64e(0xeaa8d4f7be1a3921), C64e(0x5cf47e094c232751),
C64e(0x26a32453ba323cd2), C64e(0x44a3174a6da6d5ad),
C64e(0xb51d3ea6aff2c908), C64e(0x83593d98916b3c56),
C64e(0x4cf87ca17286604d), C64e(0x46e23ecc086ec7f6),
C64e(0x2f9833b3b1bc765e), C64e(0x2bd666a5efc4e62a),
C64e(0x06f4b6e8bec1d436), C64e(0x74ee8215bcef2163),
C64e(0xfdc14e0df453c969), C64e(0xa77d5ac406585826),
C64e(0x7ec1141606e0fa16), C64e(0x7e90af3d28639d3f),
C64e(0xd2c9f2e3009bd20c), C64e(0x5faace30b7d40c30),
C64e(0x742a5116f2e03298), C64e(0x0deb30d8e3cef89a),
C64e(0x4bc59e7bb5f17992), C64e(0xff51e66e048668d3),
C64e(0x9b234d57e6966731), C64e(0xcce6a6f3170a7505),
C64e(0xb17681d913326cce), C64e(0x3c175284f805a262),
C64e(0xf42bcbb378471547), C64e(0xff46548223936a48),
C64e(0x38df58074e5e6565), C64e(0xf2fc7c89fc86508e),
C64e(0x31702e44d00bca86), C64e(0xf04009a23078474e),
C64e(0x65a0ee39d1f73883), C64e(0xf75ee937e42c3abd),
C64e(0x2197b2260113f86f), C64e(0xa344edd1ef9fdee7),
C64e(0x8ba0df15762592d9), C64e(0x3c85f7f612dc42be),
C64e(0xd8a7ec7cab27b07e), C64e(0x538d7ddaaa3ea8de),
C64e(0xaa25ce93bd0269d8), C64e(0x5af643fd1a7308f9),
C64e(0xc05fefda174a19a5), C64e(0x974d66334cfd216a),
C64e(0x35b49831db411570), C64e(0xea1e0fbbedcd549b),
C64e(0x9ad063a151974072), C64e(0xf6759dbf91476fe2)
};
#define Ceven_hi(r) (C[((r) << 2) + 0])
#define Ceven_lo(r) (C[((r) << 2) + 1])
#define Codd_hi(r) (C[((r) << 2) + 2])
#define Codd_lo(r) (C[((r) << 2) + 3])
#define S(x0, x1, x2, x3, cb, r) do { \
Sb(x0 ## h, x1 ## h, x2 ## h, x3 ## h, cb ## hi(r)); \
Sb(x0 ## l, x1 ## l, x2 ## l, x3 ## l, cb ## lo(r)); \
} while (0)
#define L(x0, x1, x2, x3, x4, x5, x6, x7) do { \
Lb(x0 ## h, x1 ## h, x2 ## h, x3 ## h, \
x4 ## h, x5 ## h, x6 ## h, x7 ## h); \
Lb(x0 ## l, x1 ## l, x2 ## l, x3 ## l, \
x4 ## l, x5 ## l, x6 ## l, x7 ## l); \
} while (0)
#define Wz(x, c, n) do { \
sph_u64 t = (x ## h & (c)) << (n); \
x ## h = ((x ## h >> (n)) & (c)) | t; \
t = (x ## l & (c)) << (n); \
x ## l = ((x ## l >> (n)) & (c)) | t; \
} while (0)
#define W0(x) Wz(x, SPH_C64(0x5555555555555555), 1)
#define W1(x) Wz(x, SPH_C64(0x3333333333333333), 2)
#define W2(x) Wz(x, SPH_C64(0x0F0F0F0F0F0F0F0F), 4)
#define W3(x) Wz(x, SPH_C64(0x00FF00FF00FF00FF), 8)
#define W4(x) Wz(x, SPH_C64(0x0000FFFF0000FFFF), 16)
#define W5(x) Wz(x, SPH_C64(0x00000000FFFFFFFF), 32)
#define W6(x) do { \
sph_u64 t = x ## h; \
x ## h = x ## l; \
x ## l = t; \
} while (0)
__constant static const sph_u64 JH_IV512[] = {
C64e(0x6fd14b963e00aa17), C64e(0x636a2e057a15d543),
C64e(0x8a225e8d0c97ef0b), C64e(0xe9341259f2b3c361),
C64e(0x891da0c1536f801e), C64e(0x2aa9056bea2b6d80),
C64e(0x588eccdb2075baa6), C64e(0xa90f3a76baf83bf7),
C64e(0x0169e60541e34a69), C64e(0x46b58a8e2e6fe65a),
C64e(0x1047a7d0c1843c24), C64e(0x3b6e71b12d5ac199),
C64e(0xcf57f6ec9db1f856), C64e(0xa706887c5716b156),
C64e(0xe3c2fcdfe68517fb), C64e(0x545a4678cc8cdd4b)
};
#define SL(ro) SLu(r + ro, ro)
#define SLu(r, ro) do { \
S(h0, h2, h4, h6, Ceven_, r); \
S(h1, h3, h5, h7, Codd_, r); \
L(h0, h2, h4, h6, h1, h3, h5, h7); \
W ## ro(h1); \
W ## ro(h3); \
W ## ro(h5); \
W ## ro(h7); \
} while (0)
#if SPH_SMALL_FOOTPRINT_JH
/*
* The "small footprint" 64-bit version just uses a partially unrolled
* loop.
*/
#define E8 do { \
unsigned r; \
for (r = 0; r < 42; r += 7) { \
SL(0); \
SL(1); \
SL(2); \
SL(3); \
SL(4); \
SL(5); \
SL(6); \
} \
} while (0)
#else
/*
* On a "true 64-bit" architecture, we can unroll at will.
*/
#define E8 do { \
SLu( 0, 0); \
SLu( 1, 1); \
SLu( 2, 2); \
SLu( 3, 3); \
SLu( 4, 4); \
SLu( 5, 5); \
SLu( 6, 6); \
SLu( 7, 0); \
SLu( 8, 1); \
SLu( 9, 2); \
SLu(10, 3); \
SLu(11, 4); \
SLu(12, 5); \
SLu(13, 6); \
SLu(14, 0); \
SLu(15, 1); \
SLu(16, 2); \
SLu(17, 3); \
SLu(18, 4); \
SLu(19, 5); \
SLu(20, 6); \
SLu(21, 0); \
SLu(22, 1); \
SLu(23, 2); \
SLu(24, 3); \
SLu(25, 4); \
SLu(26, 5); \
SLu(27, 6); \
SLu(28, 0); \
SLu(29, 1); \
SLu(30, 2); \
SLu(31, 3); \
SLu(32, 4); \
SLu(33, 5); \
SLu(34, 6); \
SLu(35, 0); \
SLu(36, 1); \
SLu(37, 2); \
SLu(38, 3); \
SLu(39, 4); \
SLu(40, 5); \
SLu(41, 6); \
} while (0)
#endif

683
kernel/keccak.cl
Unescape View File

@ -0,0 +1,683 @@ @@ -0,0 +1,683 @@
/* $Id: keccak.c 259 2011-07-19 22:11:27Z tp $ */
/*
* Keccak implementation.
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2007-2010 Projet RNRT SAPHIR
*
* 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 Thomas Pornin <thomas.pornin@cryptolog.com>
*/
#ifdef __cplusplus
extern "C"{
#endif
/*
* Parameters:
*
* SPH_KECCAK_64 use a 64-bit type
* SPH_KECCAK_UNROLL number of loops to unroll (0/undef for full unroll)
* SPH_KECCAK_INTERLEAVE use bit-interleaving (32-bit type only)
* SPH_KECCAK_NOCOPY do not copy the state into local variables
*
* If there is no usable 64-bit type, the code automatically switches
* back to the 32-bit implementation.
*
* Some tests on an Intel Core2 Q6600 (both 64-bit and 32-bit, 32 kB L1
* code cache), a PowerPC (G3, 32 kB L1 code cache), an ARM920T core
* (16 kB L1 code cache), and a small MIPS-compatible CPU (Broadcom BCM3302,
* 8 kB L1 code cache), seem to show that the following are optimal:
*
* -- x86, 64-bit: use the 64-bit implementation, unroll 8 rounds,
* do not copy the state; unrolling 2, 6 or all rounds also provides
* near-optimal performance.
* -- x86, 32-bit: use the 32-bit implementation, unroll 6 rounds,
* interleave, do not copy the state. Unrolling 1, 2, 4 or 8 rounds
* also provides near-optimal performance.
* -- PowerPC: use the 64-bit implementation, unroll 8 rounds,
* copy the state. Unrolling 4 or 6 rounds is near-optimal.
* -- ARM: use the 64-bit implementation, unroll 2 or 4 rounds,
* copy the state.
* -- MIPS: use the 64-bit implementation, unroll 2 rounds, copy
* the state. Unrolling only 1 round is also near-optimal.
*
* Also, interleaving does not always yield actual improvements when
* using a 32-bit implementation; in particular when the architecture
* does not offer a native rotation opcode (interleaving replaces one
* 64-bit rotation with two 32-bit rotations, which is a gain only if
* there is a native 32-bit rotation opcode and not a native 64-bit
* rotation opcode; also, interleaving implies a small overhead when
* processing input words).
*
* To sum up:
* -- when possible, use the 64-bit code
* -- exception: on 32-bit x86, use 32-bit code
* -- when using 32-bit code, use interleaving
* -- copy the state, except on x86
* -- unroll 8 rounds on "big" machine, 2 rounds on "small" machines
*/
/*
* Unroll 8 rounds on big systems, 2 rounds on small systems.
*/
#ifndef SPH_KECCAK_UNROLL
#if SPH_SMALL_FOOTPRINT_KECCAK
#define SPH_KECCAK_UNROLL 2
#else
#define SPH_KECCAK_UNROLL 8
#endif
#endif
__constant static const sph_u64 RC[] = {
SPH_C64(0x0000000000000001), SPH_C64(0x0000000000008082),
SPH_C64(0x800000000000808A), SPH_C64(0x8000000080008000),
SPH_C64(0x000000000000808B), SPH_C64(0x0000000080000001),
SPH_C64(0x8000000080008081), SPH_C64(0x8000000000008009),
SPH_C64(0x000000000000008A), SPH_C64(0x0000000000000088),
SPH_C64(0x0000000080008009), SPH_C64(0x000000008000000A),
SPH_C64(0x000000008000808B), SPH_C64(0x800000000000008B),
SPH_C64(0x8000000000008089), SPH_C64(0x8000000000008003),
SPH_C64(0x8000000000008002), SPH_C64(0x8000000000000080),
SPH_C64(0x000000000000800A), SPH_C64(0x800000008000000A),
SPH_C64(0x8000000080008081), SPH_C64(0x8000000000008080),
SPH_C64(0x0000000080000001), SPH_C64(0x8000000080008008)
};
#define DECL64(x) sph_u64 x
#define MOV64(d, s) (d = s)
#define XOR64(d, a, b) (d = a ^ b)
#define AND64(d, a, b) (d = a & b)
#define OR64(d, a, b) (d = a | b)
#define NOT64(d, s) (d = SPH_T64(~s))
#define ROL64(d, v, n) (d = SPH_ROTL64(v, n))
#define XOR64_IOTA XOR64
#define TH_ELT(t, c0, c1, c2, c3, c4, d0, d1, d2, d3, d4) do { \
DECL64(tt0); \
DECL64(tt1); \
DECL64(tt2); \
DECL64(tt3); \
XOR64(tt0, d0, d1); \
XOR64(tt1, d2, d3); \
XOR64(tt0, tt0, d4); \
XOR64(tt0, tt0, tt1); \
ROL64(tt0, tt0, 1); \
XOR64(tt2, c0, c1); \
XOR64(tt3, c2, c3); \
XOR64(tt0, tt0, c4); \
XOR64(tt2, tt2, tt3); \
XOR64(t, tt0, tt2); \
} while (0)
#define THETA(b00, b01, b02, b03, b04, b10, b11, b12, b13, b14, \
b20, b21, b22, b23, b24, b30, b31, b32, b33, b34, \
b40, b41, b42, b43, b44) \
do { \
DECL64(t0); \
DECL64(t1); \
DECL64(t2); \
DECL64(t3); \
DECL64(t4); \
TH_ELT(t0, b40, b41, b42, b43, b44, b10, b11, b12, b13, b14); \
TH_ELT(t1, b00, b01, b02, b03, b04, b20, b21, b22, b23, b24); \
TH_ELT(t2, b10, b11, b12, b13, b14, b30, b31, b32, b33, b34); \
TH_ELT(t3, b20, b21, b22, b23, b24, b40, b41, b42, b43, b44); \
TH_ELT(t4, b30, b31, b32, b33, b34, b00, b01, b02, b03, b04); \
XOR64(b00, b00, t0); \
XOR64(b01, b01, t0); \
XOR64(b02, b02, t0); \
XOR64(b03, b03, t0); \
XOR64(b04, b04, t0); \
XOR64(b10, b10, t1); \
XOR64(b11, b11, t1); \
XOR64(b12, b12, t1); \
XOR64(b13, b13, t1); \
XOR64(b14, b14, t1); \
XOR64(b20, b20, t2); \
XOR64(b21, b21, t2); \
XOR64(b22, b22, t2); \
XOR64(b23, b23, t2); \
XOR64(b24, b24, t2); \
XOR64(b30, b30, t3); \
XOR64(b31, b31, t3); \
XOR64(b32, b32, t3); \
XOR64(b33, b33, t3); \
XOR64(b34, b34, t3); \
XOR64(b40, b40, t4); \
XOR64(b41, b41, t4); \
XOR64(b42, b42, t4); \
XOR64(b43, b43, t4); \
XOR64(b44, b44, t4); \
} while (0)
#define RHO(b00, b01, b02, b03, b04, b10, b11, b12, b13, b14, \
b20, b21, b22, b23, b24, b30, b31, b32, b33, b34, \
b40, b41, b42, b43, b44) \
do { \
/* ROL64(b00, b00, 0); */ \
ROL64(b01, b01, 36); \
ROL64(b02, b02, 3); \
ROL64(b03, b03, 41); \
ROL64(b04, b04, 18); \
ROL64(b10, b10, 1); \
ROL64(b11, b11, 44); \
ROL64(b12, b12, 10); \
ROL64(b13, b13, 45); \
ROL64(b14, b14, 2); \
ROL64(b20, b20, 62); \
ROL64(b21, b21, 6); \
ROL64(b22, b22, 43); \
ROL64(b23, b23, 15); \
ROL64(b24, b24, 61); \
ROL64(b30, b30, 28); \
ROL64(b31, b31, 55); \
ROL64(b32, b32, 25); \
ROL64(b33, b33, 21); \
ROL64(b34, b34, 56); \
ROL64(b40, b40, 27); \
ROL64(b41, b41, 20); \
ROL64(b42, b42, 39); \
ROL64(b43, b43, 8); \
ROL64(b44, b44, 14); \
} while (0)
/*
* The KHI macro integrates the "lane complement" optimization. On input,
* some words are complemented:
* a00 a01 a02 a04 a13 a20 a21 a22 a30 a33 a34 a43
* On output, the following words are complemented:
* a04 a10 a20 a22 a23 a31
*
* The (implicit) permutation and the theta expansion will bring back
* the input mask for the next round.
*/
#define KHI_XO(d, a, b, c) do { \
DECL64(kt); \
OR64(kt, b, c); \
XOR64(d, a, kt); \
} while (0)
#define KHI_XA(d, a, b, c) do { \
DECL64(kt); \
AND64(kt, b, c); \
XOR64(d, a, kt); \
} while (0)
#define KHI(b00, b01, b02, b03, b04, b10, b11, b12, b13, b14, \
b20, b21, b22, b23, b24, b30, b31, b32, b33, b34, \
b40, b41, b42, b43, b44) \
do { \
DECL64(c0); \
DECL64(c1); \
DECL64(c2); \
DECL64(c3); \
DECL64(c4); \
DECL64(bnn); \
NOT64(bnn, b20); \
KHI_XO(c0, b00, b10, b20); \
KHI_XO(c1, b10, bnn, b30); \
KHI_XA(c2, b20, b30, b40); \
KHI_XO(c3, b30, b40, b00); \
KHI_XA(c4, b40, b00, b10); \
MOV64(b00, c0); \
MOV64(b10, c1); \
MOV64(b20, c2); \
MOV64(b30, c3); \
MOV64(b40, c4); \
NOT64(bnn, b41); \
KHI_XO(c0, b01, b11, b21); \
KHI_XA(c1, b11, b21, b31); \
KHI_XO(c2, b21, b31, bnn); \
KHI_XO(c3, b31, b41, b01); \
KHI_XA(c4, b41, b01, b11); \
MOV64(b01, c0); \
MOV64(b11, c1); \
MOV64(b21, c2); \
MOV64(b31, c3); \
MOV64(b41, c4); \
NOT64(bnn, b32); \
KHI_XO(c0, b02, b12, b22); \
KHI_XA(c1, b12, b22, b32); \
KHI_XA(c2, b22, bnn, b42); \
KHI_XO(c3, bnn, b42, b02); \
KHI_XA(c4, b42, b02, b12); \
MOV64(b02, c0); \
MOV64(b12, c1); \
MOV64(b22, c2); \
MOV64(b32, c3); \
MOV64(b42, c4); \
NOT64(bnn, b33); \
KHI_XA(c0, b03, b13, b23); \
KHI_XO(c1, b13, b23, b33); \
KHI_XO(c2, b23, bnn, b43); \
KHI_XA(c3, bnn, b43, b03); \
KHI_XO(c4, b43, b03, b13); \
MOV64(b03, c0); \
MOV64(b13, c1); \
MOV64(b23, c2); \
MOV64(b33, c3); \
MOV64(b43, c4); \
NOT64(bnn, b14); \
KHI_XA(c0, b04, bnn, b24); \
KHI_XO(c1, bnn, b24, b34); \
KHI_XA(c2, b24, b34, b44); \
KHI_XO(c3, b34, b44, b04); \
KHI_XA(c4, b44, b04, b14); \
MOV64(b04, c0); \
MOV64(b14, c1); \
MOV64(b24, c2); \
MOV64(b34, c3); \
MOV64(b44, c4); \
} while (0)
#define IOTA(r) XOR64_IOTA(a00, a00, r)
#define P0 a00, a01, a02, a03, a04, a10, a11, a12, a13, a14, a20, a21, \
a22, a23, a24, a30, a31, a32, a33, a34, a40, a41, a42, a43, a44
#define P1 a00, a30, a10, a40, a20, a11, a41, a21, a01, a31, a22, a02, \
a32, a12, a42, a33, a13, a43, a23, a03, a44, a24, a04, a34, a14
#define P2 a00, a33, a11, a44, a22, a41, a24, a02, a30, a13, a32, a10, \
a43, a21, a04, a23, a01, a34, a12, a40, a14, a42, a20, a03, a31
#define P3 a00, a23, a41, a14, a32, a24, a42, a10, a33, a01, a43, a11, \
a34, a02, a20, a12, a30, a03, a21, a44, a31, a04, a22, a40, a13
#define P4 a00, a12, a24, a31, a43, a42, a04, a11, a23, a30, a34, a41, \
a03, a10, a22, a21, a33, a40, a02, a14, a13, a20, a32, a44, a01
#define P5 a00, a21, a42, a13, a34, a04, a20, a41, a12, a33, a03, a24, \
a40, a11, a32, a02, a23, a44, a10, a31, a01, a22, a43, a14, a30
#define P6 a00, a02, a04, a01, a03, a20, a22, a24, a21, a23, a40, a42, \
a44, a41, a43, a10, a12, a14, a11, a13, a30, a32, a34, a31, a33
#define P7 a00, a10, a20, a30, a40, a22, a32, a42, a02, a12, a44, a04, \
a14, a24, a34, a11, a21, a31, a41, a01, a33, a43, a03, a13, a23
#define P8 a00, a11, a22, a33, a44, a32, a43, a04, a10, a21, a14, a20, \
a31, a42, a03, a41, a02, a13, a24, a30, a23, a34, a40, a01, a12
#define P9 a00, a41, a32, a23, a14, a43, a34, a20, a11, a02, a31, a22, \
a13, a04, a40, a24, a10, a01, a42, a33, a12, a03, a44, a30, a21
#define P10 a00, a24, a43, a12, a31, a34, a03, a22, a41, a10, a13, a32, \
a01, a20, a44, a42, a11, a30, a04, a23, a21, a40, a14, a33, a02
#define P11 a00, a42, a34, a21, a13, a03, a40, a32, a24, a11, a01, a43, \
a30, a22, a14, a04, a41, a33, a20, a12, a02, a44, a31, a23, a10
#define P12 a00, a04, a03, a02, a01, a40, a44, a43, a42, a41, a30, a34, \
a33, a32, a31, a20, a24, a23, a22, a21, a10, a14, a13, a12, a11
#define P13 a00, a20, a40, a10, a30, a44, a14, a34, a04, a24, a33, a03, \
a23, a43, a13, a22, a42, a12, a32, a02, a11, a31, a01, a21, a41
#define P14 a00, a22, a44, a11, a33, a14, a31, a03, a20, a42, a23, a40, \
a12, a34, a01, a32, a04, a21, a43, a10, a41, a13, a30, a02, a24
#define P15 a00, a32, a14, a41, a23, a31, a13, a40, a22, a04, a12, a44, \
a21, a03, a30, a43, a20, a02, a34, a11, a24, a01, a33, a10, a42
#define P16 a00, a43, a31, a24, a12, a13, a01, a44, a32, a20, a21, a14, \
a02, a40, a33, a34, a22, a10, a03, a41, a42, a30, a23, a11, a04
#define P17 a00, a34, a13, a42, a21, a01, a30, a14, a43, a22, a02, a31, \
a10, a44, a23, a03, a32, a11, a40, a24, a04, a33, a12, a41, a20
#define P18 a00, a03, a01, a04, a02, a30, a33, a31, a34, a32, a10, a13, \
a11, a14, a12, a40, a43, a41, a44, a42, a20, a23, a21, a24, a22
#define P19 a00, a40, a30, a20, a10, a33, a23, a13, a03, a43, a11, a01, \
a41, a31, a21, a44, a34, a24, a14, a04, a22, a12, a02, a42, a32
#define P20 a00, a44, a33, a22, a11, a23, a12, a01, a40, a34, a41, a30, \
a24, a13, a02, a14, a03, a42, a31, a20, a32, a21, a10, a04, a43
#define P21 a00, a14, a23, a32, a41, a12, a21, a30, a44, a03, a24, a33, \
a42, a01, a10, a31, a40, a04, a13, a22, a43, a02, a11, a20, a34
#define P22 a00, a31, a12, a43, a24, a21, a02, a33, a14, a40, a42, a23, \
a04, a30, a11, a13, a44, a20, a01, a32, a34, a10, a41, a22, a03
#define P23 a00, a13, a21, a34, a42, a02, a10, a23, a31, a44, a04, a12, \
a20, a33, a41, a01, a14, a22, a30, a43, a03, a11, a24, a32, a40
#define P1_TO_P0 do { \
DECL64(t); \
MOV64(t, a01); \
MOV64(a01, a30); \
MOV64(a30, a33); \
MOV64(a33, a23); \
MOV64(a23, a12); \
MOV64(a12, a21); \
MOV64(a21, a02); \
MOV64(a02, a10); \
MOV64(a10, a11); \
MOV64(a11, a41); \
MOV64(a41, a24); \
MOV64(a24, a42); \
MOV64(a42, a04); \
MOV64(a04, a20); \
MOV64(a20, a22); \
MOV64(a22, a32); \
MOV64(a32, a43); \
MOV64(a43, a34); \
MOV64(a34, a03); \
MOV64(a03, a40); \
MOV64(a40, a44); \
MOV64(a44, a14); \
MOV64(a14, a31); \
MOV64(a31, a13); \
MOV64(a13, t); \
} while (0)
#define P2_TO_P0 do { \
DECL64(t); \
MOV64(t, a01); \
MOV64(a01, a33); \
MOV64(a33, a12); \
MOV64(a12, a02); \
MOV64(a02, a11); \
MOV64(a11, a24); \
MOV64(a24, a04); \
MOV64(a04, a22); \
MOV64(a22, a43); \
MOV64(a43, a03); \
MOV64(a03, a44); \
MOV64(a44, a31); \
MOV64(a31, t); \
MOV64(t, a10); \
MOV64(a10, a41); \
MOV64(a41, a42); \
MOV64(a42, a20); \
MOV64(a20, a32); \
MOV64(a32, a34); \
MOV64(a34, a40); \
MOV64(a40, a14); \
MOV64(a14, a13); \
MOV64(a13, a30); \
MOV64(a30, a23); \
MOV64(a23, a21); \
MOV64(a21, t); \
} while (0)
#define P4_TO_P0 do { \
DECL64(t); \
MOV64(t, a01); \
MOV64(a01, a12); \
MOV64(a12, a11); \
MOV64(a11, a04); \
MOV64(a04, a43); \
MOV64(a43, a44); \
MOV64(a44, t); \
MOV64(t, a02); \
MOV64(a02, a24); \
MOV64(a24, a22); \
MOV64(a22, a03); \
MOV64(a03, a31); \
MOV64(a31, a33); \
MOV64(a33, t); \
MOV64(t, a10); \
MOV64(a10, a42); \
MOV64(a42, a32); \
MOV64(a32, a40); \
MOV64(a40, a13); \
MOV64(a13, a23); \
MOV64(a23, t); \
MOV64(t, a14); \
MOV64(a14, a30); \
MOV64(a30, a21); \
MOV64(a21, a41); \
MOV64(a41, a20); \
MOV64(a20, a34); \
MOV64(a34, t); \
} while (0)
#define P6_TO_P0 do { \
DECL64(t); \
MOV64(t, a01); \
MOV64(a01, a02); \
MOV64(a02, a04); \
MOV64(a04, a03); \
MOV64(a03, t); \
MOV64(t, a10); \
MOV64(a10, a20); \
MOV64(a20, a40); \
MOV64(a40, a30); \
MOV64(a30, t); \
MOV64(t, a11); \
MOV64(a11, a22); \
MOV64(a22, a44); \
MOV64(a44, a33); \
MOV64(a33, t); \
MOV64(t, a12); \
MOV64(a12, a24); \
MOV64(a24, a43); \
MOV64(a43, a31); \
MOV64(a31, t); \
MOV64(t, a13); \
MOV64(a13, a21); \
MOV64(a21, a42); \
MOV64(a42, a34); \
MOV64(a34, t); \
MOV64(t, a14); \
MOV64(a14, a23); \
MOV64(a23, a41); \
MOV64(a41, a32); \
MOV64(a32, t); \
} while (0)
#define P8_TO_P0 do { \
DECL64(t); \
MOV64(t, a01); \
MOV64(a01, a11); \
MOV64(a11, a43); \
MOV64(a43, t); \
MOV64(t, a02); \
MOV64(a02, a22); \
MOV64(a22, a31); \
MOV64(a31, t); \
MOV64(t, a03); \
MOV64(a03, a33); \
MOV64(a33, a24); \
MOV64(a24, t); \
MOV64(t, a04); \
MOV64(a04, a44); \
MOV64(a44, a12); \
MOV64(a12, t); \
MOV64(t, a10); \
MOV64(a10, a32); \
MOV64(a32, a13); \
MOV64(a13, t); \
MOV64(t, a14); \
MOV64(a14, a21); \
MOV64(a21, a20); \
MOV64(a20, t); \
MOV64(t, a23); \
MOV64(a23, a42); \
MOV64(a42, a40); \
MOV64(a40, t); \
MOV64(t, a30); \
MOV64(a30, a41); \
MOV64(a41, a34); \
MOV64(a34, t); \
} while (0)
#define P12_TO_P0 do { \
DECL64(t); \
MOV64(t, a01); \
MOV64(a01, a04); \
MOV64(a04, t); \
MOV64(t, a02); \
MOV64(a02, a03); \
MOV64(a03, t); \
MOV64(t, a10); \
MOV64(a10, a40); \
MOV64(a40, t); \
MOV64(t, a11); \
MOV64(a11, a44); \
MOV64(a44, t); \
MOV64(t, a12); \
MOV64(a12, a43); \
MOV64(a43, t); \
MOV64(t, a13); \
MOV64(a13, a42); \
MOV64(a42, t); \
MOV64(t, a14); \
MOV64(a14, a41); \
MOV64(a41, t); \
MOV64(t, a20); \
MOV64(a20, a30); \
MOV64(a30, t); \
MOV64(t, a21); \
MOV64(a21, a34); \
MOV64(a34, t); \
MOV64(t, a22); \
MOV64(a22, a33); \
MOV64(a33, t); \
MOV64(t, a23); \
MOV64(a23, a32); \
MOV64(a32, t); \
MOV64(t, a24); \
MOV64(a24, a31); \
MOV64(a31, t); \
} while (0)
#define LPAR (
#define RPAR )
#define KF_ELT(r, s, k) do { \
THETA LPAR P ## r RPAR; \
RHO LPAR P ## r RPAR; \
KHI LPAR P ## s RPAR; \
IOTA(k); \
} while (0)
#define DO(x) x
#define KECCAK_F_1600 DO(KECCAK_F_1600_)
#if SPH_KECCAK_UNROLL == 1
#define KECCAK_F_1600_ do { \
int j; \
for (j = 0; j < 24; j ++) { \
KF_ELT( 0, 1, RC[j + 0]); \
P1_TO_P0; \
} \
} while (0)
#elif SPH_KECCAK_UNROLL == 2
#define KECCAK_F_1600_ do { \
int j; \
for (j = 0; j < 24; j += 2) { \
KF_ELT( 0, 1, RC[j + 0]); \
KF_ELT( 1, 2, RC[j + 1]); \
P2_TO_P0; \
} \
} while (0)
#elif SPH_KECCAK_UNROLL == 4
#define KECCAK_F_1600_ do { \
int j; \
for (j = 0; j < 24; j += 4) { \
KF_ELT( 0, 1, RC[j + 0]); \
KF_ELT( 1, 2, RC[j + 1]); \
KF_ELT( 2, 3, RC[j + 2]); \
KF_ELT( 3, 4, RC[j + 3]); \
P4_TO_P0; \
} \
} while (0)
#elif SPH_KECCAK_UNROLL == 6
#define KECCAK_F_1600_ do { \
int j; \
for (j = 0; j < 24; j += 6) { \
KF_ELT( 0, 1, RC[j + 0]); \
KF_ELT( 1, 2, RC[j + 1]); \
KF_ELT( 2, 3, RC[j + 2]); \
KF_ELT( 3, 4, RC[j + 3]); \
KF_ELT( 4, 5, RC[j + 4]); \
KF_ELT( 5, 6, RC[j + 5]); \
P6_TO_P0; \
} \
} while (0)
#elif SPH_KECCAK_UNROLL == 8
#define KECCAK_F_1600_ do { \
int j; \
for (j = 0; j < 24; j += 8) { \
KF_ELT( 0, 1, RC[j + 0]); \
KF_ELT( 1, 2, RC[j + 1]); \
KF_ELT( 2, 3, RC[j + 2]); \
KF_ELT( 3, 4, RC[j + 3]); \
KF_ELT( 4, 5, RC[j + 4]); \
KF_ELT( 5, 6, RC[j + 5]); \
KF_ELT( 6, 7, RC[j + 6]); \
KF_ELT( 7, 8, RC[j + 7]); \
P8_TO_P0; \
} \
} while (0)
#elif SPH_KECCAK_UNROLL == 12
#define KECCAK_F_1600_ do { \
int j; \
for (j = 0; j < 24; j += 12) { \
KF_ELT( 0, 1, RC[j + 0]); \
KF_ELT( 1, 2, RC[j + 1]); \
KF_ELT( 2, 3, RC[j + 2]); \
KF_ELT( 3, 4, RC[j + 3]); \
KF_ELT( 4, 5, RC[j + 4]); \
KF_ELT( 5, 6, RC[j + 5]); \
KF_ELT( 6, 7, RC[j + 6]); \
KF_ELT( 7, 8, RC[j + 7]); \
KF_ELT( 8, 9, RC[j + 8]); \
KF_ELT( 9, 10, RC[j + 9]); \
KF_ELT(10, 11, RC[j + 10]); \
KF_ELT(11, 12, RC[j + 11]); \
P12_TO_P0; \
} \
} while (0)
#elif SPH_KECCAK_UNROLL == 0
#define KECCAK_F_1600_ do { \
KF_ELT( 0, 1, RC[ 0]); \
KF_ELT( 1, 2, RC[ 1]); \
KF_ELT( 2, 3, RC[ 2]); \
KF_ELT( 3, 4, RC[ 3]); \
KF_ELT( 4, 5, RC[ 4]); \
KF_ELT( 5, 6, RC[ 5]); \
KF_ELT( 6, 7, RC[ 6]); \
KF_ELT( 7, 8, RC[ 7]); \
KF_ELT( 8, 9, RC[ 8]); \
KF_ELT( 9, 10, RC[ 9]); \
KF_ELT(10, 11, RC[10]); \
KF_ELT(11, 12, RC[11]); \
KF_ELT(12, 13, RC[12]); \
KF_ELT(13, 14, RC[13]); \
KF_ELT(14, 15, RC[14]); \
KF_ELT(15, 16, RC[15]); \
KF_ELT(16, 17, RC[16]); \
KF_ELT(17, 18, RC[17]); \
KF_ELT(18, 19, RC[18]); \
KF_ELT(19, 20, RC[19]); \
KF_ELT(20, 21, RC[20]); \
KF_ELT(21, 22, RC[21]); \
KF_ELT(22, 23, RC[22]); \
KF_ELT(23, 0, RC[23]); \
} while (0)
#else
#error Unimplemented unroll count for Keccak.
#endif

488
kernel/luffa.cl
Unescape View File

@ -0,0 +1,488 @@ @@ -0,0 +1,488 @@
/* $Id: luffa.c 219 2010-06-08 17:24:41Z tp $ */
/*
* Luffa implementation.
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2007-2010 Projet RNRT SAPHIR
*
* 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 Thomas Pornin <thomas.pornin@cryptolog.com>
*/
#ifdef __cplusplus
extern "C"{
#endif
#if SPH_64_TRUE && !defined SPH_LUFFA_PARALLEL
#define SPH_LUFFA_PARALLEL 1
#endif
__constant static const sph_u32 V_INIT[5][8] = {
{
SPH_C32(0x6d251e69), SPH_C32(0x44b051e0),
SPH_C32(0x4eaa6fb4), SPH_C32(0xdbf78465),
SPH_C32(0x6e292011), SPH_C32(0x90152df4),
SPH_C32(0xee058139), SPH_C32(0xdef610bb)
}, {
SPH_C32(0xc3b44b95), SPH_C32(0xd9d2f256),
SPH_C32(0x70eee9a0), SPH_C32(0xde099fa3),
SPH_C32(0x5d9b0557), SPH_C32(0x8fc944b3),
SPH_C32(0xcf1ccf0e), SPH_C32(0x746cd581)
}, {
SPH_C32(0xf7efc89d), SPH_C32(0x5dba5781),
SPH_C32(0x04016ce5), SPH_C32(0xad659c05),
SPH_C32(0x0306194f), SPH_C32(0x666d1836),
SPH_C32(0x24aa230a), SPH_C32(0x8b264ae7)
}, {
SPH_C32(0x858075d5), SPH_C32(0x36d79cce),
SPH_C32(0xe571f7d7), SPH_C32(0x204b1f67),
SPH_C32(0x35870c6a), SPH_C32(0x57e9e923),
SPH_C32(0x14bcb808), SPH_C32(0x7cde72ce)
}, {
SPH_C32(0x6c68e9be), SPH_C32(0x5ec41e22),
SPH_C32(0xc825b7c7), SPH_C32(0xaffb4363),
SPH_C32(0xf5df3999), SPH_C32(0x0fc688f1),
SPH_C32(0xb07224cc), SPH_C32(0x03e86cea)
}
};
__constant static const sph_u32 RC00[8] = {
SPH_C32(0x303994a6), SPH_C32(0xc0e65299),
SPH_C32(0x6cc33a12), SPH_C32(0xdc56983e),
SPH_C32(0x1e00108f), SPH_C32(0x7800423d),
SPH_C32(0x8f5b7882), SPH_C32(0x96e1db12)
};
__constant static const sph_u32 RC04[8] = {
SPH_C32(0xe0337818), SPH_C32(0x441ba90d),
SPH_C32(0x7f34d442), SPH_C32(0x9389217f),
SPH_C32(0xe5a8bce6), SPH_C32(0x5274baf4),
SPH_C32(0x26889ba7), SPH_C32(0x9a226e9d)
};
__constant static const sph_u32 RC10[8] = {
SPH_C32(0xb6de10ed), SPH_C32(0x70f47aae),
SPH_C32(0x0707a3d4), SPH_C32(0x1c1e8f51),
SPH_C32(0x707a3d45), SPH_C32(0xaeb28562),
SPH_C32(0xbaca1589), SPH_C32(0x40a46f3e)
};
__constant static const sph_u32 RC14[8] = {
SPH_C32(0x01685f3d), SPH_C32(0x05a17cf4),
SPH_C32(0xbd09caca), SPH_C32(0xf4272b28),
SPH_C32(0x144ae5cc), SPH_C32(0xfaa7ae2b),
SPH_C32(0x2e48f1c1), SPH_C32(0xb923c704)
};
#if SPH_LUFFA_PARALLEL
__constant static const sph_u64 RCW010[8] = {
SPH_C64(0xb6de10ed303994a6), SPH_C64(0x70f47aaec0e65299),
SPH_C64(0x0707a3d46cc33a12), SPH_C64(0x1c1e8f51dc56983e),
SPH_C64(0x707a3d451e00108f), SPH_C64(0xaeb285627800423d),
SPH_C64(0xbaca15898f5b7882), SPH_C64(0x40a46f3e96e1db12)
};
__constant static const sph_u64 RCW014[8] = {
SPH_C64(0x01685f3de0337818), SPH_C64(0x05a17cf4441ba90d),
SPH_C64(0xbd09caca7f34d442), SPH_C64(0xf4272b289389217f),
SPH_C64(0x144ae5cce5a8bce6), SPH_C64(0xfaa7ae2b5274baf4),
SPH_C64(0x2e48f1c126889ba7), SPH_C64(0xb923c7049a226e9d)
};
#endif
__constant static const sph_u32 RC20[8] = {
SPH_C32(0xfc20d9d2), SPH_C32(0x34552e25),
SPH_C32(0x7ad8818f), SPH_C32(0x8438764a),
SPH_C32(0xbb6de032), SPH_C32(0xedb780c8),
SPH_C32(0xd9847356), SPH_C32(0xa2c78434)
};
__constant static const sph_u32 RC24[8] = {
SPH_C32(0xe25e72c1), SPH_C32(0xe623bb72),
SPH_C32(0x5c58a4a4), SPH_C32(0x1e38e2e7),
SPH_C32(0x78e38b9d), SPH_C32(0x27586719),
SPH_C32(0x36eda57f), SPH_C32(0x703aace7)
};
__constant static const sph_u32 RC30[8] = {
SPH_C32(0xb213afa5), SPH_C32(0xc84ebe95),
SPH_C32(0x4e608a22), SPH_C32(0x56d858fe),
SPH_C32(0x343b138f), SPH_C32(0xd0ec4e3d),
SPH_C32(0x2ceb4882), SPH_C32(0xb3ad2208)
};
__constant static const sph_u32 RC34[8] = {
SPH_C32(0xe028c9bf), SPH_C32(0x44756f91),
SPH_C32(0x7e8fce32), SPH_C32(0x956548be),
SPH_C32(0xfe191be2), SPH_C32(0x3cb226e5),
SPH_C32(0x5944a28e), SPH_C32(0xa1c4c355)
};
#if SPH_LUFFA_PARALLEL
__constant static const sph_u64 RCW230[8] = {
SPH_C64(0xb213afa5fc20d9d2), SPH_C64(0xc84ebe9534552e25),
SPH_C64(0x4e608a227ad8818f), SPH_C64(0x56d858fe8438764a),
SPH_C64(0x343b138fbb6de032), SPH_C64(0xd0ec4e3dedb780c8),
SPH_C64(0x2ceb4882d9847356), SPH_C64(0xb3ad2208a2c78434)
};
__constant static const sph_u64 RCW234[8] = {
SPH_C64(0xe028c9bfe25e72c1), SPH_C64(0x44756f91e623bb72),
SPH_C64(0x7e8fce325c58a4a4), SPH_C64(0x956548be1e38e2e7),
SPH_C64(0xfe191be278e38b9d), SPH_C64(0x3cb226e527586719),
SPH_C64(0x5944a28e36eda57f), SPH_C64(0xa1c4c355703aace7)
};
#endif
__constant static const sph_u32 RC40[8] = {
SPH_C32(0xf0d2e9e3), SPH_C32(0xac11d7fa),
SPH_C32(0x1bcb66f2), SPH_C32(0x6f2d9bc9),
SPH_C32(0x78602649), SPH_C32(0x8edae952),
SPH_C32(0x3b6ba548), SPH_C32(0xedae9520)
};
__constant static const sph_u32 RC44[8] = {
SPH_C32(0x5090d577), SPH_C32(0x2d1925ab),
SPH_C32(0xb46496ac), SPH_C32(0xd1925ab0),
SPH_C32(0x29131ab6), SPH_C32(0x0fc053c3),
SPH_C32(0x3f014f0c), SPH_C32(0xfc053c31)
};
#define DECL_TMP8(w) \
sph_u32 w ## 0, w ## 1, w ## 2, w ## 3, w ## 4, w ## 5, w ## 6, w ## 7;
#define M2(d, s) do { \
sph_u32 tmp = s ## 7; \
d ## 7 = s ## 6; \
d ## 6 = s ## 5; \
d ## 5 = s ## 4; \
d ## 4 = s ## 3 ^ tmp; \
d ## 3 = s ## 2 ^ tmp; \
d ## 2 = s ## 1; \
d ## 1 = s ## 0 ^ tmp; \
d ## 0 = tmp; \
} while (0)
#define XOR(d, s1, s2) do { \
d ## 0 = s1 ## 0 ^ s2 ## 0; \
d ## 1 = s1 ## 1 ^ s2 ## 1; \
d ## 2 = s1 ## 2 ^ s2 ## 2; \
d ## 3 = s1 ## 3 ^ s2 ## 3; \
d ## 4 = s1 ## 4 ^ s2 ## 4; \
d ## 5 = s1 ## 5 ^ s2 ## 5; \
d ## 6 = s1 ## 6 ^ s2 ## 6; \
d ## 7 = s1 ## 7 ^ s2 ## 7; \
} while (0)
#if SPH_LUFFA_PARALLEL
#define SUB_CRUMB_GEN(a0, a1, a2, a3, width) do { \
sph_u ## width tmp; \
tmp = (a0); \
(a0) |= (a1); \
(a2) ^= (a3); \
(a1) = SPH_T ## width(~(a1)); \
(a0) ^= (a3); \
(a3) &= tmp; \
(a1) ^= (a3); \
(a3) ^= (a2); \
(a2) &= (a0); \
(a0) = SPH_T ## width(~(a0)); \
(a2) ^= (a1); \
(a1) |= (a3); \
tmp ^= (a1); \
(a3) ^= (a2); \
(a2) &= (a1); \
(a1) ^= (a0); \
(a0) = tmp; \
} while (0)
#define SUB_CRUMB(a0, a1, a2, a3) SUB_CRUMB_GEN(a0, a1, a2, a3, 32)
#define SUB_CRUMBW(a0, a1, a2, a3) SUB_CRUMB_GEN(a0, a1, a2, a3, 64)
#define MIX_WORDW(u, v) do { \
sph_u32 ul, uh, vl, vh; \
(v) ^= (u); \
ul = SPH_T32((sph_u32)(u)); \
uh = SPH_T32((sph_u32)((u) >> 32)); \
vl = SPH_T32((sph_u32)(v)); \
vh = SPH_T32((sph_u32)((v) >> 32)); \
ul = SPH_ROTL32(ul, 2) ^ vl; \
vl = SPH_ROTL32(vl, 14) ^ ul; \
ul = SPH_ROTL32(ul, 10) ^ vl; \
vl = SPH_ROTL32(vl, 1); \
uh = SPH_ROTL32(uh, 2) ^ vh; \
vh = SPH_ROTL32(vh, 14) ^ uh; \
uh = SPH_ROTL32(uh, 10) ^ vh; \
vh = SPH_ROTL32(vh, 1); \
(u) = (sph_u64)ul | ((sph_u64)uh << 32); \
(v) = (sph_u64)vl | ((sph_u64)vh << 32); \
} while (0)
#else
#define SUB_CRUMB(a0, a1, a2, a3) do { \
sph_u32 tmp; \
tmp = (a0); \
(a0) |= (a1); \
(a2) ^= (a3); \
(a1) = SPH_T32(~(a1)); \
(a0) ^= (a3); \
(a3) &= tmp; \
(a1) ^= (a3); \
(a3) ^= (a2); \
(a2) &= (a0); \
(a0) = SPH_T32(~(a0)); \
(a2) ^= (a1); \
(a1) |= (a3); \
tmp ^= (a1); \
(a3) ^= (a2); \
(a2) &= (a1); \
(a1) ^= (a0); \
(a0) = tmp; \
} while (0)
#endif
#define MIX_WORD(u, v) do { \
(v) ^= (u); \
(u) = SPH_ROTL32((u), 2) ^ (v); \
(v) = SPH_ROTL32((v), 14) ^ (u); \
(u) = SPH_ROTL32((u), 10) ^ (v); \
(v) = SPH_ROTL32((v), 1); \
} while (0)
#define MI5 do { \
DECL_TMP8(a) \
DECL_TMP8(b) \
XOR(a, V0, V1); \
XOR(b, V2, V3); \
XOR(a, a, b); \
XOR(a, a, V4); \
M2(a, a); \
XOR(V0, a, V0); \
XOR(V1, a, V1); \
XOR(V2, a, V2); \
XOR(V3, a, V3); \
XOR(V4, a, V4); \
M2(b, V0); \
XOR(b, b, V1); \
M2(V1, V1); \
XOR(V1, V1, V2); \
M2(V2, V2); \
XOR(V2, V2, V3); \
M2(V3, V3); \
XOR(V3, V3, V4); \
M2(V4, V4); \
XOR(V4, V4, V0); \
M2(V0, b); \
XOR(V0, V0, V4); \
M2(V4, V4); \
XOR(V4, V4, V3); \
M2(V3, V3); \
XOR(V3, V3, V2); \
M2(V2, V2); \
XOR(V2, V2, V1); \
M2(V1, V1); \
XOR(V1, V1, b); \
XOR(V0, V0, M); \
M2(M, M); \
XOR(V1, V1, M); \
M2(M, M); \
XOR(V2, V2, M); \
M2(M, M); \
XOR(V3, V3, M); \
M2(M, M); \
XOR(V4, V4, M); \
} while (0)
#define TWEAK5 do { \
V14 = SPH_ROTL32(V14, 1); \
V15 = SPH_ROTL32(V15, 1); \
V16 = SPH_ROTL32(V16, 1); \
V17 = SPH_ROTL32(V17, 1); \
V24 = SPH_ROTL32(V24, 2); \
V25 = SPH_ROTL32(V25, 2); \
V26 = SPH_ROTL32(V26, 2); \
V27 = SPH_ROTL32(V27, 2); \
V34 = SPH_ROTL32(V34, 3); \
V35 = SPH_ROTL32(V35, 3); \
V36 = SPH_ROTL32(V36, 3); \
V37 = SPH_ROTL32(V37, 3); \
V44 = SPH_ROTL32(V44, 4); \
V45 = SPH_ROTL32(V45, 4); \
V46 = SPH_ROTL32(V46, 4); \
V47 = SPH_ROTL32(V47, 4); \
} while (0)
#if SPH_LUFFA_PARALLEL
#define LUFFA_P5 do { \
int r; \
sph_u64 W0, W1, W2, W3, W4, W5, W6, W7; \
TWEAK5; \
W0 = (sph_u64)V00 | ((sph_u64)V10 << 32); \
W1 = (sph_u64)V01 | ((sph_u64)V11 << 32); \
W2 = (sph_u64)V02 | ((sph_u64)V12 << 32); \
W3 = (sph_u64)V03 | ((sph_u64)V13 << 32); \
W4 = (sph_u64)V04 | ((sph_u64)V14 << 32); \
W5 = (sph_u64)V05 | ((sph_u64)V15 << 32); \
W6 = (sph_u64)V06 | ((sph_u64)V16 << 32); \
W7 = (sph_u64)V07 | ((sph_u64)V17 << 32); \
for (r = 0; r < 8; r ++) { \
SUB_CRUMBW(W0, W1, W2, W3); \
SUB_CRUMBW(W5, W6, W7, W4); \
MIX_WORDW(W0, W4); \
MIX_WORDW(W1, W5); \
MIX_WORDW(W2, W6); \
MIX_WORDW(W3, W7); \
W0 ^= RCW010[r]; \
W4 ^= RCW014[r]; \
} \
V00 = SPH_T32((sph_u32)W0); \
V10 = SPH_T32((sph_u32)(W0 >> 32)); \
V01 = SPH_T32((sph_u32)W1); \
V11 = SPH_T32((sph_u32)(W1 >> 32)); \
V02 = SPH_T32((sph_u32)W2); \
V12 = SPH_T32((sph_u32)(W2 >> 32)); \
V03 = SPH_T32((sph_u32)W3); \
V13 = SPH_T32((sph_u32)(W3 >> 32)); \
V04 = SPH_T32((sph_u32)W4); \
V14 = SPH_T32((sph_u32)(W4 >> 32)); \
V05 = SPH_T32((sph_u32)W5); \
V15 = SPH_T32((sph_u32)(W5 >> 32)); \
V06 = SPH_T32((sph_u32)W6); \
V16 = SPH_T32((sph_u32)(W6 >> 32)); \
V07 = SPH_T32((sph_u32)W7); \
V17 = SPH_T32((sph_u32)(W7 >> 32)); \
W0 = (sph_u64)V20 | ((sph_u64)V30 << 32); \
W1 = (sph_u64)V21 | ((sph_u64)V31 << 32); \
W2 = (sph_u64)V22 | ((sph_u64)V32 << 32); \
W3 = (sph_u64)V23 | ((sph_u64)V33 << 32); \
W4 = (sph_u64)V24 | ((sph_u64)V34 << 32); \
W5 = (sph_u64)V25 | ((sph_u64)V35 << 32); \
W6 = (sph_u64)V26 | ((sph_u64)V36 << 32); \
W7 = (sph_u64)V27 | ((sph_u64)V37 << 32); \
for (r = 0; r < 8; r ++) { \
SUB_CRUMBW(W0, W1, W2, W3); \
SUB_CRUMBW(W5, W6, W7, W4); \
MIX_WORDW(W0, W4); \
MIX_WORDW(W1, W5); \
MIX_WORDW(W2, W6); \
MIX_WORDW(W3, W7); \
W0 ^= RCW230[r]; \
W4 ^= RCW234[r]; \
} \
V20 = SPH_T32((sph_u32)W0); \
V30 = SPH_T32((sph_u32)(W0 >> 32)); \
V21 = SPH_T32((sph_u32)W1); \
V31 = SPH_T32((sph_u32)(W1 >> 32)); \
V22 = SPH_T32((sph_u32)W2); \
V32 = SPH_T32((sph_u32)(W2 >> 32)); \
V23 = SPH_T32((sph_u32)W3); \
V33 = SPH_T32((sph_u32)(W3 >> 32)); \
V24 = SPH_T32((sph_u32)W4); \
V34 = SPH_T32((sph_u32)(W4 >> 32)); \
V25 = SPH_T32((sph_u32)W5); \
V35 = SPH_T32((sph_u32)(W5 >> 32)); \
V26 = SPH_T32((sph_u32)W6); \
V36 = SPH_T32((sph_u32)(W6 >> 32)); \
V27 = SPH_T32((sph_u32)W7); \
V37 = SPH_T32((sph_u32)(W7 >> 32)); \
for (r = 0; r < 8; r ++) { \
SUB_CRUMB(V40, V41, V42, V43); \
SUB_CRUMB(V45, V46, V47, V44); \
MIX_WORD(V40, V44); \
MIX_WORD(V41, V45); \
MIX_WORD(V42, V46); \
MIX_WORD(V43, V47); \
V40 ^= RC40[r]; \
V44 ^= RC44[r]; \
} \
} while (0)
#else
#define LUFFA_P5 do { \
int r; \
TWEAK5; \
for (r = 0; r < 8; r ++) { \
SUB_CRUMB(V00, V01, V02, V03); \
SUB_CRUMB(V05, V06, V07, V04); \
MIX_WORD(V00, V04); \
MIX_WORD(V01, V05); \
MIX_WORD(V02, V06); \
MIX_WORD(V03, V07); \
V00 ^= RC00[r]; \
V04 ^= RC04[r]; \
} \
for (r = 0; r < 8; r ++) { \
SUB_CRUMB(V10, V11, V12, V13); \
SUB_CRUMB(V15, V16, V17, V14); \
MIX_WORD(V10, V14); \
MIX_WORD(V11, V15); \
MIX_WORD(V12, V16); \
MIX_WORD(V13, V17); \
V10 ^= RC10[r]; \
V14 ^= RC14[r]; \
} \
for (r = 0; r < 8; r ++) { \
SUB_CRUMB(V20, V21, V22, V23); \
SUB_CRUMB(V25, V26, V27, V24); \
MIX_WORD(V20, V24); \
MIX_WORD(V21, V25); \
MIX_WORD(V22, V26); \
MIX_WORD(V23, V27); \
V20 ^= RC20[r]; \
V24 ^= RC24[r]; \
} \
for (r = 0; r < 8; r ++) { \
SUB_CRUMB(V30, V31, V32, V33); \
SUB_CRUMB(V35, V36, V37, V34); \
MIX_WORD(V30, V34); \
MIX_WORD(V31, V35); \
MIX_WORD(V32, V36); \
MIX_WORD(V33, V37); \
V30 ^= RC30[r]; \
V34 ^= RC34[r]; \
} \
for (r = 0; r < 8; r ++) { \
SUB_CRUMB(V40, V41, V42, V43); \
SUB_CRUMB(V45, V46, V47, V44); \
MIX_WORD(V40, V44); \
MIX_WORD(V41, V45); \
MIX_WORD(V42, V46); \
MIX_WORD(V43, V47); \
V40 ^= RC40[r]; \
V44 ^= RC44[r]; \
} \
} while (0)
#endif

607
kernel/shavite.cl
Unescape View File

@ -0,0 +1,607 @@ @@ -0,0 +1,607 @@
/* $Id: shavite.c 227 2010-06-16 17:28:38Z tp $ */
/*
* SHAvite-3 implementation.
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2007-2010 Projet RNRT SAPHIR
*
* 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 Thomas Pornin <thomas.pornin@cryptolog.com>
*/
/*
* As of round 2 of the SHA-3 competition, the published reference
* implementation and test vectors are wrong, because they use
* big-endian AES tables while the internal decoding uses little-endian.
* The code below follows the specification. To turn it into a code
* which follows the reference implementation (the one called "BugFix"
* on the SHAvite-3 web site, published on Nov 23rd, 2009), comment out
* the code below (from the '#define AES_BIG_ENDIAN...' to the definition
* of the AES_ROUND_NOKEY macro) and replace it with the version which
* is commented out afterwards.
*/
#define AES_BIG_ENDIAN 0
#include "aes_helper.cl"
#define AES_ROUND_NOKEY(x0, x1, x2, x3) do { \
sph_u32 t0 = (x0); \
sph_u32 t1 = (x1); \
sph_u32 t2 = (x2); \
sph_u32 t3 = (x3); \
AES_ROUND_NOKEY_LE(t0, t1, t2, t3, x0, x1, x2, x3); \
} while (0)
#define KEY_EXPAND_ELT(k0, k1, k2, k3) do { \
sph_u32 kt; \
AES_ROUND_NOKEY(k1, k2, k3, k0); \
kt = (k0); \
(k0) = (k1); \
(k1) = (k2); \
(k2) = (k3); \
(k3) = kt; \
} while (0)
/*
* This function assumes that "msg" is aligned for 32-bit access.
*/
#define c512(msg) do { \
sph_u32 p0, p1, p2, p3, p4, p5, p6, p7; \
sph_u32 p8, p9, pA, pB, pC, pD, pE, pF; \
sph_u32 x0, x1, x2, x3; \
int r; \
\
p0 = h0; \
p1 = h1; \
p2 = h2; \
p3 = h3; \
p4 = h4; \
p5 = h5; \
p6 = h6; \
p7 = h7; \
p8 = h8; \
p9 = h9; \
pA = hA; \
pB = hB; \
pC = hC; \
pD = hD; \
pE = hE; \
pF = hF; \
/* round 0 */ \
x0 = p4 ^ rk00; \
x1 = p5 ^ rk01; \
x2 = p6 ^ rk02; \
x3 = p7 ^ rk03; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
x0 ^= rk04; \
x1 ^= rk05; \
x2 ^= rk06; \
x3 ^= rk07; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
x0 ^= rk08; \
x1 ^= rk09; \
x2 ^= rk0A; \
x3 ^= rk0B; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
x0 ^= rk0C; \
x1 ^= rk0D; \
x2 ^= rk0E; \
x3 ^= rk0F; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
p0 ^= x0; \
p1 ^= x1; \
p2 ^= x2; \
p3 ^= x3; \
x0 = pC ^ rk10; \
x1 = pD ^ rk11; \
x2 = pE ^ rk12; \
x3 = pF ^ rk13; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
x0 ^= rk14; \
x1 ^= rk15; \
x2 ^= rk16; \
x3 ^= rk17; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
x0 ^= rk18; \
x1 ^= rk19; \
x2 ^= rk1A; \
x3 ^= rk1B; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
x0 ^= rk1C; \
x1 ^= rk1D; \
x2 ^= rk1E; \
x3 ^= rk1F; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
p8 ^= x0; \
p9 ^= x1; \
pA ^= x2; \
pB ^= x3; \
\
for (r = 0; r < 3; r ++) { \
/* round 1, 5, 9 */ \
KEY_EXPAND_ELT(rk00, rk01, rk02, rk03); \
rk00 ^= rk1C; \
rk01 ^= rk1D; \
rk02 ^= rk1E; \
rk03 ^= rk1F; \
if (r == 0) { \
rk00 ^= sc_count0; \
rk01 ^= sc_count1; \
rk02 ^= sc_count2; \
rk03 ^= SPH_T32(~sc_count3); \
} \
x0 = p0 ^ rk00; \
x1 = p1 ^ rk01; \
x2 = p2 ^ rk02; \
x3 = p3 ^ rk03; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk04, rk05, rk06, rk07); \
rk04 ^= rk00; \
rk05 ^= rk01; \
rk06 ^= rk02; \
rk07 ^= rk03; \
if (r == 1) { \
rk04 ^= sc_count3; \
rk05 ^= sc_count2; \
rk06 ^= sc_count1; \
rk07 ^= SPH_T32(~sc_count0); \
} \
x0 ^= rk04; \
x1 ^= rk05; \
x2 ^= rk06; \
x3 ^= rk07; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk08, rk09, rk0A, rk0B); \
rk08 ^= rk04; \
rk09 ^= rk05; \
rk0A ^= rk06; \
rk0B ^= rk07; \
x0 ^= rk08; \
x1 ^= rk09; \
x2 ^= rk0A; \
x3 ^= rk0B; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk0C, rk0D, rk0E, rk0F); \
rk0C ^= rk08; \
rk0D ^= rk09; \
rk0E ^= rk0A; \
rk0F ^= rk0B; \
x0 ^= rk0C; \
x1 ^= rk0D; \
x2 ^= rk0E; \
x3 ^= rk0F; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
pC ^= x0; \
pD ^= x1; \
pE ^= x2; \
pF ^= x3; \
KEY_EXPAND_ELT(rk10, rk11, rk12, rk13); \
rk10 ^= rk0C; \
rk11 ^= rk0D; \
rk12 ^= rk0E; \
rk13 ^= rk0F; \
x0 = p8 ^ rk10; \
x1 = p9 ^ rk11; \
x2 = pA ^ rk12; \
x3 = pB ^ rk13; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk14, rk15, rk16, rk17); \
rk14 ^= rk10; \
rk15 ^= rk11; \
rk16 ^= rk12; \
rk17 ^= rk13; \
x0 ^= rk14; \
x1 ^= rk15; \
x2 ^= rk16; \
x3 ^= rk17; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk18, rk19, rk1A, rk1B); \
rk18 ^= rk14; \
rk19 ^= rk15; \
rk1A ^= rk16; \
rk1B ^= rk17; \
x0 ^= rk18; \
x1 ^= rk19; \
x2 ^= rk1A; \
x3 ^= rk1B; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk1C, rk1D, rk1E, rk1F); \
rk1C ^= rk18; \
rk1D ^= rk19; \
rk1E ^= rk1A; \
rk1F ^= rk1B; \
if (r == 2) { \
rk1C ^= sc_count2; \
rk1D ^= sc_count3; \
rk1E ^= sc_count0; \
rk1F ^= SPH_T32(~sc_count1); \
} \
x0 ^= rk1C; \
x1 ^= rk1D; \
x2 ^= rk1E; \
x3 ^= rk1F; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
p4 ^= x0; \
p5 ^= x1; \
p6 ^= x2; \
p7 ^= x3; \
/* round 2, 6, 10 */ \
rk00 ^= rk19; \
x0 = pC ^ rk00; \
rk01 ^= rk1A; \
x1 = pD ^ rk01; \
rk02 ^= rk1B; \
x2 = pE ^ rk02; \
rk03 ^= rk1C; \
x3 = pF ^ rk03; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
rk04 ^= rk1D; \
x0 ^= rk04; \
rk05 ^= rk1E; \
x1 ^= rk05; \
rk06 ^= rk1F; \
x2 ^= rk06; \
rk07 ^= rk00; \
x3 ^= rk07; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
rk08 ^= rk01; \
x0 ^= rk08; \
rk09 ^= rk02; \
x1 ^= rk09; \
rk0A ^= rk03; \
x2 ^= rk0A; \
rk0B ^= rk04; \
x3 ^= rk0B; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
rk0C ^= rk05; \
x0 ^= rk0C; \
rk0D ^= rk06; \
x1 ^= rk0D; \
rk0E ^= rk07; \
x2 ^= rk0E; \
rk0F ^= rk08; \
x3 ^= rk0F; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
p8 ^= x0; \
p9 ^= x1; \
pA ^= x2; \
pB ^= x3; \
rk10 ^= rk09; \
x0 = p4 ^ rk10; \
rk11 ^= rk0A; \
x1 = p5 ^ rk11; \
rk12 ^= rk0B; \
x2 = p6 ^ rk12; \
rk13 ^= rk0C; \
x3 = p7 ^ rk13; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
rk14 ^= rk0D; \
x0 ^= rk14; \
rk15 ^= rk0E; \
x1 ^= rk15; \
rk16 ^= rk0F; \
x2 ^= rk16; \
rk17 ^= rk10; \
x3 ^= rk17; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
rk18 ^= rk11; \
x0 ^= rk18; \
rk19 ^= rk12; \
x1 ^= rk19; \
rk1A ^= rk13; \
x2 ^= rk1A; \
rk1B ^= rk14; \
x3 ^= rk1B; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
rk1C ^= rk15; \
x0 ^= rk1C; \
rk1D ^= rk16; \
x1 ^= rk1D; \
rk1E ^= rk17; \
x2 ^= rk1E; \
rk1F ^= rk18; \
x3 ^= rk1F; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
p0 ^= x0; \
p1 ^= x1; \
p2 ^= x2; \
p3 ^= x3; \
/* round 3, 7, 11 */ \
KEY_EXPAND_ELT(rk00, rk01, rk02, rk03); \
rk00 ^= rk1C; \
rk01 ^= rk1D; \
rk02 ^= rk1E; \
rk03 ^= rk1F; \
x0 = p8 ^ rk00; \
x1 = p9 ^ rk01; \
x2 = pA ^ rk02; \
x3 = pB ^ rk03; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk04, rk05, rk06, rk07); \
rk04 ^= rk00; \
rk05 ^= rk01; \
rk06 ^= rk02; \
rk07 ^= rk03; \
x0 ^= rk04; \
x1 ^= rk05; \
x2 ^= rk06; \
x3 ^= rk07; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk08, rk09, rk0A, rk0B); \
rk08 ^= rk04; \
rk09 ^= rk05; \
rk0A ^= rk06; \
rk0B ^= rk07; \
x0 ^= rk08; \
x1 ^= rk09; \
x2 ^= rk0A; \
x3 ^= rk0B; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk0C, rk0D, rk0E, rk0F); \
rk0C ^= rk08; \
rk0D ^= rk09; \
rk0E ^= rk0A; \
rk0F ^= rk0B; \
x0 ^= rk0C; \
x1 ^= rk0D; \
x2 ^= rk0E; \
x3 ^= rk0F; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
p4 ^= x0; \
p5 ^= x1; \
p6 ^= x2; \
p7 ^= x3; \
KEY_EXPAND_ELT(rk10, rk11, rk12, rk13); \
rk10 ^= rk0C; \
rk11 ^= rk0D; \
rk12 ^= rk0E; \
rk13 ^= rk0F; \
x0 = p0 ^ rk10; \
x1 = p1 ^ rk11; \
x2 = p2 ^ rk12; \
x3 = p3 ^ rk13; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk14, rk15, rk16, rk17); \
rk14 ^= rk10; \
rk15 ^= rk11; \
rk16 ^= rk12; \
rk17 ^= rk13; \
x0 ^= rk14; \
x1 ^= rk15; \
x2 ^= rk16; \
x3 ^= rk17; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk18, rk19, rk1A, rk1B); \
rk18 ^= rk14; \
rk19 ^= rk15; \
rk1A ^= rk16; \
rk1B ^= rk17; \
x0 ^= rk18; \
x1 ^= rk19; \
x2 ^= rk1A; \
x3 ^= rk1B; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk1C, rk1D, rk1E, rk1F); \
rk1C ^= rk18; \
rk1D ^= rk19; \
rk1E ^= rk1A; \
rk1F ^= rk1B; \
x0 ^= rk1C; \
x1 ^= rk1D; \
x2 ^= rk1E; \
x3 ^= rk1F; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
pC ^= x0; \
pD ^= x1; \
pE ^= x2; \
pF ^= x3; \
/* round 4, 8, 12 */ \
rk00 ^= rk19; \
x0 = p4 ^ rk00; \
rk01 ^= rk1A; \
x1 = p5 ^ rk01; \
rk02 ^= rk1B; \
x2 = p6 ^ rk02; \
rk03 ^= rk1C; \
x3 = p7 ^ rk03; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
rk04 ^= rk1D; \
x0 ^= rk04; \
rk05 ^= rk1E; \
x1 ^= rk05; \
rk06 ^= rk1F; \
x2 ^= rk06; \
rk07 ^= rk00; \
x3 ^= rk07; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
rk08 ^= rk01; \
x0 ^= rk08; \
rk09 ^= rk02; \
x1 ^= rk09; \
rk0A ^= rk03; \
x2 ^= rk0A; \
rk0B ^= rk04; \
x3 ^= rk0B; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
rk0C ^= rk05; \
x0 ^= rk0C; \
rk0D ^= rk06; \
x1 ^= rk0D; \
rk0E ^= rk07; \
x2 ^= rk0E; \
rk0F ^= rk08; \
x3 ^= rk0F; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
p0 ^= x0; \
p1 ^= x1; \
p2 ^= x2; \
p3 ^= x3; \
rk10 ^= rk09; \
x0 = pC ^ rk10; \
rk11 ^= rk0A; \
x1 = pD ^ rk11; \
rk12 ^= rk0B; \
x2 = pE ^ rk12; \
rk13 ^= rk0C; \
x3 = pF ^ rk13; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
rk14 ^= rk0D; \
x0 ^= rk14; \
rk15 ^= rk0E; \
x1 ^= rk15; \
rk16 ^= rk0F; \
x2 ^= rk16; \
rk17 ^= rk10; \
x3 ^= rk17; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
rk18 ^= rk11; \
x0 ^= rk18; \
rk19 ^= rk12; \
x1 ^= rk19; \
rk1A ^= rk13; \
x2 ^= rk1A; \
rk1B ^= rk14; \
x3 ^= rk1B; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
rk1C ^= rk15; \
x0 ^= rk1C; \
rk1D ^= rk16; \
x1 ^= rk1D; \
rk1E ^= rk17; \
x2 ^= rk1E; \
rk1F ^= rk18; \
x3 ^= rk1F; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
p8 ^= x0; \
p9 ^= x1; \
pA ^= x2; \
pB ^= x3; \
} \
/* round 13 */ \
KEY_EXPAND_ELT(rk00, rk01, rk02, rk03); \
rk00 ^= rk1C; \
rk01 ^= rk1D; \
rk02 ^= rk1E; \
rk03 ^= rk1F; \
x0 = p0 ^ rk00; \
x1 = p1 ^ rk01; \
x2 = p2 ^ rk02; \
x3 = p3 ^ rk03; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk04, rk05, rk06, rk07); \
rk04 ^= rk00; \
rk05 ^= rk01; \
rk06 ^= rk02; \
rk07 ^= rk03; \
x0 ^= rk04; \
x1 ^= rk05; \
x2 ^= rk06; \
x3 ^= rk07; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk08, rk09, rk0A, rk0B); \
rk08 ^= rk04; \
rk09 ^= rk05; \
rk0A ^= rk06; \
rk0B ^= rk07; \
x0 ^= rk08; \
x1 ^= rk09; \
x2 ^= rk0A; \
x3 ^= rk0B; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk0C, rk0D, rk0E, rk0F); \
rk0C ^= rk08; \
rk0D ^= rk09; \
rk0E ^= rk0A; \
rk0F ^= rk0B; \
x0 ^= rk0C; \
x1 ^= rk0D; \
x2 ^= rk0E; \
x3 ^= rk0F; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
pC ^= x0; \
pD ^= x1; \
pE ^= x2; \
pF ^= x3; \
KEY_EXPAND_ELT(rk10, rk11, rk12, rk13); \
rk10 ^= rk0C; \
rk11 ^= rk0D; \
rk12 ^= rk0E; \
rk13 ^= rk0F; \
x0 = p8 ^ rk10; \
x1 = p9 ^ rk11; \
x2 = pA ^ rk12; \
x3 = pB ^ rk13; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk14, rk15, rk16, rk17); \
rk14 ^= rk10; \
rk15 ^= rk11; \
rk16 ^= rk12; \
rk17 ^= rk13; \
x0 ^= rk14; \
x1 ^= rk15; \
x2 ^= rk16; \
x3 ^= rk17; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk18, rk19, rk1A, rk1B); \
rk18 ^= rk14 ^ sc_count1; \
rk19 ^= rk15 ^ sc_count0; \
rk1A ^= rk16 ^ sc_count3; \
rk1B ^= rk17 ^ SPH_T32(~sc_count2); \
x0 ^= rk18; \
x1 ^= rk19; \
x2 ^= rk1A; \
x3 ^= rk1B; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
KEY_EXPAND_ELT(rk1C, rk1D, rk1E, rk1F); \
rk1C ^= rk18; \
rk1D ^= rk19; \
rk1E ^= rk1A; \
rk1F ^= rk1B; \
x0 ^= rk1C; \
x1 ^= rk1D; \
x2 ^= rk1E; \
x3 ^= rk1F; \
AES_ROUND_NOKEY(x0, x1, x2, x3); \
p4 ^= x0; \
p5 ^= x1; \
p6 ^= x2; \
p7 ^= x3; \
h0 ^= p8; \
h1 ^= p9; \
h2 ^= pA; \
h3 ^= pB; \
h4 ^= pC; \
h5 ^= pD; \
h6 ^= pE; \
h7 ^= pF; \
h8 ^= p0; \
h9 ^= p1; \
hA ^= p2; \
hB ^= p3; \
hC ^= p4; \
hD ^= p5; \
hE ^= p6; \
hF ^= p7; \
} while (0)

1694
kernel/simd.cl
View File

File diff suppressed because it is too large Load Diff

350
kernel/skein.cl
Unescape View File

@ -0,0 +1,350 @@ @@ -0,0 +1,350 @@
/* $Id: skein.c 254 2011-06-07 19:38:58Z tp $ */
/*
* Skein implementation.
*
* ==========================(LICENSE BEGIN)============================
*
* Copyright (c) 2007-2010 Projet RNRT SAPHIR
*
* 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 Thomas Pornin <thomas.pornin@cryptolog.com>
*/
/*
* M9_ ## s ## _ ## i evaluates to s+i mod 9 (0 <= s <= 18, 0 <= i <= 7).
*/
#define M9_0_0 0
#define M9_0_1 1
#define M9_0_2 2
#define M9_0_3 3
#define M9_0_4 4
#define M9_0_5 5
#define M9_0_6 6
#define M9_0_7 7
#define M9_1_0 1
#define M9_1_1 2
#define M9_1_2 3
#define M9_1_3 4
#define M9_1_4 5
#define M9_1_5 6
#define M9_1_6 7
#define M9_1_7 8
#define M9_2_0 2
#define M9_2_1 3
#define M9_2_2 4
#define M9_2_3 5
#define M9_2_4 6
#define M9_2_5 7
#define M9_2_6 8
#define M9_2_7 0
#define M9_3_0 3
#define M9_3_1 4
#define M9_3_2 5
#define M9_3_3 6
#define M9_3_4 7
#define M9_3_5 8
#define M9_3_6 0
#define M9_3_7 1
#define M9_4_0 4
#define M9_4_1 5
#define M9_4_2 6
#define M9_4_3 7
#define M9_4_4 8
#define M9_4_5 0
#define M9_4_6 1
#define M9_4_7 2
#define M9_5_0 5
#define M9_5_1 6
#define M9_5_2 7
#define M9_5_3 8
#define M9_5_4 0
#define M9_5_5 1
#define M9_5_6 2
#define M9_5_7 3
#define M9_6_0 6
#define M9_6_1 7
#define M9_6_2 8
#define M9_6_3 0
#define M9_6_4 1
#define M9_6_5 2
#define M9_6_6 3
#define M9_6_7 4
#define M9_7_0 7
#define M9_7_1 8
#define M9_7_2 0
#define M9_7_3 1
#define M9_7_4 2
#define M9_7_5 3
#define M9_7_6 4
#define M9_7_7 5
#define M9_8_0 8
#define M9_8_1 0
#define M9_8_2 1
#define M9_8_3 2
#define M9_8_4 3
#define M9_8_5 4
#define M9_8_6 5
#define M9_8_7 6
#define M9_9_0 0
#define M9_9_1 1
#define M9_9_2 2
#define M9_9_3 3
#define M9_9_4 4
#define M9_9_5 5
#define M9_9_6 6
#define M9_9_7 7
#define M9_10_0 1
#define M9_10_1 2
#define M9_10_2 3
#define M9_10_3 4
#define M9_10_4 5
#define M9_10_5 6
#define M9_10_6 7
#define M9_10_7 8
#define M9_11_0 2
#define M9_11_1 3
#define M9_11_2 4
#define M9_11_3 5
#define M9_11_4 6
#define M9_11_5 7
#define M9_11_6 8
#define M9_11_7 0
#define M9_12_0 3
#define M9_12_1 4
#define M9_12_2 5
#define M9_12_3 6
#define M9_12_4 7
#define M9_12_5 8
#define M9_12_6 0
#define M9_12_7 1
#define M9_13_0 4
#define M9_13_1 5
#define M9_13_2 6
#define M9_13_3 7
#define M9_13_4 8
#define M9_13_5 0
#define M9_13_6 1
#define M9_13_7 2
#define M9_14_0 5
#define M9_14_1 6
#define M9_14_2 7
#define M9_14_3 8
#define M9_14_4 0
#define M9_14_5 1
#define M9_14_6 2
#define M9_14_7 3
#define M9_15_0 6
#define M9_15_1 7
#define M9_15_2 8
#define M9_15_3 0
#define M9_15_4 1
#define M9_15_5 2
#define M9_15_6 3
#define M9_15_7 4
#define M9_16_0 7
#define M9_16_1 8
#define M9_16_2 0
#define M9_16_3 1
#define M9_16_4 2
#define M9_16_5 3
#define M9_16_6 4
#define M9_16_7 5
#define M9_17_0 8
#define M9_17_1 0
#define M9_17_2 1
#define M9_17_3 2
#define M9_17_4 3
#define M9_17_5 4
#define M9_17_6 5
#define M9_17_7 6
#define M9_18_0 0
#define M9_18_1 1
#define M9_18_2 2
#define M9_18_3 3
#define M9_18_4 4
#define M9_18_5 5
#define M9_18_6 6
#define M9_18_7 7
/*
* M3_ ## s ## _ ## i evaluates to s+i mod 3 (0 <= s <= 18, 0 <= i <= 1).
*/
#define M3_0_0 0
#define M3_0_1 1
#define M3_1_0 1
#define M3_1_1 2
#define M3_2_0 2
#define M3_2_1 0
#define M3_3_0 0
#define M3_3_1 1
#define M3_4_0 1
#define M3_4_1 2
#define M3_5_0 2
#define M3_5_1 0
#define M3_6_0 0
#define M3_6_1 1
#define M3_7_0 1
#define M3_7_1 2
#define M3_8_0 2
#define M3_8_1 0
#define M3_9_0 0
#define M3_9_1 1
#define M3_10_0 1
#define M3_10_1 2
#define M3_11_0 2
#define M3_11_1 0
#define M3_12_0 0
#define M3_12_1 1
#define M3_13_0 1
#define M3_13_1 2
#define M3_14_0 2
#define M3_14_1 0
#define M3_15_0 0
#define M3_15_1 1
#define M3_16_0 1
#define M3_16_1 2
#define M3_17_0 2
#define M3_17_1 0
#define M3_18_0 0
#define M3_18_1 1
#define XCAT(x, y) XCAT_(x, y)
#define XCAT_(x, y) x ## y
#define SKBI(k, s, i) XCAT(k, XCAT(XCAT(XCAT(M9_, s), _), i))
#define SKBT(t, s, v) XCAT(t, XCAT(XCAT(XCAT(M3_, s), _), v))
#define TFBIG_KINIT(k0, k1, k2, k3, k4, k5, k6, k7, k8, t0, t1, t2) do { \
k8 = ((k0 ^ k1) ^ (k2 ^ k3)) ^ ((k4 ^ k5) ^ (k6 ^ k7)) \
^ SPH_C64(0x1BD11BDAA9FC1A22); \
t2 = t0 ^ t1; \
} while (0)
#define TFBIG_ADDKEY(w0, w1, w2, w3, w4, w5, w6, w7, k, t, s) do { \
w0 = SPH_T64(w0 + SKBI(k, s, 0)); \
w1 = SPH_T64(w1 + SKBI(k, s, 1)); \
w2 = SPH_T64(w2 + SKBI(k, s, 2)); \
w3 = SPH_T64(w3 + SKBI(k, s, 3)); \
w4 = SPH_T64(w4 + SKBI(k, s, 4)); \
w5 = SPH_T64(w5 + SKBI(k, s, 5) + SKBT(t, s, 0)); \
w6 = SPH_T64(w6 + SKBI(k, s, 6) + SKBT(t, s, 1)); \
w7 = SPH_T64(w7 + SKBI(k, s, 7) + (sph_u64)s); \
} while (0)
#define TFBIG_MIX(x0, x1, rc) do { \
x0 = SPH_T64(x0 + x1); \
x1 = SPH_ROTL64(x1, rc) ^ x0; \
} while (0)
#define TFBIG_MIX8(w0, w1, w2, w3, w4, w5, w6, w7, rc0, rc1, rc2, rc3) do { \
TFBIG_MIX(w0, w1, rc0); \
TFBIG_MIX(w2, w3, rc1); \
TFBIG_MIX(w4, w5, rc2); \
TFBIG_MIX(w6, w7, rc3); \
} while (0)
#define TFBIG_4e(s) do { \
TFBIG_ADDKEY(p0, p1, p2, p3, p4, p5, p6, p7, h, t, s); \
TFBIG_MIX8(p0, p1, p2, p3, p4, p5, p6, p7, 46, 36, 19, 37); \
TFBIG_MIX8(p2, p1, p4, p7, p6, p5, p0, p3, 33, 27, 14, 42); \
TFBIG_MIX8(p4, p1, p6, p3, p0, p5, p2, p7, 17, 49, 36, 39); \
TFBIG_MIX8(p6, p1, p0, p7, p2, p5, p4, p3, 44, 9, 54, 56); \
} while (0)
#define TFBIG_4o(s) do { \
TFBIG_ADDKEY(p0, p1, p2, p3, p4, p5, p6, p7, h, t, s); \
TFBIG_MIX8(p0, p1, p2, p3, p4, p5, p6, p7, 39, 30, 34, 24); \
TFBIG_MIX8(p2, p1, p4, p7, p6, p5, p0, p3, 13, 50, 10, 17); \
TFBIG_MIX8(p4, p1, p6, p3, p0, p5, p2, p7, 25, 29, 39, 43); \
TFBIG_MIX8(p6, p1, p0, p7, p2, p5, p4, p3, 8, 35, 56, 22); \
} while (0)
#define UBI_BIG(etype, extra) do { \
sph_u64 h8, t0, t1, t2; \
sph_u64 p0 = m0; \
sph_u64 p1 = m1; \
sph_u64 p2 = m2; \
sph_u64 p3 = m3; \
sph_u64 p4 = m4; \
sph_u64 p5 = m5; \
sph_u64 p6 = m6; \
sph_u64 p7 = m7; \
t0 = SPH_T64(bcount << 6) + (sph_u64)(extra); \
t1 = (bcount >> 58) + ((sph_u64)(etype) << 55); \
TFBIG_KINIT(h0, h1, h2, h3, h4, h5, h6, h7, h8, t0, t1, t2); \
TFBIG_4e(0); \
TFBIG_4o(1); \
TFBIG_4e(2); \
TFBIG_4o(3); \
TFBIG_4e(4); \
TFBIG_4o(5); \
TFBIG_4e(6); \
TFBIG_4o(7); \
TFBIG_4e(8); \
TFBIG_4o(9); \
TFBIG_4e(10); \
TFBIG_4o(11); \
TFBIG_4e(12); \
TFBIG_4o(13); \
TFBIG_4e(14); \
TFBIG_4o(15); \
TFBIG_4e(16); \
TFBIG_4o(17); \
TFBIG_ADDKEY(p0, p1, p2, p3, p4, p5, p6, p7, h, t, 18); \
h0 = m0 ^ p0; \
h1 = m1 ^ p1; \
h2 = m2 ^ p2; \
h3 = m3 ^ p3; \
h4 = m4 ^ p4; \
h5 = m5 ^ p5; \
h6 = m6 ^ p6; \
h7 = m7 ^ p7; \
} while (0)
__constant static const sph_u64 SKEIN_IV512[] = {
SPH_C64(0x4903ADFF749C51CE), SPH_C64(0x0D95DE399746DF03),
SPH_C64(0x8FD1934127C79BCE), SPH_C64(0x9A255629FF352CB1),
SPH_C64(0x5DB62599DF6CA7B0), SPH_C64(0xEABE394CA9D5C3F4),
SPH_C64(0x991112C71A75B523), SPH_C64(0xAE18A40B660FCC33)
};

14
miner.h
Unescape View File

@ -381,6 +381,7 @@ enum cl_kernels { @@ -381,6 +381,7 @@ enum cl_kernels {
KL_CKOLIVAS,
KL_PSW,
KL_ZUIKKIS,
KL_DARKCOIN,
};
enum dev_reason {
@ -505,7 +506,7 @@ struct cgpu_info { @@ -505,7 +506,7 @@ struct cgpu_info {
int gpu_powertune;
float gpu_vddc;
#endif
int diff1;
double diff1;
double diff_accepted;
double diff_rejected;
int last_share_pool;
@ -1096,7 +1097,8 @@ extern double total_rolling; @@ -1096,7 +1097,8 @@ extern double total_rolling;
extern double total_mhashes_done;
extern unsigned int new_blocks;
extern unsigned int found_blocks;
extern int total_accepted, total_rejected, total_diff1;;
extern int total_accepted, total_rejected;
extern double total_diff1;
extern int total_getworks, total_stale, total_discarded;
extern double total_diff_accepted, total_diff_rejected, total_diff_stale;
extern unsigned int local_work;
@ -1106,7 +1108,7 @@ extern int opt_log_interval; @@ -1106,7 +1108,7 @@ extern int opt_log_interval;
extern unsigned long long global_hashrate;
extern char current_hash[68];
extern double current_diff;
extern uint64_t best_diff;
extern double best_diff;
extern struct timeval block_timeval;
extern char *workpadding;
@ -1177,7 +1179,7 @@ struct pool { @@ -1177,7 +1179,7 @@ struct pool {
int seq_rejects;
int seq_getfails;
int solved;
int diff1;
double diff1;
char diff[8];
int quota;
int quota_gcd;
@ -1234,7 +1236,7 @@ struct pool { @@ -1234,7 +1236,7 @@ struct pool {
time_t last_share_time;
double last_share_diff;
uint64_t best_diff;
double best_diff;
struct sgminer_stats sgminer_stats;
struct sgminer_pool_stats sgminer_pool_stats;
@ -1310,7 +1312,7 @@ struct work { @@ -1310,7 +1312,7 @@ struct work {
unsigned char device_target[32];
double device_diff;
uint64_t share_diff;
double share_diff;
int rolls;
int drv_rolllimit; /* How much the driver can roll ntime */

8
ocl.c
Unescape View File

@ -454,6 +454,11 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize) @@ -454,6 +454,11 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
/* Kernel only supports worksize 256 */
cgpu->work_size = 256;
break;
case KL_DARKCOIN:
applog(LOG_WARNING, "Kernel darkcoin is experimental.");
strcpy(filename, DARKCOIN_KERNNAME".cl");
strcpy(binaryfilename, DARKCOIN_KERNNAME);
break;
case KL_NONE: /* Shouldn't happen */
break;
}
@ -587,7 +592,8 @@ build: @@ -587,7 +592,8 @@ build:
/* create a cl program executable for all the devices specified */
char *CompilerOptions = calloc(1, 256);
sprintf(CompilerOptions, "-D LOOKUP_GAP=%d -D CONCURRENT_THREADS=%d -D WORKSIZE=%d",
sprintf(CompilerOptions, "-I \"%s\" -I \"%s\" -I \"%skernel\" -I \".\" -D LOOKUP_GAP=%d -D CONCURRENT_THREADS=%d -D WORKSIZE=%d",
opt_kernel_path, sgminer_path, sgminer_path,
cgpu->lookup_gap, (unsigned int)cgpu->thread_concurrency, (int)clState->wsize);
applog(LOG_DEBUG, "Setting worksize to %d", (int)(clState->wsize));

2
ocl.h
Unescape View File

@ -21,7 +21,7 @@ typedef struct { @@ -21,7 +21,7 @@ typedef struct {
cl_mem CLbuffer0;
cl_mem padbuffer8;
size_t padbufsize;
void * cldata;
unsigned char cldata[80];
bool hasBitAlign;
bool hasOpenCL11plus;
bool hasOpenCL12plus;

74
sgminer.c
Unescape View File

@ -55,6 +55,7 @@ char *curly = ":D"; @@ -55,6 +55,7 @@ char *curly = ":D";
#include "driver-opencl.h"
#include "bench_block.h"
#include "scrypt.h"
#include "darkcoin.h"
#if defined(unix) || defined(__APPLE__)
#include <errno.h>
@ -202,7 +203,8 @@ cglock_t control_lock; @@ -202,7 +203,8 @@ cglock_t control_lock;
pthread_mutex_t stats_lock;
int hw_errors;
int total_accepted, total_rejected, total_diff1;
int total_accepted, total_rejected;
double total_diff1;
int total_getworks, total_stale, total_discarded;
double total_diff_accepted, total_diff_rejected, total_diff_stale;
static int staged_rollable;
@ -239,7 +241,7 @@ struct timeval block_timeval; @@ -239,7 +241,7 @@ struct timeval block_timeval;
static char best_share[8] = "0";
double current_diff = 0xFFFFFFFFFFFFFFFFULL;
static char block_diff[8];
uint64_t best_diff = 0;
double best_diff = 0;
struct block {
char hash[68];
@ -297,6 +299,7 @@ struct schedtime { @@ -297,6 +299,7 @@ struct schedtime {
struct schedtime schedstart;
struct schedtime schedstop;
bool sched_paused;
bool is_scrypt = true;
static bool time_before(struct tm *tm1, struct tm *tm2)
{
@ -2073,6 +2076,18 @@ static void suffix_string(uint64_t val, char *buf, size_t bufsiz, int sigdigits) @@ -2073,6 +2076,18 @@ static void suffix_string(uint64_t val, char *buf, size_t bufsiz, int sigdigits)
}
}
/* Convert a double value into a truncated string for displaying with its
* associated suitable for Mega, Giga etc. Buf array needs to be long enough */
static void suffix_string_double(double val, char *buf, size_t bufsiz, int sigdigits)
{
if (val < 10) {
snprintf(buf, bufsiz, "%.3lf", val);
} else {
return suffix_string(val, buf, bufsiz, sigdigits);
}
}
double cgpu_runtime(struct cgpu_info *cgpu)
{
struct timeval now;
@ -2107,7 +2122,7 @@ static void get_statline(char *buf, size_t bufsiz, struct cgpu_info *cgpu) @@ -2107,7 +2122,7 @@ static void get_statline(char *buf, size_t bufsiz, struct cgpu_info *cgpu)
snprintf(buf, bufsiz, "%s%d ", cgpu->drv->name, cgpu->device_id);
cgpu->drv->get_statline_before(buf, bufsiz, cgpu);
tailsprintf(buf, bufsiz, "(%ds):%s (avg):%sh/s | A:%.0f R:%.0f HW:%d WU:%.1f/m",
tailsprintf(buf, bufsiz, "(%ds):%s (avg):%sh/s | A:%.0f R:%.0f HW:%d WU:%.3f/m",
opt_log_interval,
displayed_rolling,
displayed_hashes,
@ -2236,7 +2251,7 @@ static void curses_print_devstatus(struct cgpu_info *cgpu, int count) @@ -2236,7 +2251,7 @@ static void curses_print_devstatus(struct cgpu_info *cgpu, int count)
adj_width(cgpu->hw_errors, &hwwidth);
adj_width(wu, &wuwidth);
cg_wprintw(statuswin, "/%6sh/s | R:%*.1f%% HW:%*d WU:%*.1f/m",
cg_wprintw(statuswin, "/%6sh/s | R:%*.1f%% HW:%*d WU:%*.3f/m",
displayed_hashes,
drwidth, reject_pct,
hwwidth, cgpu->hw_errors,
@ -2519,7 +2534,7 @@ share_result(json_t *val, json_t *res, json_t *err, const struct work *work, @@ -2519,7 +2534,7 @@ share_result(json_t *val, json_t *res, json_t *err, const struct work *work,
static void show_hash(struct work *work, char *hashshow)
{
unsigned char rhash[32];
char diffdisp[16];
char diffdisp[16], wdiffdisp[16];
unsigned long h32;
uint32_t *hash32;
int intdiff, ofs;
@ -2532,8 +2547,9 @@ static void show_hash(struct work *work, char *hashshow) @@ -2532,8 +2547,9 @@ static void show_hash(struct work *work, char *hashshow)
hash32 = (uint32_t *)(rhash + ofs);
h32 = be32toh(*hash32);
intdiff = round(work->work_difficulty);
suffix_string(work->share_diff, diffdisp, sizeof (diffdisp), 0);
snprintf(hashshow, 64, "%08lx Diff %s/%d%s", h32, diffdisp, intdiff,
suffix_string_double(work->share_diff, diffdisp, sizeof (diffdisp), 0);
suffix_string_double(work->work_difficulty, wdiffdisp, sizeof (wdiffdisp), 0);
snprintf(hashshow, 64, "%08lx Diff %s/%s%s", h32, diffdisp, wdiffdisp,
work->block? " BLOCK!" : "");
}
@ -2943,7 +2959,7 @@ static void calc_diff(struct work *work, double known) @@ -2943,7 +2959,7 @@ static void calc_diff(struct work *work, double known)
else {
double d64, dcut64;
d64 = (double)65536 * truediffone;
d64 = (is_scrypt ? (double)65536 * truediffone : truediffone);
dcut64 = le256todouble(work->target);
if (unlikely(!dcut64))
@ -2953,8 +2969,7 @@ static void calc_diff(struct work *work, double known) @@ -2953,8 +2969,7 @@ static void calc_diff(struct work *work, double known)
difficulty = work->work_difficulty;
pool_stats->last_diff = difficulty;
uintdiff = round(difficulty);
suffix_string(uintdiff, work->pool->diff, sizeof(work->pool->diff), 0);
suffix_string_double(difficulty, work->pool->diff, sizeof(work->pool->diff), 0);
if (difficulty == pool_stats->min_diff)
pool_stats->min_diff_count++;
@ -3555,24 +3570,24 @@ static bool stale_work(struct work *work, bool share) @@ -3555,24 +3570,24 @@ static bool stale_work(struct work *work, bool share)
return false;
}
static uint64_t share_diff(const struct work *work)
static double share_diff(const struct work *work)
{
bool new_best = false;
double d64, s64;
uint64_t ret;
double ret;
d64 = (double)65536 * truediffone;
d64 = (is_scrypt ? (double)65536 * truediffone : truediffone);
s64 = le256todouble(work->hash);
if (unlikely(!s64))
s64 = 0;
ret = round(d64 / s64);
ret = d64 / s64;
cg_wlock(&control_lock);
if (unlikely(ret > best_diff)) {
new_best = true;
best_diff = ret;
suffix_string(best_diff, best_share, sizeof(best_share), 0);
suffix_string_double(best_diff, best_share, sizeof(best_share), 0);
}
if (unlikely(ret > work->pool->best_diff))
work->pool->best_diff = ret;
@ -3884,7 +3899,7 @@ static void set_blockdiff(const struct work *work) @@ -3884,7 +3899,7 @@ static void set_blockdiff(const struct work *work)
uint8_t pow = work->data[72];
int powdiff = (8 * (0x1d - 3)) - (8 * (pow - 3));
uint32_t diff32 = be32toh(*((uint32_t *)(work->data + 72))) & 0x00FFFFFF;
double numerator = 0xFFFFFFFFULL << powdiff;
double numerator = (is_scrypt ? 0xFFFFFFFFULL : 0xFFFFULL) << powdiff;
double ddiff = numerator / (double)diff32;
if (unlikely(current_diff != ddiff)) {
@ -4209,6 +4224,9 @@ void write_config(FILE *fcfg) @@ -4209,6 +4224,9 @@ void write_config(FILE *fcfg)
case KL_ZUIKKIS:
fprintf(fcfg, ZUIKKIS_KERNNAME);
break;
case KL_DARKCOIN:
fprintf(fcfg, DARKCOIN_KERNNAME);
break;
}
}
@ -4369,7 +4387,7 @@ void zero_bestshare(void) @@ -4369,7 +4387,7 @@ void zero_bestshare(void)
best_diff = 0;
memset(best_share, 0, 8);
suffix_string(best_diff, best_share, sizeof(best_share), 0);
suffix_string_double(best_diff, best_share, sizeof(best_share), 0);
for (i = 0; i < total_pools; i++) {
struct pool *pool = pools[i];
@ -5003,7 +5021,7 @@ static void hashmeter(int thr_id, struct timeval *diff, @@ -5003,7 +5021,7 @@ static void hashmeter(int thr_id, struct timeval *diff,
suffix_string(dr64, displayed_rolling, sizeof(displayed_rolling), 4);
snprintf(statusline, sizeof(statusline),
"%s(%ds):%s (avg):%sh/s | A:%.0f R:%.0f HW:%d WU:%.1f/m",
"%s(%ds):%s (avg):%sh/s | A:%.0f R:%.0f HW:%d WU:%.3f/m",
want_per_device_stats ? "ALL " : "",
opt_log_interval, displayed_rolling, displayed_hashes,
total_diff_accepted, total_diff_rejected, hw_errors,
@ -5797,7 +5815,7 @@ void set_target(unsigned char *dest_target, double diff) @@ -5797,7 +5815,7 @@ void set_target(unsigned char *dest_target, double diff)
}
// FIXME: is target set right?
d64 = (double)65536 * truediffone;
d64 = (is_scrypt ? (double)65536 * truediffone : truediffone);
d64 /= diff;
dcut64 = d64 / bits192;
@ -6012,7 +6030,14 @@ static void rebuild_nonce(struct work *work, uint32_t nonce) @@ -6012,7 +6030,14 @@ static void rebuild_nonce(struct work *work, uint32_t nonce)
*work_nonce = htole32(nonce);
scrypt_regenhash(work);
switch (gpus[0].kernel) {
case KL_DARKCOIN:
darkcoin_regenhash(work);
break;
default:
scrypt_regenhash(work);
break;
}
}
/* For testing a nonce against diff 1 */
@ -6033,7 +6058,7 @@ bool test_nonce_diff(struct work *work, uint32_t nonce, double diff) @@ -6033,7 +6058,7 @@ bool test_nonce_diff(struct work *work, uint32_t nonce, double diff)
uint64_t *hash64 = (uint64_t *)(work->hash + 24), diff64;
rebuild_nonce(work, nonce);
diff64 = 0x0000ffff00000000ULL;
diff64 = (is_scrypt ? 0x0000ffff00000000ULL : 0x00000000ffff0000ULL);
diff64 /= diff;
return (le64toh(*hash64) <= diff64);
@ -6042,11 +6067,13 @@ bool test_nonce_diff(struct work *work, uint32_t nonce, double diff) @@ -6042,11 +6067,13 @@ bool test_nonce_diff(struct work *work, uint32_t nonce, double diff)
static void update_work_stats(struct thr_info *thr, struct work *work)
{
double test_diff = current_diff;
test_diff *= 65536;
if (is_scrypt)
test_diff *= 65536;
work->share_diff = share_diff(work);
test_diff *= 65536;
if (is_scrypt)
test_diff *= 65536;
if (unlikely(work->share_diff >= test_diff)) {
work->block = true;
@ -6058,6 +6085,7 @@ static void update_work_stats(struct thr_info *thr, struct work *work) @@ -6058,6 +6085,7 @@ static void update_work_stats(struct thr_info *thr, struct work *work)
mutex_lock(&stats_lock);
total_diff1 += work->device_diff;
applog(LOG_DEBUG, "total_diff1: %lf\n", total_diff1);
thr->cgpu->diff1 += work->device_diff;
work->pool->diff1 += work->device_diff;
thr->cgpu->last_device_valid_work = time(NULL);

2
util.c
Unescape View File

@ -1657,7 +1657,7 @@ static bool parse_diff(struct pool *pool, json_t *val) @@ -1657,7 +1657,7 @@ static bool parse_diff(struct pool *pool, json_t *val)
if ((double)idiff == diff)
applog(LOG_NOTICE, "%s difficulty changed to %d", pool->poolname ,idiff);
else
applog(LOG_NOTICE, "%s difficulty changed to %.1f", pool->poolname, diff);
applog(LOG_NOTICE, "%s difficulty changed to %.3f", pool->poolname, diff);
} else
applog(LOG_DEBUG, "%s difficulty set to %f", pool->poolname, diff);

Write Preview
Loading…
Cancel
Save