Skip to content

Commit

Permalink
Merge branch 'master' into feature/#228
Browse files Browse the repository at this point in the history
  • Loading branch information
fancyIX authored and fancyIX committed Mar 23, 2021
2 parents e6c9e05 + 8e7448a commit 7c484e1
Show file tree
Hide file tree
Showing 12 changed files with 2,848 additions and 23 deletions.
3 changes: 3 additions & 0 deletions algorithm.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -2547,6 +2547,7 @@ static algorithm_settings_t algos[] = {

{ "phi", ALGO_PHI, "", 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 5, 8 * 16 * 4194304, 0, phi_regenhash, NULL, NULL, queue_phi_kernel, gen_hash, append_x11_compiler_options },
{ "phi2", ALGO_PHI2, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 11, 8 * 16 * 4194304, 0, phi2_regenhash, NULL, NULL, queue_phi2_kernel, gen_hash, append_x11_compiler_options },
{ "phi2_navi", ALGO_PHI2_NAVI, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 11, 8 * 16 * 4194304, 0, phi2_regenhash, NULL, NULL, queue_phi2_kernel, gen_hash, append_x11_compiler_options },

{ "fresh", ALGO_FRESH, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 4, 4 * 16 * 4194304, 0, fresh_regenhash, NULL, NULL, queue_fresh_kernel, gen_hash, NULL },

Expand All @@ -2555,9 +2556,11 @@ static algorithm_settings_t algos[] = {
{ "lyra2rev3", ALGO_LYRA2REV3, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 8, -1, 0, lyra2rev3_regenhash, blake256_midstate, blake256_prepare_work, queue_lyra2rev3_kernel, gen_hash, append_neoscrypt_compiler_options },
{ "lyra2rev3_navi", ALGO_LYRA2REV3, "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 8, -1, 0, lyra2rev3_regenhash, blake256_midstate, blake256_prepare_work, queue_lyra2rev3_kernel, gen_hash, append_neoscrypt_compiler_options },
{ "lyra2Z" , ALGO_LYRA2Z , "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 4, -1, 0, lyra2Z_regenhash, blake256_midstate, blake256_prepare_work, queue_lyra2z_kernel, gen_hash, NULL },
{ "lyra2Z_navi" , ALGO_LYRA2Z_NAVI , "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 4, -1, 0, lyra2Z_regenhash, blake256_midstate, blake256_prepare_work, queue_lyra2z_kernel, gen_hash, NULL },
{ "lyra2Zz" , ALGO_LYRA2ZZ , "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 4, -1, 0, lyra2Zz_regenhash, blake256_midstate_112, blake256_prepare_work_112, queue_lyra2zz_kernel, gen_hash, NULL },
{ "lyra2h" , ALGO_LYRA2H , "", 1, 256, 256, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 1, -1, 0, lyra2h_regenhash, blake256_midstate, blake256_prepare_work, queue_lyra2h_kernel, gen_hash, NULL },
{ "allium", ALGO_ALLIUM, "", 1, 128, 128, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 10, 2 * 8 * 4194304, 0, allium_regenhash, blake256_midstate, blake256_prepare_work, queue_allium_kernel, gen_hash, NULL },
{ "allium_navi", ALGO_ALLIUM_NAVI, "", 1, 128, 128, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 10, 2 * 8 * 4194304, 0, allium_regenhash, blake256_midstate, blake256_prepare_work, queue_allium_kernel, gen_hash, NULL },
{ "mtp" , ALGO_MTP , "", 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 1, 0,0, mtp_regenhash , NULL, NULL, queue_mtp_kernel , gen_hash, NULL },
{ "mtp_vega" , ALGO_MTP , "", 1, 1, 1, 0, 0, 0xFF, 0xFFFFULL, 0x0000ffffUL, 1, 0,0, mtp_regenhash , NULL, NULL, queue_mtp_kernel , gen_hash, NULL },

Expand Down
3 changes: 3 additions & 0 deletions algorithm.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -15,6 +15,7 @@ typedef enum {
ALGO_MTP,
ALGO_UNK,
ALGO_ALLIUM,
ALGO_ALLIUM_NAVI,
ALGO_CRE,
ALGO_SCRYPT,
ALGO_NSCRYPT,
Expand All @@ -41,6 +42,7 @@ typedef enum {
ALGO_LYRA2REV2,
ALGO_LYRA2REV3,
ALGO_LYRA2Z,
ALGO_LYRA2Z_NAVI,
ALGO_LYRA2ZZ,
ALGO_LYRA2H,
ALGO_PLUCK,
Expand All @@ -56,6 +58,7 @@ typedef enum {
ALGO_LBRY,
ALGO_PHI,
ALGO_PHI2,
ALGO_PHI2_NAVI,
ALGO_SIBCOIN,
ALGO_ARGON2D
} algorithm_type_t;
Expand Down
4 changes: 2 additions & 2 deletions configure.ac
View file
Original file line number Diff line number Diff line change
Expand Up @@ -2,8 +2,8 @@
##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##
m4_define([v_maj], [0])
m4_define([v_min], [7])
m4_define([v_mic], [2])
m4_define([v_rev], [1])
m4_define([v_mic], [3])
m4_define([v_rev], [0])
##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##--##
m4_ifdef([v_rev], [m4_define([v_ver], [v_maj.v_min.v_mic-v_rev])], [m4_define([v_ver], [v_maj.v_min.v_mic])])
m4_define([lt_rev], m4_eval(v_maj + v_min))
Expand Down
14 changes: 7 additions & 7 deletions driver-opencl.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1459,8 +1459,8 @@ if (gpu->algorithm.type != ALGO_MTP && gpu->algorithm.type != ALGO_YESCRYPT_NAVI
}

for (i = 0; i < clState->n_extra_kernels; i++) {
if (gpu->algorithm.type == ALGO_PHI2 && i == 1) {
if (clState->prebuilt) {
if ((gpu->algorithm.type == ALGO_PHI2 || gpu->algorithm.type == ALGO_PHI2_NAVI) && i == 1) {
if (clState->prebuilt || gpu->algorithm.type == ALGO_PHI2_NAVI) {
const size_t off2[] = { 0, 0, *p_global_work_offset };
const size_t gws[] = { 4, 4, globalThreads[0] };
const size_t expand[] = { 4, 4, 16 };
Expand All @@ -1472,15 +1472,15 @@ if (gpu->algorithm.type != ALGO_MTP && gpu->algorithm.type != ALGO_YESCRYPT_NAVI
status = clEnqueueNDRangeKernel(clState->commandQueue, clState->extra_kernels[i], 2, off2, gws, expand, 0, NULL, NULL); // lyra 4w monolithic
}

} else if (gpu->algorithm.type == ALGO_PHI2 && (i == 0 || i == 2)) {
} else if ((gpu->algorithm.type == ALGO_PHI2 || gpu->algorithm.type == ALGO_PHI2_NAVI) && (i == 0 || i == 2)) {
size_t globalThreads2[1];
size_t localThreads2[1];
globalThreads2[0] = globalThreads[0] * 2;
localThreads2[0] = localThreads[0];
status = clEnqueueNDRangeKernel(clState->commandQueue, clState->extra_kernels[i], 1, p_global_work_offset,
globalThreads2, localThreads2, 0, NULL, NULL);
} else if (gpu->algorithm.type == ALGO_LYRA2Z && i == 1) {
if (clState->prebuilt) {
} else if ((gpu->algorithm.type == ALGO_LYRA2Z || gpu->algorithm.type == ALGO_LYRA2Z_NAVI) && i == 1) {
if (clState->prebuilt || gpu->algorithm.type == ALGO_LYRA2Z_NAVI) {
const size_t off2[] = { 0, 0, *p_global_work_offset };
const size_t gws[] = { 4, 4, globalThreads[0] / 2 };
const size_t expand[] = { 4, 4, 16 };
Expand All @@ -1503,8 +1503,8 @@ if (gpu->algorithm.type != ALGO_MTP && gpu->algorithm.type != ALGO_YESCRYPT_NAVI
const size_t expand[] = { 4, 5 };
status = clEnqueueNDRangeKernel(clState->commandQueue, clState->extra_kernels[i], 2, off2, gws, expand, 0, NULL, NULL); // lyra 4w monolithic
}
} else if (gpu->algorithm.type == ALGO_ALLIUM && (i == 2 || i == 6)) {
if (clState->prebuilt) {
} else if ((gpu->algorithm.type == ALGO_ALLIUM || gpu->algorithm.type == ALGO_ALLIUM_NAVI) && (i == 2 || i == 6)) {
if (clState->prebuilt || gpu->algorithm.type == ALGO_ALLIUM_NAVI) {
const size_t off2[] = { 0, 0, *p_global_work_offset };
const size_t gws[] = { 4, 4, globalThreads[0] / 2 };
const size_t expand[] = { 4, 4, 16 };
Expand Down
Loading

0 comments on commit 7c484e1

Please sign in to comment.