mirror of https://github.com/GOSTSec/ccminer
Tanguy Pruvot
7 years ago
12 changed files with 761 additions and 0 deletions
@ -0,0 +1,139 @@
@@ -0,0 +1,139 @@
|
||||
#include <stdio.h> |
||||
#include <stdint.h> |
||||
#include <memory.h> |
||||
|
||||
#include <cuda_helper.h> |
||||
#include <miner.h> |
||||
|
||||
#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 <<<grid, block>>>(threads, g_hash); |
||||
//MyStreamSynchronize(NULL, order, thr_id); |
||||
} |
||||
@ -0,0 +1,265 @@
@@ -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; |
||||
} |
||||
@ -0,0 +1,220 @@
@@ -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 <string.h> |
||||
|
||||
#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)); |
||||
} |
||||
@ -0,0 +1,109 @@
@@ -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 <sys/types.h> |
||||
#include <stdint.h> |
||||
#include <string.h> |
||||
|
||||
#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 |
||||
Loading…
Reference in new issue