GOSTSec
/
ccminer
mirror of https://github.com/GOSTSec/ccminer
1
0
Fork 0
Code Issues Releases Wiki Activity
GOSTcoin support for ccminer CUDA miner project, compatible with most nvidia cards
You can not select more than 25 topics Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
1029 Commits
2 Branches
52 Tags
66 MiB
 
 
 
 
 
Branch: master
Branches Tags
${ item.name }
Create tag ${ searchTerm }
Create branch ${ searchTerm }
from 'master'
${ noResults }
ccminer/Algo256/cuda_bmw.cu

380 lines
13 KiB
Raw Permalink Blame History

/**
* BMW-256 CUDA Implementation - tpruvot 2015
*
* Not optimal but close to the sph version and easier to adapt.
*/
#include <stdio.h>
#include <memory.h>
#define SPH_64 1
#define USE_MIDSTATE
extern "C" {
#include "sph/sph_bmw.h"
}
#include "cuda_helper.h"
__constant__ uint64_t c_PaddedMessage80[16]; // padded message (80 bytes + padding)
#ifndef USE_MIDSTATE
__constant__ static sph_u32 IV256[16] = {
0x40414243, 0x44454647, 0x48494A4B, 0x4C4D4E4F,
0x50515253, 0x54555657, 0x58595A5B, 0x5C5D5E5F,
0x60616263, 0x64656667, 0x68696A6B, 0x6C6D6E6F,
0x70717273, 0x74757677, 0x78797A7B, 0x7C7D7E7F
};
#endif
__constant__ static sph_u32 final_s[16] = {
0xaaaaaaa0, 0xaaaaaaa1, 0xaaaaaaa2, 0xaaaaaaa3,
0xaaaaaaa4, 0xaaaaaaa5, 0xaaaaaaa6, 0xaaaaaaa7,
0xaaaaaaa8, 0xaaaaaaa9, 0xaaaaaaaa, 0xaaaaaaab,
0xaaaaaaac, 0xaaaaaaad, 0xaaaaaaae, 0xaaaaaaaf
};
static sph_bmw_small_context* d_midstate[MAX_GPUS];
#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) ^ ((x) << 3) ^ ROTL32(x, 4) ^ ROTL32(x, 19))
#define ss1(x) (((x) >> 1) ^ ((x) << 2) ^ ROTL32(x, 8) ^ ROTL32(x, 23))
#define ss2(x) (((x) >> 2) ^ ((x) << 1) ^ ROTL32(x, 12) ^ ROTL32(x, 25))
#define ss3(x) (((x) >> 2) ^ ((x) << 2) ^ ROTL32(x, 15) ^ ROTL32(x, 29))
#define ss4(x) (((x) >> 1) ^ (x))
#define ss5(x) (((x) >> 2) ^ (x))
#define rs1(x) ROTL32(x, 3)
#define rs2(x) ROTL32(x, 7)
#define rs3(x) ROTL32(x, 13)
#define rs4(x) ROTL32(x, 16)
#define rs5(x) ROTL32(x, 19)
#define rs6(x) ROTL32(x, 23)
#define rs7(x) ROTL32(x, 27)
#define MAKE_W(tt, i0, op01, i1, op12, i2, op23, i3, op34, i4) \
tt((data[i0] ^ h[i0]) op01 (data[i1] ^ h[i1]) op12 (data[i2] ^ h[i2]) op23 (data[i3] ^ h[i3]) op34 (data[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)
__device__
static void gpu_compress_small(const sph_u32 *data, const sph_u32 h[16], sph_u32 dh[16])
{
// FOLD MAKE_Qas;
sph_u32 dx[16];
for (int i=0; i<16; i++)
dx[i] = data[i] ^ h[i];
sph_u32 qt[32];
qt[ 0] = dx[ 5] - dx[7] + dx[10] + dx[13] + dx[14]; // Ws0
qt[ 1] = dx[ 6] - dx[8] + dx[11] + dx[14] - dx[15]; // Ws1
qt[ 2] = dx[ 0] + dx[7] + dx[ 9] - dx[12] + dx[15]; // Ws2
qt[ 3] = dx[ 0] - dx[1] + dx[ 8] - dx[10] + dx[13]; // Ws3
qt[ 4] = dx[ 1] + dx[2] + dx[ 9] - dx[11] - dx[14]; // Ws4;
qt[ 5] = dx[ 3] - dx[2] + dx[10] - dx[12] + dx[15]; // Ws5;
qt[ 6] = dx[ 4] - dx[0] - dx[ 3] - dx[11] + dx[13]; // Ws6;
qt[ 7] = dx[ 1] - dx[4] - dx[ 5] - dx[12] - dx[14]; // Ws7;
qt[ 8] = dx[ 2] - dx[5] - dx[ 6] + dx[13] - dx[15]; // Ws8;
qt[ 9] = dx[ 0] - dx[3] + dx[ 6] - dx[ 7] + dx[14]; // Ws9;
qt[10] = dx[ 8] - dx[1] - dx[ 4] - dx[ 7] + dx[15]; // Ws10;
qt[11] = dx[ 8] - dx[0] - dx[ 2] - dx[ 5] + dx[ 9]; // Ws11;
qt[12] = dx[ 1] + dx[3] - dx[ 6] - dx[ 9] + dx[10]; // Ws12;
qt[13] = dx[ 2] + dx[4] + dx[ 7] + dx[10] + dx[11]; // Ws13;
qt[14] = dx[ 3] - dx[5] + dx[ 8] - dx[11] - dx[12]; // Ws14;
qt[15] = dx[12] - dx[4] - dx[ 6] - dx[ 9] + dx[13]; // Ws15;
qt[ 0] = ss0(qt[ 0]) + h[ 1];
qt[ 1] = ss1(qt[ 1]) + h[ 2];
qt[ 2] = ss2(qt[ 2]) + h[ 3];
qt[ 3] = ss3(qt[ 3]) + h[ 4];
qt[ 4] = ss4(qt[ 4]) + h[ 5];
qt[ 5] = ss0(qt[ 5]) + h[ 6];
qt[ 6] = ss1(qt[ 6]) + h[ 7];
qt[ 7] = ss2(qt[ 7]) + h[ 8];
qt[ 8] = ss3(qt[ 8]) + h[ 9];
qt[ 9] = ss4(qt[ 9]) + h[10];
qt[10] = ss0(qt[10]) + h[11];
qt[11] = ss1(qt[11]) + h[12];
qt[12] = ss2(qt[12]) + h[13];
qt[13] = ss3(qt[13]) + h[14];
qt[14] = ss4(qt[14]) + h[15];
qt[15] = ss0(qt[15]) + h[ 0];
//MAKE_Qbs;
#define Ks(j) ((sph_u32)(0x05555555UL * j))
#define Qs(j) (qt[j])
#define expand1s_in(i16, \
i0, i1, i2, i3, i4, i5, i6, i7, i8, i9, i10, i11, i12, i13, i14, i15, \
i0m, i1m, i3m, i4m, i7m, i10m, i11m) \
(ss1(qt[i0]) + ss2(qt[i1]) + ss3(qt[i2]) + ss0(qt[i3]) + ss1(qt[i4]) + ss2(qt[i5]) + ss3(qt[i6]) + ss0(qt[i7]) \
+ ss1(qt[i8]) + ss2(qt[i9]) + ss3(qt[i10]) + ss0(qt[i11]) + ss1(qt[i12]) + ss2(qt[i13]) + ss3(qt[i14]) + ss0(qt[i15]) \
+ ((ROTL32(data[i0m], i1m) + ROTL32(data[i3m], i4m) - ROTL32(data[i10m], i11m) + Ks(i16)) ^ h[i7m]))
qt[16] = expand1s_in(16, 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 0, 1, 3, 4, 7, 10, 11);
qt[17] = expand1s_in(17, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 1, 2, 4, 5, 8, 11, 12);
#define expand2s_inner(qf, i16, \
i0, i1, i2, i3, i4, i5, i6, i7, i8, i9, i10, i11, i12, i13, i14, i15, \
i0m, i1m, i3m, i4m, i7m, i10m, i11m) \
(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)) \
+ ((ROTL32(data[i0m], i1m) + ROTL32(data[i3m], i4m) - ROTL32(data[i10m], i11m) + Ks(i16)) ^ h[i7m]))
#ifdef _MSC_VER
#define LPAR (
#define expand2s(i16) \
expand2s_(Qs, i16, I16_ ## i16, M16_ ## i16)
#define expand2s_(qf, i16, ix, iy) \
expand2s_inner LPAR qf, i16, ix, iy)
#else
#define expand2s_(i16, ix, iy) \
expand2s_inner(Qs, i16, ix, iy)
#define expand2s(i16) \
expand2s_(i16, I16_ ## i16, M16_ ## i16)
#endif
qt[18] = expand2s(18);
qt[19] = expand2s(19);
qt[20] = expand2s(20);
qt[21] = expand2s(21);
qt[22] = expand2s(22);
qt[23] = expand2s(23);
qt[24] = expand2s(24);
qt[25] = expand2s(25);
qt[26] = expand2s(26);
qt[27] = expand2s(27);
qt[28] = expand2s(28);
qt[29] = expand2s(29);
qt[30] = expand2s(30);
qt[31] = expand2s(31);
sph_u32 xl, xh;
xl = Qs(16) ^ Qs(17) ^ Qs(18) ^ Qs(19) ^ Qs(20) ^ Qs(21) ^ Qs(22) ^ Qs(23);
xh = xl ^ Qs(24) ^ Qs(25) ^ Qs(26) ^ Qs(27) ^ Qs(28) ^ Qs(29) ^ Qs(30) ^ Qs(31);
dh[ 0] = ((xh << 5) ^ (Qs(16) >> 5) ^ data[ 0]) + (xl ^ Qs(24) ^ Qs(0));
dh[ 1] = ((xh >> 7) ^ (Qs(17) << 8) ^ data[ 1]) + (xl ^ Qs(25) ^ Qs(1));
dh[ 2] = ((xh >> 5) ^ (Qs(18) << 5) ^ data[ 2]) + (xl ^ Qs(26) ^ Qs(2));
dh[ 3] = ((xh >> 1) ^ (Qs(19) << 5) ^ data[ 3]) + (xl ^ Qs(27) ^ Qs(3));
dh[ 4] = ((xh >> 3) ^ (Qs(20) << 0) ^ data[ 4]) + (xl ^ Qs(28) ^ Qs(4));
dh[ 5] = ((xh << 6) ^ (Qs(21) >> 6) ^ data[ 5]) + (xl ^ Qs(29) ^ Qs(5));
dh[ 6] = ((xh >> 4) ^ (Qs(22) << 6) ^ data[ 6]) + (xl ^ Qs(30) ^ Qs(6));
dh[ 7] = ((xh >> 11) ^ (Qs(23) << 2) ^ data[ 7]) + (xl ^ Qs(31) ^ Qs(7));
dh[ 8] = ROTL32(dh[4], 9) + (xh ^ Qs(24) ^ data[ 8]) + ((xl << 8) ^ Qs(23) ^ Qs( 8));
dh[ 9] = ROTL32(dh[5], 10) + (xh ^ Qs(25) ^ data[ 9]) + ((xl >> 6) ^ Qs(16) ^ Qs( 9));
dh[10] = ROTL32(dh[6], 11) + (xh ^ Qs(26) ^ data[10]) + ((xl << 6) ^ Qs(17) ^ Qs(10));
dh[11] = ROTL32(dh[7], 12) + (xh ^ Qs(27) ^ data[11]) + ((xl << 4) ^ Qs(18) ^ Qs(11));
dh[12] = ROTL32(dh[0], 13) + (xh ^ Qs(28) ^ data[12]) + ((xl >> 3) ^ Qs(19) ^ Qs(12));
dh[13] = ROTL32(dh[1], 14) + (xh ^ Qs(29) ^ data[13]) + ((xl >> 4) ^ Qs(20) ^ Qs(13));
dh[14] = ROTL32(dh[2], 15) + (xh ^ Qs(30) ^ data[14]) + ((xl >> 7) ^ Qs(21) ^ Qs(14));
dh[15] = ROTL32(dh[3], 16) + (xh ^ Qs(31) ^ data[15]) + ((xl >> 2) ^ Qs(22) ^ Qs(15));
}
#ifndef USE_MIDSTATE
__device__
static void gpu_bmw256_init(sph_bmw_small_context *sc)
{
memcpy(sc->H, IV256, sizeof sc->H);
sc->ptr = 0;
sc->bit_count = 0;
}
__device__
static void gpu_bmw256(sph_bmw_small_context *sc, const void *data, size_t len)
{
sph_u32 htmp[16];
sph_u32 *h1, *h2;
unsigned char *buf = sc->buf;
size_t ptr = sc->ptr;
sc->bit_count += (sph_u64)len << 3;
h1 = sc->H;
h2 = htmp;
while (len > 0) {
size_t clen;
clen = (sizeof sc->buf) - ptr;
if (clen > len)
clen = len;
memcpy(buf + ptr, data, clen);
data = (const unsigned char *)data + clen;
len -= clen;
ptr += clen;
if (ptr == sizeof sc->buf) {
sph_u32 *ht;
gpu_compress_small((sph_u32 *) buf, h1, h2);
ht = h1;
h1 = h2;
h2 = ht;
ptr = 0;
}
}
sc->ptr = ptr;
if (h1 != sc->H)
memcpy(sc->H, h1, sizeof sc->H);
}
#endif
#define sph_enc64le(ptr, x) \
*((uint64_t*)(ptr)) = x
#define sph_enc64le_aligned sph_enc64le
__device__
static void gpu_bmw256_close(sph_bmw_small_context *sc, uint2 *out)
{
unsigned char *buf = sc->buf;
size_t ptr = sc->ptr;
buf[ptr ++] = 0x80;
sph_u32 *h = sc->H;
sph_u32 h1[16];
if (ptr > (sizeof sc->buf) - 8) {
memset(buf + ptr, 0, (sizeof sc->buf) - ptr);
gpu_compress_small((sph_u32 *) buf, h, h1);
ptr = 0;
h = h1;
}
memset(buf + ptr, 0, sizeof(sc->buf) - 8 - ptr);
sph_enc64le_aligned(buf + sizeof(sc->buf) - 8, SPH_T64(sc->bit_count));
sph_u32 h2[16];
gpu_compress_small((sph_u32 *) buf, h, h2);
gpu_compress_small(h2, final_s, h1);
uint64_t* h64 = (uint64_t*) (&h1[8]);
#pragma unroll
for (int i = 0; i < 4; i++) {
out[i] = vectorize(h64[i]);
}
}
__global__ /* __launch_bounds__(256, 3) */
void bmw256_gpu_hash_80(uint32_t threads, uint32_t startNonce, uint64_t *g_hash, sph_bmw256_context *d_midstate, int swap)
{
uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
if (thread < threads)
{
uint32_t nonce = startNonce + thread;
nonce = swap ? cuda_swab32(nonce): nonce;
#ifndef USE_MIDSTATE
uint2 hash[10];
#pragma unroll
for(int i=0;i<9;i++)
hash[i] = vectorize(c_PaddedMessage80[i]);
hash[9] = make_uint2(c_PaddedMessage80[9], nonce);
sph_bmw256_context ctx;
gpu_bmw256_init(&ctx);
gpu_bmw256(&ctx, (void*) hash, 80);
#else
sph_bmw256_context ctx;
ctx.ptr = 16; ctx.bit_count = 640;
uint2 *buf = (uint2 *) ctx.buf;
buf[0] = vectorize(c_PaddedMessage80[8]);
buf[1] = make_uint2(c_PaddedMessage80[9], nonce);
#pragma unroll
for(int i=0;i<16;i++)
ctx.H[i] = d_midstate->H[i];
#endif
gpu_bmw256_close(&ctx, (uint2*) &g_hash[thread << 2]);
}
}
__host__
void bmw256_cpu_hash_80(int thr_id, uint32_t threads, uint32_t startNonce, uint32_t *d_outputHash, int swap)
{
const uint32_t threadsperblock = 256;
dim3 grid((threads + threadsperblock-1)/threadsperblock);
dim3 block(threadsperblock);
bmw256_gpu_hash_80<<<grid, block>>>(threads, startNonce, (uint64_t*)d_outputHash, d_midstate[thr_id], swap);
}
__host__
void bmw256_setBlock_80(int thr_id, void *pdata)
{
uint64_t PaddedMessage[16];
memcpy(PaddedMessage, pdata, 80);
memset(&PaddedMessage[10], 0, 48);
CUDA_SAFE_CALL(cudaMemcpyToSymbol(c_PaddedMessage80, PaddedMessage, sizeof(PaddedMessage), 0, cudaMemcpyHostToDevice));
sph_bmw256_context ctx;
sph_bmw256_init(&ctx);
sph_bmw256(&ctx, (void*) PaddedMessage, 80);
CUDA_SAFE_CALL(cudaMemcpy(d_midstate[thr_id], &ctx, sizeof(sph_bmw256_context), cudaMemcpyHostToDevice));
}
__host__
void bmw256_midstate_init(int thr_id, uint32_t threads)
{
cudaMalloc(&d_midstate[thr_id], sizeof(sph_bmw256_context));
}
__host__
void bmw256_midstate_free(int thr_id)
{
cudaFree(d_midstate[thr_id]);
}