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.
1726 lines
50 KiB
1726 lines
50 KiB
// |
|
// Experimental Kernel for Kepler (Compute 3.5) devices |
|
// code submitted by nVidia performance engineer Alexey Panteleev |
|
// with modifications by Christian Buchner |
|
// |
|
// for Compute 3.5 |
|
// NOTE: compile this .cu module for compute_35,sm_35 with --maxrregcount=80 |
|
// for Compute 3.0 |
|
// NOTE: compile this .cu module for compute_30,sm_30 with --maxrregcount=63 |
|
// |
|
|
|
#include <map> |
|
|
|
#include "cuda_runtime.h" |
|
#include "miner.h" |
|
|
|
#include "salsa_kernel.h" |
|
#include "nv_kernel2.h" |
|
|
|
#define THREADS_PER_WU 1 // single thread per hash |
|
|
|
#if __CUDA_ARCH__ < 350 |
|
// Kepler (Compute 3.0) |
|
#define __ldg(x) (*(x)) |
|
#endif |
|
|
|
// grab lane ID |
|
static __device__ __inline__ unsigned int __laneId() { unsigned int laneId; asm( "mov.u32 %0, %%laneid;" : "=r"( laneId ) ); return laneId; } |
|
|
|
// forward references |
|
template <int ALGO> __global__ void nv2_scrypt_core_kernelA(uint32_t *g_idata, int begin, int end); |
|
template <int ALGO> __global__ void nv2_scrypt_core_kernelB(uint32_t *g_odata, int begin, int end); |
|
template <int ALGO> __global__ void nv2_scrypt_core_kernelA_LG(uint32_t *g_idata, int begin, int end, unsigned int LOOKUP_GAP); |
|
template <int ALGO> __global__ void nv2_scrypt_core_kernelB_LG(uint32_t *g_odata, int begin, int end, unsigned int LOOKUP_GAP); |
|
|
|
// scratchbuf constants (pointers to scratch buffer for each work unit) |
|
__constant__ uint32_t* c_V[TOTAL_WARP_LIMIT]; |
|
|
|
// iteration count N |
|
__constant__ uint32_t c_N; |
|
__constant__ uint32_t c_N_1; // N - 1 |
|
__constant__ uint32_t c_spacing; // (N+LOOKUP_GAP-1)/LOOKUP_GAP |
|
|
|
|
|
NV2Kernel::NV2Kernel() : KernelInterface() |
|
{ |
|
} |
|
|
|
void NV2Kernel::set_scratchbuf_constants(int MAXWARPS, uint32_t** h_V) |
|
{ |
|
checkCudaErrors(cudaMemcpyToSymbol(c_V, h_V, MAXWARPS*sizeof(uint32_t*), 0, cudaMemcpyHostToDevice)); |
|
} |
|
|
|
bool NV2Kernel::run_kernel(dim3 grid, dim3 threads, int WARPS_PER_BLOCK, int thr_id, cudaStream_t stream, uint32_t* d_idata, uint32_t* d_odata, unsigned int N, unsigned int LOOKUP_GAP, bool interactive, bool benchmark, int texture_cache) |
|
{ |
|
bool success = true; |
|
|
|
// make some constants available to kernel, update only initially and when changing |
|
static uint32_t prev_N[MAX_GPUS] = { 0 }; |
|
|
|
if (N != prev_N[thr_id]) { |
|
uint32_t h_N = N; |
|
uint32_t h_N_1 = N-1; |
|
uint32_t h_spacing = (N+LOOKUP_GAP-1)/LOOKUP_GAP; |
|
|
|
cudaMemcpyToSymbolAsync(c_N, &h_N, sizeof(uint32_t), 0, cudaMemcpyHostToDevice, stream); |
|
cudaMemcpyToSymbolAsync(c_N_1, &h_N_1, sizeof(uint32_t), 0, cudaMemcpyHostToDevice, stream); |
|
cudaMemcpyToSymbolAsync(c_spacing, &h_spacing, sizeof(uint32_t), 0, cudaMemcpyHostToDevice, stream); |
|
|
|
prev_N[thr_id] = N; |
|
} |
|
|
|
// First phase: Sequential writes to scratchpad. |
|
const int batch = device_batchsize[thr_id]; |
|
unsigned int pos = 0; |
|
|
|
do |
|
{ |
|
if (LOOKUP_GAP == 1) { |
|
if (IS_SCRYPT()) nv2_scrypt_core_kernelA<A_SCRYPT> <<< grid, threads, 0, stream >>>(d_idata, pos, min(pos+batch, N)); |
|
if (IS_SCRYPT_JANE()) nv2_scrypt_core_kernelA<A_SCRYPT_JANE><<< grid, threads, 0, stream >>>(d_idata, pos, min(pos+batch, N)); |
|
} else { |
|
if (IS_SCRYPT()) nv2_scrypt_core_kernelA_LG<A_SCRYPT> <<< grid, threads, 0, stream >>>(d_idata, pos, min(pos+batch, N), LOOKUP_GAP); |
|
if (IS_SCRYPT_JANE()) nv2_scrypt_core_kernelA_LG<A_SCRYPT_JANE><<< grid, threads, 0, stream >>>(d_idata, pos, min(pos+batch, N), LOOKUP_GAP); |
|
} |
|
pos += batch; |
|
} while (pos < N); |
|
|
|
// Second phase: Random read access from scratchpad. |
|
pos = 0; |
|
do |
|
{ |
|
if (LOOKUP_GAP == 1) { |
|
if (IS_SCRYPT()) nv2_scrypt_core_kernelB<A_SCRYPT > <<< grid, threads, 0, stream >>>(d_odata, pos, min(pos+batch, N)); |
|
if (IS_SCRYPT_JANE()) nv2_scrypt_core_kernelB<A_SCRYPT_JANE> <<< grid, threads, 0, stream >>>(d_odata, pos, min(pos+batch, N)); |
|
} else { |
|
if (IS_SCRYPT()) nv2_scrypt_core_kernelB_LG<A_SCRYPT > <<< grid, threads, 0, stream >>>(d_odata, pos, min(pos+batch, N), LOOKUP_GAP); |
|
if (IS_SCRYPT_JANE()) nv2_scrypt_core_kernelB_LG<A_SCRYPT_JANE> <<< grid, threads, 0, stream >>>(d_odata, pos, min(pos+batch, N), LOOKUP_GAP); |
|
} |
|
|
|
pos += batch; |
|
} while (pos < N); |
|
|
|
return success; |
|
} |
|
|
|
static __device__ uint4& operator^=(uint4& left, const uint4& right) |
|
{ |
|
left.x ^= right.x; |
|
left.y ^= right.y; |
|
left.z ^= right.z; |
|
left.w ^= right.w; |
|
return left; |
|
} |
|
|
|
__device__ __forceinline__ uint4 __shfl(const uint4 val, unsigned int lane, unsigned int width) |
|
{ |
|
return make_uint4( |
|
(unsigned int)__shfl((int)val.x, lane, width), |
|
(unsigned int)__shfl((int)val.y, lane, width), |
|
(unsigned int)__shfl((int)val.z, lane, width), |
|
(unsigned int)__shfl((int)val.w, lane, width)); |
|
} |
|
|
|
__device__ __forceinline__ void __transposed_write_BC(uint4 (&B)[4], uint4 (&C)[4], uint4 *D, int spacing) |
|
{ |
|
unsigned int laneId = __laneId(); |
|
|
|
unsigned int lane8 = laneId%8; |
|
unsigned int tile = laneId/8; |
|
|
|
uint4 T1[8], T2[8]; |
|
|
|
/* Source matrix, A-H are threads, 0-7 are data items, thread A is marked with `*`: |
|
|
|
*A0 B0 C0 D0 E0 F0 G0 H0 |
|
*A1 B1 C1 D1 E1 F1 G1 H1 |
|
*A2 B2 C2 D2 E2 F2 G2 H2 |
|
*A3 B3 C3 D3 E3 F3 G3 H3 |
|
*A4 B4 C4 D4 E4 F4 G4 H4 |
|
*A5 B5 C5 D5 E5 F5 G5 H5 |
|
*A6 B6 C6 D6 E6 F6 G6 H6 |
|
*A7 B7 C7 D7 E7 F7 G7 H7 |
|
*/ |
|
|
|
// rotate rows |
|
T1[0] = B[0]; |
|
T1[1] = __shfl(B[1], lane8 + 7, 8); |
|
T1[2] = __shfl(B[2], lane8 + 6, 8); |
|
T1[3] = __shfl(B[3], lane8 + 5, 8); |
|
T1[4] = __shfl(C[0], lane8 + 4, 8); |
|
T1[5] = __shfl(C[1], lane8 + 3, 8); |
|
T1[6] = __shfl(C[2], lane8 + 2, 8); |
|
T1[7] = __shfl(C[3], lane8 + 1, 8); |
|
|
|
/* Matrix after row rotates: |
|
|
|
*A0 B0 C0 D0 E0 F0 G0 H0 |
|
H1 *A1 B1 C1 D1 E1 F1 G1 |
|
G2 H2 *A2 B2 C2 D2 E2 F2 |
|
F3 G3 H3 *A3 B3 C3 D3 E3 |
|
E4 F4 G4 H4 *A4 B4 C4 D4 |
|
D5 E5 F5 G5 H5 *A5 B5 C5 |
|
C6 D6 E6 F6 G6 H6 *A6 B6 |
|
B7 C7 D7 E7 F7 G7 H7 *A7 |
|
*/ |
|
|
|
// rotate columns up using a barrel shifter simulation |
|
// column X is rotated up by (X+1) items |
|
#pragma unroll 8 |
|
for(int n = 0; n < 8; n++) T2[n] = ((lane8+1) & 1) ? T1[(n+1) % 8] : T1[n]; |
|
#pragma unroll 8 |
|
for(int n = 0; n < 8; n++) T1[n] = ((lane8+1) & 2) ? T2[(n+2) % 8] : T2[n]; |
|
#pragma unroll 8 |
|
for(int n = 0; n < 8; n++) T2[n] = ((lane8+1) & 4) ? T1[(n+4) % 8] : T1[n]; |
|
|
|
/* Matrix after column rotates: |
|
|
|
H1 H2 H3 H4 H5 H6 H7 H0 |
|
G2 G3 G4 G5 G6 G7 G0 G1 |
|
F3 F4 F5 F6 F7 F0 F1 F2 |
|
E4 E5 E6 E7 E0 E1 E2 E3 |
|
D5 D6 D7 D0 D1 D2 D3 D4 |
|
C6 C7 C0 C1 C2 C3 C4 C5 |
|
B7 B0 B1 B2 B3 B4 B5 B6 |
|
*A0 *A1 *A2 *A3 *A4 *A5 *A6 *A7 |
|
*/ |
|
|
|
// rotate rows again using address math and write to D, in reverse row order |
|
D[spacing*2*(32*tile )+ lane8 ] = T2[7]; |
|
D[spacing*2*(32*tile+4 )+(lane8+7)%8] = T2[6]; |
|
D[spacing*2*(32*tile+8 )+(lane8+6)%8] = T2[5]; |
|
D[spacing*2*(32*tile+12)+(lane8+5)%8] = T2[4]; |
|
D[spacing*2*(32*tile+16)+(lane8+4)%8] = T2[3]; |
|
D[spacing*2*(32*tile+20)+(lane8+3)%8] = T2[2]; |
|
D[spacing*2*(32*tile+24)+(lane8+2)%8] = T2[1]; |
|
D[spacing*2*(32*tile+28)+(lane8+1)%8] = T2[0]; |
|
} |
|
|
|
__device__ __forceinline__ void __transposed_read_BC(const uint4 *S, uint4 (&B)[4], uint4 (&C)[4], int spacing, int row) |
|
{ |
|
unsigned int laneId = __laneId(); |
|
|
|
unsigned int lane8 = laneId%8; |
|
unsigned int tile = laneId/8; |
|
|
|
// Perform the same transposition as in __transposed_write_BC, but in reverse order. |
|
// See the illustrations in comments for __transposed_write_BC. |
|
|
|
// read and rotate rows, in reverse row order |
|
uint4 T1[8], T2[8]; |
|
T1[7] = __ldg(&S[(spacing*2*(32*tile ) + lane8 + 8*__shfl(row, 0, 8))]); |
|
T1[6] = __ldg(&S[(spacing*2*(32*tile+4 ) + (lane8+7)%8 + 8*__shfl(row, 1, 8))]); |
|
T1[5] = __ldg(&S[(spacing*2*(32*tile+8 ) + (lane8+6)%8 + 8*__shfl(row, 2, 8))]); |
|
T1[4] = __ldg(&S[(spacing*2*(32*tile+12) + (lane8+5)%8 + 8*__shfl(row, 3, 8))]); |
|
T1[3] = __ldg(&S[(spacing*2*(32*tile+16) + (lane8+4)%8 + 8*__shfl(row, 4, 8))]); |
|
T1[2] = __ldg(&S[(spacing*2*(32*tile+20) + (lane8+3)%8 + 8*__shfl(row, 5, 8))]); |
|
T1[1] = __ldg(&S[(spacing*2*(32*tile+24) + (lane8+2)%8 + 8*__shfl(row, 6, 8))]); |
|
T1[0] = __ldg(&S[(spacing*2*(32*tile+28) + (lane8+1)%8 + 8*__shfl(row, 7, 8))]); |
|
|
|
// rotate columns down using a barrel shifter simulation |
|
// column X is rotated down by (X+1) items, or up by (8-(X+1)) = (7-X) items |
|
#pragma unroll 8 |
|
for(int n = 0; n < 8; n++) T2[n] = ((7-lane8) & 1) ? T1[(n+1) % 8] : T1[n]; |
|
#pragma unroll 8 |
|
for(int n = 0; n < 8; n++) T1[n] = ((7-lane8) & 2) ? T2[(n+2) % 8] : T2[n]; |
|
#pragma unroll 8 |
|
for(int n = 0; n < 8; n++) T2[n] = ((7-lane8) & 4) ? T1[(n+4) % 8] : T1[n]; |
|
|
|
// rotate rows |
|
B[0] = T2[0]; |
|
B[1] = __shfl(T2[1], lane8 + 1, 8); |
|
B[2] = __shfl(T2[2], lane8 + 2, 8); |
|
B[3] = __shfl(T2[3], lane8 + 3, 8); |
|
C[0] = __shfl(T2[4], lane8 + 4, 8); |
|
C[1] = __shfl(T2[5], lane8 + 5, 8); |
|
C[2] = __shfl(T2[6], lane8 + 6, 8); |
|
C[3] = __shfl(T2[7], lane8 + 7, 8); |
|
|
|
} |
|
|
|
__device__ __forceinline__ void __transposed_xor_BC(const uint4 *S, uint4 (&B)[4], uint4 (&C)[4], int spacing, int row) |
|
{ |
|
uint4 BT[4], CT[4]; |
|
__transposed_read_BC(S, BT, CT, spacing, row); |
|
|
|
#pragma unroll 4 |
|
for(int n = 0; n < 4; n++) |
|
{ |
|
B[n] ^= BT[n]; |
|
C[n] ^= CT[n]; |
|
} |
|
} |
|
|
|
#if __CUDA_ARCH__ < 350 |
|
// Kepler (Compute 3.0) |
|
#define ROTL(a, b) ((a)<<(b))|((a)>>(32-(b))) |
|
#else |
|
// Kepler (Compute 3.5) |
|
#define ROTL(a, b) __funnelshift_l( a, a, b ); |
|
#endif |
|
|
|
|
|
|
|
#if 0 |
|
|
|
#define QUARTER(a,b,c,d) \ |
|
a += b; d ^= a; d = ROTL(d,16); \ |
|
c += d; b ^= c; b = ROTL(b,12); \ |
|
a += b; d ^= a; d = ROTL(d,8); \ |
|
c += d; b ^= c; b = ROTL(b,7); |
|
|
|
static __device__ void xor_chacha8(uint4 *B, uint4 *C) |
|
{ |
|
uint32_t x[16]; |
|
x[0]=(B[0].x ^= C[0].x); |
|
x[1]=(B[0].y ^= C[0].y); |
|
x[2]=(B[0].z ^= C[0].z); |
|
x[3]=(B[0].w ^= C[0].w); |
|
x[4]=(B[1].x ^= C[1].x); |
|
x[5]=(B[1].y ^= C[1].y); |
|
x[6]=(B[1].z ^= C[1].z); |
|
x[7]=(B[1].w ^= C[1].w); |
|
x[8]=(B[2].x ^= C[2].x); |
|
x[9]=(B[2].y ^= C[2].y); |
|
x[10]=(B[2].z ^= C[2].z); |
|
x[11]=(B[2].w ^= C[2].w); |
|
x[12]=(B[3].x ^= C[3].x); |
|
x[13]=(B[3].y ^= C[3].y); |
|
x[14]=(B[3].z ^= C[3].z); |
|
x[15]=(B[3].w ^= C[3].w); |
|
|
|
/* Operate on columns. */ |
|
QUARTER( x[0], x[4], x[ 8], x[12] ) |
|
QUARTER( x[1], x[5], x[ 9], x[13] ) |
|
QUARTER( x[2], x[6], x[10], x[14] ) |
|
QUARTER( x[3], x[7], x[11], x[15] ) |
|
|
|
/* Operate on diagonals */ |
|
QUARTER( x[0], x[5], x[10], x[15] ) |
|
QUARTER( x[1], x[6], x[11], x[12] ) |
|
QUARTER( x[2], x[7], x[ 8], x[13] ) |
|
QUARTER( x[3], x[4], x[ 9], x[14] ) |
|
|
|
/* Operate on columns. */ |
|
QUARTER( x[0], x[4], x[ 8], x[12] ) |
|
QUARTER( x[1], x[5], x[ 9], x[13] ) |
|
QUARTER( x[2], x[6], x[10], x[14] ) |
|
QUARTER( x[3], x[7], x[11], x[15] ) |
|
|
|
/* Operate on diagonals */ |
|
QUARTER( x[0], x[5], x[10], x[15] ) |
|
QUARTER( x[1], x[6], x[11], x[12] ) |
|
QUARTER( x[2], x[7], x[ 8], x[13] ) |
|
QUARTER( x[3], x[4], x[ 9], x[14] ) |
|
|
|
/* Operate on columns. */ |
|
QUARTER( x[0], x[4], x[ 8], x[12] ) |
|
QUARTER( x[1], x[5], x[ 9], x[13] ) |
|
QUARTER( x[2], x[6], x[10], x[14] ) |
|
QUARTER( x[3], x[7], x[11], x[15] ) |
|
|
|
/* Operate on diagonals */ |
|
QUARTER( x[0], x[5], x[10], x[15] ) |
|
QUARTER( x[1], x[6], x[11], x[12] ) |
|
QUARTER( x[2], x[7], x[ 8], x[13] ) |
|
QUARTER( x[3], x[4], x[ 9], x[14] ) |
|
|
|
/* Operate on columns. */ |
|
QUARTER( x[0], x[4], x[ 8], x[12] ) |
|
QUARTER( x[1], x[5], x[ 9], x[13] ) |
|
QUARTER( x[2], x[6], x[10], x[14] ) |
|
QUARTER( x[3], x[7], x[11], x[15] ) |
|
|
|
/* Operate on diagonals */ |
|
QUARTER( x[0], x[5], x[10], x[15] ) |
|
QUARTER( x[1], x[6], x[11], x[12] ) |
|
QUARTER( x[2], x[7], x[ 8], x[13] ) |
|
QUARTER( x[3], x[4], x[ 9], x[14] ) |
|
|
|
B[0].x += x[0]; B[0].y += x[1]; B[0].z += x[2]; B[0].w += x[3]; B[1].x += x[4]; B[1].y += x[5]; B[1].z += x[6]; B[1].w += x[7]; |
|
B[2].x += x[8]; B[2].y += x[9]; B[2].z += x[10]; B[2].w += x[11]; B[3].x += x[12]; B[3].y += x[13]; B[3].z += x[14]; B[3].w += x[15]; |
|
} |
|
|
|
#else |
|
|
|
#define ADD4(d1,d2,d3,d4,s1,s2,s3,s4) \ |
|
d1 += s1; d2 += s2; d3 += s3; d4 += s4; |
|
|
|
#define XOR4(d1,d2,d3,d4,s1,s2,s3,s4) \ |
|
d1 ^= s1; d2 ^= s2; d3 ^= s3; d4 ^= s4; |
|
|
|
#define ROTL4(d1,d2,d3,d4,amt) \ |
|
d1 = ROTL(d1, amt); d2 = ROTL(d2, amt); d3 = ROTL(d3, amt); d4 = ROTL(d4, amt); |
|
|
|
#define QROUND(a1,a2,a3,a4, b1,b2,b3,b4, c1,c2,c3,c4, amt) \ |
|
ADD4 (a1,a2,a3,a4, c1,c2,c3,c4) \ |
|
XOR4 (b1,b2,b3,b4, a1,a2,a3,a4) \ |
|
ROTL4(b1,b2,b3,b4, amt) |
|
|
|
static __device__ void xor_chacha8(uint4 *B, uint4 *C) |
|
{ |
|
uint32_t x[16]; |
|
x[0]=(B[0].x ^= C[0].x); |
|
x[1]=(B[0].y ^= C[0].y); |
|
x[2]=(B[0].z ^= C[0].z); |
|
x[3]=(B[0].w ^= C[0].w); |
|
x[4]=(B[1].x ^= C[1].x); |
|
x[5]=(B[1].y ^= C[1].y); |
|
x[6]=(B[1].z ^= C[1].z); |
|
x[7]=(B[1].w ^= C[1].w); |
|
x[8]=(B[2].x ^= C[2].x); |
|
x[9]=(B[2].y ^= C[2].y); |
|
x[10]=(B[2].z ^= C[2].z); |
|
x[11]=(B[2].w ^= C[2].w); |
|
x[12]=(B[3].x ^= C[3].x); |
|
x[13]=(B[3].y ^= C[3].y); |
|
x[14]=(B[3].z ^= C[3].z); |
|
x[15]=(B[3].w ^= C[3].w); |
|
|
|
/* Operate on columns. */ |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 16); |
|
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 12); |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 8); |
|
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 7); |
|
|
|
/* Operate on diagonals */ |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 16); |
|
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 12); |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 8); |
|
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 7); |
|
|
|
/* Operate on columns. */ |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 16); |
|
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 12); |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 8); |
|
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 7); |
|
|
|
/* Operate on diagonals */ |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 16); |
|
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 12); |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 8); |
|
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 7); |
|
|
|
/* Operate on columns. */ |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 16); |
|
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 12); |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 8); |
|
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 7); |
|
|
|
/* Operate on diagonals */ |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 16); |
|
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 12); |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 8); |
|
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 7); |
|
|
|
/* Operate on columns. */ |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 16); |
|
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 12); |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[12],x[13],x[14],x[15], x[ 4],x[ 5],x[ 6],x[ 7], 8); |
|
QROUND(x[ 8],x[ 9],x[10],x[11], x[ 4],x[ 5],x[ 6],x[ 7], x[12],x[13],x[14],x[15], 7); |
|
|
|
/* Operate on diagonals */ |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 16); |
|
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 12); |
|
QROUND(x[ 0],x[ 1],x[ 2],x[ 3], x[15],x[12],x[13],x[14], x[ 5],x[ 6],x[ 7],x[ 4], 8); |
|
QROUND(x[10],x[11],x[ 8],x[ 9], x[ 5],x[ 6],x[ 7],x[ 4], x[15],x[12],x[13],x[14], 7); |
|
|
|
B[0].x += x[0]; B[0].y += x[1]; B[0].z += x[2]; B[0].w += x[3]; B[1].x += x[4]; B[1].y += x[5]; B[1].z += x[6]; B[1].w += x[7]; |
|
B[2].x += x[8]; B[2].y += x[9]; B[2].z += x[10]; B[2].w += x[11]; B[3].x += x[12]; B[3].y += x[13]; B[3].z += x[14]; B[3].w += x[15]; |
|
} |
|
|
|
#endif |
|
|
|
|
|
#define ROTL7(a0,a1,a2,a3,a00,a10,a20,a30){\ |
|
a0^=ROTL(a00, 7); a1^=ROTL(a10, 7); a2^=ROTL(a20, 7); a3^=ROTL(a30, 7);\ |
|
};\ |
|
|
|
#define ROTL9(a0,a1,a2,a3,a00,a10,a20,a30){\ |
|
a0^=ROTL(a00, 9); a1^=ROTL(a10, 9); a2^=ROTL(a20, 9); a3^=ROTL(a30, 9);\ |
|
};\ |
|
|
|
#define ROTL13(a0,a1,a2,a3,a00,a10,a20,a30){\ |
|
a0^=ROTL(a00, 13); a1^=ROTL(a10, 13); a2^=ROTL(a20, 13); a3^=ROTL(a30, 13);\ |
|
};\ |
|
|
|
#define ROTL18(a0,a1,a2,a3,a00,a10,a20,a30){\ |
|
a0^=ROTL(a00, 18); a1^=ROTL(a10, 18); a2^=ROTL(a20, 18); a3^=ROTL(a30, 18);\ |
|
};\ |
|
|
|
static __device__ void xor_salsa8(uint4 *B, uint4 *C) |
|
{ |
|
uint32_t x[16]; |
|
x[0]=(B[0].x ^= C[0].x); |
|
x[1]=(B[0].y ^= C[0].y); |
|
x[2]=(B[0].z ^= C[0].z); |
|
x[3]=(B[0].w ^= C[0].w); |
|
x[4]=(B[1].x ^= C[1].x); |
|
x[5]=(B[1].y ^= C[1].y); |
|
x[6]=(B[1].z ^= C[1].z); |
|
x[7]=(B[1].w ^= C[1].w); |
|
x[8]=(B[2].x ^= C[2].x); |
|
x[9]=(B[2].y ^= C[2].y); |
|
x[10]=(B[2].z ^= C[2].z); |
|
x[11]=(B[2].w ^= C[2].w); |
|
x[12]=(B[3].x ^= C[3].x); |
|
x[13]=(B[3].y ^= C[3].y); |
|
x[14]=(B[3].z ^= C[3].z); |
|
x[15]=(B[3].w ^= C[3].w); |
|
|
|
/* Operate on columns. */ |
|
ROTL7(x[4],x[9],x[14],x[3],x[0]+x[12],x[1]+x[5],x[6]+x[10],x[11]+x[15]); |
|
ROTL9(x[8],x[13],x[2],x[7],x[0]+x[4],x[5]+x[9],x[10]+x[14],x[3]+x[15]); |
|
ROTL13(x[12],x[1],x[6],x[11],x[4]+x[8],x[9]+x[13],x[2]+x[14],x[3]+x[7]); |
|
ROTL18(x[0],x[5],x[10],x[15],x[8]+x[12],x[1]+x[13],x[2]+x[6],x[7]+x[11]); |
|
|
|
/* Operate on rows. */ |
|
ROTL7(x[1],x[6],x[11],x[12],x[0]+x[3],x[4]+x[5],x[9]+x[10],x[14]+x[15]); |
|
ROTL9(x[2],x[7],x[8],x[13],x[0]+x[1],x[5]+x[6],x[10]+x[11],x[12]+x[15]); |
|
ROTL13(x[3],x[4],x[9],x[14],x[1]+x[2],x[6]+x[7],x[8]+x[11],x[12]+x[13]); |
|
ROTL18(x[0],x[5],x[10],x[15],x[2]+x[3],x[4]+x[7],x[8]+x[9],x[13]+x[14]); |
|
|
|
/* Operate on columns. */ |
|
ROTL7(x[4],x[9],x[14],x[3],x[0]+x[12],x[1]+x[5],x[6]+x[10],x[11]+x[15]); |
|
ROTL9(x[8],x[13],x[2],x[7],x[0]+x[4],x[5]+x[9],x[10]+x[14],x[3]+x[15]); |
|
ROTL13(x[12],x[1],x[6],x[11],x[4]+x[8],x[9]+x[13],x[2]+x[14],x[3]+x[7]); |
|
ROTL18(x[0],x[5],x[10],x[15],x[8]+x[12],x[1]+x[13],x[2]+x[6],x[7]+x[11]); |
|
|
|
/* Operate on rows. */ |
|
ROTL7(x[1],x[6],x[11],x[12],x[0]+x[3],x[4]+x[5],x[9]+x[10],x[14]+x[15]); |
|
ROTL9(x[2],x[7],x[8],x[13],x[0]+x[1],x[5]+x[6],x[10]+x[11],x[12]+x[15]); |
|
ROTL13(x[3],x[4],x[9],x[14],x[1]+x[2],x[6]+x[7],x[8]+x[11],x[12]+x[13]); |
|
ROTL18(x[0],x[5],x[10],x[15],x[2]+x[3],x[4]+x[7],x[8]+x[9],x[13]+x[14]); |
|
|
|
/* Operate on columns. */ |
|
ROTL7(x[4],x[9],x[14],x[3],x[0]+x[12],x[1]+x[5],x[6]+x[10],x[11]+x[15]); |
|
ROTL9(x[8],x[13],x[2],x[7],x[0]+x[4],x[5]+x[9],x[10]+x[14],x[3]+x[15]); |
|
ROTL13(x[12],x[1],x[6],x[11],x[4]+x[8],x[9]+x[13],x[2]+x[14],x[3]+x[7]); |
|
ROTL18(x[0],x[5],x[10],x[15],x[8]+x[12],x[1]+x[13],x[2]+x[6],x[7]+x[11]); |
|
|
|
/* Operate on rows. */ |
|
ROTL7(x[1],x[6],x[11],x[12],x[0]+x[3],x[4]+x[5],x[9]+x[10],x[14]+x[15]); |
|
ROTL9(x[2],x[7],x[8],x[13],x[0]+x[1],x[5]+x[6],x[10]+x[11],x[12]+x[15]); |
|
ROTL13(x[3],x[4],x[9],x[14],x[1]+x[2],x[6]+x[7],x[8]+x[11],x[12]+x[13]); |
|
ROTL18(x[0],x[5],x[10],x[15],x[2]+x[3],x[4]+x[7],x[8]+x[9],x[13]+x[14]); |
|
|
|
/* Operate on columns. */ |
|
ROTL7(x[4],x[9],x[14],x[3],x[0]+x[12],x[1]+x[5],x[6]+x[10],x[11]+x[15]); |
|
ROTL9(x[8],x[13],x[2],x[7],x[0]+x[4],x[5]+x[9],x[10]+x[14],x[3]+x[15]); |
|
ROTL13(x[12],x[1],x[6],x[11],x[4]+x[8],x[9]+x[13],x[2]+x[14],x[3]+x[7]); |
|
ROTL18(x[0],x[5],x[10],x[15],x[8]+x[12],x[1]+x[13],x[2]+x[6],x[7]+x[11]); |
|
|
|
/* Operate on rows. */ |
|
ROTL7(x[1],x[6],x[11],x[12],x[0]+x[3],x[4]+x[5],x[9]+x[10],x[14]+x[15]); |
|
ROTL9(x[2],x[7],x[8],x[13],x[0]+x[1],x[5]+x[6],x[10]+x[11],x[12]+x[15]); |
|
ROTL13(x[3],x[4],x[9],x[14],x[1]+x[2],x[6]+x[7],x[8]+x[11],x[12]+x[13]); |
|
ROTL18(x[0],x[5],x[10],x[15],x[2]+x[3],x[4]+x[7],x[8]+x[9],x[13]+x[14]); |
|
|
|
B[0].x += x[0]; B[0].y += x[1]; B[0].z += x[2]; B[0].w += x[3]; B[1].x += x[4]; B[1].y += x[5]; B[1].z += x[6]; B[1].w += x[7]; |
|
B[2].x += x[8]; B[2].y += x[9]; B[2].z += x[10]; B[2].w += x[11]; B[3].x += x[12]; B[3].y += x[13]; B[3].z += x[14]; B[3].w += x[15]; |
|
} |
|
|
|
|
|
template <int ALGO> static __device__ void block_mixer(uint4 *B, uint4 *C) |
|
{ |
|
switch (ALGO) |
|
{ |
|
case A_SCRYPT: xor_salsa8(B, C); break; |
|
case A_SCRYPT_JANE: xor_chacha8(B, C); break; |
|
} |
|
} |
|
|
|
//////////////////////////////////////////////////////////////////////////////// |
|
//! Experimental Scrypt core kernel for Titan devices. |
|
//! @param g_idata input data in global memory |
|
//! @param g_odata output data in global memory |
|
//////////////////////////////////////////////////////////////////////////////// |
|
template <int ALGO> __global__ void nv2_scrypt_core_kernelA(uint32_t *g_idata, int begin, int end) |
|
{ |
|
int offset = blockIdx.x * blockDim.x + threadIdx.x / warpSize * warpSize; |
|
g_idata += 32 * offset; |
|
uint32_t * V = c_V[offset / warpSize]; |
|
uint4 B[4], C[4]; |
|
int i = begin; |
|
|
|
if(i == 0) { |
|
__transposed_read_BC((uint4*)g_idata, B, C, 1, 0); |
|
__transposed_write_BC(B, C, (uint4*)V, c_N); |
|
++i; |
|
} else |
|
__transposed_read_BC((uint4*)(V + (i-1)*32), B, C, c_N, 0); |
|
|
|
while(i < end) { |
|
block_mixer<ALGO>(B, C); block_mixer<ALGO>(C, B); |
|
__transposed_write_BC(B, C, (uint4*)(V + i*32), c_N); |
|
++i; |
|
} |
|
} |
|
|
|
template <int ALGO> __global__ void nv2_scrypt_core_kernelA_LG(uint32_t *g_idata, int begin, int end, unsigned int LOOKUP_GAP) |
|
{ |
|
int offset = blockIdx.x * blockDim.x + threadIdx.x / warpSize * warpSize; |
|
g_idata += 32 * offset; |
|
uint32_t * V = c_V[offset / warpSize]; |
|
uint4 B[4], C[4]; |
|
int i = begin; |
|
|
|
if(i == 0) { |
|
__transposed_read_BC((uint4*)g_idata, B, C, 1, 0); |
|
__transposed_write_BC(B, C, (uint4*)V, c_spacing); |
|
++i; |
|
} else { |
|
int pos = (i-1)/LOOKUP_GAP, loop = (i-1)-pos*LOOKUP_GAP; |
|
__transposed_read_BC((uint4*)(V + pos*32), B, C, c_spacing, 0); |
|
while(loop--) { block_mixer<ALGO>(B, C); block_mixer<ALGO>(C, B); } |
|
} |
|
|
|
while(i < end) { |
|
block_mixer<ALGO>(B, C); block_mixer<ALGO>(C, B); |
|
if (i % LOOKUP_GAP == 0) |
|
__transposed_write_BC(B, C, (uint4*)(V + (i/LOOKUP_GAP)*32), c_spacing); |
|
++i; |
|
} |
|
} |
|
|
|
template <int ALGO> __global__ void nv2_scrypt_core_kernelB(uint32_t *g_odata, int begin, int end) |
|
{ |
|
int offset = blockIdx.x * blockDim.x + threadIdx.x / warpSize * warpSize; |
|
g_odata += 32 * offset; |
|
uint32_t * V = c_V[offset / warpSize]; |
|
uint4 B[4], C[4]; |
|
|
|
if(begin == 0) { |
|
__transposed_read_BC((uint4*)V, B, C, c_N, c_N_1); |
|
block_mixer<ALGO>(B, C); block_mixer<ALGO>(C, B); |
|
} else |
|
__transposed_read_BC((uint4*)g_odata, B, C, 1, 0); |
|
|
|
for (int i = begin; i < end; i++) { |
|
int slot = C[0].x & c_N_1; |
|
__transposed_xor_BC((uint4*)(V), B, C, c_N, slot); |
|
block_mixer<ALGO>(B, C); block_mixer<ALGO>(C, B); |
|
} |
|
|
|
__transposed_write_BC(B, C, (uint4*)(g_odata), 1); |
|
} |
|
|
|
template <int ALGO> __global__ void nv2_scrypt_core_kernelB_LG(uint32_t *g_odata, int begin, int end, unsigned int LOOKUP_GAP) |
|
{ |
|
int offset = blockIdx.x * blockDim.x + threadIdx.x / warpSize * warpSize; |
|
g_odata += 32 * offset; |
|
uint32_t * V = c_V[offset / warpSize]; |
|
uint4 B[4], C[4]; |
|
|
|
if(begin == 0) { |
|
int pos = c_N_1/LOOKUP_GAP, loop = 1 + (c_N_1-pos*LOOKUP_GAP); |
|
__transposed_read_BC((uint4*)V, B, C, c_spacing, pos); |
|
while(loop--) { block_mixer<ALGO>(B, C); block_mixer<ALGO>(C, B); } |
|
} else { |
|
__transposed_read_BC((uint4*)g_odata, B, C, 1, 0); |
|
} |
|
|
|
for (int i = begin; i < end; i++) { |
|
int slot = C[0].x & c_N_1; |
|
int pos = slot/LOOKUP_GAP, loop = slot-pos*LOOKUP_GAP; |
|
uint4 b[4], c[4]; __transposed_read_BC((uint4*)(V), b, c, c_spacing, pos); |
|
while(loop--) { block_mixer<ALGO>(b, c); block_mixer<ALGO>(c, b); } |
|
#pragma unroll 4 |
|
for(int n = 0; n < 4; n++) { B[n] ^= b[n]; C[n] ^= c[n]; } |
|
block_mixer<ALGO>(B, C); block_mixer<ALGO>(C, B); |
|
} |
|
|
|
__transposed_write_BC(B, C, (uint4*)(g_odata), 1); |
|
} |
|
|
|
|
|
// |
|
// Maxcoin related Keccak implementation (Keccak256) |
|
// |
|
|
|
// from salsa_kernel.cu |
|
extern std::map<int, int> context_blocks; |
|
extern std::map<int, int> context_wpb; |
|
extern std::map<int, KernelInterface *> context_kernel; |
|
extern std::map<int, cudaStream_t> context_streams[2]; |
|
extern std::map<int, uint32_t *> context_hash[2]; |
|
|
|
__constant__ uint64_t ptarget64[4]; |
|
|
|
// ROL macro replaced with the inline assembly code below to work around a performance issue |
|
//#define ROL(a, offset) ((((uint64_t)a) << ((offset) % 64)) ^ (((uint64_t)a) >> (64-((offset) % 64)))) |
|
__inline__ __device__ uint2 ROL(const uint2 a, const int offset) { |
|
uint2 result; |
|
if(offset >= 32) { |
|
asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(a.x), "r"(a.y), "r"(offset)); |
|
asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(a.y), "r"(a.x), "r"(offset)); |
|
} else { |
|
asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.x) : "r"(a.y), "r"(a.x), "r"(offset)); |
|
asm("shf.l.wrap.b32 %0, %1, %2, %3;" : "=r"(result.y) : "r"(a.x), "r"(a.y), "r"(offset)); |
|
} |
|
return result; |
|
} |
|
#define ROL_mult8(a, offset) ROL(a, offset) |
|
|
|
__inline__ __device__ uint64_t devectorize(uint2 v) { return __double_as_longlong(__hiloint2double(v.y, v.x)); } |
|
__inline__ __device__ uint2 vectorize(uint64_t v) { return make_uint2(__double2loint(__longlong_as_double(v)), __double2hiint(__longlong_as_double(v))); } |
|
__inline__ __device__ uint2 operator^ (uint2 a, uint2 b) { return make_uint2(a.x ^ b.x, a.y ^ b.y); } |
|
__inline__ __device__ uint2 operator& (uint2 a, uint2 b) { return make_uint2(a.x & b.x, a.y & b.y); } |
|
__inline__ __device__ uint2 operator| (uint2 a, uint2 b) { return make_uint2(a.x | b.x, a.y | b.y); } |
|
__inline__ __device__ uint2 operator~ (uint2 a) { return make_uint2(~a.x, ~a.y); } |
|
__inline__ __device__ void operator^= (uint2 &a, uint2 b) { a = a ^ b; } |
|
|
|
__constant__ uint64_t KeccakF_RoundConstants[24]; |
|
|
|
static uint64_t host_KeccakF_RoundConstants[24] = |
|
{ |
|
(uint64_t)0x0000000000000001ULL, |
|
(uint64_t)0x0000000000008082ULL, |
|
(uint64_t)0x800000000000808aULL, |
|
(uint64_t)0x8000000080008000ULL, |
|
(uint64_t)0x000000000000808bULL, |
|
(uint64_t)0x0000000080000001ULL, |
|
(uint64_t)0x8000000080008081ULL, |
|
(uint64_t)0x8000000000008009ULL, |
|
(uint64_t)0x000000000000008aULL, |
|
(uint64_t)0x0000000000000088ULL, |
|
(uint64_t)0x0000000080008009ULL, |
|
(uint64_t)0x000000008000000aULL, |
|
(uint64_t)0x000000008000808bULL, |
|
(uint64_t)0x800000000000008bULL, |
|
(uint64_t)0x8000000000008089ULL, |
|
(uint64_t)0x8000000000008003ULL, |
|
(uint64_t)0x8000000000008002ULL, |
|
(uint64_t)0x8000000000000080ULL, |
|
(uint64_t)0x000000000000800aULL, |
|
(uint64_t)0x800000008000000aULL, |
|
(uint64_t)0x8000000080008081ULL, |
|
(uint64_t)0x8000000000008080ULL, |
|
(uint64_t)0x0000000080000001ULL, |
|
(uint64_t)0x8000000080008008ULL |
|
}; |
|
|
|
__constant__ uint64_t pdata64[10]; |
|
|
|
static __device__ uint32_t cuda_swab32(uint32_t x) |
|
{ |
|
return (((x << 24) & 0xff000000u) | ((x << 8) & 0x00ff0000u) |
|
| ((x >> 8) & 0x0000ff00u) | ((x >> 24) & 0x000000ffu)); |
|
} |
|
|
|
// in this implementation the first and last iteration of the for() loop were explicitly |
|
// unrolled and redundant operations were removed (e.g. operations on zero inputs, and |
|
// computation of unnecessary outputs) |
|
__global__ void titan_crypto_hash( uint64_t *g_out, uint32_t nonce, uint32_t *g_good, bool validate ) |
|
{ |
|
uint2 Aba, Abe, Abi, Abo, Abu; |
|
uint2 Aga, Age, Agi, Ago, Agu; |
|
uint2 Aka, Ake, Aki, Ako, Aku; |
|
uint2 Ama, Ame, Ami, Amo, Amu; |
|
uint2 Asa, Ase, Asi, Aso, Asu; |
|
uint2 BCa, BCe, BCi, BCo, BCu; |
|
uint2 Da, De, Di, Do, Du; |
|
uint2 Eba, Ebe, Ebi, Ebo, Ebu; |
|
uint2 Ega, Ege, Egi, Ego, Egu; |
|
uint2 Eka, Eke, Eki, Eko, Eku; |
|
uint2 Ema, Eme, Emi, Emo, Emu; |
|
uint2 Esa, Ese, Esi, Eso, Esu; |
|
|
|
// embed unique nonce into source data stream in pdata[] |
|
Agu = vectorize((pdata64[9] & 0x00000000FFFFFFFFULL) | (((uint64_t)cuda_swab32(nonce + ((blockIdx.x * blockDim.x) + threadIdx.x))) << 32)); |
|
|
|
// prepareTheta |
|
BCa = vectorize(pdata64[0]^pdata64[5]^0x0000000000000001ULL); |
|
BCe = vectorize(pdata64[1]^pdata64[6]^0x8000000000000000ULL); |
|
BCi = vectorize(pdata64[2]^pdata64[7]); |
|
BCo = vectorize(pdata64[3]^pdata64[8]); |
|
BCu = vectorize(pdata64[4])^Agu; |
|
|
|
//thetaRhoPiChiIotaPrepareTheta(round , A, E) |
|
Da = BCu^ROL(BCe, 1); |
|
De = BCa^ROL(BCi, 1); |
|
Di = BCe^ROL(BCo, 1); |
|
Do = BCi^ROL(BCu, 1); |
|
Du = BCo^ROL(BCa, 1); |
|
|
|
Aba = vectorize(pdata64[0]) ^ Da; |
|
BCa = Aba; |
|
Age = vectorize(pdata64[6]) ^ De; |
|
BCe = ROL(Age, 44); |
|
Aki = Di; |
|
BCi = ROL(Aki, 43); |
|
Amo = Do; |
|
BCo = ROL(Amo, 21); |
|
Asu = Du; |
|
BCu = ROL(Asu, 14); |
|
Eba = BCa ^((~BCe)& BCi ); |
|
Eba ^= vectorize((uint64_t)KeccakF_RoundConstants[0]); |
|
Ebe = BCe ^((~BCi)& BCo ); |
|
Ebi = BCi ^((~BCo)& BCu ); |
|
Ebo = BCo ^((~BCu)& BCa ); |
|
Ebu = BCu ^((~BCa)& BCe ); |
|
|
|
Abo = vectorize(pdata64[3]) ^ Do; |
|
BCa = ROL(Abo, 28); |
|
Agu ^= Du; |
|
BCe = ROL(Agu, 20); |
|
Aka = vectorize(0x0000000000000001ULL) ^ Da; |
|
BCi = ROL(Aka, 3); |
|
Ame = vectorize(0x8000000000000000ULL) ^ De; |
|
BCo = ROL(Ame, 45); |
|
Asi = Di; |
|
BCu = ROL(Asi, 61); |
|
Ega = BCa ^((~BCe)& BCi ); |
|
Ege = BCe ^((~BCi)& BCo ); |
|
Egi = BCi ^((~BCo)& BCu ); |
|
Ego = BCo ^((~BCu)& BCa ); |
|
Egu = BCu ^((~BCa)& BCe ); |
|
|
|
Abe = vectorize(pdata64[1]) ^ De; |
|
BCa = ROL(Abe, 1); |
|
Agi = vectorize(pdata64[7]) ^ Di; |
|
BCe = ROL(Agi, 6); |
|
Ako = Do; |
|
BCi = ROL(Ako, 25); |
|
Amu = Du; |
|
BCo = ROL(Amu, 8); |
|
Asa = Da; |
|
BCu = ROL(Asa, 18); |
|
Eka = BCa ^((~BCe)& BCi ); |
|
Eke = BCe ^((~BCi)& BCo ); |
|
Eki = BCi ^((~BCo)& BCu ); |
|
Eko = BCo ^((~BCu)& BCa ); |
|
Eku = BCu ^((~BCa)& BCe ); |
|
|
|
Abu = vectorize(pdata64[4]) ^ Du; |
|
BCa = ROL(Abu, 27); |
|
Aga = vectorize(pdata64[5]) ^ Da; |
|
BCe = ROL(Aga, 36); |
|
Ake = De; |
|
BCi = ROL(Ake, 10); |
|
Ami = Di; |
|
BCo = ROL(Ami, 15); |
|
Aso = Do; |
|
BCu = ROL(Aso, 56); |
|
Ema = BCa ^((~BCe)& BCi ); |
|
Eme = BCe ^((~BCi)& BCo ); |
|
Emi = BCi ^((~BCo)& BCu ); |
|
Emo = BCo ^((~BCu)& BCa ); |
|
Emu = BCu ^((~BCa)& BCe ); |
|
|
|
Abi = vectorize(pdata64[2]) ^ Di; |
|
BCa = ROL(Abi, 62); |
|
Ago = vectorize(pdata64[8]) ^ Do; |
|
BCe = ROL(Ago, 55); |
|
Aku = Du; |
|
BCi = ROL(Aku, 39); |
|
Ama = Da; |
|
BCo = ROL(Ama, 41); |
|
Ase = De; |
|
BCu = ROL(Ase, 2); |
|
Esa = BCa ^((~BCe)& BCi ); |
|
Ese = BCe ^((~BCi)& BCo ); |
|
Esi = BCi ^((~BCo)& BCu ); |
|
Eso = BCo ^((~BCu)& BCa ); |
|
Esu = BCu ^((~BCa)& BCe ); |
|
|
|
// prepareTheta |
|
BCa = Eba^Ega^Eka^Ema^Esa; |
|
BCe = Ebe^Ege^Eke^Eme^Ese; |
|
BCi = Ebi^Egi^Eki^Emi^Esi; |
|
BCo = Ebo^Ego^Eko^Emo^Eso; |
|
BCu = Ebu^Egu^Eku^Emu^Esu; |
|
|
|
//thetaRhoPiChiIotaPrepareTheta(round+1, E, A) |
|
Da = BCu^ROL(BCe, 1); |
|
De = BCa^ROL(BCi, 1); |
|
Di = BCe^ROL(BCo, 1); |
|
Do = BCi^ROL(BCu, 1); |
|
Du = BCo^ROL(BCa, 1); |
|
|
|
Eba ^= Da; |
|
BCa = Eba; |
|
Ege ^= De; |
|
BCe = ROL(Ege, 44); |
|
Eki ^= Di; |
|
BCi = ROL(Eki, 43); |
|
Emo ^= Do; |
|
BCo = ROL(Emo, 21); |
|
Esu ^= Du; |
|
BCu = ROL(Esu, 14); |
|
Aba = BCa ^((~BCe)& BCi ); |
|
Aba ^= vectorize((uint64_t)KeccakF_RoundConstants[1]); |
|
Abe = BCe ^((~BCi)& BCo ); |
|
Abi = BCi ^((~BCo)& BCu ); |
|
Abo = BCo ^((~BCu)& BCa ); |
|
Abu = BCu ^((~BCa)& BCe ); |
|
|
|
Ebo ^= Do; |
|
BCa = ROL(Ebo, 28); |
|
Egu ^= Du; |
|
BCe = ROL(Egu, 20); |
|
Eka ^= Da; |
|
BCi = ROL(Eka, 3); |
|
Eme ^= De; |
|
BCo = ROL(Eme, 45); |
|
Esi ^= Di; |
|
BCu = ROL(Esi, 61); |
|
Aga = BCa ^((~BCe)& BCi ); |
|
Age = BCe ^((~BCi)& BCo ); |
|
Agi = BCi ^((~BCo)& BCu ); |
|
Ago = BCo ^((~BCu)& BCa ); |
|
Agu = BCu ^((~BCa)& BCe ); |
|
|
|
Ebe ^= De; |
|
BCa = ROL(Ebe, 1); |
|
Egi ^= Di; |
|
BCe = ROL(Egi, 6); |
|
Eko ^= Do; |
|
BCi = ROL(Eko, 25); |
|
Emu ^= Du; |
|
BCo = ROL(Emu, 8); |
|
Esa ^= Da; |
|
BCu = ROL(Esa, 18); |
|
Aka = BCa ^((~BCe)& BCi ); |
|
Ake = BCe ^((~BCi)& BCo ); |
|
Aki = BCi ^((~BCo)& BCu ); |
|
Ako = BCo ^((~BCu)& BCa ); |
|
Aku = BCu ^((~BCa)& BCe ); |
|
|
|
Ebu ^= Du; |
|
BCa = ROL(Ebu, 27); |
|
Ega ^= Da; |
|
BCe = ROL(Ega, 36); |
|
Eke ^= De; |
|
BCi = ROL(Eke, 10); |
|
Emi ^= Di; |
|
BCo = ROL(Emi, 15); |
|
Eso ^= Do; |
|
BCu = ROL(Eso, 56); |
|
Ama = BCa ^((~BCe)& BCi ); |
|
Ame = BCe ^((~BCi)& BCo ); |
|
Ami = BCi ^((~BCo)& BCu ); |
|
Amo = BCo ^((~BCu)& BCa ); |
|
Amu = BCu ^((~BCa)& BCe ); |
|
|
|
Ebi ^= Di; |
|
BCa = ROL(Ebi, 62); |
|
Ego ^= Do; |
|
BCe = ROL(Ego, 55); |
|
Eku ^= Du; |
|
BCi = ROL(Eku, 39); |
|
Ema ^= Da; |
|
BCo = ROL(Ema, 41); |
|
Ese ^= De; |
|
BCu = ROL(Ese, 2); |
|
Asa = BCa ^((~BCe)& BCi ); |
|
Ase = BCe ^((~BCi)& BCo ); |
|
Asi = BCi ^((~BCo)& BCu ); |
|
Aso = BCo ^((~BCu)& BCa ); |
|
Asu = BCu ^((~BCa)& BCe ); |
|
|
|
//#pragma unroll 10 |
|
for( int laneCount = 2; laneCount < 22; laneCount += 2 ) |
|
{ |
|
// prepareTheta |
|
BCa = Aba^Aga^Aka^Ama^Asa; |
|
BCe = Abe^Age^Ake^Ame^Ase; |
|
BCi = Abi^Agi^Aki^Ami^Asi; |
|
BCo = Abo^Ago^Ako^Amo^Aso; |
|
BCu = Abu^Agu^Aku^Amu^Asu; |
|
|
|
//thetaRhoPiChiIotaPrepareTheta(round , A, E) |
|
Da = BCu^ROL(BCe, 1); |
|
De = BCa^ROL(BCi, 1); |
|
Di = BCe^ROL(BCo, 1); |
|
Do = BCi^ROL(BCu, 1); |
|
Du = BCo^ROL(BCa, 1); |
|
|
|
Aba ^= Da; |
|
BCa = Aba; |
|
Age ^= De; |
|
BCe = ROL(Age, 44); |
|
Aki ^= Di; |
|
BCi = ROL(Aki, 43); |
|
Amo ^= Do; |
|
BCo = ROL(Amo, 21); |
|
Asu ^= Du; |
|
BCu = ROL(Asu, 14); |
|
Eba = BCa ^((~BCe)& BCi ); |
|
Eba ^= vectorize((uint64_t)KeccakF_RoundConstants[laneCount]); |
|
Ebe = BCe ^((~BCi)& BCo ); |
|
Ebi = BCi ^((~BCo)& BCu ); |
|
Ebo = BCo ^((~BCu)& BCa ); |
|
Ebu = BCu ^((~BCa)& BCe ); |
|
|
|
Abo ^= Do; |
|
BCa = ROL(Abo, 28); |
|
Agu ^= Du; |
|
BCe = ROL(Agu, 20); |
|
Aka ^= Da; |
|
BCi = ROL(Aka, 3); |
|
Ame ^= De; |
|
BCo = ROL(Ame, 45); |
|
Asi ^= Di; |
|
BCu = ROL(Asi, 61); |
|
Ega = BCa ^((~BCe)& BCi ); |
|
Ege = BCe ^((~BCi)& BCo ); |
|
Egi = BCi ^((~BCo)& BCu ); |
|
Ego = BCo ^((~BCu)& BCa ); |
|
Egu = BCu ^((~BCa)& BCe ); |
|
|
|
Abe ^= De; |
|
BCa = ROL(Abe, 1); |
|
Agi ^= Di; |
|
BCe = ROL(Agi, 6); |
|
Ako ^= Do; |
|
BCi = ROL(Ako, 25); |
|
Amu ^= Du; |
|
BCo = ROL(Amu, 8); |
|
Asa ^= Da; |
|
BCu = ROL(Asa, 18); |
|
Eka = BCa ^((~BCe)& BCi ); |
|
Eke = BCe ^((~BCi)& BCo ); |
|
Eki = BCi ^((~BCo)& BCu ); |
|
Eko = BCo ^((~BCu)& BCa ); |
|
Eku = BCu ^((~BCa)& BCe ); |
|
|
|
Abu ^= Du; |
|
BCa = ROL(Abu, 27); |
|
Aga ^= Da; |
|
BCe = ROL(Aga, 36); |
|
Ake ^= De; |
|
BCi = ROL(Ake, 10); |
|
Ami ^= Di; |
|
BCo = ROL(Ami, 15); |
|
Aso ^= Do; |
|
BCu = ROL(Aso, 56); |
|
Ema = BCa ^((~BCe)& BCi ); |
|
Eme = BCe ^((~BCi)& BCo ); |
|
Emi = BCi ^((~BCo)& BCu ); |
|
Emo = BCo ^((~BCu)& BCa ); |
|
Emu = BCu ^((~BCa)& BCe ); |
|
|
|
Abi ^= Di; |
|
BCa = ROL(Abi, 62); |
|
Ago ^= Do; |
|
BCe = ROL(Ago, 55); |
|
Aku ^= Du; |
|
BCi = ROL(Aku, 39); |
|
Ama ^= Da; |
|
BCo = ROL(Ama, 41); |
|
Ase ^= De; |
|
BCu = ROL(Ase, 2); |
|
Esa = BCa ^((~BCe)& BCi ); |
|
Ese = BCe ^((~BCi)& BCo ); |
|
Esi = BCi ^((~BCo)& BCu ); |
|
Eso = BCo ^((~BCu)& BCa ); |
|
Esu = BCu ^((~BCa)& BCe ); |
|
|
|
// prepareTheta |
|
BCa = Eba^Ega^Eka^Ema^Esa; |
|
BCe = Ebe^Ege^Eke^Eme^Ese; |
|
BCi = Ebi^Egi^Eki^Emi^Esi; |
|
BCo = Ebo^Ego^Eko^Emo^Eso; |
|
BCu = Ebu^Egu^Eku^Emu^Esu; |
|
|
|
//thetaRhoPiChiIotaPrepareTheta(round+1, E, A) |
|
Da = BCu^ROL(BCe, 1); |
|
De = BCa^ROL(BCi, 1); |
|
Di = BCe^ROL(BCo, 1); |
|
Do = BCi^ROL(BCu, 1); |
|
Du = BCo^ROL(BCa, 1); |
|
|
|
Eba ^= Da; |
|
BCa = Eba; |
|
Ege ^= De; |
|
BCe = ROL(Ege, 44); |
|
Eki ^= Di; |
|
BCi = ROL(Eki, 43); |
|
Emo ^= Do; |
|
BCo = ROL(Emo, 21); |
|
Esu ^= Du; |
|
BCu = ROL(Esu, 14); |
|
Aba = BCa ^((~BCe)& BCi ); |
|
Aba ^= vectorize((uint64_t)KeccakF_RoundConstants[laneCount+1]); |
|
Abe = BCe ^((~BCi)& BCo ); |
|
Abi = BCi ^((~BCo)& BCu ); |
|
Abo = BCo ^((~BCu)& BCa ); |
|
Abu = BCu ^((~BCa)& BCe ); |
|
|
|
Ebo ^= Do; |
|
BCa = ROL(Ebo, 28); |
|
Egu ^= Du; |
|
BCe = ROL(Egu, 20); |
|
Eka ^= Da; |
|
BCi = ROL(Eka, 3); |
|
Eme ^= De; |
|
BCo = ROL(Eme, 45); |
|
Esi ^= Di; |
|
BCu = ROL(Esi, 61); |
|
Aga = BCa ^((~BCe)& BCi ); |
|
Age = BCe ^((~BCi)& BCo ); |
|
Agi = BCi ^((~BCo)& BCu ); |
|
Ago = BCo ^((~BCu)& BCa ); |
|
Agu = BCu ^((~BCa)& BCe ); |
|
|
|
Ebe ^= De; |
|
BCa = ROL(Ebe, 1); |
|
Egi ^= Di; |
|
BCe = ROL(Egi, 6); |
|
Eko ^= Do; |
|
BCi = ROL(Eko, 25); |
|
Emu ^= Du; |
|
BCo = ROL(Emu, 8); |
|
Esa ^= Da; |
|
BCu = ROL(Esa, 18); |
|
Aka = BCa ^((~BCe)& BCi ); |
|
Ake = BCe ^((~BCi)& BCo ); |
|
Aki = BCi ^((~BCo)& BCu ); |
|
Ako = BCo ^((~BCu)& BCa ); |
|
Aku = BCu ^((~BCa)& BCe ); |
|
|
|
Ebu ^= Du; |
|
BCa = ROL(Ebu, 27); |
|
Ega ^= Da; |
|
BCe = ROL(Ega, 36); |
|
Eke ^= De; |
|
BCi = ROL(Eke, 10); |
|
Emi ^= Di; |
|
BCo = ROL(Emi, 15); |
|
Eso ^= Do; |
|
BCu = ROL(Eso, 56); |
|
Ama = BCa ^((~BCe)& BCi ); |
|
Ame = BCe ^((~BCi)& BCo ); |
|
Ami = BCi ^((~BCo)& BCu ); |
|
Amo = BCo ^((~BCu)& BCa ); |
|
Amu = BCu ^((~BCa)& BCe ); |
|
|
|
Ebi ^= Di; |
|
BCa = ROL(Ebi, 62); |
|
Ego ^= Do; |
|
BCe = ROL(Ego, 55); |
|
Eku ^= Du; |
|
BCi = ROL(Eku, 39); |
|
Ema ^= Da; |
|
BCo = ROL(Ema, 41); |
|
Ese ^= De; |
|
BCu = ROL(Ese, 2); |
|
Asa = BCa ^((~BCe)& BCi ); |
|
Ase = BCe ^((~BCi)& BCo ); |
|
Asi = BCi ^((~BCo)& BCu ); |
|
Aso = BCo ^((~BCu)& BCa ); |
|
Asu = BCu ^((~BCa)& BCe ); |
|
} |
|
|
|
// prepareTheta |
|
BCa = Aba^Aga^Aka^Ama^Asa; |
|
BCe = Abe^Age^Ake^Ame^Ase; |
|
BCi = Abi^Agi^Aki^Ami^Asi; |
|
BCo = Abo^Ago^Ako^Amo^Aso; |
|
BCu = Abu^Agu^Aku^Amu^Asu; |
|
|
|
//thetaRhoPiChiIotaPrepareTheta(round , A, E) |
|
Da = BCu^ROL(BCe, 1); |
|
De = BCa^ROL(BCi, 1); |
|
Di = BCe^ROL(BCo, 1); |
|
Do = BCi^ROL(BCu, 1); |
|
Du = BCo^ROL(BCa, 1); |
|
|
|
Aba ^= Da; |
|
BCa = Aba; |
|
Age ^= De; |
|
BCe = ROL(Age, 44); |
|
Aki ^= Di; |
|
BCi = ROL(Aki, 43); |
|
Amo ^= Do; |
|
BCo = ROL(Amo, 21); |
|
Asu ^= Du; |
|
BCu = ROL(Asu, 14); |
|
Eba = BCa ^((~BCe)& BCi ); |
|
Eba ^= vectorize((uint64_t)KeccakF_RoundConstants[22]); |
|
Ebe = BCe ^((~BCi)& BCo ); |
|
Ebi = BCi ^((~BCo)& BCu ); |
|
Ebo = BCo ^((~BCu)& BCa ); |
|
Ebu = BCu ^((~BCa)& BCe ); |
|
|
|
Abo ^= Do; |
|
BCa = ROL(Abo, 28); |
|
Agu ^= Du; |
|
BCe = ROL(Agu, 20); |
|
Aka ^= Da; |
|
BCi = ROL(Aka, 3); |
|
Ame ^= De; |
|
BCo = ROL(Ame, 45); |
|
Asi ^= Di; |
|
BCu = ROL(Asi, 61); |
|
Ega = BCa ^((~BCe)& BCi ); |
|
Ege = BCe ^((~BCi)& BCo ); |
|
Egi = BCi ^((~BCo)& BCu ); |
|
Ego = BCo ^((~BCu)& BCa ); |
|
Egu = BCu ^((~BCa)& BCe ); |
|
|
|
Abe ^= De; |
|
BCa = ROL(Abe, 1); |
|
Agi ^= Di; |
|
BCe = ROL(Agi, 6); |
|
Ako ^= Do; |
|
BCi = ROL(Ako, 25); |
|
Amu ^= Du; |
|
BCo = ROL(Amu, 8); |
|
Asa ^= Da; |
|
BCu = ROL(Asa, 18); |
|
Eka = BCa ^((~BCe)& BCi ); |
|
Eke = BCe ^((~BCi)& BCo ); |
|
Eki = BCi ^((~BCo)& BCu ); |
|
Eko = BCo ^((~BCu)& BCa ); |
|
Eku = BCu ^((~BCa)& BCe ); |
|
|
|
Abu ^= Du; |
|
BCa = ROL(Abu, 27); |
|
Aga ^= Da; |
|
BCe = ROL(Aga, 36); |
|
Ake ^= De; |
|
BCi = ROL(Ake, 10); |
|
Ami ^= Di; |
|
BCo = ROL(Ami, 15); |
|
Aso ^= Do; |
|
BCu = ROL(Aso, 56); |
|
Ema = BCa ^((~BCe)& BCi ); |
|
Eme = BCe ^((~BCi)& BCo ); |
|
Emi = BCi ^((~BCo)& BCu ); |
|
Emo = BCo ^((~BCu)& BCa ); |
|
Emu = BCu ^((~BCa)& BCe ); |
|
|
|
Abi ^= Di; |
|
BCa = ROL(Abi, 62); |
|
Ago ^= Do; |
|
BCe = ROL(Ago, 55); |
|
Aku ^= Du; |
|
BCi = ROL(Aku, 39); |
|
Ama ^= Da; |
|
BCo = ROL(Ama, 41); |
|
Ase ^= De; |
|
BCu = ROL(Ase, 2); |
|
Esa = BCa ^((~BCe)& BCi ); |
|
Ese = BCe ^((~BCi)& BCo ); |
|
Esi = BCi ^((~BCo)& BCu ); |
|
Eso = BCo ^((~BCu)& BCa ); |
|
Esu = BCu ^((~BCa)& BCe ); |
|
|
|
// prepareTheta |
|
BCa = Eba^Ega^Eka^Ema^Esa; |
|
BCe = Ebe^Ege^Eke^Eme^Ese; |
|
BCi = Ebi^Egi^Eki^Emi^Esi; |
|
BCo = Ebo^Ego^Eko^Emo^Eso; |
|
BCu = Ebu^Egu^Eku^Emu^Esu; |
|
|
|
//thetaRhoPiChiIotaPrepareTheta(round+1, E, A) |
|
Da = BCu^ROL(BCe, 1); |
|
De = BCa^ROL(BCi, 1); |
|
Di = BCe^ROL(BCo, 1); |
|
Do = BCi^ROL(BCu, 1); |
|
Du = BCo^ROL(BCa, 1); |
|
|
|
Eba ^= Da; |
|
BCa = Eba; |
|
Ege ^= De; |
|
BCe = ROL(Ege, 44); |
|
Eki ^= Di; |
|
BCi = ROL(Eki, 43); |
|
Emo ^= Do; |
|
BCo = ROL(Emo, 21); |
|
Esu ^= Du; |
|
BCu = ROL(Esu, 14); |
|
Aba = BCa ^((~BCe)& BCi ); |
|
Aba ^= vectorize((uint64_t)KeccakF_RoundConstants[23]); |
|
Abe = BCe ^((~BCi)& BCo ); |
|
Abi = BCi ^((~BCo)& BCu ); |
|
Abo = BCo ^((~BCu)& BCa ); |
|
|
|
if (validate) { |
|
g_out += 4 * ((blockIdx.x * blockDim.x) + threadIdx.x); |
|
g_out[3] = devectorize(Abo); |
|
g_out[2] = devectorize(Abi); |
|
g_out[1] = devectorize(Abe); |
|
g_out[0] = devectorize(Aba); |
|
} |
|
|
|
// the likelyhood of meeting the hashing target is so low, that we're not guarding this |
|
// with atomic writes, locks or similar... |
|
uint64_t *g_good64 = (uint64_t*)g_good; |
|
if (devectorize(Abo) <= ptarget64[3]) { |
|
if (devectorize(Abo) < g_good64[3]) { |
|
g_good64[3] = devectorize(Abo); |
|
g_good64[2] = devectorize(Abi); |
|
g_good64[1] = devectorize(Abe); |
|
g_good64[0] = devectorize(Aba); |
|
g_good[8] = nonce + ((blockIdx.x * blockDim.x) + threadIdx.x); |
|
} |
|
} |
|
} |
|
|
|
static std::map<int, uint32_t *> context_good[2]; |
|
|
|
bool NV2Kernel::prepare_keccak256(int thr_id, const uint32_t host_pdata[20], const uint32_t host_ptarget[8]) |
|
{ |
|
static bool init[MAX_GPUS] = { 0 }; |
|
|
|
if (!init[thr_id]) |
|
{ |
|
checkCudaErrors(cudaMemcpyToSymbol(KeccakF_RoundConstants, host_KeccakF_RoundConstants, sizeof(host_KeccakF_RoundConstants), 0, cudaMemcpyHostToDevice)); |
|
|
|
// allocate pinned host memory for good hashes |
|
uint32_t *tmp; |
|
checkCudaErrors(cudaMalloc((void **) &tmp, 9*sizeof(uint32_t))); context_good[0][thr_id] = tmp; |
|
checkCudaErrors(cudaMalloc((void **) &tmp, 9*sizeof(uint32_t))); context_good[1][thr_id] = tmp; |
|
|
|
init[thr_id] = true; |
|
} |
|
checkCudaErrors(cudaMemcpyToSymbol(pdata64, host_pdata, 20*sizeof(uint32_t), 0, cudaMemcpyHostToDevice)); |
|
checkCudaErrors(cudaMemcpyToSymbol(ptarget64, host_ptarget, 8*sizeof(uint32_t), 0, cudaMemcpyHostToDevice)); |
|
|
|
return context_good[0][thr_id] && context_good[1][thr_id]; |
|
} |
|
|
|
void NV2Kernel::do_keccak256(dim3 grid, dim3 threads, int thr_id, int stream, uint32_t *hash, uint32_t nonce, int throughput, bool do_d2h) |
|
{ |
|
checkCudaErrors(cudaMemsetAsync(context_good[stream][thr_id], 0xff, 9 * sizeof(uint32_t), context_streams[stream][thr_id])); |
|
|
|
titan_crypto_hash<<<grid, threads, 0, context_streams[stream][thr_id]>>>((uint64_t*)context_hash[stream][thr_id], nonce, context_good[stream][thr_id], do_d2h); |
|
|
|
// copy hashes from device memory to host (ALL hashes, lots of data...) |
|
if (do_d2h && hash != NULL) { |
|
size_t mem_size = throughput * sizeof(uint32_t) * 8; |
|
checkCudaErrors(cudaMemcpyAsync(hash, context_hash[stream][thr_id], mem_size, |
|
cudaMemcpyDeviceToHost, context_streams[stream][thr_id])); |
|
} |
|
else if (hash != NULL) { |
|
// asynchronous copy of winning nonce (just 4 bytes...) |
|
checkCudaErrors(cudaMemcpyAsync(hash, context_good[stream][thr_id]+8, sizeof(uint32_t), |
|
cudaMemcpyDeviceToHost, context_streams[stream][thr_id])); |
|
} |
|
} |
|
|
|
|
|
// |
|
// Blakecoin related Keccak implementation (Keccak256) |
|
// |
|
|
|
typedef uint32_t sph_u32; |
|
#define SPH_C32(x) ((sph_u32)(x)) |
|
#define SPH_T32(x) ((x) & SPH_C32(0xFFFFFFFF)) |
|
#if __CUDA_ARCH__ < 350 |
|
// Kepler (Compute 3.0) |
|
#define SPH_ROTL32(a, b) ((a)<<(b))|((a)>>(32-(b))) |
|
#else |
|
// Kepler (Compute 3.5) |
|
#define SPH_ROTL32(a, b) __funnelshift_l( a, a, b ); |
|
#endif |
|
#define SPH_ROTR32(x, n) SPH_ROTL32(x, (32 - (n))) |
|
|
|
__constant__ uint32_t pdata[20]; |
|
|
|
#ifdef _MSC_VER |
|
#pragma warning (disable: 4146) |
|
#endif |
|
|
|
static __device__ sph_u32 cuda_sph_bswap32(sph_u32 x) |
|
{ |
|
return (((x << 24) & 0xff000000u) | ((x << 8) & 0x00ff0000u) |
|
| ((x >> 8) & 0x0000ff00u) | ((x >> 24) & 0x000000ffu)); |
|
} |
|
|
|
/** |
|
* Encode a 32-bit value into the provided buffer (big endian convention). |
|
* |
|
* @param dst the destination buffer |
|
* @param val the 32-bit value to encode |
|
*/ |
|
static __device__ void |
|
cuda_sph_enc32be(void *dst, sph_u32 val) |
|
{ |
|
*(sph_u32 *)dst = cuda_sph_bswap32(val); |
|
} |
|
|
|
#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) |
|
|
|
#define GS(m0, m1, c0, c1, a, b, c, d) do { \ |
|
a = SPH_T32(a + b + (m0 ^ c1)); \ |
|
d = SPH_ROTR32(d ^ a, 16); \ |
|
c = SPH_T32(c + d); \ |
|
b = SPH_ROTR32(b ^ c, 12); \ |
|
a = SPH_T32(a + b + (m1 ^ c0)); \ |
|
d = SPH_ROTR32(d ^ a, 8); \ |
|
c = SPH_T32(c + d); \ |
|
b = SPH_ROTR32(b ^ c, 7); \ |
|
} while (0) |
|
|
|
#define ROUND_S(r) do { \ |
|
GS(Mx(r, 0), Mx(r, 1), CSx(r, 0), CSx(r, 1), V0, V4, V8, VC); \ |
|
GS(Mx(r, 2), Mx(r, 3), CSx(r, 2), CSx(r, 3), V1, V5, V9, VD); \ |
|
GS(Mx(r, 4), Mx(r, 5), CSx(r, 4), CSx(r, 5), V2, V6, VA, VE); \ |
|
GS(Mx(r, 6), Mx(r, 7), CSx(r, 6), CSx(r, 7), V3, V7, VB, VF); \ |
|
GS(Mx(r, 8), Mx(r, 9), CSx(r, 8), CSx(r, 9), V0, V5, VA, VF); \ |
|
GS(Mx(r, A), Mx(r, B), CSx(r, A), CSx(r, B), V1, V6, VB, VC); \ |
|
GS(Mx(r, C), Mx(r, D), CSx(r, C), CSx(r, D), V2, V7, V8, VD); \ |
|
GS(Mx(r, E), Mx(r, F), CSx(r, E), CSx(r, F), V3, V4, V9, VE); \ |
|
} while (0) |
|
|
|
#define COMPRESS32 do { \ |
|
sph_u32 M0, M1, M2, M3, M4, M5, M6, M7; \ |
|
sph_u32 M8, M9, MA, MB, MC, MD, ME, MF; \ |
|
sph_u32 V0, V1, V2, V3, V4, V5, V6, V7; \ |
|
sph_u32 V8, V9, VA, VB, VC, VD, VE, VF; \ |
|
V0 = H0; \ |
|
V1 = H1; \ |
|
V2 = H2; \ |
|
V3 = H3; \ |
|
V4 = H4; \ |
|
V5 = H5; \ |
|
V6 = H6; \ |
|
V7 = H7; \ |
|
V8 = S0 ^ CS0; \ |
|
V9 = S1 ^ CS1; \ |
|
VA = S2 ^ CS2; \ |
|
VB = S3 ^ CS3; \ |
|
VC = T0 ^ CS4; \ |
|
VD = T0 ^ CS5; \ |
|
VE = T1 ^ CS6; \ |
|
VF = T1 ^ CS7; \ |
|
M0 = input[0]; \ |
|
M1 = input[1]; \ |
|
M2 = input[2]; \ |
|
M3 = input[3]; \ |
|
M4 = input[4]; \ |
|
M5 = input[5]; \ |
|
M6 = input[6]; \ |
|
M7 = input[7]; \ |
|
M8 = input[8]; \ |
|
M9 = input[9]; \ |
|
MA = input[10]; \ |
|
MB = input[11]; \ |
|
MC = input[12]; \ |
|
MD = input[13]; \ |
|
ME = input[14]; \ |
|
MF = input[15]; \ |
|
ROUND_S(0); \ |
|
ROUND_S(1); \ |
|
ROUND_S(2); \ |
|
ROUND_S(3); \ |
|
ROUND_S(4); \ |
|
ROUND_S(5); \ |
|
ROUND_S(6); \ |
|
ROUND_S(7); \ |
|
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) |
|
|
|
|
|
__global__ void titan_blake256_hash( uint64_t *g_out, uint32_t nonce, uint32_t *g_good, bool validate ) |
|
{ |
|
uint32_t input[16]; |
|
uint64_t output[4]; |
|
|
|
#pragma unroll 16 |
|
for (int i=0; i < 16; ++i) input[i] = pdata[i]; |
|
|
|
sph_u32 H0 = 0x6A09E667; |
|
sph_u32 H1 = 0xBB67AE85; |
|
sph_u32 H2 = 0x3C6EF372; |
|
sph_u32 H3 = 0xA54FF53A; |
|
sph_u32 H4 = 0x510E527F; |
|
sph_u32 H5 = 0x9B05688C; |
|
sph_u32 H6 = 0x1F83D9AB; |
|
sph_u32 H7 = 0x5BE0CD19; |
|
sph_u32 S0 = 0; |
|
sph_u32 S1 = 0; |
|
sph_u32 S2 = 0; |
|
sph_u32 S3 = 0; |
|
sph_u32 T0 = 0; |
|
sph_u32 T1 = 0; |
|
T0 = SPH_T32(T0 + 512); |
|
COMPRESS32; |
|
|
|
#pragma unroll 3 |
|
for (int i=0; i < 3; ++i) input[i] = pdata[16+i]; |
|
input[3] = nonce + ((blockIdx.x * blockDim.x) + threadIdx.x); |
|
input[4] = 0x80000000; |
|
#pragma unroll 8 |
|
for (int i=5; i < 13; ++i) input[i] = 0; |
|
input[13] = 0x00000001; |
|
input[14] = T1; |
|
input[15] = T0 + 128; |
|
|
|
T0 = SPH_T32(T0 + 128); |
|
COMPRESS32; |
|
|
|
cuda_sph_enc32be((unsigned char*)output + 4*6, H6); |
|
cuda_sph_enc32be((unsigned char*)output + 4*7, H7); |
|
if (validate || output[3] <= ptarget64[3]) |
|
{ |
|
// this data is only needed when we actually need to save the hashes |
|
cuda_sph_enc32be((unsigned char*)output + 4*0, H0); |
|
cuda_sph_enc32be((unsigned char*)output + 4*1, H1); |
|
cuda_sph_enc32be((unsigned char*)output + 4*2, H2); |
|
cuda_sph_enc32be((unsigned char*)output + 4*3, H3); |
|
cuda_sph_enc32be((unsigned char*)output + 4*4, H4); |
|
cuda_sph_enc32be((unsigned char*)output + 4*5, H5); |
|
} |
|
|
|
if (validate) |
|
{ |
|
g_out += 4 * ((blockIdx.x * blockDim.x) + threadIdx.x); |
|
#pragma unroll 4 |
|
for (int i=0; i < 4; ++i) g_out[i] = output[i]; |
|
} |
|
|
|
if (output[3] <= ptarget64[3]) { |
|
uint64_t *g_good64 = (uint64_t*)g_good; |
|
if (output[3] < g_good64[3]) { |
|
g_good64[3] = output[3]; |
|
g_good64[2] = output[2]; |
|
g_good64[1] = output[1]; |
|
g_good64[0] = output[0]; |
|
g_good[8] = nonce + ((blockIdx.x * blockDim.x) + threadIdx.x); |
|
} |
|
} |
|
} |
|
|
|
bool NV2Kernel::prepare_blake256(int thr_id, const uint32_t host_pdata[20], const uint32_t host_ptarget[8]) |
|
{ |
|
static bool init[MAX_GPUS] = { 0 }; |
|
|
|
if (!init[thr_id]) |
|
{ |
|
// allocate pinned host memory for good hashes |
|
uint32_t *tmp; |
|
checkCudaErrors(cudaMalloc((void **) &tmp, 9*sizeof(uint32_t))); context_good[0][thr_id] = tmp; |
|
checkCudaErrors(cudaMalloc((void **) &tmp, 9*sizeof(uint32_t))); context_good[1][thr_id] = tmp; |
|
|
|
init[thr_id] = true; |
|
} |
|
checkCudaErrors(cudaMemcpyToSymbol(pdata, host_pdata, 20*sizeof(uint32_t), 0, cudaMemcpyHostToDevice)); |
|
checkCudaErrors(cudaMemcpyToSymbol(ptarget64, host_ptarget, 8*sizeof(uint32_t), 0, cudaMemcpyHostToDevice)); |
|
|
|
return context_good[0][thr_id] && context_good[1][thr_id]; |
|
} |
|
|
|
void NV2Kernel::do_blake256(dim3 grid, dim3 threads, int thr_id, int stream, uint32_t *hash, uint32_t nonce, int throughput, bool do_d2h) |
|
{ |
|
checkCudaErrors(cudaMemsetAsync(context_good[stream][thr_id], 0xff, 9 * sizeof(uint32_t), context_streams[stream][thr_id])); |
|
|
|
titan_blake256_hash<<<grid, threads, 0, context_streams[stream][thr_id]>>>((uint64_t*)context_hash[stream][thr_id], nonce, context_good[stream][thr_id], do_d2h); |
|
|
|
// copy hashes from device memory to host (ALL hashes, lots of data...) |
|
if (do_d2h && hash != NULL) { |
|
size_t mem_size = throughput * sizeof(uint32_t) * 8; |
|
checkCudaErrors(cudaMemcpyAsync(hash, context_hash[stream][thr_id], mem_size, |
|
cudaMemcpyDeviceToHost, context_streams[stream][thr_id])); |
|
} |
|
else if (hash != NULL) { |
|
// asynchronous copy of winning nonce (just 4 bytes...) |
|
checkCudaErrors(cudaMemcpyAsync(hash, context_good[stream][thr_id]+8, sizeof(uint32_t), |
|
cudaMemcpyDeviceToHost, context_streams[stream][thr_id])); |
|
} |
|
}
|
|
|