Skip to content

Commit 41543b5

Browse files
tpruvottpruvot
committed
lyra2v2: add support for SM 2.1 devices
and improve a bit SM 3 perf
1 parent e4476a5 commit 41543b5

5 files changed

+133
-22
lines changed

‎bench.cpp

-1
Original file line numberDiff line numberDiff line change
@@ -103,7 +103,6 @@ bool bench_algo_switch_next(int thr_id)
103103
if (algo == ALGO_GROESTL) algo++;
104104
if (algo == ALGO_MYR_GR) algo++;
105105
if (algo == ALGO_JACKPOT) algo++; // compact shuffle
106-
if (algo == ALGO_LYRA2v2) algo++;
107106
if (algo == ALGO_NEOSCRYPT) algo++;
108107
if (algo == ALGO_WHIRLPOOLX) algo++;
109108
}

‎lyra2/cuda_lyra2_vectors.h

+12
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,10 @@
1212

1313
#include "cuda_helper.h"
1414

15+
#if __CUDA_ARCH__ < 300
16+
#define __shfl(x, y) (x)
17+
#endif
18+
1519
#if __CUDA_ARCH__ < 320 && !defined(__ldg4)
1620
#define __ldg4(x) (*(x))
1721
#endif
@@ -545,6 +549,7 @@ static __forceinline__ __device__ uint16 swapvec(const uint16 &buf)
545549

546550
static __device__ __forceinline__ uint28 shuffle4(const uint28 &var, int lane)
547551
{
552+
#if __CUDA_ARCH__ >= 300
548553
uint28 res;
549554
res.x.x = __shfl(var.x.x, lane);
550555
res.x.y = __shfl(var.x.y, lane);
@@ -555,10 +560,14 @@ static __device__ __forceinline__ uint28 shuffle4(const uint28 &var, int lane)
555560
res.w.x = __shfl(var.w.x, lane);
556561
res.w.y = __shfl(var.w.y, lane);
557562
return res;
563+
#else
564+
return var;
565+
#endif
558566
}
559567

560568
static __device__ __forceinline__ ulonglong4 shuffle4(ulonglong4 var, int lane)
561569
{
570+
#if __CUDA_ARCH__ >= 300
562571
ulonglong4 res;
563572
uint2 temp;
564573
temp = vectorize(var.x);
@@ -578,6 +587,9 @@ static __device__ __forceinline__ ulonglong4 shuffle4(ulonglong4 var, int lane)
578587
temp.y = __shfl(temp.y, lane);
579588
res.w = devectorize(temp);
580589
return res;
590+
#else
591+
return var;
592+
#endif
581593
}
582594

583595
#endif // #ifndef CUDA_LYRA_VECTOR_H

‎lyra2/cuda_lyra2v2.cu

+5-4
Original file line numberDiff line numberDiff line change
@@ -342,7 +342,7 @@ void lyra2v2_gpu_hash_32(const uint32_t threads, uint32_t startNounce, uint2 *g_
342342
}
343343
#else
344344
#include "cuda_helper.h"
345-
#if __CUDA_ARCH__ < 300
345+
#if __CUDA_ARCH__ < 200
346346
__device__ void* DMatrix;
347347
#endif
348348
__global__ void lyra2v2_gpu_hash_32(const uint32_t threads, uint32_t startNounce, uint2 *g_hash) {}
@@ -362,9 +362,10 @@ void lyra2v2_cpu_hash_32(int thr_id, uint32_t threads, uint32_t startNounce, uin
362362
int dev_id = device_map[thr_id % MAX_GPUS];
363363
uint32_t tpb = TPB52;
364364

365-
if (device_sm[dev_id] == 500 || cuda_arch[dev_id] == 500) tpb = TPB50;
366-
else if (device_sm[dev_id] == 350 || cuda_arch[dev_id] == 350) tpb = TPB35;
367-
else if (device_sm[dev_id] < 350 || cuda_arch[dev_id] < 350) tpb = TPB30;
365+
if (cuda_arch[dev_id] == 500) tpb = TPB50;
366+
else if (cuda_arch[dev_id] >= 350) tpb = TPB35;
367+
else if (cuda_arch[dev_id] >= 300) tpb = TPB30;
368+
else if (cuda_arch[dev_id] >= 200) tpb = TPB20;
368369

369370
dim3 grid((threads + tpb - 1) / tpb);
370371
dim3 block(tpb);

‎lyra2/cuda_lyra2v2_sm3.cuh

+94-10
Original file line numberDiff line numberDiff line change
@@ -1,15 +1,16 @@
1-
/* SM 3/3.5 Variant for lyra2REv2 */
1+
/* SM 2/3/3.5 Variant for lyra2REv2 */
22

33
#ifdef __INTELLISENSE__
44
/* just for vstudio code colors */
55
#undef __CUDA_ARCH__
66
#define __CUDA_ARCH__ 350
77
#endif
88

9-
#define TPB30 16
9+
#define TPB20 64
10+
#define TPB30 64
1011
#define TPB35 64
1112

12-
#if __CUDA_ARCH__ >= 300 && __CUDA_ARCH__ < 500
13+
#if __CUDA_ARCH__ >= 200 && __CUDA_ARCH__ < 500
1314

1415
#include "cuda_lyra2_vectors.h"
1516

@@ -165,6 +166,7 @@ void reduceDuplexRowtV3(const int rowIn, const int rowInOut, const int rowOut, v
165166
}
166167
}
167168

