diff --git a/Makefile.am b/Makefile.am
index 4423354..6277e9f 100644
--- a/Makefile.am
+++ b/Makefile.am
@@ -71,6 +71,7 @@ ccminer_SOURCES = elist.h miner.h compat.h \
x11/cuda_x11_shavite512.cu x11/cuda_x11_simd512.cu x11/cuda_x11_echo.cu \
x11/cuda_x11_luffa512_Cubehash.cu x11/x11evo.cu x11/timetravel.cu x11/bitcore.cu \
x13/x13.cu x13/cuda_x13_hamsi512.cu x13/cuda_x13_fugue512.cu \
+ x13/hsr.cu x13/cuda_hsr_sm3.cu x13/sm3.c \
x15/x14.cu x15/x15.cu x15/cuda_x14_shabal512.cu x15/cuda_x15_whirlpool.cu \
x15/whirlpool.cu x15/cuda_x15_whirlpool_sm3.cu \
x17/x17.cu x17/hmq17.cu x17/cuda_x17_haval256.cu x17/cuda_x17_sha512.cu \
diff --git a/algos.h b/algos.h
index 7febd71..3c1528b 100644
--- a/algos.h
+++ b/algos.h
@@ -22,6 +22,7 @@ enum sha_algos {
ALGO_GROESTL,
ALGO_HEAVY, /* Heavycoin hash */
ALGO_HMQ1725,
+ ALGO_HSR,
ALGO_KECCAK,
ALGO_JACKPOT,
ALGO_JHA,
@@ -88,6 +89,7 @@ static const char *algo_names[] = {
"groestl",
"heavy",
"hmq1725",
+ "hsr",
"keccak",
"jackpot",
"jha",
@@ -163,6 +165,8 @@ static inline int algo_to_int(char* arg)
i = ALGO_LUFFA;
else if (!strcasecmp("hmq17", arg))
i = ALGO_HMQ1725;
+ else if (!strcasecmp("hshare", arg))
+ i = ALGO_HSR;
//else if (!strcasecmp("jackpot", arg))
// i = ALGO_JHA;
else if (!strcasecmp("lyra2re", arg))
diff --git a/bench.cpp b/bench.cpp
index a91b4c3..271345f 100644
--- a/bench.cpp
+++ b/bench.cpp
@@ -66,6 +66,7 @@ void algo_free_all(int thr_id)
free_groestlcoin(thr_id);
free_heavy(thr_id);
free_hmq17(thr_id);
+ free_hsr(thr_id);
free_jackpot(thr_id);
free_jha(thr_id);
free_lbry(thr_id);
diff --git a/ccminer.cpp b/ccminer.cpp
index d41a1dc..ae7a073 100644
--- a/ccminer.cpp
+++ b/ccminer.cpp
@@ -2227,6 +2227,7 @@ static void *miner_thread(void *userdata)
case ALGO_HEAVY:
case ALGO_JACKPOT:
case ALGO_JHA:
+ case ALGO_HSR:
case ALGO_LYRA2v2:
case ALGO_PHI:
case ALGO_S3:
@@ -2361,6 +2362,9 @@ static void *miner_thread(void *userdata)
case ALGO_HMQ1725:
rc = scanhash_hmq17(thr_id, &work, max_nonce, &hashes_done);
break;
+ case ALGO_HSR:
+ rc = scanhash_hsr(thr_id, &work, max_nonce, &hashes_done);
+ break;
case ALGO_HEAVY:
rc = scanhash_heavy(thr_id, &work, max_nonce, &hashes_done, work.maxvote, HEAVYCOIN_BLKHDR_SZ);
diff --git a/ccminer.vcxproj b/ccminer.vcxproj
index ea0a22d..5554061 100644
--- a/ccminer.vcxproj
+++ b/ccminer.vcxproj
@@ -331,6 +331,7 @@
+
@@ -576,6 +577,8 @@
+
+
diff --git a/ccminer.vcxproj.filters b/ccminer.vcxproj.filters
index c1579f9..11fb230 100644
--- a/ccminer.vcxproj.filters
+++ b/ccminer.vcxproj.filters
@@ -225,6 +225,9 @@
Source Files\sph
+
+ Source Files\sph
+
Source Files
@@ -727,6 +730,12 @@
Source Files\CUDA\x13
+
+ Source Files\CUDA\x13
+
+
+ Source Files\CUDA\x13
+
Source Files\CUDA
diff --git a/miner.h b/miner.h
index 0d58eec..02a8e7c 100644
--- a/miner.h
+++ b/miner.h
@@ -289,6 +289,7 @@ extern int scanhash_fugue256(int thr_id, struct work* work, uint32_t max_nonce,
extern int scanhash_groestlcoin(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_hmq17(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_heavy(int thr_id,struct work *work, uint32_t max_nonce, unsigned long *hashes_done, uint32_t maxvote, int blocklen);
+extern int scanhash_hsr(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_jha(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done);
extern int scanhash_jackpot(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done); // quark method
extern int scanhash_lbry(int thr_id, struct work *work, uint32_t max_nonce, unsigned long *hashes_done);
@@ -351,6 +352,7 @@ extern void free_fugue256(int thr_id);
extern void free_groestlcoin(int thr_id);
extern void free_heavy(int thr_id);
extern void free_hmq17(int thr_id);
+extern void free_hsr(int thr_id);
extern void free_jackpot(int thr_id);
extern void free_jha(int thr_id);
extern void free_lbry(int thr_id);
@@ -892,6 +894,7 @@ void fresh_hash(void *state, const void *input);
void fugue256_hash(unsigned char* output, const unsigned char* input, int len);
void heavycoin_hash(unsigned char* output, const unsigned char* input, int len);
void hmq17hash(void *output, const void *input);
+void hsr_hash(void *output, const void *input);
void keccak256_hash(void *state, const void *input);
void jackpothash(void *state, const void *input);
void groestlhash(void *state, const void *input);
diff --git a/util.cpp b/util.cpp
index e5fe5b6..e8f3833 100644
--- a/util.cpp
+++ b/util.cpp
@@ -2206,6 +2206,9 @@ void print_hash_tests(void)
hmq17hash(&hash[0], &buf[0]);
printpfx("hmq1725", hash);
+ hsr_hash(&hash[0], &buf[0]);
+ printpfx("hsr", hash);
+
jha_hash(&hash[0], &buf[0]);
printpfx("jha", hash);
diff --git a/x13/cuda_hsr_sm3.cu b/x13/cuda_hsr_sm3.cu
new file mode 100644
index 0000000..5ce0186
--- /dev/null
+++ b/x13/cuda_hsr_sm3.cu
@@ -0,0 +1,139 @@
+#include
+#include
+#include
+
+#include
+#include
+
+#define F(x, y, z) (((x) ^ (y) ^ (z)))
+#define FF(x, y, z) (((x) & (y)) | ((x) & (z)) | ((y) & (z)))
+#define GG(x, y, z) ((z) ^ ((x) & ((y) ^ (z))))
+
+#define P0(x) x ^ ROTL32(x, 9) ^ ROTL32(x, 17)
+#define P1(x) x ^ ROTL32(x, 15) ^ ROTL32(x, 23)
+
+static __forceinline__ __device__
+void sm3_compress2(uint32_t digest[8], const uint32_t pblock[16])
+{
+ uint32_t tt1, tt2, i, t, ss1, ss2, x, y;
+ uint32_t w[68];
+ uint32_t a = digest[0];
+ uint32_t b = digest[1];
+ uint32_t c = digest[2];
+ uint32_t d = digest[3];
+ uint32_t e = digest[4];
+ uint32_t f = digest[5];
+ uint32_t g = digest[6];
+ uint32_t h = digest[7];
+
+ #pragma unroll
+ for (i = 0; i<16; i++) {
+ w[i] = cuda_swab32(pblock[i]);
+ }
+
+ for (i = 16; i<68; i++) {
+ x = ROTL32(w[i - 3], 15);
+ y = ROTL32(w[i - 13], 7);
+
+ x ^= w[i - 16];
+ x ^= w[i - 9];
+ y ^= w[i - 6];
+
+ w[i] = P1(x) ^ y;
+ }
+
+ for (i = 0; i<64; i++) {
+
+ t = (i < 16) ? 0x79cc4519 : 0x7a879d8a;
+
+ ss2 = ROTL32(a, 12);
+ ss1 = ROTL32(ss2 + e + ROTL32(t, i), 7);
+ ss2 ^= ss1;
+
+ tt1 = d + ss2 + (w[i] ^ w[i + 4]);
+ tt2 = h + ss1 + w[i];
+
+ if (i < 16) {
+ tt1 += F(a, b, c);
+ tt2 += F(e, f, g);
+ }
+ else {
+ tt1 += FF(a, b, c);
+ tt2 += GG(e, f, g);
+ }
+ d = c;
+ c = ROTL32(b, 9);
+ b = a;
+ a = tt1;
+ h = g;
+ g = ROTL32(f, 19);
+ f = e;
+ e = P0(tt2);
+ }
+
+ digest[0] ^= a;
+ digest[1] ^= b;
+ digest[2] ^= c;
+ digest[3] ^= d;
+ digest[4] ^= e;
+ digest[5] ^= f;
+ digest[6] ^= g;
+ digest[7] ^= h;
+}
+
+/***************************************************/
+// GPU Hash Function
+__global__
+void sm3_gpu_hash_64(const uint32_t threads, uint32_t *g_hash)
+{
+ const uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
+
+ if (thread < threads)
+ {
+ const size_t hashPosition = thread;
+
+ uint32_t digest[8];
+ digest[0] = 0x7380166F;
+ digest[1] = 0x4914B2B9;
+ digest[2] = 0x172442D7;
+ digest[3] = 0xDA8A0600;
+ digest[4] = 0xA96F30BC;
+ digest[5] = 0x163138AA;
+ digest[6] = 0xE38DEE4D;
+ digest[7] = 0xB0FB0E4E;
+
+ uint32_t *pHash = &g_hash[hashPosition << 4];
+ sm3_compress2(digest, pHash);
+
+ uint32_t block[16];
+ block[0] = 0x80;
+
+ #pragma unroll
+ for (int i = 1; i < 14; i++)
+ block[i] = 0;
+
+ // count
+ block[14] = cuda_swab32(1 >> 23);
+ block[15] = cuda_swab32((1 << 9) + (0 << 3));
+
+ sm3_compress2(digest, block);
+
+ for (int i = 0; i < 8; i++)
+ pHash[i] = cuda_swab32(digest[i]);
+
+ for (int i = 8; i < 16; i++)
+ pHash[i] = 0;
+ }
+}
+
+__host__
+void sm3_cuda_hash_64(int thr_id, uint32_t threads, uint32_t *g_hash, int order)
+{
+ const uint32_t threadsperblock = 256;
+
+ dim3 grid((threads + threadsperblock - 1) / threadsperblock);
+ dim3 block(threadsperblock);
+
+ sm3_gpu_hash_64 <<>>(threads, g_hash);
+ //MyStreamSynchronize(NULL, order, thr_id);
+}
diff --git a/x13/hsr.cu b/x13/hsr.cu
new file mode 100644
index 0000000..e864446
--- /dev/null
+++ b/x13/hsr.cu
@@ -0,0 +1,265 @@
+/*
+ * X13 algorithm
+ */
+extern "C"
+{
+#include "sph/sph_blake.h"
+#include "sph/sph_bmw.h"
+#include "sph/sph_groestl.h"
+#include "sph/sph_skein.h"
+#include "sph/sph_jh.h"
+#include "sph/sph_keccak.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"
+
+#include "sph/sph_hamsi.h"
+#include "sph/sph_fugue.h"
+}
+#include "sm3.h"
+
+#include "miner.h"
+
+#include "cuda_helper.h"
+#include "x11/cuda_x11.h"
+
+static uint32_t *d_hash[MAX_GPUS];
+
+extern void sm3_cuda_hash_64(int thr_id, uint32_t threads, uint32_t *d_hash, int order);
+
+extern void x13_hamsi512_cpu_init(int thr_id, uint32_t threads);
+extern void x13_hamsi512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+
+extern void x13_fugue512_cpu_init(int thr_id, uint32_t threads);
+extern void x13_fugue512_cpu_hash_64(int thr_id, uint32_t threads, uint32_t startNounce, uint32_t *d_nonceVector, uint32_t *d_hash, int order);
+extern void x13_fugue512_cpu_free(int thr_id);
+
+// HSR CPU Hash
+extern "C" void hsr_hash(void *output, const void *input)
+{
+ // blake1-bmw2-grs3-skein4-jh5-keccak6-luffa7-cubehash8-shavite9-simd10-echo11-hamsi12-fugue13
+
+ sph_blake512_context ctx_blake;
+ sph_bmw512_context ctx_bmw;
+ sph_groestl512_context ctx_groestl;
+ sph_jh512_context ctx_jh;
+ sph_keccak512_context ctx_keccak;
+ sph_skein512_context ctx_skein;
+ sph_luffa512_context ctx_luffa;
+ sph_cubehash512_context ctx_cubehash;
+ sph_shavite512_context ctx_shavite;
+ sph_simd512_context ctx_simd;
+ sph_echo512_context ctx_echo;
+ sm3_ctx_t ctx_sm3;
+ sph_hamsi512_context ctx_hamsi;
+ sph_fugue512_context ctx_fugue;
+
+ uint32_t hash[32];
+ memset(hash, 0, sizeof hash);
+
+ sph_blake512_init(&ctx_blake);
+ sph_blake512(&ctx_blake, input, 80);
+ sph_blake512_close(&ctx_blake, (void*) hash);
+
+ sph_bmw512_init(&ctx_bmw);
+ sph_bmw512(&ctx_bmw, (const void*) hash, 64);
+ sph_bmw512_close(&ctx_bmw, (void*) hash);
+
+ sph_groestl512_init(&ctx_groestl);
+ sph_groestl512(&ctx_groestl, (const void*) hash, 64);
+ sph_groestl512_close(&ctx_groestl, (void*) hash);
+
+ sph_skein512_init(&ctx_skein);
+ sph_skein512(&ctx_skein, (const void*) hash, 64);
+ sph_skein512_close(&ctx_skein, (void*) hash);
+
+ sph_jh512_init(&ctx_jh);
+ sph_jh512(&ctx_jh, (const void*) hash, 64);
+ sph_jh512_close(&ctx_jh, (void*) hash);
+
+ sph_keccak512_init(&ctx_keccak);
+ sph_keccak512(&ctx_keccak, (const void*) hash, 64);
+ sph_keccak512_close(&ctx_keccak, (void*) hash);
+
+ sph_luffa512_init(&ctx_luffa);
+ sph_luffa512(&ctx_luffa, (const void*) hash, 64);
+ sph_luffa512_close (&ctx_luffa, (void*) hash);
+
+ sph_cubehash512_init(&ctx_cubehash);
+ sph_cubehash512(&ctx_cubehash, (const void*) hash, 64);
+ sph_cubehash512_close(&ctx_cubehash, (void*) hash);
+
+ sph_shavite512_init(&ctx_shavite);
+ sph_shavite512(&ctx_shavite, (const void*) hash, 64);
+ sph_shavite512_close(&ctx_shavite, (void*) hash);
+
+ sph_simd512_init(&ctx_simd);
+ sph_simd512(&ctx_simd, (const void*) hash, 64);
+ sph_simd512_close(&ctx_simd, (void*) hash);
+
+ sph_echo512_init(&ctx_echo);
+ sph_echo512(&ctx_echo, (const void*) hash, 64);
+ sph_echo512_close(&ctx_echo, (void*) hash);
+
+ sm3_init(&ctx_sm3);
+ sm3_update(&ctx_sm3, (const unsigned char*) hash, 64);
+ memset(hash, 0, sizeof hash);
+ sm3_close(&ctx_sm3, (void*) hash);
+
+ sph_hamsi512_init(&ctx_hamsi);
+ sph_hamsi512(&ctx_hamsi, (const void*) hash, 64);
+ sph_hamsi512_close(&ctx_hamsi, (void*) hash);
+
+ sph_fugue512_init(&ctx_fugue);
+ sph_fugue512(&ctx_fugue, (const void*) hash, 64);
+ sph_fugue512_close(&ctx_fugue, (void*) hash);
+
+ memcpy(output, hash, 32);
+}
+
+static bool init[MAX_GPUS] = { 0 };
+
+extern "C" int scanhash_hsr(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done)
+{
+ uint32_t *pdata = work->data;
+ uint32_t *ptarget = work->target;
+ const uint32_t first_nonce = pdata[19];
+ int intensity = 19; // (device_sm[device_map[thr_id]] > 500 && !is_windows()) ? 20 : 19;
+ uint32_t throughput = cuda_default_throughput(thr_id, 1 << intensity); // 19=256*256*8;
+ //if (init[thr_id]) throughput = min(throughput, max_nonce - first_nonce);
+
+ if (opt_benchmark)
+ ((uint32_t*)ptarget)[7] = 0x000f;
+
+ if (!init[thr_id])
+ {
+ cudaSetDevice(device_map[thr_id]);
+ if (opt_cudaschedule == -1 && gpu_threads == 1) {
+ cudaDeviceReset();
+ // reduce cpu usage
+ cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
+ CUDA_LOG_ERROR();
+ }
+ gpulog(LOG_INFO, thr_id, "Intensity set to %g, %u cuda threads", throughput2intensity(throughput), throughput);
+
+ quark_blake512_cpu_init(thr_id, throughput);
+ quark_groestl512_cpu_init(thr_id, throughput);
+ quark_skein512_cpu_init(thr_id, throughput);
+ quark_bmw512_cpu_init(thr_id, throughput);
+ quark_keccak512_cpu_init(thr_id, throughput);
+ quark_jh512_cpu_init(thr_id, throughput);
+ x11_luffaCubehash512_cpu_init(thr_id, throughput);
+ x11_shavite512_cpu_init(thr_id, throughput);
+ if (x11_simd512_cpu_init(thr_id, throughput) != 0) {
+ return 0;
+ }
+ x11_echo512_cpu_init(thr_id, throughput);
+ x13_hamsi512_cpu_init(thr_id, throughput);
+ x13_fugue512_cpu_init(thr_id, throughput);
+
+ CUDA_CALL_OR_RET_X(cudaMalloc(&d_hash[thr_id], 16 * sizeof(uint32_t) * throughput), 0);
+
+ cuda_check_cpu_init(thr_id, throughput);
+
+ init[thr_id] = true;
+ }
+
+ uint32_t endiandata[20];
+ for (int k=0; k < 20; k++)
+ be32enc(&endiandata[k], pdata[k]);
+
+ quark_blake512_cpu_setBlock_80(thr_id, endiandata);
+ cuda_check_cpu_setTarget(ptarget);
+
+ do {
+ int order = 0;
+
+ quark_blake512_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id]); order++;
+ quark_bmw512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
+ quark_groestl512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
+ quark_skein512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
+ quark_jh512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
+ quark_keccak512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
+ x11_luffaCubehash512_cpu_hash_64(thr_id, throughput, d_hash[thr_id], order++);
+ x11_shavite512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
+ x11_simd512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
+ x11_echo512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
+ sm3_cuda_hash_64(thr_id, throughput, d_hash[thr_id], order++);
+ x13_hamsi512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
+ x13_fugue512_cpu_hash_64(thr_id, throughput, pdata[19], NULL, d_hash[thr_id], order++);
+
+ *hashes_done = pdata[19] - first_nonce + throughput;
+
+ CUDA_LOG_ERROR();
+
+ work->nonces[0] = cuda_check_hash(thr_id, throughput, pdata[19], d_hash[thr_id]);
+ if (work->nonces[0] != UINT32_MAX)
+ {
+ const uint32_t Htarg = ptarget[7];
+ uint32_t _ALIGN(64) vhash[8];
+ be32enc(&endiandata[19], work->nonces[0]);
+ hsr_hash(vhash, endiandata);
+
+ if (vhash[7] <= ptarget[7] && fulltest(vhash, ptarget)) {
+ work->valid_nonces = 1;
+ work->nonces[1] = cuda_check_hash_suppl(thr_id, throughput, pdata[19], d_hash[thr_id], 1);
+ work_set_target_ratio(work, vhash);
+ if (work->nonces[1] != 0) {
+ be32enc(&endiandata[19], work->nonces[1]);
+ hsr_hash(vhash, endiandata);
+ bn_set_target_ratio(work, vhash, 1);
+ work->valid_nonces++;
+ pdata[19] = max(work->nonces[0], work->nonces[1]) + 1;
+ } else {
+ pdata[19] = work->nonces[0] + 1; // cursor
+ }
+ return work->valid_nonces;
+ }
+ else if (vhash[7] > Htarg) {
+ gpu_increment_reject(thr_id);
+ if (!opt_quiet)
+ gpulog(LOG_WARNING, thr_id, "result for %08x does not validate on CPU!", work->nonces[0]);
+ pdata[19] = work->nonces[0] + 1;
+ continue;
+ }
+ }
+
+ if ((uint64_t)throughput + pdata[19] >= max_nonce) {
+ pdata[19] = max_nonce;
+ break;
+ }
+ pdata[19] += throughput;
+
+ } while (!work_restart[thr_id].restart);
+
+ *hashes_done = pdata[19] - first_nonce;
+
+ CUDA_LOG_ERROR();
+
+ return 0;
+}
+
+// cleanup
+extern "C" void free_hsr(int thr_id)
+{
+ if (!init[thr_id])
+ return;
+
+ cudaThreadSynchronize();
+
+ cudaFree(d_hash[thr_id]);
+
+ quark_blake512_cpu_free(thr_id);
+ quark_groestl512_cpu_free(thr_id);
+ x11_simd512_cpu_free(thr_id);
+ x13_fugue512_cpu_free(thr_id);
+
+ cuda_check_cpu_free(thr_id);
+ CUDA_LOG_ERROR();
+
+ cudaDeviceSynchronize();
+ init[thr_id] = false;
+}
diff --git a/x13/sm3.c b/x13/sm3.c
new file mode 100644
index 0000000..295ba15
--- /dev/null
+++ b/x13/sm3.c
@@ -0,0 +1,220 @@
+/* ====================================================================
+ * Copyright (c) 2014 - 2017 The GmSSL Project. 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.
+ *
+ * 3. All advertising materials mentioning features or use of this
+ * software must display the following acknowledgment:
+ * "This product includes software developed by the GmSSL Project.
+ * (http://gmssl.org/)"
+ *
+ * 4. The name "GmSSL Project" must not be used to endorse or promote
+ * products derived from this software without prior written
+ * permission. For written permission, please contact
+ * guanzhi1980@gmail.com.
+ *
+ * 5. Products derived from this software may not be called "GmSSL"
+ * nor may "GmSSL" appear in their names without prior written
+ * permission of the GmSSL Project.
+ *
+ * 6. Redistributions of any form whatsoever must retain the following
+ * acknowledgment:
+ * "This product includes software developed by the GmSSL Project
+ * (http://gmssl.org/)"
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE GmSSL PROJECT ``AS IS'' AND ANY
+ * EXPRESSED 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 GmSSL PROJECT OR
+ * ITS 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.
+ * ====================================================================
+ */
+
+#include
+
+#include "sm3.h"
+
+void sm3_init(sm3_ctx_t *ctx)
+{
+ ctx->digest[0] = 0x7380166F;
+ ctx->digest[1] = 0x4914B2B9;
+ ctx->digest[2] = 0x172442D7;
+ ctx->digest[3] = 0xDA8A0600;
+ ctx->digest[4] = 0xA96F30BC;
+ ctx->digest[5] = 0x163138AA;
+ ctx->digest[6] = 0xE38DEE4D;
+ ctx->digest[7] = 0xB0FB0E4E;
+
+ ctx->nblocks = 0;
+ ctx->num = 0;
+}
+
+void sm3_update(sm3_ctx_t *ctx, const unsigned char* data, size_t data_len)
+{
+ if (ctx->num) {
+ unsigned int left = SM3_BLOCK_SIZE - ctx->num;
+ if (data_len < left) {
+ memcpy(ctx->block + ctx->num, data, data_len);
+ ctx->num += data_len;
+ return;
+ } else {
+ memcpy(ctx->block + ctx->num, data, left);
+ sm3_compress(ctx->digest, ctx->block);
+ ctx->nblocks++;
+ data += left;
+ data_len -= left;
+ }
+ }
+ while (data_len >= SM3_BLOCK_SIZE) {
+ sm3_compress(ctx->digest, data);
+ ctx->nblocks++;
+ data += SM3_BLOCK_SIZE;
+ data_len -= SM3_BLOCK_SIZE;
+ }
+ ctx->num = data_len;
+ if (data_len) {
+ memcpy(ctx->block, data, data_len);
+ }
+}
+
+void sm3_close(void *cc, void *dst)
+{
+ sm3_final(cc, dst);
+ memset(cc, 0, sizeof(sm3_ctx_t));
+}
+
+void sm3_final(sm3_ctx_t *ctx, unsigned char *digest)
+{
+ int i;
+ uint32_t *pdigest = (uint32_t *)digest;
+ uint32_t *count = (uint32_t *)(ctx->block + SM3_BLOCK_SIZE - 8);
+
+ ctx->block[ctx->num] = 0x80;
+
+ if (ctx->num + 9 <= SM3_BLOCK_SIZE) {
+ memset(ctx->block + ctx->num + 1, 0, SM3_BLOCK_SIZE - ctx->num - 9);
+ } else {
+ memset(ctx->block + ctx->num + 1, 0, SM3_BLOCK_SIZE - ctx->num - 1);
+ sm3_compress(ctx->digest, ctx->block);
+ memset(ctx->block, 0, SM3_BLOCK_SIZE - 8);
+ }
+
+ count[0] = cpu_to_be32((ctx->nblocks) >> 23);
+ count[1] = cpu_to_be32((ctx->nblocks << 9) + (ctx->num << 3));
+
+ sm3_compress(ctx->digest, ctx->block);
+ for (i = 0; i < sizeof(ctx->digest)/sizeof(ctx->digest[0]); i++) {
+ pdigest[i] = cpu_to_be32(ctx->digest[i]);
+ }
+}
+
+#define ROTATELEFT(X,n) (((X)<<(n)) | ((X)>>(32-(n))))
+
+#define P0(x) ((x) ^ ROTATELEFT((x),9) ^ ROTATELEFT((x),17))
+#define P1(x) ((x) ^ ROTATELEFT((x),15) ^ ROTATELEFT((x),23))
+
+#define FF0(x,y,z) ( (x) ^ (y) ^ (z))
+#define FF1(x,y,z) (((x) & (y)) | ( (x) & (z)) | ( (y) & (z)))
+
+#define GG0(x,y,z) ( (x) ^ (y) ^ (z))
+#define GG1(x,y,z) (((x) & (y)) | ( (~(x)) & (z)) )
+
+
+void sm3_compress(uint32_t digest[8], const unsigned char block[64])
+{
+ int j;
+ uint32_t W[68], W1[64];
+ const uint32_t *pblock = (const uint32_t *)block;
+
+ uint32_t A = digest[0];
+ uint32_t B = digest[1];
+ uint32_t C = digest[2];
+ uint32_t D = digest[3];
+ uint32_t E = digest[4];
+ uint32_t F = digest[5];
+ uint32_t G = digest[6];
+ uint32_t H = digest[7];
+ uint32_t SS1,SS2,TT1,TT2,T[64];
+
+ for (j = 0; j < 16; j++) {
+ W[j] = cpu_to_be32(pblock[j]);
+ }
+ for (j = 16; j < 68; j++) {
+ W[j] = P1( W[j-16] ^ W[j-9] ^ ROTATELEFT(W[j-3],15)) ^ ROTATELEFT(W[j - 13],7 ) ^ W[j-6];;
+ }
+ for( j = 0; j < 64; j++) {
+ W1[j] = W[j] ^ W[j+4];
+ }
+
+ for(j =0; j < 16; j++) {
+
+ T[j] = 0x79CC4519;
+ SS1 = ROTATELEFT((ROTATELEFT(A,12) + E + ROTATELEFT(T[j],j)), 7);
+ SS2 = SS1 ^ ROTATELEFT(A,12);
+ TT1 = FF0(A,B,C) + D + SS2 + W1[j];
+ TT2 = GG0(E,F,G) + H + SS1 + W[j];
+ D = C;
+ C = ROTATELEFT(B,9);
+ B = A;
+ A = TT1;
+ H = G;
+ G = ROTATELEFT(F,19);
+ F = E;
+ E = P0(TT2);
+ }
+
+ for(j =16; j < 64; j++) {
+
+ T[j] = 0x7A879D8A;
+ SS1 = ROTATELEFT((ROTATELEFT(A,12) + E + ROTATELEFT(T[j],j)), 7);
+ SS2 = SS1 ^ ROTATELEFT(A,12);
+ TT1 = FF1(A,B,C) + D + SS2 + W1[j];
+ TT2 = GG1(E,F,G) + H + SS1 + W[j];
+ D = C;
+ C = ROTATELEFT(B,9);
+ B = A;
+ A = TT1;
+ H = G;
+ G = ROTATELEFT(F,19);
+ F = E;
+ E = P0(TT2);
+ }
+
+ digest[0] ^= A;
+ digest[1] ^= B;
+ digest[2] ^= C;
+ digest[3] ^= D;
+ digest[4] ^= E;
+ digest[5] ^= F;
+ digest[6] ^= G;
+ digest[7] ^= H;
+}
+
+void sm3(const unsigned char *msg, size_t msglen,
+ unsigned char dgst[SM3_DIGEST_LENGTH])
+{
+ sm3_ctx_t ctx;
+
+ sm3_init(&ctx);
+ sm3_update(&ctx, msg, msglen);
+ sm3_final(&ctx, dgst);
+
+ memset(&ctx, 0, sizeof(sm3_ctx_t));
+}
diff --git a/x13/sm3.h b/x13/sm3.h
new file mode 100644
index 0000000..05c6595
--- /dev/null
+++ b/x13/sm3.h
@@ -0,0 +1,109 @@
+/* ====================================================================
+ * Copyright (c) 2014 - 2016 The GmSSL Project. All rights reserved.
+ * Copyright (c) 2017 - YiiMP (cleaned hmac dead stuff)
+ *
+ * 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.
+ *
+ * 3. All advertising materials mentioning features or use of this
+ * software must display the following acknowledgment:
+ * "This product includes software developed by the GmSSL Project.
+ * (http://gmssl.org/)"
+ *
+ * 4. The name "GmSSL Project" must not be used to endorse or promote
+ * products derived from this software without prior written
+ * permission. For written permission, please contact
+ * guanzhi1980@gmail.com.
+ *
+ * 5. Products derived from this software may not be called "GmSSL"
+ * nor may "GmSSL" appear in their names without prior written
+ * permission of the GmSSL Project.
+ *
+ * 6. Redistributions of any form whatsoever must retain the following
+ * acknowledgment:
+ * "This product includes software developed by the GmSSL Project
+ * (http://gmssl.org/)"
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE GmSSL PROJECT ``AS IS'' AND ANY
+ * EXPRESSED 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 GmSSL PROJECT OR
+ * ITS 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.
+ * ====================================================================
+ */
+
+#ifndef _SM3_H
+#define _SM3_H
+
+#define SM3_DIGEST_LENGTH 32
+#define SM3_BLOCK_SIZE 64
+#define SM3_CBLOCK (SM3_BLOCK_SIZE)
+#define SM3_HMAC_SIZE (SM3_DIGEST_LENGTH)
+
+
+#include
+#include
+#include
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+
+typedef struct {
+ uint32_t digest[8];
+ int nblocks;
+ unsigned char block[64];
+ int num;
+} sm3_ctx_t;
+
+void sm3_init(sm3_ctx_t *ctx);
+void sm3_update(sm3_ctx_t *ctx, const unsigned char* data, size_t data_len);
+void sm3_close(void *cc, void *dst);
+
+void sm3_final(sm3_ctx_t *ctx, unsigned char digest[SM3_DIGEST_LENGTH]);
+void sm3_compress(uint32_t digest[8], const unsigned char block[SM3_BLOCK_SIZE]);
+void sm3(const unsigned char *data, size_t datalen,
+ unsigned char digest[SM3_DIGEST_LENGTH]);
+
+#ifdef CPU_BIGENDIAN
+
+#define cpu_to_be16(v) (v)
+#define cpu_to_be32(v) (v)
+#define be16_to_cpu(v) (v)
+#define be32_to_cpu(v) (v)
+
+#else
+
+#define cpu_to_le16(v) (v)
+#define cpu_to_le32(v) (v)
+#define le16_to_cpu(v) (v)
+#define le32_to_cpu(v) (v)
+
+#define cpu_to_be16(v) (((v)<< 8) | ((v)>>8))
+#define cpu_to_be32(v) (((v)>>24) | (((v)>>8)&0xff00) | (((v)<<8)&0xff0000) | ((v)<<24))
+#define be16_to_cpu(v) cpu_to_be16(v)
+#define be32_to_cpu(v) cpu_to_be32(v)
+
+#endif
+
+#ifdef __cplusplus
+}
+#endif
+#endif