169+
#if __CUDA_ARCH__ >= 300
168170
__global__ __launch_bounds__(TPB35, 1)
169171
void lyra2v2_gpu_hash_32_v3(uint32_t threads, uint32_t startNounce, uint2 *outputHash)
170172
{
@@ -177,14 +179,14 @@ void lyra2v2_gpu_hash_32_v3(uint32_t threads, uint32_t startNounce, uint2 *outpu
177179
if (threadIdx.x == 0) {
178180

179181
((uint16*)blake2b_IV)[0] = make_uint16(
180-
0xf3bcc908, 0x6a09e667 , 0x84caa73b, 0xbb67ae85 ,
181-
0xfe94f82b, 0x3c6ef372 , 0x5f1d36f1, 0xa54ff53a ,
182-
0xade682d1, 0x510e527f , 0x2b3e6c1f, 0x9b05688c ,
182+
0xf3bcc908, 0x6a09e667 , 0x84caa73b, 0xbb67ae85,
183+
0xfe94f82b, 0x3c6ef372 , 0x5f1d36f1, 0xa54ff53a,
184+
0xade682d1, 0x510e527f , 0x2b3e6c1f, 0x9b05688c,
183185
0xfb41bd6b, 0x1f83d9ab , 0x137e2179, 0x5be0cd19
184186
);
185187
((uint16*)padding)[0] = make_uint16(
186-
0x20, 0x0 , 0x20, 0x0 , 0x20, 0x0 , 0x01, 0x0 ,
187-
0x04, 0x0 , 0x04, 0x0 , 0x80, 0x0 , 0x0, 0x01000000
188+
0x20, 0x0 , 0x20, 0x0 , 0x20, 0x0 , 0x01, 0x0,
189+
0x04, 0x0 , 0x04, 0x0 , 0x80, 0x0 , 0x0, 0x01000000
188190
);
189191
}
190192

@@ -194,6 +196,7 @@ void lyra2v2_gpu_hash_32_v3(uint32_t threads, uint32_t startNounce, uint2 *outpu
194196
((uint2*)state)[1] = __ldg(&outputHash[thread + threads]);
195197
((uint2*)state)[2] = __ldg(&outputHash[thread + 2 * threads]);
196198
((uint2*)state)[3] = __ldg(&outputHash[thread + 3 * threads]);
199+
197200
state[1] = state[0];
198201
state[2] = shuffle4(((vectype*)blake2b_IV)[0], 0);
199202
state[3] = shuffle4(((vectype*)blake2b_IV)[1], 0);
@@ -246,9 +249,90 @@ void lyra2v2_gpu_hash_32_v3(uint32_t threads, uint32_t startNounce, uint2 *outpu
246249

247250
} //thread
248251
}
252+
#elif __CUDA_ARCH__ >= 200
253+
__global__ __launch_bounds__(TPB20, 1)
254+
void lyra2v2_gpu_hash_32_v3(uint32_t threads, uint32_t startNounce, uint2 *outputHash)
255+
{
256+
uint32_t thread = (blockDim.x * blockIdx.x + threadIdx.x);
257+
258+
vectype state[4];
259+
vectype blake2b_IV[2];
260+
vectype padding[2];
261+
262+
((uint16*)blake2b_IV)[0] = make_uint16(
263+
0xf3bcc908, 0x6a09e667, 0x84caa73b, 0xbb67ae85,
264+
0xfe94f82b, 0x3c6ef372, 0x5f1d36f1, 0xa54ff53a,
265+
0xade682d1, 0x510e527f, 0x2b3e6c1f, 0x9b05688c,
266+
0xfb41bd6b, 0x1f83d9ab, 0x137e2179, 0x5be0cd19
267+
);
268+
((uint16*)padding)[0] = make_uint16(
269+
0x20, 0x0, 0x20, 0x0, 0x20, 0x0, 0x01, 0x0,
270+
0x04, 0x0, 0x04, 0x0, 0x80, 0x0, 0x0, 0x01000000
271+
);
272+
273+
if (thread < threads)
274+
{
275+
276+
((uint2*)state)[0] = outputHash[thread];
277+
((uint2*)state)[1] = outputHash[thread + threads];
278+
((uint2*)state)[2] = outputHash[thread + 2 * threads];
279+
((uint2*)state)[3] = outputHash[thread + 3 * threads];
280+
281+
state[1] = state[0];
282+
state[2] = ((vectype*)blake2b_IV)[0];
283+
state[3] = ((vectype*)blake2b_IV)[1];
284+
285+
for (int i = 0; i<12; i++)
286+
round_lyra_v35(state);
287+
288+
state[0] ^= ((vectype*)padding)[0];
289+
state[1] ^= ((vectype*)padding)[1];
290+
291+
for (int i = 0; i<12; i++)
292+
round_lyra_v35(state);
293+
294+
uint32_t ps1 = (4 * memshift * 3 + 16 * memshift * thread);
295+
296+
//#pragma unroll 4
297+
for (int i = 0; i < 4; i++)
298+
{
299+
uint32_t s1 = ps1 - 4 * memshift * i;
300+
for (int j = 0; j < 3; j++)
301+
(DMatrix + s1)[j] = (state)[j];
302+
303+
round_lyra_v35(state);
304+
}
305+
306+
reduceDuplexV3(state, thread);
307+
reduceDuplexRowSetupV3(1, 0, 2, state, thread);
308+
reduceDuplexRowSetupV3(2, 1, 3, state, thread);
309+
310+
uint32_t rowa;
311+
int prev = 3;
312+
for (int i = 0; i < 4; i++)
313+
{
314+
rowa = ((uint2*)state)[0].x & 3; reduceDuplexRowtV3(prev, rowa, i, state, thread);
315+
prev = i;
316+
}
317+
318+
uint32_t shift = (memshift * rowa + 16 * memshift * thread);
319+
320+
for (int j = 0; j < 3; j++)
321+
state[j] ^= __ldg4(&(DMatrix + shift)[j]);
322+
323+
for (int i = 0; i < 12; i++)
324+
round_lyra_v35(state);
325+
326+
outputHash[thread] = ((uint2*)state)[0];
327+
outputHash[thread + threads] = ((uint2*)state)[1];
328+
outputHash[thread + 2 * threads] = ((uint2*)state)[2];
329+
outputHash[thread + 3 * threads] = ((uint2*)state)[3];
330+
331+
} //thread
332+
}
333+
#endif
249334

250335
#else
251-
/* if __CUDA_ARCH__ < 300 .. */
336+
/* host & sm5+ */
252337
__global__ void lyra2v2_gpu_hash_32_v3(uint32_t threads, uint32_t startNounce, uint2 *outputHash) {}
253338
#endif
254-

‎lyra2/lyra2REv2.cu

+22-7
Original file line numberDiff line numberDiff line change
@@ -73,6 +73,21 @@ void lyra2v2_hash(void *state, const void *input)
7373
memcpy(state, hashA, 32);
7474
}
7575

76+
#ifdef _DEBUG
77+
#define TRACE(algo) { \
78+
if (max_nonce == 1 && pdata[19] <= 1) { \
79+
uint32_t* debugbuf = NULL; \
80+
cudaMallocHost(&debugbuf, 32); \
81+
cudaMemcpy(debugbuf, d_hash[thr_id], 32, cudaMemcpyDeviceToHost); \
82+
printf("lyra2 %s %08x %08x %08x %08x...%08x... \n", algo, swab32(debugbuf[0]), swab32(debugbuf[1]), \
83+
swab32(debugbuf[2]), swab32(debugbuf[3]), swab32(debugbuf[7])); \
84+
cudaFreeHost(debugbuf); \
85+
} \
86+
}
87+
#else
88+
#define TRACE(algo) {}
89+
#endif
90+
7691
static bool init[MAX_GPUS] = { 0 };
7792

7893
extern "C" int scanhash_lyra2v2(int thr_id, struct work* work, uint32_t max_nonce, unsigned long *hashes_done)
@@ -113,12 +128,6 @@ extern "C" int scanhash_lyra2v2(int thr_id, struct work* work, uint32_t max_nonc
113128

114129
CUDA_SAFE_CALL(cudaMalloc(&d_hash[thr_id], (size_t)32 * throughput));
115130

116-
if (device_sm[dev_id] < 300) {
117-
gpulog(LOG_ERR, thr_id, "Device SM 3.0 or more recent required!");
118-
proper_exit(1);
119-
return -1;
120-
}
121-
122131
api_set_throughput(thr_id, throughput);
123132
init[thr_id] = true;
124133
}
@@ -135,11 +144,17 @@ extern "C" int scanhash_lyra2v2(int thr_id, struct work* work, uint32_t max_nonc
135144
uint32_t foundNonces[2] = { 0, 0 };
136145

137146
blake256_cpu_hash_80(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
147+
TRACE("blake :");
138148
keccak256_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
149+
TRACE("keccak :");
139150
cubehash256_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
151+
TRACE("cube :");
140152
lyra2v2_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
153+
TRACE("lyra2 :");
141154
skein256_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], order++);
155+
TRACE("skein :");
142156
cubehash256_cpu_hash_32(thr_id, throughput,pdata[19], d_hash[thr_id], order++);
157+
TRACE("cube :");
143158

144159
bmw256_cpu_hash_32(thr_id, throughput, pdata[19], d_hash[thr_id], foundNonces);
145160

@@ -181,7 +196,7 @@ extern "C" int scanhash_lyra2v2(int thr_id, struct work* work, uint32_t max_nonc
181196
}
182197
pdata[19] += throughput;
183198

184-
} while (!work_restart[thr_id].restart);
199+
} while (!work_restart[thr_id].restart && !abort_flag);
185200

186201
*hashes_done = pdata[19] - first_nonce;
187202
return 0;

0 commit comments

Comments
 (0)