mirror of
https://github.com/GOSTSec/sgminer
synced 2025-03-13 06:01:03 +00:00
Merge branch 'master' into pools
This commit is contained in:
Noel Maersk
commit
1ebcb07302
@ -1,6 +1,6 @@
|
||||
# Authors
|
||||
|
||||
## Current
|
||||
## Core
|
||||
|
||||
* Scrypt-only refactor and maintenance: Noel Maersk <veox@wemakethings.net> LfxRFgXzA13TSTxgFGUFEtumv5ShGzAmLw
|
||||
* Core: Martin Danielsen <kalroth@gmail.com> 1DNBcSEENBwDKrcTyTW61ezWhzsPy5imkn
|
||||
@ -9,6 +9,13 @@
|
||||
* API: Andrew Smith <kan0i {at} kano-kun [dot] net> 1Jjk2LmktEQKnv8r2cZ9MvLiZwZ9gxabKm
|
||||
|
||||
|
||||
## OpenCL kernels
|
||||
|
||||
* scrypt: Colin Percival
|
||||
* zuikkis: Zuikkis LeXck7EYgxyjw13zNDxZFmmgmWffFvhmSh
|
||||
* alexkarnew/alexkarold: Alexey Karimov LMqRcHdwnZtTMH6c2kWoxSoKM5KySfaP5C
|
||||
|
||||
|
||||
## Bug fixes
|
||||
|
||||
* Gabriel Devenyi
|
||||
|
||||
@ -37,7 +37,7 @@ sgminer_SOURCES += ocl.c ocl.h
|
||||
sgminer_SOURCES += findnonce.c findnonce.h
|
||||
sgminer_SOURCES += adl.c adl.h adl_functions.h
|
||||
sgminer_SOURCES += scrypt.c scrypt.h
|
||||
sgminer_SOURCES += *.cl
|
||||
sgminer_SOURCES += kernel/*.cl
|
||||
|
||||
bin_SCRIPTS = $(top_srcdir)/*.cl
|
||||
bin_SCRIPTS = $(top_srcdir)/kernel/*.cl
|
||||
|
||||
|
||||
@ -18,9 +18,10 @@ License: GPLv3. See `COPYING` for details.
|
||||
Documentation is available in directory `doc`. For details on several topics, see:
|
||||
|
||||
* `API` for the RPC API specification;
|
||||
* `SCRYPT` for how to find the right balance in GPU configuration to mine Scrypt-based coins effectively;
|
||||
* `FAQ` for frequently asked questions;
|
||||
* `GPU` for semi-obsolete information on GPU configuration options and mining SHA256d-based coins;
|
||||
* `KERNEL.md` for OpenCL kernel-related information;
|
||||
* `MINING` for how to find the right balance in GPU configuration to mine Scrypt-based coins effectively;
|
||||
* `windows-build.txt` for information on how to build on Windows.
|
||||
|
||||
Note that **most of the documentation is outdated**. If you want to contribute, fork this repository, update as needed, and submit a pull request.
|
||||
|
||||
137
adl.c
137
adl.c
@ -63,7 +63,7 @@ static void * __stdcall ADL_Main_Memory_Alloc(int iSize)
|
||||
// Optional Memory de-allocation function
|
||||
static void __stdcall ADL_Main_Memory_Free (void **lpBuffer)
|
||||
{
|
||||
if (*lpBuffer) {
|
||||
if (*lpBuffer != NULL) {
|
||||
free (*lpBuffer);
|
||||
*lpBuffer = NULL;
|
||||
}
|
||||
@ -77,40 +77,42 @@ static void *GetProcAddress(void *pLibrary, const char *name)
|
||||
}
|
||||
#endif
|
||||
|
||||
static ADL_MAIN_CONTROL_CREATE ADL_Main_Control_Create;
|
||||
static ADL_MAIN_CONTROL_DESTROY ADL_Main_Control_Destroy;
|
||||
static ADL_MAIN_CONTROL_CREATE ADL_Main_Control_Create;
|
||||
static ADL_MAIN_CONTROL_DESTROY ADL_Main_Control_Destroy;
|
||||
static ADL_ADAPTER_NUMBEROFADAPTERS_GET ADL_Adapter_NumberOfAdapters_Get;
|
||||
static ADL_ADAPTER_ADAPTERINFO_GET ADL_Adapter_AdapterInfo_Get;
|
||||
static ADL_ADAPTER_ID_GET ADL_Adapter_ID_Get;
|
||||
static ADL_MAIN_CONTROL_REFRESH ADL_Main_Control_Refresh;
|
||||
static ADL_ADAPTER_VIDEOBIOSINFO_GET ADL_Adapter_VideoBiosInfo_Get;
|
||||
static ADL_ADAPTER_ADAPTERINFO_GET ADL_Adapter_AdapterInfo_Get;
|
||||
static ADL_ADAPTER_ID_GET ADL_Adapter_ID_Get;
|
||||
static ADL_MAIN_CONTROL_REFRESH ADL_Main_Control_Refresh;
|
||||
static ADL_ADAPTER_VIDEOBIOSINFO_GET ADL_Adapter_VideoBiosInfo_Get;
|
||||
static ADL_DISPLAY_DISPLAYINFO_GET ADL_Display_DisplayInfo_Get;
|
||||
static ADL_ADAPTER_ACCESSIBILITY_GET ADL_Adapter_Accessibility_Get;
|
||||
|
||||
static ADL_OVERDRIVE_CAPS ADL_Overdrive_Caps;
|
||||
static ADL_OVERDRIVE_CAPS ADL_Overdrive_Caps;
|
||||
|
||||
static ADL_OVERDRIVE5_TEMPERATURE_GET ADL_Overdrive5_Temperature_Get;
|
||||
static ADL_OVERDRIVE5_TEMPERATURE_GET ADL_Overdrive5_Temperature_Get;
|
||||
static ADL_OVERDRIVE5_CURRENTACTIVITY_GET ADL_Overdrive5_CurrentActivity_Get;
|
||||
static ADL_OVERDRIVE5_ODPARAMETERS_GET ADL_Overdrive5_ODParameters_Get;
|
||||
static ADL_OVERDRIVE5_FANSPEEDINFO_GET ADL_Overdrive5_FanSpeedInfo_Get;
|
||||
static ADL_OVERDRIVE5_FANSPEED_GET ADL_Overdrive5_FanSpeed_Get;
|
||||
static ADL_OVERDRIVE5_FANSPEED_SET ADL_Overdrive5_FanSpeed_Set;
|
||||
static ADL_OVERDRIVE5_ODPARAMETERS_GET ADL_Overdrive5_ODParameters_Get;
|
||||
static ADL_OVERDRIVE5_FANSPEEDINFO_GET ADL_Overdrive5_FanSpeedInfo_Get;
|
||||
static ADL_OVERDRIVE5_FANSPEED_GET ADL_Overdrive5_FanSpeed_Get;
|
||||
static ADL_OVERDRIVE5_FANSPEED_SET ADL_Overdrive5_FanSpeed_Set;
|
||||
static ADL_OVERDRIVE5_ODPERFORMANCELEVELS_GET ADL_Overdrive5_ODPerformanceLevels_Get;
|
||||
static ADL_OVERDRIVE5_ODPERFORMANCELEVELS_SET ADL_Overdrive5_ODPerformanceLevels_Set;
|
||||
static ADL_OVERDRIVE5_POWERCONTROL_GET ADL_Overdrive5_PowerControl_Get;
|
||||
static ADL_OVERDRIVE5_POWERCONTROL_SET ADL_Overdrive5_PowerControl_Set;
|
||||
static ADL_OVERDRIVE5_POWERCONTROL_GET ADL_Overdrive5_PowerControl_Get;
|
||||
static ADL_OVERDRIVE5_POWERCONTROL_SET ADL_Overdrive5_PowerControl_Set;
|
||||
static ADL_OVERDRIVE5_FANSPEEDTODEFAULT_SET ADL_Overdrive5_FanSpeedToDefault_Set;
|
||||
|
||||
static ADL_OVERDRIVE6_CAPABILITIES_GET ADL_Overdrive6_Capabilities_Get;
|
||||
static ADL_OVERDRIVE6_FANSPEED_GET ADL_Overdrive6_FanSpeed_Get;
|
||||
static ADL_OVERDRIVE6_THERMALCONTROLLER_CAPS ADL_Overdrive6_ThermalController_Caps;
|
||||
static ADL_OVERDRIVE6_TEMPERATURE_GET ADL_Overdrive6_Temperature_Get;
|
||||
static ADL_OVERDRIVE6_STATEINFO_GET ADL_Overdrive6_StateInfo_Get;
|
||||
static ADL_OVERDRIVE6_CURRENTSTATUS_GET ADL_Overdrive6_CurrentStatus_Get;
|
||||
static ADL_OVERDRIVE6_POWERCONTROL_CAPS ADL_Overdrive6_PowerControl_Caps;
|
||||
static ADL_OVERDRIVE6_POWERCONTROLINFO_GET ADL_Overdrive6_PowerControlInfo_Get;
|
||||
static ADL_OVERDRIVE6_POWERCONTROL_GET ADL_Overdrive6_PowerControl_Get;
|
||||
static ADL_OVERDRIVE6_FANSPEED_SET ADL_Overdrive6_FanSpeed_Set;
|
||||
static ADL_OVERDRIVE6_STATE_SET ADL_Overdrive6_State_Set;
|
||||
static ADL_OVERDRIVE6_POWERCONTROL_SET ADL_Overdrive6_PowerControl_Set;
|
||||
static ADL_OVERDRIVE6_CAPABILITIES_GET ADL_Overdrive6_Capabilities_Get;
|
||||
static ADL_OVERDRIVE6_FANSPEED_GET ADL_Overdrive6_FanSpeed_Get;
|
||||
static ADL_OVERDRIVE6_THERMALCONTROLLER_CAPS ADL_Overdrive6_ThermalController_Caps;
|
||||
static ADL_OVERDRIVE6_TEMPERATURE_GET ADL_Overdrive6_Temperature_Get;
|
||||
static ADL_OVERDRIVE6_STATEINFO_GET ADL_Overdrive6_StateInfo_Get;
|
||||
static ADL_OVERDRIVE6_CURRENTSTATUS_GET ADL_Overdrive6_CurrentStatus_Get;
|
||||
static ADL_OVERDRIVE6_POWERCONTROL_CAPS ADL_Overdrive6_PowerControl_Caps;
|
||||
static ADL_OVERDRIVE6_POWERCONTROLINFO_GET ADL_Overdrive6_PowerControlInfo_Get;
|
||||
static ADL_OVERDRIVE6_POWERCONTROL_GET ADL_Overdrive6_PowerControl_Get;
|
||||
static ADL_OVERDRIVE6_FANSPEED_SET ADL_Overdrive6_FanSpeed_Set;
|
||||
static ADL_OVERDRIVE6_STATE_SET ADL_Overdrive6_State_Set;
|
||||
static ADL_OVERDRIVE6_POWERCONTROL_SET ADL_Overdrive6_PowerControl_Set;
|
||||
|
||||
#if defined (LINUX)
|
||||
static void *hDLL; // Handle to .so library
|
||||
@ -119,6 +121,7 @@ static ADL_OVERDRIVE6_POWERCONTROL_SET ADL_Overdrive6_PowerControl_Set;
|
||||
#endif
|
||||
static int iNumberAdapters;
|
||||
static LPAdapterInfo lpInfo = NULL;
|
||||
static LPADLDisplayInfo lpAdlDisplayInfo = NULL;
|
||||
|
||||
int set_fanspeed(int gpu, int iFanSpeed);
|
||||
static float __gpu_temp(struct gpu_adl *ga);
|
||||
@ -166,8 +169,10 @@ static bool init_overdrive5()
|
||||
!ADL_Overdrive5_ODPerformanceLevels_Get || !ADL_Overdrive5_ODPerformanceLevels_Set ||
|
||||
!ADL_Overdrive5_PowerControl_Get || !ADL_Overdrive5_PowerControl_Set ||
|
||||
!ADL_Overdrive5_FanSpeedToDefault_Set) {
|
||||
applog(LOG_WARNING, "ATI ADL Overdrive5's API is missing");
|
||||
return false;
|
||||
applog(LOG_WARNING, "ATI ADL Overdrive5's API is missing or broken.");
|
||||
return false;
|
||||
} else {
|
||||
applog(LOG_INFO, "ATI ADL Overdrive5 API found.");
|
||||
}
|
||||
|
||||
return true;
|
||||
@ -194,8 +199,10 @@ static bool init_overdrive6()
|
||||
!ADL_Overdrive6_PowerControl_Caps || !ADL_Overdrive6_PowerControlInfo_Get ||
|
||||
!ADL_Overdrive6_PowerControl_Get || !ADL_Overdrive6_FanSpeed_Set ||
|
||||
!ADL_Overdrive6_State_Set || !ADL_Overdrive6_PowerControl_Set) {
|
||||
applog(LOG_WARNING, "ATI ADL Overdrive6's API is missing");
|
||||
applog(LOG_WARNING, "ATI ADL Overdrive6's API is missing or broken.");
|
||||
return false;
|
||||
} else {
|
||||
applog(LOG_INFO, "ATI ADL Overdrive6 API found.");
|
||||
}
|
||||
|
||||
return true;
|
||||
@ -215,40 +222,46 @@ static bool prepare_adl(void)
|
||||
hDLL = LoadLibrary("atiadlxy.dll");
|
||||
#endif
|
||||
if (hDLL == NULL) {
|
||||
applog(LOG_INFO, "Unable to load ati adl library");
|
||||
applog(LOG_INFO, "Unable to load ATI ADL library.");
|
||||
return false;
|
||||
}
|
||||
ADL_Main_Control_Create = (ADL_MAIN_CONTROL_CREATE) GetProcAddress(hDLL,"ADL_Main_Control_Create");
|
||||
ADL_Main_Control_Destroy = (ADL_MAIN_CONTROL_DESTROY) GetProcAddress(hDLL,"ADL_Main_Control_Destroy");
|
||||
ADL_Adapter_NumberOfAdapters_Get = (ADL_ADAPTER_NUMBEROFADAPTERS_GET) GetProcAddress(hDLL,"ADL_Adapter_NumberOfAdapters_Get");
|
||||
ADL_Adapter_AdapterInfo_Get = (ADL_ADAPTER_ADAPTERINFO_GET) GetProcAddress(hDLL,"ADL_Adapter_AdapterInfo_Get");
|
||||
ADL_Display_DisplayInfo_Get = (ADL_DISPLAY_DISPLAYINFO_GET) GetProcAddress(hDLL,"ADL_Display_DisplayInfo_Get");
|
||||
ADL_Adapter_ID_Get = (ADL_ADAPTER_ID_GET) GetProcAddress(hDLL,"ADL_Adapter_ID_Get");
|
||||
ADL_Main_Control_Refresh = (ADL_MAIN_CONTROL_REFRESH) GetProcAddress(hDLL, "ADL_Main_Control_Refresh");
|
||||
ADL_Adapter_VideoBiosInfo_Get = (ADL_ADAPTER_VIDEOBIOSINFO_GET)GetProcAddress(hDLL,"ADL_Adapter_VideoBiosInfo_Get");
|
||||
ADL_Overdrive_Caps = (ADL_OVERDRIVE_CAPS)GetProcAddress(hDLL, "ADL_Overdrive_Caps");
|
||||
|
||||
ADL_Adapter_Accessibility_Get = (ADL_ADAPTER_ACCESSIBILITY_GET)GetProcAddress(hDLL, "ADL_Adapter_Accessibility_Get");
|
||||
|
||||
if (!ADL_Main_Control_Create || !ADL_Main_Control_Destroy ||
|
||||
!ADL_Adapter_NumberOfAdapters_Get || !ADL_Adapter_AdapterInfo_Get ||
|
||||
!ADL_Display_DisplayInfo_Get ||
|
||||
!ADL_Adapter_ID_Get || !ADL_Main_Control_Refresh ||
|
||||
!ADL_Adapter_VideoBiosInfo_Get || !ADL_Overdrive_Caps) {
|
||||
applog(LOG_WARNING, "ATI ADL's API is missing");
|
||||
applog(LOG_WARNING, "ATI ADL API is missing or broken.");
|
||||
return false;
|
||||
}
|
||||
|
||||
// Initialise ADL. The second parameter is 1, which means:
|
||||
// retrieve adapter information only for adapters that are physically present and enabled in the system
|
||||
result = ADL_Main_Control_Create (ADL_Main_Memory_Alloc, 1);
|
||||
result = ADL_Main_Control_Create(ADL_Main_Memory_Alloc, 1);
|
||||
if (result != ADL_OK) {
|
||||
applog(LOG_INFO, "ADL Initialisation Error! Error %d!", result);
|
||||
applog(LOG_INFO, "ADL initialisation error: %d!", result);
|
||||
return false;
|
||||
}
|
||||
|
||||
result = ADL_Main_Control_Refresh();
|
||||
if (result != ADL_OK) {
|
||||
applog(LOG_INFO, "ADL Refresh Error! Error %d!", result);
|
||||
applog(LOG_INFO, "ADL refresh error: %d!", result);
|
||||
return false;
|
||||
}
|
||||
|
||||
init_overdrive5();
|
||||
init_overdrive6(); // FIXME: don't if ADL6 is not present
|
||||
|
||||
return true;
|
||||
}
|
||||
@ -258,6 +271,10 @@ void init_adl(int nDevs)
|
||||
int result, i, j, devices = 0, last_adapter = -1, gpu = 0, dummy = 0;
|
||||
struct gpu_adapters adapters[MAX_GPUDEVICES], vadapters[MAX_GPUDEVICES];
|
||||
bool devs_match = true;
|
||||
ADLBiosInfo BiosInfo;
|
||||
int iNumDisplays;
|
||||
|
||||
applog(LOG_INFO, "Number of ADL devices: %d", nDevs);
|
||||
|
||||
if (unlikely(pthread_mutex_init(&adl_lock, NULL))) {
|
||||
applog(LOG_ERR, "Failed to init adl_lock in init_adl");
|
||||
@ -290,38 +307,53 @@ void init_adl(int nDevs)
|
||||
return;
|
||||
}
|
||||
|
||||
applog(LOG_INFO, "Found %d ADL adapters", iNumberAdapters);
|
||||
|
||||
/* Iterate over iNumberAdapters and find the lpAdapterID of real devices */
|
||||
for (i = 0; i < iNumberAdapters; i++) {
|
||||
int iAdapterIndex;
|
||||
int lpAdapterID;
|
||||
|
||||
iAdapterIndex = lpInfo[i].iAdapterIndex;
|
||||
|
||||
/* Get unique identifier of the adapter, 0 means not AMD */
|
||||
result = ADL_Adapter_ID_Get(iAdapterIndex, &lpAdapterID);
|
||||
|
||||
if (ADL_Adapter_VideoBiosInfo_Get(iAdapterIndex, &BiosInfo) == ADL_ERR) {
|
||||
applog(LOG_INFO, "ADL index %d, id %d - FAILED to get BIOS info", iAdapterIndex, lpAdapterID);
|
||||
} else {
|
||||
applog(LOG_INFO, "ADL index %d, id %d - BIOS partno.: %s, version: %s, date: %s", iAdapterIndex, lpAdapterID, BiosInfo.strPartNumber, BiosInfo.strVersion, BiosInfo.strDate);
|
||||
}
|
||||
|
||||
if (result != ADL_OK) {
|
||||
applog(LOG_INFO, "Failed to ADL_Adapter_ID_Get. Error %d", result);
|
||||
if (result == -10)
|
||||
applog(LOG_INFO, "This error says the device is not enabled");
|
||||
applog(LOG_INFO, "(Device is not enabled.)");
|
||||
continue;
|
||||
}
|
||||
|
||||
/* Each adapter may have multiple entries */
|
||||
if (lpAdapterID == last_adapter)
|
||||
if (lpAdapterID == last_adapter) {
|
||||
continue;
|
||||
}
|
||||
|
||||
applog(LOG_DEBUG, "GPU %d "
|
||||
"iAdapterIndex %d "
|
||||
"strUDID %s "
|
||||
"iBusNumber %d "
|
||||
"iDeviceNumber %d "
|
||||
"iFunctionNumber %d "
|
||||
"iVendorID %d "
|
||||
"strAdapterName %s ",
|
||||
applog(LOG_INFO, "GPU %d assigned: "
|
||||
"iAdapterIndex:%d "
|
||||
"iPresent:%d "
|
||||
"strUDID:%s "
|
||||
"iBusNumber:%d "
|
||||
"iDeviceNumber:%d "
|
||||
"iDrvIndex:%d "
|
||||
"iFunctionNumber:%d "
|
||||
"iVendorID:%d "
|
||||
"name:%s",
|
||||
devices,
|
||||
iAdapterIndex,
|
||||
lpInfo[i].iAdapterIndex,
|
||||
lpInfo[i].iPresent,
|
||||
lpInfo[i].strUDID,
|
||||
lpInfo[i].iBusNumber,
|
||||
lpInfo[i].iDeviceNumber,
|
||||
lpInfo[i].iDrvIndex,
|
||||
lpInfo[i].iFunctionNumber,
|
||||
lpInfo[i].iVendorID,
|
||||
lpInfo[i].strAdapterName);
|
||||
@ -365,8 +397,9 @@ void init_adl(int nDevs)
|
||||
if (gpus[i].mapped) {
|
||||
vadapters[gpus[i].virtual_adl].virtual_gpu = i;
|
||||
applog(LOG_INFO, "Mapping OpenCL device %d to ADL device %d", i, gpus[i].virtual_adl);
|
||||
} else
|
||||
} else {
|
||||
gpus[i].virtual_adl = i;
|
||||
}
|
||||
}
|
||||
|
||||
if (!devs_match) {
|
||||
@ -451,6 +484,8 @@ void init_adl(int nDevs)
|
||||
ga->twin = NULL;
|
||||
ga->def_fan_valid = false;
|
||||
|
||||
applog(LOG_INFO, "ADL GPU %d is Adapter index %d and maps to adapter id %d", ga->gpu, ga->iAdapterIndex, ga->lpAdapterID);
|
||||
|
||||
if (ADL_Adapter_VideoBiosInfo_Get(iAdapterIndex, &BiosInfo) != ADL_ERR)
|
||||
applog(LOG_INFO, "GPU %d BIOS partno.: %s, version: %s, date: %s", gpu, BiosInfo.strPartNumber, BiosInfo.strVersion, BiosInfo.strDate);
|
||||
|
||||
@ -545,7 +580,7 @@ void init_adl(int nDevs)
|
||||
ga->has_fanspeed = true;
|
||||
|
||||
/* Save the fanspeed values as defaults in case we reset later */
|
||||
if (ADL_Overdrive5_FanSpeed_Get(ga->iAdapterIndex, 0, &ga->DefFanSpeedValue) != ADL_OK)
|
||||
if (ADL_Overdrive5_FanSpeed_Get(iAdapterIndex, 0, &ga->DefFanSpeedValue) != ADL_OK)
|
||||
applog(LOG_INFO, "Failed to ADL_Overdrive5_FanSpeed_Get for default value");
|
||||
else
|
||||
ga->def_fan_valid = true;
|
||||
@ -556,12 +591,12 @@ void init_adl(int nDevs)
|
||||
gpus[gpu].gpu_fan = 85; /* Set a nominal upper limit of 85% */
|
||||
|
||||
/* Not fatal if powercontrol get fails */
|
||||
if (ADL_Overdrive5_PowerControl_Get(ga->iAdapterIndex, &ga->iPercentage, &dummy) != ADL_OK)
|
||||
if (ADL_Overdrive5_PowerControl_Get(iAdapterIndex, &ga->iPercentage, &dummy) != ADL_OK)
|
||||
applog(LOG_INFO, "Failed to ADL_Overdrive5_PowerControl_get");
|
||||
|
||||
if (gpus[gpu].gpu_powertune) {
|
||||
ADL_Overdrive5_PowerControl_Set(ga->iAdapterIndex, gpus[gpu].gpu_powertune);
|
||||
ADL_Overdrive5_PowerControl_Get(ga->iAdapterIndex, &ga->iPercentage, &dummy);
|
||||
ADL_Overdrive5_PowerControl_Set(iAdapterIndex, gpus[gpu].gpu_powertune);
|
||||
ADL_Overdrive5_PowerControl_Get(iAdapterIndex, &ga->iPercentage, &dummy);
|
||||
ga->managed = true;
|
||||
}
|
||||
|
||||
|
||||
@ -317,8 +317,10 @@ fi
|
||||
|
||||
AC_DEFINE_UNQUOTED([SGMINER_PREFIX], ["$prefix/bin"], [Path to sgminer install])
|
||||
|
||||
AC_DEFINE_UNQUOTED([SCRYPT_KERNNAME], ["scrypt130511"], [Filename for scrypt kernel])
|
||||
|
||||
AC_DEFINE_UNQUOTED([ALEXKARNEW_KERNNAME], ["alexkarnew"], [Filename for Alexey Karimov's optimised kernel for Catalyst >=13.4])
|
||||
AC_DEFINE_UNQUOTED([ALEXKAROLD_KERNNAME], ["alexkarold"], [Filename for Alexey Karimov's optimised kernel for Catalyst <13.4])
|
||||
AC_DEFINE_UNQUOTED([CKOLIVAS_KERNNAME], ["ckolivas"], [Filename for original scrypt kernel])
|
||||
AC_DEFINE_UNQUOTED([ZUIKKIS_KERNNAME], ["zuikkis"], [Filename for Zuikkis' optimised kernel])
|
||||
|
||||
AC_SUBST(OPENCL_LIBS)
|
||||
AC_SUBST(OPENCL_FLAGS)
|
||||
|
||||
86
doc/KERNEL.md
Normal file
86
doc/KERNEL.md
Normal file
@ -0,0 +1,86 @@
|
||||
# Kernels
|
||||
|
||||
## Available OpenCL kernels
|
||||
|
||||
See directory `kernel`.
|
||||
|
||||
|
||||
## Parameter configuration
|
||||
|
||||
### Common
|
||||
|
||||
In general, switching kernels requires reconfiguring mining parameters,
|
||||
such as (but not necessarily limited to) `thread-concurrency`, `intensity`,
|
||||
`gpu-engine` and `gpu-memclock`.
|
||||
|
||||
A description of how to do this is available in `doc/MINING`.
|
||||
|
||||
|
||||
### alexkarnew
|
||||
|
||||
Alaxey Karimov's optimised kernel, based on `ckolivas`. For Catalyst >=13.4.
|
||||
|
||||
Only supports `vectors=1`.
|
||||
|
||||
[Announcement](https://litecointalk.org/index.php?topic=4082.0).
|
||||
|
||||
|
||||
### alexkarold
|
||||
|
||||
Alaxey Karimov's optimised kernel, based on `ckolivas`. For Catalyst <13.4.
|
||||
|
||||
Only supports `vectors=1`.
|
||||
|
||||
[Announcement](https://litecointalk.org/index.php?topic=4082.0).
|
||||
|
||||
|
||||
### ckolivas
|
||||
|
||||
The original Colin Percival `scrypt` kernel, maintained for a long time by
|
||||
Con Kolivas in `cgminer` and renamed to reflect the fact.
|
||||
|
||||
Only supports `vectors=1`.
|
||||
|
||||
|
||||
### zuikkis
|
||||
|
||||
Zuikkis' optimised kernel, based on `ckolivas`.
|
||||
|
||||
Only supports `vectors=1`, `lookup-gap=2` and `worksize=256`.
|
||||
|
||||
[Announcement](https://litecointalk.org/index.php?topic=6058.msg90873#msg90873).
|
||||
|
||||
|
||||
## Submitting new kernels
|
||||
|
||||
### Requirements
|
||||
|
||||
* OpenCL source code only, licenced under GPLv3 (or later).
|
||||
* Not hard-coded for a specific GPU model or manufacturer.
|
||||
* Known limitations and any specific configuration quirks must be mentioned.
|
||||
|
||||
|
||||
### Procedure
|
||||
|
||||
1. Copy the kernel you wish to modify and commit it verbatim.
|
||||
|
||||
This way, it is easy to verify that there are no hidden changes. Note in
|
||||
the commit message which kernel is used as a base.
|
||||
|
||||
2. Make changes to the kernel. Commit them.
|
||||
|
||||
This allows to produce a diff that makes sense.
|
||||
|
||||
3. Search for KL_CKOLIVAS and CKOLIVAS_KERNNAME in the top-level source
|
||||
directory and make additions to the listed files in order to integrate
|
||||
the new kernel.
|
||||
|
||||
Now it can be selected when starting via the `--kernel` argument or
|
||||
`kernel` configuration option.
|
||||
|
||||
4. Add yourself to the "kernels" section in `AUTHORS.md`. Keep it short.
|
||||
|
||||
5. Submit a pull request on GitHub, or file it at the issue tracker,
|
||||
outlining the changes made, known limitations, and tested GPUs. List
|
||||
your git repository and branch name. The current repository and issue
|
||||
tracker links should be in `README.md`.
|
||||
@ -196,8 +196,14 @@ char *set_thread_concurrency(char *arg)
|
||||
|
||||
static enum cl_kernels select_kernel(char *arg)
|
||||
{
|
||||
if (!strcmp(arg, "scrypt"))
|
||||
return KL_SCRYPT;
|
||||
if (!strcmp(arg, ALEXKARNEW_KERNNAME))
|
||||
return KL_ALEXKARNEW;
|
||||
if (!strcmp(arg, ALEXKAROLD_KERNNAME))
|
||||
return KL_ALEXKAROLD;
|
||||
if (!strcmp(arg, CKOLIVAS_KERNNAME))
|
||||
return KL_CKOLIVAS;
|
||||
if (!strcmp(arg, ZUIKKIS_KERNNAME))
|
||||
return KL_ZUIKKIS;
|
||||
|
||||
return KL_NONE;
|
||||
}
|
||||
@ -1291,8 +1297,17 @@ static bool opencl_thread_prepare(struct thr_info *thr)
|
||||
if (!cgpu->kname)
|
||||
{
|
||||
switch (clStates[i]->chosen_kernel) {
|
||||
case KL_SCRYPT:
|
||||
cgpu->kname = "scrypt";
|
||||
case KL_ALEXKARNEW:
|
||||
cgpu->kname = ALEXKARNEW_KERNNAME;
|
||||
break;
|
||||
case KL_ALEXKAROLD:
|
||||
cgpu->kname = ALEXKAROLD_KERNNAME;
|
||||
break;
|
||||
case KL_CKOLIVAS:
|
||||
cgpu->kname = CKOLIVAS_KERNNAME;
|
||||
break;
|
||||
case KL_ZUIKKIS:
|
||||
cgpu->kname = ZUIKKIS_KERNNAME;
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
@ -1322,11 +1337,15 @@ static bool opencl_thread_init(struct thr_info *thr)
|
||||
}
|
||||
|
||||
switch (clState->chosen_kernel) {
|
||||
case KL_SCRYPT:
|
||||
thrdata->queue_kernel_parameters = &queue_scrypt_kernel;
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
case KL_ALEXKARNEW:
|
||||
case KL_ALEXKAROLD:
|
||||
case KL_CKOLIVAS:
|
||||
case KL_ZUIKKIS:
|
||||
thrdata->queue_kernel_parameters = &queue_scrypt_kernel;
|
||||
break;
|
||||
default:
|
||||
applog(LOG_ERR, "Failed to choose kernel in opencl_thread_init");
|
||||
break;
|
||||
}
|
||||
|
||||
thrdata->res = calloc(buffersize, 1);
|
||||
|
||||
858
kernel/alexkarnew.cl
Normal file
858
kernel/alexkarnew.cl
Normal file
@ -0,0 +1,858 @@
|
||||
/*-
|
||||
* Copyright 2009 Colin Percival, 2011 ArtForz, 2011 pooler, 2012 mtrlt,
|
||||
* 2012-2013 Con Kolivas, 2013 Alexey Karimov.
|
||||
* 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.
|
||||
*
|
||||
* THIS SOFTWARE IS PROVIDED BY THE AUTHOR AND CONTRIBUTORS ``AS IS'' AND
|
||||
* ANY EXPRESS 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 AUTHOR OR 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.
|
||||
*
|
||||
* This file was originally written by Colin Percival as part of the Tarsnap
|
||||
* online backup system.
|
||||
*/
|
||||
|
||||
__constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 };
|
||||
__constant uint K[] = {
|
||||
0x428a2f98U,
|
||||
0x71374491U,
|
||||
0xb5c0fbcfU,
|
||||
0xe9b5dba5U,
|
||||
0x3956c25bU,
|
||||
0x59f111f1U,
|
||||
0x923f82a4U,
|
||||
0xab1c5ed5U,
|
||||
0xd807aa98U,
|
||||
0x12835b01U,
|
||||
0x243185beU, // 10
|
||||
0x550c7dc3U,
|
||||
0x72be5d74U,
|
||||
0x80deb1feU,
|
||||
0x9bdc06a7U,
|
||||
0xe49b69c1U,
|
||||
0xefbe4786U,
|
||||
0x0fc19dc6U,
|
||||
0x240ca1ccU,
|
||||
0x2de92c6fU,
|
||||
0x4a7484aaU, // 20
|
||||
0x5cb0a9dcU,
|
||||
0x76f988daU,
|
||||
0x983e5152U,
|
||||
0xa831c66dU,
|
||||
0xb00327c8U,
|
||||
0xbf597fc7U,
|
||||
0xc6e00bf3U,
|
||||
0xd5a79147U,
|
||||
0x06ca6351U,
|
||||
0x14292967U, // 30
|
||||
0x27b70a85U,
|
||||
0x2e1b2138U,
|
||||
0x4d2c6dfcU,
|
||||
0x53380d13U,
|
||||
0x650a7354U,
|
||||
0x766a0abbU,
|
||||
0x81c2c92eU,
|
||||
0x92722c85U,
|
||||
0xa2bfe8a1U,
|
||||
0xa81a664bU, // 40
|
||||
0xc24b8b70U,
|
||||
0xc76c51a3U,
|
||||
0xd192e819U,
|
||||
0xd6990624U,
|
||||
0xf40e3585U,
|
||||
0x106aa070U,
|
||||
0x19a4c116U,
|
||||
0x1e376c08U,
|
||||
0x2748774cU,
|
||||
0x34b0bcb5U, // 50
|
||||
0x391c0cb3U,
|
||||
0x4ed8aa4aU,
|
||||
0x5b9cca4fU,
|
||||
0x682e6ff3U,
|
||||
0x748f82eeU,
|
||||
0x78a5636fU,
|
||||
0x84c87814U,
|
||||
0x8cc70208U,
|
||||
0x90befffaU,
|
||||
0xa4506cebU, // 60
|
||||
0xbef9a3f7U,
|
||||
0xc67178f2U,
|
||||
0x98c7e2a2U,
|
||||
0xfc08884dU,
|
||||
0xcd2a11aeU,
|
||||
0x510e527fU,
|
||||
0x9b05688cU,
|
||||
0xC3910C8EU,
|
||||
0xfb6feee7U,
|
||||
0x2a01a605U, // 70
|
||||
0x0c2e12e0U,
|
||||
0x4498517BU,
|
||||
0x6a09e667U,
|
||||
0xa4ce148bU,
|
||||
0x95F61999U,
|
||||
0xc19bf174U,
|
||||
0xBB67AE85U,
|
||||
0x3C6EF372U,
|
||||
0xA54FF53AU,
|
||||
0x1F83D9ABU, // 80
|
||||
0x5BE0CD19U,
|
||||
0x5C5C5C5CU,
|
||||
0x36363636U,
|
||||
0x80000000U,
|
||||
0x000003FFU,
|
||||
0x00000280U,
|
||||
0x000004a0U,
|
||||
0x00000300U
|
||||
};
|
||||
|
||||
#define rotl(x,y) rotate(x,y)
|
||||
#define Ch(x,y,z) bitselect(z,y,x)
|
||||
#define Maj(x,y,z) Ch((x^z),y,z)
|
||||
|
||||
#define EndianSwap(n) (rotl(n & ES[0], 24U)|rotl(n & ES[1], 8U))
|
||||
|
||||
#define Tr2(x) (rotl(x, 30U) ^ rotl(x, 19U) ^ rotl(x, 10U))
|
||||
#define Tr1(x) (rotl(x, 26U) ^ rotl(x, 21U) ^ rotl(x, 7U))
|
||||
#define Wr2(x) (rotl(x, 25U) ^ rotl(x, 14U) ^ (x>>3U))
|
||||
#define Wr1(x) (rotl(x, 15U) ^ rotl(x, 13U) ^ (x>>10U))
|
||||
|
||||
#define RND(a, b, c, d, e, f, g, h, k) \
|
||||
h += Tr1(e); \
|
||||
h += Ch(e, f, g); \
|
||||
h += k; \
|
||||
d += h; \
|
||||
h += Tr2(a); \
|
||||
h += Maj(a, b, c);
|
||||
|
||||
void SHA256(uint4*restrict state0,uint4*restrict state1, const uint4 block0, const uint4 block1, const uint4 block2, const uint4 block3)
|
||||
{
|
||||
uint4 S0 = *state0;
|
||||
uint4 S1 = *state1;
|
||||
|
||||
#define A S0.x
|
||||
#define B S0.y
|
||||
#define C S0.z
|
||||
#define D S0.w
|
||||
#define E S1.x
|
||||
#define F S1.y
|
||||
#define G S1.z
|
||||
#define H S1.w
|
||||
|
||||
uint4 W[4];
|
||||
|
||||
W[ 0].x = block0.x;
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[0]);
|
||||
W[ 0].y = block0.y;
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[1]);
|
||||
W[ 0].z = block0.z;
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[2]);
|
||||
W[ 0].w = block0.w;
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[3]);
|
||||
|
||||
W[ 1].x = block1.x;
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[4]);
|
||||
W[ 1].y = block1.y;
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[5]);
|
||||
W[ 1].z = block1.z;
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[6]);
|
||||
W[ 1].w = block1.w;
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[7]);
|
||||
|
||||
W[ 2].x = block2.x;
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[8]);
|
||||
W[ 2].y = block2.y;
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[9]);
|
||||
W[ 2].z = block2.z;
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[10]);
|
||||
W[ 2].w = block2.w;
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[11]);
|
||||
|
||||
W[ 3].x = block3.x;
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[12]);
|
||||
W[ 3].y = block3.y;
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[13]);
|
||||
W[ 3].z = block3.z;
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[14]);
|
||||
W[ 3].w = block3.w;
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[76]);
|
||||
|
||||
W[ 0].x += Wr1(W[ 3].z) + W[ 2].y + Wr2(W[ 0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[15]);
|
||||
|
||||
W[ 0].y += Wr1(W[ 3].w) + W[ 2].z + Wr2(W[ 0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[16]);
|
||||
|
||||
W[ 0].z += Wr1(W[ 0].x) + W[ 2].w + Wr2(W[ 0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[17]);
|
||||
|
||||
W[ 0].w += Wr1(W[ 0].y) + W[ 3].x + Wr2(W[ 1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[18]);
|
||||
|
||||
W[ 1].x += Wr1(W[ 0].z) + W[ 3].y + Wr2(W[ 1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[19]);
|
||||
|
||||
W[ 1].y += Wr1(W[ 0].w) + W[ 3].z + Wr2(W[ 1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[20]);
|
||||
|
||||
W[ 1].z += Wr1(W[ 1].x) + W[ 3].w + Wr2(W[ 1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[21]);
|
||||
|
||||
W[ 1].w += Wr1(W[ 1].y) + W[ 0].x + Wr2(W[ 2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[22]);
|
||||
|
||||
W[ 2].x += Wr1(W[ 1].z) + W[ 0].y + Wr2(W[ 2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[23]);
|
||||
|
||||
W[ 2].y += Wr1(W[ 1].w) + W[ 0].z + Wr2(W[ 2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[24]);
|
||||
|
||||
W[ 2].z += Wr1(W[ 2].x) + W[ 0].w + Wr2(W[ 2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[25]);
|
||||
|
||||
W[ 2].w += Wr1(W[ 2].y) + W[ 1].x + Wr2(W[ 3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[26]);
|
||||
|
||||
W[ 3].x += Wr1(W[ 2].z) + W[ 1].y + Wr2(W[ 3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[27]);
|
||||
|
||||
W[ 3].y += Wr1(W[ 2].w) + W[ 1].z + Wr2(W[ 3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[28]);
|
||||
|
||||
W[ 3].z += Wr1(W[ 3].x) + W[ 1].w + Wr2(W[ 3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[29]);
|
||||
|
||||
W[ 3].w += Wr1(W[ 3].y) + W[ 2].x + Wr2(W[ 0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[30]);
|
||||
|
||||
W[ 0].x += Wr1(W[ 3].z) + W[ 2].y + Wr2(W[ 0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[31]);
|
||||
|
||||
W[ 0].y += Wr1(W[ 3].w) + W[ 2].z + Wr2(W[ 0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[32]);
|
||||
|
||||
W[ 0].z += Wr1(W[ 0].x) + W[ 2].w + Wr2(W[ 0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[33]);
|
||||
|
||||
W[ 0].w += Wr1(W[ 0].y) + W[ 3].x + Wr2(W[ 1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[34]);
|
||||
|
||||
W[ 1].x += Wr1(W[ 0].z) + W[ 3].y + Wr2(W[ 1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[35]);
|
||||
|
||||
W[ 1].y += Wr1(W[ 0].w) + W[ 3].z + Wr2(W[ 1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[36]);
|
||||
|
||||
W[ 1].z += Wr1(W[ 1].x) + W[ 3].w + Wr2(W[ 1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[37]);
|
||||
|
||||
W[ 1].w += Wr1(W[ 1].y) + W[ 0].x + Wr2(W[ 2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[38]);
|
||||
|
||||
W[ 2].x += Wr1(W[ 1].z) + W[ 0].y + Wr2(W[ 2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[39]);
|
||||
|
||||
W[ 2].y += Wr1(W[ 1].w) + W[ 0].z + Wr2(W[ 2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[40]);
|
||||
|
||||
W[ 2].z += Wr1(W[ 2].x) + W[ 0].w + Wr2(W[ 2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[41]);
|
||||
|
||||
W[ 2].w += Wr1(W[ 2].y) + W[ 1].x + Wr2(W[ 3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[42]);
|
||||
|
||||
W[ 3].x += Wr1(W[ 2].z) + W[ 1].y + Wr2(W[ 3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[43]);
|
||||
|
||||
W[ 3].y += Wr1(W[ 2].w) + W[ 1].z + Wr2(W[ 3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[44]);
|
||||
|
||||
W[ 3].z += Wr1(W[ 3].x) + W[ 1].w + Wr2(W[ 3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[45]);
|
||||
|
||||
W[ 3].w += Wr1(W[ 3].y) + W[ 2].x + Wr2(W[ 0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[46]);
|
||||
|
||||
W[ 0].x += Wr1(W[ 3].z) + W[ 2].y + Wr2(W[ 0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[47]);
|
||||
|
||||
W[ 0].y += Wr1(W[ 3].w) + W[ 2].z + Wr2(W[ 0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[48]);
|
||||
|
||||
W[ 0].z += Wr1(W[ 0].x) + W[ 2].w + Wr2(W[ 0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[49]);
|
||||
|
||||
W[ 0].w += Wr1(W[ 0].y) + W[ 3].x + Wr2(W[ 1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[50]);
|
||||
|
||||
W[ 1].x += Wr1(W[ 0].z) + W[ 3].y + Wr2(W[ 1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[51]);
|
||||
|
||||
W[ 1].y += Wr1(W[ 0].w) + W[ 3].z + Wr2(W[ 1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[52]);
|
||||
|
||||
W[ 1].z += Wr1(W[ 1].x) + W[ 3].w + Wr2(W[ 1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[53]);
|
||||
|
||||
W[ 1].w += Wr1(W[ 1].y) + W[ 0].x + Wr2(W[ 2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[54]);
|
||||
|
||||
W[ 2].x += Wr1(W[ 1].z) + W[ 0].y + Wr2(W[ 2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[55]);
|
||||
|
||||
W[ 2].y += Wr1(W[ 1].w) + W[ 0].z + Wr2(W[ 2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[56]);
|
||||
|
||||
W[ 2].z += Wr1(W[ 2].x) + W[ 0].w + Wr2(W[ 2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[57]);
|
||||
|
||||
W[ 2].w += Wr1(W[ 2].y) + W[ 1].x + Wr2(W[ 3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[58]);
|
||||
|
||||
W[ 3].x += Wr1(W[ 2].z) + W[ 1].y + Wr2(W[ 3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[59]);
|
||||
|
||||
W[ 3].y += Wr1(W[ 2].w) + W[ 1].z + Wr2(W[ 3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[60]);
|
||||
|
||||
W[ 3].z += Wr1(W[ 3].x) + W[ 1].w + Wr2(W[ 3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[61]);
|
||||
|
||||
W[ 3].w += Wr1(W[ 3].y) + W[ 2].x + Wr2(W[ 0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[62]);
|
||||
|
||||
#undef A
|
||||
#undef B
|
||||
#undef C
|
||||
#undef D
|
||||
#undef E
|
||||
#undef F
|
||||
#undef G
|
||||
#undef H
|
||||
|
||||
*state0 += S0;
|
||||
*state1 += S1;
|
||||
}
|
||||
|
||||
void SHA256_fresh(uint4*restrict state0,uint4*restrict state1, const uint4 block0, const uint4 block1, const uint4 block2, const uint4 block3)
|
||||
{
|
||||
#define A (*state0).x
|
||||
#define B (*state0).y
|
||||
#define C (*state0).z
|
||||
#define D (*state0).w
|
||||
#define E (*state1).x
|
||||
#define F (*state1).y
|
||||
#define G (*state1).z
|
||||
#define H (*state1).w
|
||||
|
||||
uint4 W[4];
|
||||
|
||||
W[0].x = block0.x;
|
||||
D= K[63] +W[0].x;
|
||||
H= K[64] +W[0].x;
|
||||
|
||||
W[0].y = block0.y;
|
||||
C= K[65] +Tr1(D)+Ch(D, K[66], K[67])+W[0].y;
|
||||
G= K[68] +C+Tr2(H)+Ch(H, K[69] ,K[70]);
|
||||
|
||||
W[0].z = block0.z;
|
||||
B= K[71] +Tr1(C)+Ch(C,D,K[66])+W[0].z;
|
||||
F= K[72] +B+Tr2(G)+Maj(G,H, K[73]);
|
||||
|
||||
W[0].w = block0.w;
|
||||
A= K[74] +Tr1(B)+Ch(B,C,D)+W[0].w;
|
||||
E= K[75] +A+Tr2(F)+Maj(F,G,H);
|
||||
|
||||
W[1].x = block1.x;
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[4]);
|
||||
W[1].y = block1.y;
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[5]);
|
||||
W[1].z = block1.z;
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[6]);
|
||||
W[1].w = block1.w;
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[7]);
|
||||
|
||||
W[2].x = block2.x;
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[8]);
|
||||
W[2].y = block2.y;
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[9]);
|
||||
W[2].z = block2.z;
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[10]);
|
||||
W[2].w = block2.w;
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[11]);
|
||||
|
||||
W[3].x = block3.x;
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[12]);
|
||||
W[3].y = block3.y;
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[13]);
|
||||
W[3].z = block3.z;
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[14]);
|
||||
W[3].w = block3.w;
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[76]);
|
||||
|
||||
W[0].x += Wr1(W[3].z) + W[2].y + Wr2(W[0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[15]);
|
||||
|
||||
W[0].y += Wr1(W[3].w) + W[2].z + Wr2(W[0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[16]);
|
||||
|
||||
W[0].z += Wr1(W[0].x) + W[2].w + Wr2(W[0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[17]);
|
||||
|
||||
W[0].w += Wr1(W[0].y) + W[3].x + Wr2(W[1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[18]);
|
||||
|
||||
W[1].x += Wr1(W[0].z) + W[3].y + Wr2(W[1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[19]);
|
||||
|
||||
W[1].y += Wr1(W[0].w) + W[3].z + Wr2(W[1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[20]);
|
||||
|
||||
W[1].z += Wr1(W[1].x) + W[3].w + Wr2(W[1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[21]);
|
||||
|
||||
W[1].w += Wr1(W[1].y) + W[0].x + Wr2(W[2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[22]);
|
||||
|
||||
W[2].x += Wr1(W[1].z) + W[0].y + Wr2(W[2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[23]);
|
||||
|
||||
W[2].y += Wr1(W[1].w) + W[0].z + Wr2(W[2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[24]);
|
||||
|
||||
W[2].z += Wr1(W[2].x) + W[0].w + Wr2(W[2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[25]);
|
||||
|
||||
W[2].w += Wr1(W[2].y) + W[1].x + Wr2(W[3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[26]);
|
||||
|
||||
W[3].x += Wr1(W[2].z) + W[1].y + Wr2(W[3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[27]);
|
||||
|
||||
W[3].y += Wr1(W[2].w) + W[1].z + Wr2(W[3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[28]);
|
||||
|
||||
W[3].z += Wr1(W[3].x) + W[1].w + Wr2(W[3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[29]);
|
||||
|
||||
W[3].w += Wr1(W[3].y) + W[2].x + Wr2(W[0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[30]);
|
||||
|
||||
W[0].x += Wr1(W[3].z) + W[2].y + Wr2(W[0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[31]);
|
||||
|
||||
W[0].y += Wr1(W[3].w) + W[2].z + Wr2(W[0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[32]);
|
||||
|
||||
W[0].z += Wr1(W[0].x) + W[2].w + Wr2(W[0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[33]);
|
||||
|
||||
W[0].w += Wr1(W[0].y) + W[3].x + Wr2(W[1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[34]);
|
||||
|
||||
W[1].x += Wr1(W[0].z) + W[3].y + Wr2(W[1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[35]);
|
||||
|
||||
W[1].y += Wr1(W[0].w) + W[3].z + Wr2(W[1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[36]);
|
||||
|
||||
W[1].z += Wr1(W[1].x) + W[3].w + Wr2(W[1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[37]);
|
||||
|
||||
W[1].w += Wr1(W[1].y) + W[0].x + Wr2(W[2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[38]);
|
||||
|
||||
W[2].x += Wr1(W[1].z) + W[0].y + Wr2(W[2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[39]);
|
||||
|
||||
W[2].y += Wr1(W[1].w) + W[0].z + Wr2(W[2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[40]);
|
||||
|
||||
W[2].z += Wr1(W[2].x) + W[0].w + Wr2(W[2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[41]);
|
||||
|
||||
W[2].w += Wr1(W[2].y) + W[1].x + Wr2(W[3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[42]);
|
||||
|
||||
W[3].x += Wr1(W[2].z) + W[1].y + Wr2(W[3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[43]);
|
||||
|
||||
W[3].y += Wr1(W[2].w) + W[1].z + Wr2(W[3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[44]);
|
||||
|
||||
W[3].z += Wr1(W[3].x) + W[1].w + Wr2(W[3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[45]);
|
||||
|
||||
W[3].w += Wr1(W[3].y) + W[2].x + Wr2(W[0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[46]);
|
||||
|
||||
W[0].x += Wr1(W[3].z) + W[2].y + Wr2(W[0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[47]);
|
||||
|
||||
W[0].y += Wr1(W[3].w) + W[2].z + Wr2(W[0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[48]);
|
||||
|
||||
W[0].z += Wr1(W[0].x) + W[2].w + Wr2(W[0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[49]);
|
||||
|
||||
W[0].w += Wr1(W[0].y) + W[3].x + Wr2(W[1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[50]);
|
||||
|
||||
W[1].x += Wr1(W[0].z) + W[3].y + Wr2(W[1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[51]);
|
||||
|
||||
W[1].y += Wr1(W[0].w) + W[3].z + Wr2(W[1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[52]);
|
||||
|
||||
W[1].z += Wr1(W[1].x) + W[3].w + Wr2(W[1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[53]);
|
||||
|
||||
W[1].w += Wr1(W[1].y) + W[0].x + Wr2(W[2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[54]);
|
||||
|
||||
W[2].x += Wr1(W[1].z) + W[0].y + Wr2(W[2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[55]);
|
||||
|
||||
W[2].y += Wr1(W[1].w) + W[0].z + Wr2(W[2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[56]);
|
||||
|
||||
W[2].z += Wr1(W[2].x) + W[0].w + Wr2(W[2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[57]);
|
||||
|
||||
W[2].w += Wr1(W[2].y) + W[1].x + Wr2(W[3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[58]);
|
||||
|
||||
W[3].x += Wr1(W[2].z) + W[1].y + Wr2(W[3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[59]);
|
||||
|
||||
W[3].y += Wr1(W[2].w) + W[1].z + Wr2(W[3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[60]);
|
||||
|
||||
W[3].z += Wr1(W[3].x) + W[1].w + Wr2(W[3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[61]);
|
||||
|
||||
W[3].w += Wr1(W[3].y) + W[2].x + Wr2(W[0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[62]);
|
||||
|
||||
#undef A
|
||||
#undef B
|
||||
#undef C
|
||||
#undef D
|
||||
#undef E
|
||||
#undef F
|
||||
#undef G
|
||||
#undef H
|
||||
|
||||
*state0 += (uint4)(K[73], K[77], K[78], K[79]);
|
||||
*state1 += (uint4)(K[66], K[67], K[80], K[81]);
|
||||
}
|
||||
|
||||
__constant uint fixedW[64] =
|
||||
{
|
||||
0x428a2f99,0xf1374491,0xb5c0fbcf,0xe9b5dba5,0x3956c25b,0x59f111f1,0x923f82a4,0xab1c5ed5,
|
||||
0xd807aa98,0x12835b01,0x243185be,0x550c7dc3,0x72be5d74,0x80deb1fe,0x9bdc06a7,0xc19bf794,
|
||||
0xf59b89c2,0x73924787,0x23c6886e,0xa42ca65c,0x15ed3627,0x4d6edcbf,0xe28217fc,0xef02488f,
|
||||
0xb707775c,0x0468c23f,0xe7e72b4c,0x49e1f1a2,0x4b99c816,0x926d1570,0xaa0fc072,0xadb36e2c,
|
||||
0xad87a3ea,0xbcb1d3a3,0x7b993186,0x562b9420,0xbff3ca0c,0xda4b0c23,0x6cd8711a,0x8f337caa,
|
||||
0xc91b1417,0xc359dce1,0xa83253a7,0x3b13c12d,0x9d3d725d,0xd9031a84,0xb1a03340,0x16f58012,
|
||||
0xe64fb6a2,0xe84d923a,0xe93a5730,0x09837686,0x078ff753,0x29833341,0xd5de0b7e,0x6948ccf4,
|
||||
0xe0a1adbe,0x7c728e11,0x511c78e4,0x315b45bd,0xfca71413,0xea28f96a,0x79703128,0x4e1ef848,
|
||||
};
|
||||
|
||||
void SHA256_fixed(uint4*restrict state0,uint4*restrict state1)
|
||||
{
|
||||
uint4 S0 = *state0;
|
||||
uint4 S1 = *state1;
|
||||
|
||||
#define A S0.x
|
||||
#define B S0.y
|
||||
#define C S0.z
|
||||
#define D S0.w
|
||||
#define E S1.x
|
||||
#define F S1.y
|
||||
#define G S1.z
|
||||
#define H S1.w
|
||||
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[0]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[1]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[2]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[3]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[4]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[5]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[6]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[7]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[8]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[9]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[10]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[11]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[12]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[13]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[14]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[15]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[16]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[17]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[18]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[19]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[20]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[21]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[22]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[23]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[24]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[25]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[26]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[27]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[28]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[29]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[30]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[31]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[32]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[33]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[34]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[35]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[36]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[37]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[38]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[39]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[40]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[41]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[42]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[43]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[44]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[45]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[46]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[47]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[48]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[49]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[50]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[51]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[52]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[53]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[54]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[55]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[56]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[57]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[58]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[59]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[60]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[61]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[62]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[63]);
|
||||
|
||||
#undef A
|
||||
#undef B
|
||||
#undef C
|
||||
#undef D
|
||||
#undef E
|
||||
#undef F
|
||||
#undef G
|
||||
#undef H
|
||||
*state0 += S0;
|
||||
*state1 += S1;
|
||||
}
|
||||
|
||||
void shittify(uint4 B[8])
|
||||
{
|
||||
uint4 tmp[4];
|
||||
tmp[0] = (uint4)(B[1].x,B[2].y,B[3].z,B[0].w);
|
||||
tmp[1] = (uint4)(B[2].x,B[3].y,B[0].z,B[1].w);
|
||||
tmp[2] = (uint4)(B[3].x,B[0].y,B[1].z,B[2].w);
|
||||
tmp[3] = (uint4)(B[0].x,B[1].y,B[2].z,B[3].w);
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i] = EndianSwap(tmp[i]);
|
||||
|
||||
tmp[0] = (uint4)(B[5].x,B[6].y,B[7].z,B[4].w);
|
||||
tmp[1] = (uint4)(B[6].x,B[7].y,B[4].z,B[5].w);
|
||||
tmp[2] = (uint4)(B[7].x,B[4].y,B[5].z,B[6].w);
|
||||
tmp[3] = (uint4)(B[4].x,B[5].y,B[6].z,B[7].w);
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i+4] = EndianSwap(tmp[i]);
|
||||
}
|
||||
|
||||
void unshittify(uint4 B[8])
|
||||
{
|
||||
uint4 tmp[4];
|
||||
tmp[0] = (uint4)(B[3].x,B[2].y,B[1].z,B[0].w);
|
||||
tmp[1] = (uint4)(B[0].x,B[3].y,B[2].z,B[1].w);
|
||||
tmp[2] = (uint4)(B[1].x,B[0].y,B[3].z,B[2].w);
|
||||
tmp[3] = (uint4)(B[2].x,B[1].y,B[0].z,B[3].w);
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i] = EndianSwap(tmp[i]);
|
||||
|
||||
tmp[0] = (uint4)(B[7].x,B[6].y,B[5].z,B[4].w);
|
||||
tmp[1] = (uint4)(B[4].x,B[7].y,B[6].z,B[5].w);
|
||||
tmp[2] = (uint4)(B[5].x,B[4].y,B[7].z,B[6].w);
|
||||
tmp[3] = (uint4)(B[6].x,B[5].y,B[4].z,B[7].w);
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i+4] = EndianSwap(tmp[i]);
|
||||
}
|
||||
|
||||
void salsa(uint4 B[8])
|
||||
{
|
||||
uint4 w[4];
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
w[i] = (B[i]^=B[i+4]);
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
{
|
||||
w[0] ^= rotl(w[3] +w[2] , 7U);
|
||||
w[1] ^= rotl(w[0] +w[3] , 9U);
|
||||
w[2] ^= rotl(w[1] +w[0] ,13U);
|
||||
w[3] ^= rotl(w[2] +w[1] ,18U);
|
||||
w[2] ^= rotl(w[3].wxyz+w[0].zwxy, 7U);
|
||||
w[1] ^= rotl(w[2].wxyz+w[3].zwxy, 9U);
|
||||
w[0] ^= rotl(w[1].wxyz+w[2].zwxy,13U);
|
||||
w[3] ^= rotl(w[0].wxyz+w[1].zwxy,18U);
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
w[i] = (B[i+4]^=(B[i]+=w[i]));
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
{
|
||||
w[0] ^= rotl(w[3] +w[2] , 7U);
|
||||
w[1] ^= rotl(w[0] +w[3] , 9U);
|
||||
w[2] ^= rotl(w[1] +w[0] ,13U);
|
||||
w[3] ^= rotl(w[2] +w[1] ,18U);
|
||||
w[2] ^= rotl(w[3].wxyz+w[0].zwxy, 7U);
|
||||
w[1] ^= rotl(w[2].wxyz+w[3].zwxy, 9U);
|
||||
w[0] ^= rotl(w[1].wxyz+w[2].zwxy,13U);
|
||||
w[3] ^= rotl(w[0].wxyz+w[1].zwxy,18U);
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i+4] += w[i];
|
||||
}
|
||||
|
||||
void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
|
||||
{
|
||||
shittify(X);
|
||||
const uint zSIZE = 8;
|
||||
const uint xSIZE = CONCURRENT_THREADS;
|
||||
uint x = get_global_id(0)%xSIZE;
|
||||
uint CO=rotl(x,3U);
|
||||
uint CO_tmp=rotl(xSIZE,3U);
|
||||
|
||||
for(uint y=0; y<1024/LOOKUP_GAP; ++y, CO+=CO_tmp)
|
||||
{
|
||||
uint CO_reg=CO;
|
||||
#pragma unroll
|
||||
for(uint z=0; z<zSIZE; ++z, ++CO_reg)
|
||||
lookup[CO_reg] = X[z];
|
||||
for(uint i=0; i<LOOKUP_GAP; ++i)
|
||||
salsa(X);
|
||||
}
|
||||
|
||||
CO_tmp=rotl(x,3U);
|
||||
|
||||
#if (LOOKUP_GAP != 1) && (LOOKUP_GAP != 2) && (LOOKUP_GAP != 4) && (LOOKUP_GAP != 8)
|
||||
{
|
||||
uint y = (1024/LOOKUP_GAP);
|
||||
CO=CO_tmp+rotl(y*xSIZE,3U);
|
||||
#pragma unroll
|
||||
for(uint z=0; z<zSIZE; ++z, ++CO)
|
||||
lookup[CO] = X[z];
|
||||
for(uint i=0; i<1024%LOOKUP_GAP; ++i)
|
||||
salsa(X);
|
||||
}
|
||||
#endif
|
||||
|
||||
for (uint i=0; i<1024; ++i)
|
||||
{
|
||||
uint4 V[8];
|
||||
uint j = X[7].x & K[85];
|
||||
uint y = (j/LOOKUP_GAP);
|
||||
uint CO_reg=CO_tmp+rotl(xSIZE*y,3U);
|
||||
|
||||
for(uint z=0; z<zSIZE; ++z, ++CO_reg)
|
||||
V[z] = lookup[CO_reg];
|
||||
|
||||
#if (LOOKUP_GAP == 1)
|
||||
#elif (LOOKUP_GAP == 2)
|
||||
if (j&1)
|
||||
salsa(V);
|
||||
#else
|
||||
uint val = j%LOOKUP_GAP;
|
||||
for (uint z=0; z<val; ++z)
|
||||
salsa(V);
|
||||
#endif
|
||||
|
||||
#pragma unroll
|
||||
for(uint z=0; z<zSIZE; ++z)
|
||||
X[z] ^= V[z];
|
||||
salsa(X);
|
||||
}
|
||||
unshittify(X);
|
||||
}
|
||||
|
||||
#define FOUND (0xFF)
|
||||
#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
|
||||
|
||||
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
|
||||
__kernel void search(__global const uint4 * restrict input,
|
||||
volatile __global uint*restrict output, __global uint4*restrict padcache,
|
||||
const uint4 midstate0, const uint4 midstate16, const uint target)
|
||||
{
|
||||
uint gid = get_global_id(0);
|
||||
uint4 X[8];
|
||||
uint4 tstate0, tstate1, ostate0, ostate1, tmp0, tmp1;
|
||||
uint4 data = (uint4)(input[4].x,input[4].y,input[4].z,gid);
|
||||
uint4 pad0 = midstate0, pad1 = midstate16;
|
||||
|
||||
SHA256(&pad0,&pad1, data, (uint4)(K[84],0,0,0), (uint4)(0,0,0,0), (uint4)(0,0,0, K[86]));
|
||||
SHA256_fresh(&ostate0,&ostate1, pad0^ K[82], pad1^ K[82], K[82], K[82]);
|
||||
SHA256_fresh(&tstate0,&tstate1, pad0^ K[83], pad1^ K[83], K[83], K[83]);
|
||||
|
||||
tmp0 = tstate0;
|
||||
tmp1 = tstate1;
|
||||
SHA256(&tstate0, &tstate1, input[0],input[1],input[2],input[3]);
|
||||
|
||||
#pragma unroll
|
||||
for (uint i=0; i<4; i++)
|
||||
{
|
||||
pad0 = tstate0;
|
||||
pad1 = tstate1;
|
||||
X[rotl(i,1U) ] = ostate0;
|
||||
X[rotl(i,1U)+1] = ostate1;
|
||||
|
||||
SHA256(&pad0,&pad1, data, (uint4)(i+1,K[84],0,0), (uint4)(0,0,0,0), (uint4)(0,0,0, K[87]));
|
||||
SHA256(X+rotl(i,1U),X+rotl(i,1U)+1, pad0, pad1, (uint4)(K[84], 0U, 0U, 0U), (uint4)(0U, 0U, 0U, K[88]));
|
||||
}
|
||||
scrypt_core(X,padcache);
|
||||
SHA256(&tmp0,&tmp1, X[0], X[1], X[2], X[3]);
|
||||
SHA256(&tmp0,&tmp1, X[4], X[5], X[6], X[7]);
|
||||
SHA256_fixed(&tmp0,&tmp1);
|
||||
SHA256(&ostate0,&ostate1, tmp0, tmp1, (uint4)(K[84], 0U, 0U, 0U), (uint4)(0U, 0U, 0U, K[88]));
|
||||
|
||||
bool result = (EndianSwap(ostate1.w) <= target);
|
||||
if (result)
|
||||
SETFOUND(gid);
|
||||
}
|
||||
855
kernel/alexkarold.cl
Normal file
855
kernel/alexkarold.cl
Normal file
@ -0,0 +1,855 @@
|
||||
/*-
|
||||
* Copyright 2009 Colin Percival, 2011 ArtForz, 2011 pooler, 2012 mtrlt,
|
||||
* 2012-2013 Con Kolivas, 2013 Alexey Karimov.
|
||||
* 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.
|
||||
*
|
||||
* THIS SOFTWARE IS PROVIDED BY THE AUTHOR AND CONTRIBUTORS ``AS IS'' AND
|
||||
* ANY EXPRESS 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 AUTHOR OR 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.
|
||||
*
|
||||
* This file was originally written by Colin Percival as part of the Tarsnap
|
||||
* online backup system.
|
||||
*/
|
||||
|
||||
__constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 };
|
||||
__constant uint K[] = {
|
||||
0x428a2f98U,
|
||||
0x71374491U,
|
||||
0xb5c0fbcfU,
|
||||
0xe9b5dba5U,
|
||||
0x3956c25bU,
|
||||
0x59f111f1U,
|
||||
0x923f82a4U,
|
||||
0xab1c5ed5U,
|
||||
0xd807aa98U,
|
||||
0x12835b01U,
|
||||
0x243185beU, // 10
|
||||
0x550c7dc3U,
|
||||
0x72be5d74U,
|
||||
0x80deb1feU,
|
||||
0x9bdc06a7U,
|
||||
0xe49b69c1U,
|
||||
0xefbe4786U,
|
||||
0x0fc19dc6U,
|
||||
0x240ca1ccU,
|
||||
0x2de92c6fU,
|
||||
0x4a7484aaU, // 20
|
||||
0x5cb0a9dcU,
|
||||
0x76f988daU,
|
||||
0x983e5152U,
|
||||
0xa831c66dU,
|
||||
0xb00327c8U,
|
||||
0xbf597fc7U,
|
||||
0xc6e00bf3U,
|
||||
0xd5a79147U,
|
||||
0x06ca6351U,
|
||||
0x14292967U, // 30
|
||||
0x27b70a85U,
|
||||
0x2e1b2138U,
|
||||
0x4d2c6dfcU,
|
||||
0x53380d13U,
|
||||
0x650a7354U,
|
||||
0x766a0abbU,
|
||||
0x81c2c92eU,
|
||||
0x92722c85U,
|
||||
0xa2bfe8a1U,
|
||||
0xa81a664bU, // 40
|
||||
0xc24b8b70U,
|
||||
0xc76c51a3U,
|
||||
0xd192e819U,
|
||||
0xd6990624U,
|
||||
0xf40e3585U,
|
||||
0x106aa070U,
|
||||
0x19a4c116U,
|
||||
0x1e376c08U,
|
||||
0x2748774cU,
|
||||
0x34b0bcb5U, // 50
|
||||
0x391c0cb3U,
|
||||
0x4ed8aa4aU,
|
||||
0x5b9cca4fU,
|
||||
0x682e6ff3U,
|
||||
0x748f82eeU,
|
||||
0x78a5636fU,
|
||||
0x84c87814U,
|
||||
0x8cc70208U,
|
||||
0x90befffaU,
|
||||
0xa4506cebU, // 60
|
||||
0xbef9a3f7U,
|
||||
0xc67178f2U,
|
||||
0x98c7e2a2U,
|
||||
0xfc08884dU,
|
||||
0xcd2a11aeU,
|
||||
0x510e527fU,
|
||||
0x9b05688cU,
|
||||
0xC3910C8EU,
|
||||
0xfb6feee7U,
|
||||
0x2a01a605U, // 70
|
||||
0x0c2e12e0U,
|
||||
0x4498517BU,
|
||||
0x6a09e667U,
|
||||
0xa4ce148bU,
|
||||
0x95F61999U,
|
||||
0xc19bf174U,
|
||||
0xBB67AE85U,
|
||||
0x3C6EF372U,
|
||||
0xA54FF53AU,
|
||||
0x1F83D9ABU, // 80
|
||||
0x5BE0CD19U,
|
||||
0x5C5C5C5CU,
|
||||
0x36363636U,
|
||||
0x80000000U,
|
||||
0x000003FFU,
|
||||
0x00000280U,
|
||||
0x000004a0U,
|
||||
0x00000300U
|
||||
};
|
||||
|
||||
#define rotl(x,y) rotate(x,y)
|
||||
#define Ch(x,y,z) bitselect(z,y,x)
|
||||
#define Maj(x,y,z) Ch((x^z),y,z)
|
||||
|
||||
#define EndianSwap(n) (rotl(n & ES[0], 24U)|rotl(n & ES[1], 8U))
|
||||
|
||||
#define Tr2(x) (rotl(x, 30U) ^ rotl(x, 19U) ^ rotl(x, 10U))
|
||||
#define Tr1(x) (rotl(x, 26U) ^ rotl(x, 21U) ^ rotl(x, 7U))
|
||||
#define Wr2(x) (rotl(x, 25U) ^ rotl(x, 14U) ^ (x>>3U))
|
||||
#define Wr1(x) (rotl(x, 15U) ^ rotl(x, 13U) ^ (x>>10U))
|
||||
|
||||
#define RND(a, b, c, d, e, f, g, h, k) \
|
||||
h += Tr1(e); \
|
||||
h += Ch(e, f, g); \
|
||||
h += k; \
|
||||
d += h; \
|
||||
h += Tr2(a); \
|
||||
h += Maj(a, b, c);
|
||||
|
||||
void SHA256(uint4*restrict state0,uint4*restrict state1, const uint4 block0, const uint4 block1, const uint4 block2, const uint4 block3)
|
||||
{
|
||||
uint4 S0 = *state0;
|
||||
uint4 S1 = *state1;
|
||||
|
||||
#define A S0.x
|
||||
#define B S0.y
|
||||
#define C S0.z
|
||||
#define D S0.w
|
||||
#define E S1.x
|
||||
#define F S1.y
|
||||
#define G S1.z
|
||||
#define H S1.w
|
||||
|
||||
uint4 W[4];
|
||||
|
||||
W[ 0].x = block0.x;
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[0]);
|
||||
W[ 0].y = block0.y;
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[1]);
|
||||
W[ 0].z = block0.z;
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[2]);
|
||||
W[ 0].w = block0.w;
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[3]);
|
||||
|
||||
W[ 1].x = block1.x;
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[4]);
|
||||
W[ 1].y = block1.y;
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[5]);
|
||||
W[ 1].z = block1.z;
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[6]);
|
||||
W[ 1].w = block1.w;
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[7]);
|
||||
|
||||
W[ 2].x = block2.x;
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[8]);
|
||||
W[ 2].y = block2.y;
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[9]);
|
||||
W[ 2].z = block2.z;
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[10]);
|
||||
W[ 2].w = block2.w;
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[11]);
|
||||
|
||||
W[ 3].x = block3.x;
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[12]);
|
||||
W[ 3].y = block3.y;
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[13]);
|
||||
W[ 3].z = block3.z;
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[14]);
|
||||
W[ 3].w = block3.w;
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[76]);
|
||||
|
||||
W[ 0].x += Wr1(W[ 3].z) + W[ 2].y + Wr2(W[ 0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[15]);
|
||||
|
||||
W[ 0].y += Wr1(W[ 3].w) + W[ 2].z + Wr2(W[ 0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[16]);
|
||||
|
||||
W[ 0].z += Wr1(W[ 0].x) + W[ 2].w + Wr2(W[ 0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[17]);
|
||||
|
||||
W[ 0].w += Wr1(W[ 0].y) + W[ 3].x + Wr2(W[ 1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[18]);
|
||||
|
||||
W[ 1].x += Wr1(W[ 0].z) + W[ 3].y + Wr2(W[ 1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[19]);
|
||||
|
||||
W[ 1].y += Wr1(W[ 0].w) + W[ 3].z + Wr2(W[ 1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[20]);
|
||||
|
||||
W[ 1].z += Wr1(W[ 1].x) + W[ 3].w + Wr2(W[ 1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[21]);
|
||||
|
||||
W[ 1].w += Wr1(W[ 1].y) + W[ 0].x + Wr2(W[ 2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[22]);
|
||||
|
||||
W[ 2].x += Wr1(W[ 1].z) + W[ 0].y + Wr2(W[ 2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[23]);
|
||||
|
||||
W[ 2].y += Wr1(W[ 1].w) + W[ 0].z + Wr2(W[ 2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[24]);
|
||||
|
||||
W[ 2].z += Wr1(W[ 2].x) + W[ 0].w + Wr2(W[ 2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[25]);
|
||||
|
||||
W[ 2].w += Wr1(W[ 2].y) + W[ 1].x + Wr2(W[ 3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[26]);
|
||||
|
||||
W[ 3].x += Wr1(W[ 2].z) + W[ 1].y + Wr2(W[ 3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[27]);
|
||||
|
||||
W[ 3].y += Wr1(W[ 2].w) + W[ 1].z + Wr2(W[ 3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[28]);
|
||||
|
||||
W[ 3].z += Wr1(W[ 3].x) + W[ 1].w + Wr2(W[ 3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[29]);
|
||||
|
||||
W[ 3].w += Wr1(W[ 3].y) + W[ 2].x + Wr2(W[ 0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[30]);
|
||||
|
||||
W[ 0].x += Wr1(W[ 3].z) + W[ 2].y + Wr2(W[ 0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[31]);
|
||||
|
||||
W[ 0].y += Wr1(W[ 3].w) + W[ 2].z + Wr2(W[ 0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[32]);
|
||||
|
||||
W[ 0].z += Wr1(W[ 0].x) + W[ 2].w + Wr2(W[ 0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[33]);
|
||||
|
||||
W[ 0].w += Wr1(W[ 0].y) + W[ 3].x + Wr2(W[ 1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[34]);
|
||||
|
||||
W[ 1].x += Wr1(W[ 0].z) + W[ 3].y + Wr2(W[ 1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[35]);
|
||||
|
||||
W[ 1].y += Wr1(W[ 0].w) + W[ 3].z + Wr2(W[ 1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[36]);
|
||||
|
||||
W[ 1].z += Wr1(W[ 1].x) + W[ 3].w + Wr2(W[ 1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[37]);
|
||||
|
||||
W[ 1].w += Wr1(W[ 1].y) + W[ 0].x + Wr2(W[ 2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[38]);
|
||||
|
||||
W[ 2].x += Wr1(W[ 1].z) + W[ 0].y + Wr2(W[ 2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[39]);
|
||||
|
||||
W[ 2].y += Wr1(W[ 1].w) + W[ 0].z + Wr2(W[ 2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[40]);
|
||||
|
||||
W[ 2].z += Wr1(W[ 2].x) + W[ 0].w + Wr2(W[ 2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[41]);
|
||||
|
||||
W[ 2].w += Wr1(W[ 2].y) + W[ 1].x + Wr2(W[ 3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[42]);
|
||||
|
||||
W[ 3].x += Wr1(W[ 2].z) + W[ 1].y + Wr2(W[ 3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[43]);
|
||||
|
||||
W[ 3].y += Wr1(W[ 2].w) + W[ 1].z + Wr2(W[ 3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[44]);
|
||||
|
||||
W[ 3].z += Wr1(W[ 3].x) + W[ 1].w + Wr2(W[ 3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[45]);
|
||||
|
||||
W[ 3].w += Wr1(W[ 3].y) + W[ 2].x + Wr2(W[ 0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[46]);
|
||||
|
||||
W[ 0].x += Wr1(W[ 3].z) + W[ 2].y + Wr2(W[ 0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[47]);
|
||||
|
||||
W[ 0].y += Wr1(W[ 3].w) + W[ 2].z + Wr2(W[ 0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[48]);
|
||||
|
||||
W[ 0].z += Wr1(W[ 0].x) + W[ 2].w + Wr2(W[ 0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[49]);
|
||||
|
||||
W[ 0].w += Wr1(W[ 0].y) + W[ 3].x + Wr2(W[ 1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[50]);
|
||||
|
||||
W[ 1].x += Wr1(W[ 0].z) + W[ 3].y + Wr2(W[ 1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[51]);
|
||||
|
||||
W[ 1].y += Wr1(W[ 0].w) + W[ 3].z + Wr2(W[ 1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[52]);
|
||||
|
||||
W[ 1].z += Wr1(W[ 1].x) + W[ 3].w + Wr2(W[ 1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[53]);
|
||||
|
||||
W[ 1].w += Wr1(W[ 1].y) + W[ 0].x + Wr2(W[ 2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[54]);
|
||||
|
||||
W[ 2].x += Wr1(W[ 1].z) + W[ 0].y + Wr2(W[ 2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[55]);
|
||||
|
||||
W[ 2].y += Wr1(W[ 1].w) + W[ 0].z + Wr2(W[ 2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[56]);
|
||||
|
||||
W[ 2].z += Wr1(W[ 2].x) + W[ 0].w + Wr2(W[ 2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[57]);
|
||||
|
||||
W[ 2].w += Wr1(W[ 2].y) + W[ 1].x + Wr2(W[ 3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[58]);
|
||||
|
||||
W[ 3].x += Wr1(W[ 2].z) + W[ 1].y + Wr2(W[ 3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[59]);
|
||||
|
||||
W[ 3].y += Wr1(W[ 2].w) + W[ 1].z + Wr2(W[ 3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[60]);
|
||||
|
||||
W[ 3].z += Wr1(W[ 3].x) + W[ 1].w + Wr2(W[ 3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[61]);
|
||||
|
||||
W[ 3].w += Wr1(W[ 3].y) + W[ 2].x + Wr2(W[ 0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[62]);
|
||||
|
||||
#undef A
|
||||
#undef B
|
||||
#undef C
|
||||
#undef D
|
||||
#undef E
|
||||
#undef F
|
||||
#undef G
|
||||
#undef H
|
||||
|
||||
*state0 += S0;
|
||||
*state1 += S1;
|
||||
}
|
||||
|
||||
void SHA256_fresh(uint4*restrict state0,uint4*restrict state1, const uint4 block0, const uint4 block1, const uint4 block2, const uint4 block3)
|
||||
{
|
||||
#define A (*state0).x
|
||||
#define B (*state0).y
|
||||
#define C (*state0).z
|
||||
#define D (*state0).w
|
||||
#define E (*state1).x
|
||||
#define F (*state1).y
|
||||
#define G (*state1).z
|
||||
#define H (*state1).w
|
||||
|
||||
uint4 W[4];
|
||||
|
||||
W[0].x = block0.x;
|
||||
D= K[63] +W[0].x;
|
||||
H= K[64] +W[0].x;
|
||||
|
||||
W[0].y = block0.y;
|
||||
C= K[65] +Tr1(D)+Ch(D, K[66], K[67])+W[0].y;
|
||||
G= K[68] +C+Tr2(H)+Ch(H, K[69] ,K[70]);
|
||||
|
||||
W[0].z = block0.z;
|
||||
B= K[71] +Tr1(C)+Ch(C,D,K[66])+W[0].z;
|
||||
F= K[72] +B+Tr2(G)+Maj(G,H, K[73]);
|
||||
|
||||
W[0].w = block0.w;
|
||||
A= K[74] +Tr1(B)+Ch(B,C,D)+W[0].w;
|
||||
E= K[75] +A+Tr2(F)+Maj(F,G,H);
|
||||
|
||||
W[1].x = block1.x;
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[4]);
|
||||
W[1].y = block1.y;
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[5]);
|
||||
W[1].z = block1.z;
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[6]);
|
||||
W[1].w = block1.w;
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[7]);
|
||||
|
||||
W[2].x = block2.x;
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[8]);
|
||||
W[2].y = block2.y;
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[9]);
|
||||
W[2].z = block2.z;
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[10]);
|
||||
W[2].w = block2.w;
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[11]);
|
||||
|
||||
W[3].x = block3.x;
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[12]);
|
||||
W[3].y = block3.y;
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[13]);
|
||||
W[3].z = block3.z;
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[14]);
|
||||
W[3].w = block3.w;
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[76]);
|
||||
|
||||
W[0].x += Wr1(W[3].z) + W[2].y + Wr2(W[0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[15]);
|
||||
|
||||
W[0].y += Wr1(W[3].w) + W[2].z + Wr2(W[0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[16]);
|
||||
|
||||
W[0].z += Wr1(W[0].x) + W[2].w + Wr2(W[0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[17]);
|
||||
|
||||
W[0].w += Wr1(W[0].y) + W[3].x + Wr2(W[1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[18]);
|
||||
|
||||
W[1].x += Wr1(W[0].z) + W[3].y + Wr2(W[1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[19]);
|
||||
|
||||
W[1].y += Wr1(W[0].w) + W[3].z + Wr2(W[1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[20]);
|
||||
|
||||
W[1].z += Wr1(W[1].x) + W[3].w + Wr2(W[1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[21]);
|
||||
|
||||
W[1].w += Wr1(W[1].y) + W[0].x + Wr2(W[2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[22]);
|
||||
|
||||
W[2].x += Wr1(W[1].z) + W[0].y + Wr2(W[2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[23]);
|
||||
|
||||
W[2].y += Wr1(W[1].w) + W[0].z + Wr2(W[2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[24]);
|
||||
|
||||
W[2].z += Wr1(W[2].x) + W[0].w + Wr2(W[2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[25]);
|
||||
|
||||
W[2].w += Wr1(W[2].y) + W[1].x + Wr2(W[3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[26]);
|
||||
|
||||
W[3].x += Wr1(W[2].z) + W[1].y + Wr2(W[3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[27]);
|
||||
|
||||
W[3].y += Wr1(W[2].w) + W[1].z + Wr2(W[3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[28]);
|
||||
|
||||
W[3].z += Wr1(W[3].x) + W[1].w + Wr2(W[3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[29]);
|
||||
|
||||
W[3].w += Wr1(W[3].y) + W[2].x + Wr2(W[0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[30]);
|
||||
|
||||
W[0].x += Wr1(W[3].z) + W[2].y + Wr2(W[0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[31]);
|
||||
|
||||
W[0].y += Wr1(W[3].w) + W[2].z + Wr2(W[0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[32]);
|
||||
|
||||
W[0].z += Wr1(W[0].x) + W[2].w + Wr2(W[0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[33]);
|
||||
|
||||
W[0].w += Wr1(W[0].y) + W[3].x + Wr2(W[1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[34]);
|
||||
|
||||
W[1].x += Wr1(W[0].z) + W[3].y + Wr2(W[1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[35]);
|
||||
|
||||
W[1].y += Wr1(W[0].w) + W[3].z + Wr2(W[1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[36]);
|
||||
|
||||
W[1].z += Wr1(W[1].x) + W[3].w + Wr2(W[1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[37]);
|
||||
|
||||
W[1].w += Wr1(W[1].y) + W[0].x + Wr2(W[2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[38]);
|
||||
|
||||
W[2].x += Wr1(W[1].z) + W[0].y + Wr2(W[2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[39]);
|
||||
|
||||
W[2].y += Wr1(W[1].w) + W[0].z + Wr2(W[2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[40]);
|
||||
|
||||
W[2].z += Wr1(W[2].x) + W[0].w + Wr2(W[2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[41]);
|
||||
|
||||
W[2].w += Wr1(W[2].y) + W[1].x + Wr2(W[3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[42]);
|
||||
|
||||
W[3].x += Wr1(W[2].z) + W[1].y + Wr2(W[3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[43]);
|
||||
|
||||
W[3].y += Wr1(W[2].w) + W[1].z + Wr2(W[3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[44]);
|
||||
|
||||
W[3].z += Wr1(W[3].x) + W[1].w + Wr2(W[3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[45]);
|
||||
|
||||
W[3].w += Wr1(W[3].y) + W[2].x + Wr2(W[0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[46]);
|
||||
|
||||
W[0].x += Wr1(W[3].z) + W[2].y + Wr2(W[0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[47]);
|
||||
|
||||
W[0].y += Wr1(W[3].w) + W[2].z + Wr2(W[0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[48]);
|
||||
|
||||
W[0].z += Wr1(W[0].x) + W[2].w + Wr2(W[0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[49]);
|
||||
|
||||
W[0].w += Wr1(W[0].y) + W[3].x + Wr2(W[1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[50]);
|
||||
|
||||
W[1].x += Wr1(W[0].z) + W[3].y + Wr2(W[1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[51]);
|
||||
|
||||
W[1].y += Wr1(W[0].w) + W[3].z + Wr2(W[1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[52]);
|
||||
|
||||
W[1].z += Wr1(W[1].x) + W[3].w + Wr2(W[1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[53]);
|
||||
|
||||
W[1].w += Wr1(W[1].y) + W[0].x + Wr2(W[2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[54]);
|
||||
|
||||
W[2].x += Wr1(W[1].z) + W[0].y + Wr2(W[2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[55]);
|
||||
|
||||
W[2].y += Wr1(W[1].w) + W[0].z + Wr2(W[2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[56]);
|
||||
|
||||
W[2].z += Wr1(W[2].x) + W[0].w + Wr2(W[2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[57]);
|
||||
|
||||
W[2].w += Wr1(W[2].y) + W[1].x + Wr2(W[3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[58]);
|
||||
|
||||
W[3].x += Wr1(W[2].z) + W[1].y + Wr2(W[3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[59]);
|
||||
|
||||
W[3].y += Wr1(W[2].w) + W[1].z + Wr2(W[3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[60]);
|
||||
|
||||
W[3].z += Wr1(W[3].x) + W[1].w + Wr2(W[3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[61]);
|
||||
|
||||
W[3].w += Wr1(W[3].y) + W[2].x + Wr2(W[0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[62]);
|
||||
|
||||
#undef A
|
||||
#undef B
|
||||
#undef C
|
||||
#undef D
|
||||
#undef E
|
||||
#undef F
|
||||
#undef G
|
||||
#undef H
|
||||
|
||||
*state0 += (uint4)(K[73], K[77], K[78], K[79]);
|
||||
*state1 += (uint4)(K[66], K[67], K[80], K[81]);
|
||||
}
|
||||
|
||||
__constant uint fixedW[64] =
|
||||
{
|
||||
0x428a2f99,0xf1374491,0xb5c0fbcf,0xe9b5dba5,0x3956c25b,0x59f111f1,0x923f82a4,0xab1c5ed5,
|
||||
0xd807aa98,0x12835b01,0x243185be,0x550c7dc3,0x72be5d74,0x80deb1fe,0x9bdc06a7,0xc19bf794,
|
||||
0xf59b89c2,0x73924787,0x23c6886e,0xa42ca65c,0x15ed3627,0x4d6edcbf,0xe28217fc,0xef02488f,
|
||||
0xb707775c,0x0468c23f,0xe7e72b4c,0x49e1f1a2,0x4b99c816,0x926d1570,0xaa0fc072,0xadb36e2c,
|
||||
0xad87a3ea,0xbcb1d3a3,0x7b993186,0x562b9420,0xbff3ca0c,0xda4b0c23,0x6cd8711a,0x8f337caa,
|
||||
0xc91b1417,0xc359dce1,0xa83253a7,0x3b13c12d,0x9d3d725d,0xd9031a84,0xb1a03340,0x16f58012,
|
||||
0xe64fb6a2,0xe84d923a,0xe93a5730,0x09837686,0x078ff753,0x29833341,0xd5de0b7e,0x6948ccf4,
|
||||
0xe0a1adbe,0x7c728e11,0x511c78e4,0x315b45bd,0xfca71413,0xea28f96a,0x79703128,0x4e1ef848,
|
||||
};
|
||||
|
||||
void SHA256_fixed(uint4*restrict state0,uint4*restrict state1)
|
||||
{
|
||||
uint4 S0 = *state0;
|
||||
uint4 S1 = *state1;
|
||||
|
||||
#define A S0.x
|
||||
#define B S0.y
|
||||
#define C S0.z
|
||||
#define D S0.w
|
||||
#define E S1.x
|
||||
#define F S1.y
|
||||
#define G S1.z
|
||||
#define H S1.w
|
||||
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[0]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[1]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[2]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[3]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[4]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[5]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[6]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[7]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[8]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[9]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[10]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[11]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[12]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[13]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[14]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[15]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[16]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[17]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[18]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[19]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[20]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[21]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[22]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[23]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[24]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[25]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[26]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[27]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[28]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[29]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[30]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[31]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[32]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[33]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[34]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[35]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[36]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[37]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[38]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[39]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[40]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[41]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[42]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[43]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[44]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[45]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[46]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[47]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[48]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[49]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[50]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[51]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[52]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[53]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[54]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[55]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[56]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[57]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[58]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[59]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[60]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[61]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[62]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[63]);
|
||||
|
||||
#undef A
|
||||
#undef B
|
||||
#undef C
|
||||
#undef D
|
||||
#undef E
|
||||
#undef F
|
||||
#undef G
|
||||
#undef H
|
||||
*state0 += S0;
|
||||
*state1 += S1;
|
||||
}
|
||||
|
||||
void shittify(uint4 B[8])
|
||||
{
|
||||
uint4 tmp[4];
|
||||
tmp[0] = (uint4)(B[1].x,B[2].y,B[3].z,B[0].w);
|
||||
tmp[1] = (uint4)(B[2].x,B[3].y,B[0].z,B[1].w);
|
||||
tmp[2] = (uint4)(B[3].x,B[0].y,B[1].z,B[2].w);
|
||||
tmp[3] = (uint4)(B[0].x,B[1].y,B[2].z,B[3].w);
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i] = EndianSwap(tmp[i]);
|
||||
|
||||
tmp[0] = (uint4)(B[5].x,B[6].y,B[7].z,B[4].w);
|
||||
tmp[1] = (uint4)(B[6].x,B[7].y,B[4].z,B[5].w);
|
||||
tmp[2] = (uint4)(B[7].x,B[4].y,B[5].z,B[6].w);
|
||||
tmp[3] = (uint4)(B[4].x,B[5].y,B[6].z,B[7].w);
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i+4] = EndianSwap(tmp[i]);
|
||||
}
|
||||
|
||||
void unshittify(uint4 B[8])
|
||||
{
|
||||
uint4 tmp[4];
|
||||
tmp[0] = (uint4)(B[3].x,B[2].y,B[1].z,B[0].w);
|
||||
tmp[1] = (uint4)(B[0].x,B[3].y,B[2].z,B[1].w);
|
||||
tmp[2] = (uint4)(B[1].x,B[0].y,B[3].z,B[2].w);
|
||||
tmp[3] = (uint4)(B[2].x,B[1].y,B[0].z,B[3].w);
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i] = EndianSwap(tmp[i]);
|
||||
|
||||
tmp[0] = (uint4)(B[7].x,B[6].y,B[5].z,B[4].w);
|
||||
tmp[1] = (uint4)(B[4].x,B[7].y,B[6].z,B[5].w);
|
||||
tmp[2] = (uint4)(B[5].x,B[4].y,B[7].z,B[6].w);
|
||||
tmp[3] = (uint4)(B[6].x,B[5].y,B[4].z,B[7].w);
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i+4] = EndianSwap(tmp[i]);
|
||||
}
|
||||
|
||||
void salsa(uint4 B[8])
|
||||
{
|
||||
uint4 w[4];
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
w[i] = (B[i]^=B[i+4]);
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
{
|
||||
w[0] ^= rotl(w[3] +w[2] , 7U);
|
||||
w[1] ^= rotl(w[0] +w[3] , 9U);
|
||||
w[2] ^= rotl(w[1] +w[0] ,13U);
|
||||
w[3] ^= rotl(w[2] +w[1] ,18U);
|
||||
w[2] ^= rotl(w[3].wxyz+w[0].zwxy, 7U);
|
||||
w[1] ^= rotl(w[2].wxyz+w[3].zwxy, 9U);
|
||||
w[0] ^= rotl(w[1].wxyz+w[2].zwxy,13U);
|
||||
w[3] ^= rotl(w[0].wxyz+w[1].zwxy,18U);
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
w[i] = (B[i+4]^=(B[i]+=w[i]));
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
{
|
||||
w[0] ^= rotl(w[3] +w[2] , 7U);
|
||||
w[1] ^= rotl(w[0] +w[3] , 9U);
|
||||
w[2] ^= rotl(w[1] +w[0] ,13U);
|
||||
w[3] ^= rotl(w[2] +w[1] ,18U);
|
||||
w[2] ^= rotl(w[3].wxyz+w[0].zwxy, 7U);
|
||||
w[1] ^= rotl(w[2].wxyz+w[3].zwxy, 9U);
|
||||
w[0] ^= rotl(w[1].wxyz+w[2].zwxy,13U);
|
||||
w[3] ^= rotl(w[0].wxyz+w[1].zwxy,18U);
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i+4] += w[i];
|
||||
}
|
||||
|
||||
void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
|
||||
{
|
||||
shittify(X);
|
||||
const uint zSIZE = 8;
|
||||
const uint xSIZE = CONCURRENT_THREADS;
|
||||
uint x = get_global_id(0)%xSIZE;
|
||||
uint CO_tmp=xSIZE<<3U;
|
||||
uint CO_tmp2=x<<3U;
|
||||
|
||||
for(uint y=0; y<1024/LOOKUP_GAP; ++y)
|
||||
{
|
||||
uint CO=y*CO_tmp+CO_tmp2;
|
||||
#pragma unroll
|
||||
for(uint z=0; z<zSIZE; ++z,++CO)
|
||||
lookup[CO] = X[z];
|
||||
for(uint i=0; i<LOOKUP_GAP; ++i)
|
||||
salsa(X);
|
||||
}
|
||||
|
||||
#if (LOOKUP_GAP != 1) && (LOOKUP_GAP != 2) && (LOOKUP_GAP != 4) && (LOOKUP_GAP != 8)
|
||||
{
|
||||
uint y = (1024/LOOKUP_GAP);
|
||||
uint CO=y*CO_tmp+CO_tmp2;
|
||||
#pragma unroll
|
||||
for(uint z=0; z<zSIZE; ++z)
|
||||
lookup[CO] = X[z];
|
||||
for(uint i=0; i<1024%LOOKUP_GAP; ++i)
|
||||
salsa(X);
|
||||
}
|
||||
#endif
|
||||
for (uint i=0; i<1024; ++i)
|
||||
{
|
||||
uint4 V[8];
|
||||
uint j = X[7].x & K[85];
|
||||
uint y = (j/LOOKUP_GAP);
|
||||
uint CO=y*CO_tmp+CO_tmp2;
|
||||
#pragma unroll
|
||||
for(uint z=0; z<zSIZE; ++z)
|
||||
V[z] = lookup[CO+z];
|
||||
|
||||
#if (LOOKUP_GAP == 1)
|
||||
#elif (LOOKUP_GAP == 2)
|
||||
if (j&1)
|
||||
salsa(V);
|
||||
#else
|
||||
uint val = j%LOOKUP_GAP;
|
||||
for (uint z=0; z<val; ++z)
|
||||
salsa(V);
|
||||
#endif
|
||||
|
||||
#pragma unroll
|
||||
for(uint z=0; z<zSIZE; ++z)
|
||||
X[z] ^= V[z];
|
||||
salsa(X);
|
||||
}
|
||||
unshittify(X);
|
||||
}
|
||||
|
||||
#define FOUND (0xFF)
|
||||
#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
|
||||
|
||||
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
|
||||
__kernel void search(__global const uint4 * restrict input,
|
||||
volatile __global uint*restrict output, __global uint4*restrict padcache,
|
||||
const uint4 midstate0, const uint4 midstate16, const uint target)
|
||||
{
|
||||
uint gid = get_global_id(0);
|
||||
uint4 X[8];
|
||||
uint4 tstate0, tstate1, ostate0, ostate1, tmp0, tmp1;
|
||||
uint4 data = (uint4)(input[4].x,input[4].y,input[4].z,gid);
|
||||
uint4 pad0 = midstate0, pad1 = midstate16;
|
||||
|
||||
SHA256(&pad0,&pad1, data, (uint4)(K[84],0,0,0), (uint4)(0,0,0,0), (uint4)(0,0,0, K[86]));
|
||||
SHA256_fresh(&ostate0,&ostate1, pad0^ K[82], pad1^ K[82], K[82], K[82]);
|
||||
SHA256_fresh(&tstate0,&tstate1, pad0^ K[83], pad1^ K[83], K[83], K[83]);
|
||||
|
||||
tmp0 = tstate0;
|
||||
tmp1 = tstate1;
|
||||
SHA256(&tstate0, &tstate1, input[0],input[1],input[2],input[3]);
|
||||
|
||||
#pragma unroll
|
||||
for (uint i=0; i<4; i++)
|
||||
{
|
||||
pad0 = tstate0;
|
||||
pad1 = tstate1;
|
||||
X[i*2 ] = ostate0;
|
||||
X[i*2+1] = ostate1;
|
||||
|
||||
SHA256(&pad0,&pad1, data, (uint4)(i+1,K[84],0,0), (uint4)(0,0,0,0), (uint4)(0,0,0, K[87]));
|
||||
SHA256(X+i*2,X+i*2+1, pad0, pad1, (uint4)(K[84], 0U, 0U, 0U), (uint4)(0U, 0U, 0U, K[88]));
|
||||
}
|
||||
scrypt_core(X,padcache);
|
||||
SHA256(&tmp0,&tmp1, X[0], X[1], X[2], X[3]);
|
||||
SHA256(&tmp0,&tmp1, X[4], X[5], X[6], X[7]);
|
||||
SHA256_fixed(&tmp0,&tmp1);
|
||||
SHA256(&ostate0,&ostate1, tmp0, tmp1, (uint4)(K[84], 0U, 0U, 0U), (uint4)(0U, 0U, 0U, K[88]));
|
||||
|
||||
bool result = (EndianSwap(ostate1.w) <= target);
|
||||
if (result)
|
||||
SETFOUND(gid);
|
||||
}
|
||||
File diff suppressed because it is too large
Load Diff
841
kernel/zuikkis.cl
Normal file
841
kernel/zuikkis.cl
Normal file
@ -0,0 +1,841 @@
|
||||
/*-
|
||||
* Copyright 2009 Colin Percival, 2011 ArtForz, 2011 pooler, 2012 mtrlt,
|
||||
* 2012-2013 Con Kolivas.
|
||||
* 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.
|
||||
*
|
||||
* THIS SOFTWARE IS PROVIDED BY THE AUTHOR AND CONTRIBUTORS ``AS IS'' AND
|
||||
* ANY EXPRESS 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 AUTHOR OR 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.
|
||||
*
|
||||
* This file was originally written by Colin Percival as part of the Tarsnap
|
||||
* online backup system.
|
||||
*/
|
||||
|
||||
__constant uint ES[2] = { 0x00FF00FF, 0xFF00FF00 };
|
||||
__constant uint K[] = {
|
||||
0x428a2f98U,
|
||||
0x71374491U,
|
||||
0xb5c0fbcfU,
|
||||
0xe9b5dba5U,
|
||||
0x3956c25bU,
|
||||
0x59f111f1U,
|
||||
0x923f82a4U,
|
||||
0xab1c5ed5U,
|
||||
0xd807aa98U,
|
||||
0x12835b01U,
|
||||
0x243185beU, // 10
|
||||
0x550c7dc3U,
|
||||
0x72be5d74U,
|
||||
0x80deb1feU,
|
||||
0x9bdc06a7U,
|
||||
0xe49b69c1U,
|
||||
0xefbe4786U,
|
||||
0x0fc19dc6U,
|
||||
0x240ca1ccU,
|
||||
0x2de92c6fU,
|
||||
0x4a7484aaU, // 20
|
||||
0x5cb0a9dcU,
|
||||
0x76f988daU,
|
||||
0x983e5152U,
|
||||
0xa831c66dU,
|
||||
0xb00327c8U,
|
||||
0xbf597fc7U,
|
||||
0xc6e00bf3U,
|
||||
0xd5a79147U,
|
||||
0x06ca6351U,
|
||||
0x14292967U, // 30
|
||||
0x27b70a85U,
|
||||
0x2e1b2138U,
|
||||
0x4d2c6dfcU,
|
||||
0x53380d13U,
|
||||
0x650a7354U,
|
||||
0x766a0abbU,
|
||||
0x81c2c92eU,
|
||||
0x92722c85U,
|
||||
0xa2bfe8a1U,
|
||||
0xa81a664bU, // 40
|
||||
0xc24b8b70U,
|
||||
0xc76c51a3U,
|
||||
0xd192e819U,
|
||||
0xd6990624U,
|
||||
0xf40e3585U,
|
||||
0x106aa070U,
|
||||
0x19a4c116U,
|
||||
0x1e376c08U,
|
||||
0x2748774cU,
|
||||
0x34b0bcb5U, // 50
|
||||
0x391c0cb3U,
|
||||
0x4ed8aa4aU,
|
||||
0x5b9cca4fU,
|
||||
0x682e6ff3U,
|
||||
0x748f82eeU,
|
||||
0x78a5636fU,
|
||||
0x84c87814U,
|
||||
0x8cc70208U,
|
||||
0x90befffaU,
|
||||
0xa4506cebU, // 60
|
||||
0xbef9a3f7U,
|
||||
0xc67178f2U,
|
||||
0x98c7e2a2U,
|
||||
0xfc08884dU,
|
||||
0xcd2a11aeU,
|
||||
0x510e527fU,
|
||||
0x9b05688cU,
|
||||
0xC3910C8EU,
|
||||
0xfb6feee7U,
|
||||
0x2a01a605U, // 70
|
||||
0x0c2e12e0U,
|
||||
0x4498517BU,
|
||||
0x6a09e667U,
|
||||
0xa4ce148bU,
|
||||
0x95F61999U,
|
||||
0xc19bf174U,
|
||||
0xBB67AE85U,
|
||||
0x3C6EF372U,
|
||||
0xA54FF53AU,
|
||||
0x1F83D9ABU, // 80
|
||||
0x5BE0CD19U,
|
||||
0x5C5C5C5CU,
|
||||
0x36363636U,
|
||||
0x80000000U,
|
||||
0x000003FFU,
|
||||
0x00000280U,
|
||||
0x000004a0U,
|
||||
0x00000300U
|
||||
};
|
||||
|
||||
#define rotl(x,y) rotate(x,y)
|
||||
#define Ch(x,y,z) bitselect(z,y,x)
|
||||
#define Maj(x,y,z) Ch((x^z),y,z)
|
||||
|
||||
#define EndianSwap(n) (rotl(n & ES[0], 24U)|rotl(n & ES[1], 8U))
|
||||
|
||||
#define Tr2(x) (rotl(x, 30U) ^ rotl(x, 19U) ^ rotl(x, 10U))
|
||||
#define Tr1(x) (rotl(x, 26U) ^ rotl(x, 21U) ^ rotl(x, 7U))
|
||||
#define Wr2(x) (rotl(x, 25U) ^ rotl(x, 14U) ^ (x>>3U))
|
||||
#define Wr1(x) (rotl(x, 15U) ^ rotl(x, 13U) ^ (x>>10U))
|
||||
|
||||
#define RND(a, b, c, d, e, f, g, h, k) \
|
||||
h += Tr1(e); \
|
||||
h += Ch(e, f, g); \
|
||||
h += k; \
|
||||
d += h; \
|
||||
h += Tr2(a); \
|
||||
h += Maj(a, b, c);
|
||||
|
||||
void SHA256(uint4*restrict state0,uint4*restrict state1, const uint4 block0, const uint4 block1, const uint4 block2, const uint4 block3)
|
||||
{
|
||||
uint4 S0 = *state0;
|
||||
uint4 S1 = *state1;
|
||||
|
||||
#define A S0.x
|
||||
#define B S0.y
|
||||
#define C S0.z
|
||||
#define D S0.w
|
||||
#define E S1.x
|
||||
#define F S1.y
|
||||
#define G S1.z
|
||||
#define H S1.w
|
||||
|
||||
uint4 W[4];
|
||||
|
||||
W[ 0].x = block0.x;
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[0]);
|
||||
W[ 0].y = block0.y;
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[1]);
|
||||
W[ 0].z = block0.z;
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[2]);
|
||||
W[ 0].w = block0.w;
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[3]);
|
||||
|
||||
W[ 1].x = block1.x;
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[4]);
|
||||
W[ 1].y = block1.y;
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[5]);
|
||||
W[ 1].z = block1.z;
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[6]);
|
||||
W[ 1].w = block1.w;
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[7]);
|
||||
|
||||
W[ 2].x = block2.x;
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[8]);
|
||||
W[ 2].y = block2.y;
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[9]);
|
||||
W[ 2].z = block2.z;
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[10]);
|
||||
W[ 2].w = block2.w;
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[11]);
|
||||
|
||||
W[ 3].x = block3.x;
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[12]);
|
||||
W[ 3].y = block3.y;
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[13]);
|
||||
W[ 3].z = block3.z;
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[14]);
|
||||
W[ 3].w = block3.w;
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[76]);
|
||||
|
||||
W[ 0].x += Wr1(W[ 3].z) + W[ 2].y + Wr2(W[ 0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[15]);
|
||||
|
||||
W[ 0].y += Wr1(W[ 3].w) + W[ 2].z + Wr2(W[ 0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[16]);
|
||||
|
||||
W[ 0].z += Wr1(W[ 0].x) + W[ 2].w + Wr2(W[ 0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[17]);
|
||||
|
||||
W[ 0].w += Wr1(W[ 0].y) + W[ 3].x + Wr2(W[ 1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[18]);
|
||||
|
||||
W[ 1].x += Wr1(W[ 0].z) + W[ 3].y + Wr2(W[ 1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[19]);
|
||||
|
||||
W[ 1].y += Wr1(W[ 0].w) + W[ 3].z + Wr2(W[ 1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[20]);
|
||||
|
||||
W[ 1].z += Wr1(W[ 1].x) + W[ 3].w + Wr2(W[ 1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[21]);
|
||||
|
||||
W[ 1].w += Wr1(W[ 1].y) + W[ 0].x + Wr2(W[ 2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[22]);
|
||||
|
||||
W[ 2].x += Wr1(W[ 1].z) + W[ 0].y + Wr2(W[ 2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[23]);
|
||||
|
||||
W[ 2].y += Wr1(W[ 1].w) + W[ 0].z + Wr2(W[ 2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[24]);
|
||||
|
||||
W[ 2].z += Wr1(W[ 2].x) + W[ 0].w + Wr2(W[ 2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[25]);
|
||||
|
||||
W[ 2].w += Wr1(W[ 2].y) + W[ 1].x + Wr2(W[ 3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[26]);
|
||||
|
||||
W[ 3].x += Wr1(W[ 2].z) + W[ 1].y + Wr2(W[ 3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[27]);
|
||||
|
||||
W[ 3].y += Wr1(W[ 2].w) + W[ 1].z + Wr2(W[ 3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[28]);
|
||||
|
||||
W[ 3].z += Wr1(W[ 3].x) + W[ 1].w + Wr2(W[ 3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[29]);
|
||||
|
||||
W[ 3].w += Wr1(W[ 3].y) + W[ 2].x + Wr2(W[ 0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[30]);
|
||||
|
||||
W[ 0].x += Wr1(W[ 3].z) + W[ 2].y + Wr2(W[ 0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[31]);
|
||||
|
||||
W[ 0].y += Wr1(W[ 3].w) + W[ 2].z + Wr2(W[ 0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[32]);
|
||||
|
||||
W[ 0].z += Wr1(W[ 0].x) + W[ 2].w + Wr2(W[ 0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[33]);
|
||||
|
||||
W[ 0].w += Wr1(W[ 0].y) + W[ 3].x + Wr2(W[ 1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[34]);
|
||||
|
||||
W[ 1].x += Wr1(W[ 0].z) + W[ 3].y + Wr2(W[ 1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[35]);
|
||||
|
||||
W[ 1].y += Wr1(W[ 0].w) + W[ 3].z + Wr2(W[ 1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[36]);
|
||||
|
||||
W[ 1].z += Wr1(W[ 1].x) + W[ 3].w + Wr2(W[ 1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[37]);
|
||||
|
||||
W[ 1].w += Wr1(W[ 1].y) + W[ 0].x + Wr2(W[ 2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[38]);
|
||||
|
||||
W[ 2].x += Wr1(W[ 1].z) + W[ 0].y + Wr2(W[ 2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[39]);
|
||||
|
||||
W[ 2].y += Wr1(W[ 1].w) + W[ 0].z + Wr2(W[ 2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[40]);
|
||||
|
||||
W[ 2].z += Wr1(W[ 2].x) + W[ 0].w + Wr2(W[ 2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[41]);
|
||||
|
||||
W[ 2].w += Wr1(W[ 2].y) + W[ 1].x + Wr2(W[ 3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[42]);
|
||||
|
||||
W[ 3].x += Wr1(W[ 2].z) + W[ 1].y + Wr2(W[ 3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[43]);
|
||||
|
||||
W[ 3].y += Wr1(W[ 2].w) + W[ 1].z + Wr2(W[ 3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[44]);
|
||||
|
||||
W[ 3].z += Wr1(W[ 3].x) + W[ 1].w + Wr2(W[ 3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[45]);
|
||||
|
||||
W[ 3].w += Wr1(W[ 3].y) + W[ 2].x + Wr2(W[ 0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[46]);
|
||||
|
||||
W[ 0].x += Wr1(W[ 3].z) + W[ 2].y + Wr2(W[ 0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[47]);
|
||||
|
||||
W[ 0].y += Wr1(W[ 3].w) + W[ 2].z + Wr2(W[ 0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[48]);
|
||||
|
||||
W[ 0].z += Wr1(W[ 0].x) + W[ 2].w + Wr2(W[ 0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[49]);
|
||||
|
||||
W[ 0].w += Wr1(W[ 0].y) + W[ 3].x + Wr2(W[ 1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[50]);
|
||||
|
||||
W[ 1].x += Wr1(W[ 0].z) + W[ 3].y + Wr2(W[ 1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[51]);
|
||||
|
||||
W[ 1].y += Wr1(W[ 0].w) + W[ 3].z + Wr2(W[ 1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[52]);
|
||||
|
||||
W[ 1].z += Wr1(W[ 1].x) + W[ 3].w + Wr2(W[ 1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[53]);
|
||||
|
||||
W[ 1].w += Wr1(W[ 1].y) + W[ 0].x + Wr2(W[ 2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[54]);
|
||||
|
||||
W[ 2].x += Wr1(W[ 1].z) + W[ 0].y + Wr2(W[ 2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[55]);
|
||||
|
||||
W[ 2].y += Wr1(W[ 1].w) + W[ 0].z + Wr2(W[ 2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[56]);
|
||||
|
||||
W[ 2].z += Wr1(W[ 2].x) + W[ 0].w + Wr2(W[ 2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[57]);
|
||||
|
||||
W[ 2].w += Wr1(W[ 2].y) + W[ 1].x + Wr2(W[ 3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[58]);
|
||||
|
||||
W[ 3].x += Wr1(W[ 2].z) + W[ 1].y + Wr2(W[ 3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[59]);
|
||||
|
||||
W[ 3].y += Wr1(W[ 2].w) + W[ 1].z + Wr2(W[ 3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[60]);
|
||||
|
||||
W[ 3].z += Wr1(W[ 3].x) + W[ 1].w + Wr2(W[ 3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[61]);
|
||||
|
||||
W[ 3].w += Wr1(W[ 3].y) + W[ 2].x + Wr2(W[ 0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[62]);
|
||||
|
||||
#undef A
|
||||
#undef B
|
||||
#undef C
|
||||
#undef D
|
||||
#undef E
|
||||
#undef F
|
||||
#undef G
|
||||
#undef H
|
||||
|
||||
*state0 += S0;
|
||||
*state1 += S1;
|
||||
}
|
||||
|
||||
void SHA256_fresh(uint4*restrict state0,uint4*restrict state1, const uint4 block0, const uint4 block1, const uint4 block2, const uint4 block3)
|
||||
{
|
||||
#define A (*state0).x
|
||||
#define B (*state0).y
|
||||
#define C (*state0).z
|
||||
#define D (*state0).w
|
||||
#define E (*state1).x
|
||||
#define F (*state1).y
|
||||
#define G (*state1).z
|
||||
#define H (*state1).w
|
||||
|
||||
uint4 W[4];
|
||||
|
||||
W[0].x = block0.x;
|
||||
D= K[63] +W[0].x;
|
||||
H= K[64] +W[0].x;
|
||||
|
||||
W[0].y = block0.y;
|
||||
C= K[65] +Tr1(D)+Ch(D, K[66], K[67])+W[0].y;
|
||||
G= K[68] +C+Tr2(H)+Ch(H, K[69] ,K[70]);
|
||||
|
||||
W[0].z = block0.z;
|
||||
B= K[71] +Tr1(C)+Ch(C,D,K[66])+W[0].z;
|
||||
F= K[72] +B+Tr2(G)+Maj(G,H, K[73]);
|
||||
|
||||
W[0].w = block0.w;
|
||||
A= K[74] +Tr1(B)+Ch(B,C,D)+W[0].w;
|
||||
E= K[75] +A+Tr2(F)+Maj(F,G,H);
|
||||
|
||||
W[1].x = block1.x;
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[4]);
|
||||
W[1].y = block1.y;
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[5]);
|
||||
W[1].z = block1.z;
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[6]);
|
||||
W[1].w = block1.w;
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[7]);
|
||||
|
||||
W[2].x = block2.x;
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[8]);
|
||||
W[2].y = block2.y;
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[9]);
|
||||
W[2].z = block2.z;
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[10]);
|
||||
W[2].w = block2.w;
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[11]);
|
||||
|
||||
W[3].x = block3.x;
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[12]);
|
||||
W[3].y = block3.y;
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[13]);
|
||||
W[3].z = block3.z;
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[14]);
|
||||
W[3].w = block3.w;
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[76]);
|
||||
|
||||
W[0].x += Wr1(W[3].z) + W[2].y + Wr2(W[0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[15]);
|
||||
|
||||
W[0].y += Wr1(W[3].w) + W[2].z + Wr2(W[0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[16]);
|
||||
|
||||
W[0].z += Wr1(W[0].x) + W[2].w + Wr2(W[0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[17]);
|
||||
|
||||
W[0].w += Wr1(W[0].y) + W[3].x + Wr2(W[1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[18]);
|
||||
|
||||
W[1].x += Wr1(W[0].z) + W[3].y + Wr2(W[1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[19]);
|
||||
|
||||
W[1].y += Wr1(W[0].w) + W[3].z + Wr2(W[1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[20]);
|
||||
|
||||
W[1].z += Wr1(W[1].x) + W[3].w + Wr2(W[1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[21]);
|
||||
|
||||
W[1].w += Wr1(W[1].y) + W[0].x + Wr2(W[2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[22]);
|
||||
|
||||
W[2].x += Wr1(W[1].z) + W[0].y + Wr2(W[2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[23]);
|
||||
|
||||
W[2].y += Wr1(W[1].w) + W[0].z + Wr2(W[2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[24]);
|
||||
|
||||
W[2].z += Wr1(W[2].x) + W[0].w + Wr2(W[2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[25]);
|
||||
|
||||
W[2].w += Wr1(W[2].y) + W[1].x + Wr2(W[3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[26]);
|
||||
|
||||
W[3].x += Wr1(W[2].z) + W[1].y + Wr2(W[3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[27]);
|
||||
|
||||
W[3].y += Wr1(W[2].w) + W[1].z + Wr2(W[3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[28]);
|
||||
|
||||
W[3].z += Wr1(W[3].x) + W[1].w + Wr2(W[3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[29]);
|
||||
|
||||
W[3].w += Wr1(W[3].y) + W[2].x + Wr2(W[0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[30]);
|
||||
|
||||
W[0].x += Wr1(W[3].z) + W[2].y + Wr2(W[0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[31]);
|
||||
|
||||
W[0].y += Wr1(W[3].w) + W[2].z + Wr2(W[0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[32]);
|
||||
|
||||
W[0].z += Wr1(W[0].x) + W[2].w + Wr2(W[0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[33]);
|
||||
|
||||
W[0].w += Wr1(W[0].y) + W[3].x + Wr2(W[1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[34]);
|
||||
|
||||
W[1].x += Wr1(W[0].z) + W[3].y + Wr2(W[1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[35]);
|
||||
|
||||
W[1].y += Wr1(W[0].w) + W[3].z + Wr2(W[1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[36]);
|
||||
|
||||
W[1].z += Wr1(W[1].x) + W[3].w + Wr2(W[1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[37]);
|
||||
|
||||
W[1].w += Wr1(W[1].y) + W[0].x + Wr2(W[2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[38]);
|
||||
|
||||
W[2].x += Wr1(W[1].z) + W[0].y + Wr2(W[2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[39]);
|
||||
|
||||
W[2].y += Wr1(W[1].w) + W[0].z + Wr2(W[2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[40]);
|
||||
|
||||
W[2].z += Wr1(W[2].x) + W[0].w + Wr2(W[2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[41]);
|
||||
|
||||
W[2].w += Wr1(W[2].y) + W[1].x + Wr2(W[3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[42]);
|
||||
|
||||
W[3].x += Wr1(W[2].z) + W[1].y + Wr2(W[3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[43]);
|
||||
|
||||
W[3].y += Wr1(W[2].w) + W[1].z + Wr2(W[3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[44]);
|
||||
|
||||
W[3].z += Wr1(W[3].x) + W[1].w + Wr2(W[3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[45]);
|
||||
|
||||
W[3].w += Wr1(W[3].y) + W[2].x + Wr2(W[0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[46]);
|
||||
|
||||
W[0].x += Wr1(W[3].z) + W[2].y + Wr2(W[0].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[0].x+ K[47]);
|
||||
|
||||
W[0].y += Wr1(W[3].w) + W[2].z + Wr2(W[0].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[0].y+ K[48]);
|
||||
|
||||
W[0].z += Wr1(W[0].x) + W[2].w + Wr2(W[0].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[0].z+ K[49]);
|
||||
|
||||
W[0].w += Wr1(W[0].y) + W[3].x + Wr2(W[1].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[0].w+ K[50]);
|
||||
|
||||
W[1].x += Wr1(W[0].z) + W[3].y + Wr2(W[1].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[1].x+ K[51]);
|
||||
|
||||
W[1].y += Wr1(W[0].w) + W[3].z + Wr2(W[1].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[1].y+ K[52]);
|
||||
|
||||
W[1].z += Wr1(W[1].x) + W[3].w + Wr2(W[1].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[1].z+ K[53]);
|
||||
|
||||
W[1].w += Wr1(W[1].y) + W[0].x + Wr2(W[2].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[1].w+ K[54]);
|
||||
|
||||
W[2].x += Wr1(W[1].z) + W[0].y + Wr2(W[2].y);
|
||||
RND(A,B,C,D,E,F,G,H, W[2].x+ K[55]);
|
||||
|
||||
W[2].y += Wr1(W[1].w) + W[0].z + Wr2(W[2].z);
|
||||
RND(H,A,B,C,D,E,F,G, W[2].y+ K[56]);
|
||||
|
||||
W[2].z += Wr1(W[2].x) + W[0].w + Wr2(W[2].w);
|
||||
RND(G,H,A,B,C,D,E,F, W[2].z+ K[57]);
|
||||
|
||||
W[2].w += Wr1(W[2].y) + W[1].x + Wr2(W[3].x);
|
||||
RND(F,G,H,A,B,C,D,E, W[2].w+ K[58]);
|
||||
|
||||
W[3].x += Wr1(W[2].z) + W[1].y + Wr2(W[3].y);
|
||||
RND(E,F,G,H,A,B,C,D, W[3].x+ K[59]);
|
||||
|
||||
W[3].y += Wr1(W[2].w) + W[1].z + Wr2(W[3].z);
|
||||
RND(D,E,F,G,H,A,B,C, W[3].y+ K[60]);
|
||||
|
||||
W[3].z += Wr1(W[3].x) + W[1].w + Wr2(W[3].w);
|
||||
RND(C,D,E,F,G,H,A,B, W[3].z+ K[61]);
|
||||
|
||||
W[3].w += Wr1(W[3].y) + W[2].x + Wr2(W[0].x);
|
||||
RND(B,C,D,E,F,G,H,A, W[3].w+ K[62]);
|
||||
|
||||
#undef A
|
||||
#undef B
|
||||
#undef C
|
||||
#undef D
|
||||
#undef E
|
||||
#undef F
|
||||
#undef G
|
||||
#undef H
|
||||
|
||||
*state0 += (uint4)(K[73], K[77], K[78], K[79]);
|
||||
*state1 += (uint4)(K[66], K[67], K[80], K[81]);
|
||||
}
|
||||
|
||||
__constant uint fixedW[64] =
|
||||
{
|
||||
0x428a2f99,0xf1374491,0xb5c0fbcf,0xe9b5dba5,0x3956c25b,0x59f111f1,0x923f82a4,0xab1c5ed5,
|
||||
0xd807aa98,0x12835b01,0x243185be,0x550c7dc3,0x72be5d74,0x80deb1fe,0x9bdc06a7,0xc19bf794,
|
||||
0xf59b89c2,0x73924787,0x23c6886e,0xa42ca65c,0x15ed3627,0x4d6edcbf,0xe28217fc,0xef02488f,
|
||||
0xb707775c,0x0468c23f,0xe7e72b4c,0x49e1f1a2,0x4b99c816,0x926d1570,0xaa0fc072,0xadb36e2c,
|
||||
0xad87a3ea,0xbcb1d3a3,0x7b993186,0x562b9420,0xbff3ca0c,0xda4b0c23,0x6cd8711a,0x8f337caa,
|
||||
0xc91b1417,0xc359dce1,0xa83253a7,0x3b13c12d,0x9d3d725d,0xd9031a84,0xb1a03340,0x16f58012,
|
||||
0xe64fb6a2,0xe84d923a,0xe93a5730,0x09837686,0x078ff753,0x29833341,0xd5de0b7e,0x6948ccf4,
|
||||
0xe0a1adbe,0x7c728e11,0x511c78e4,0x315b45bd,0xfca71413,0xea28f96a,0x79703128,0x4e1ef848,
|
||||
};
|
||||
|
||||
void SHA256_fixed(uint4*restrict state0,uint4*restrict state1)
|
||||
{
|
||||
uint4 S0 = *state0;
|
||||
uint4 S1 = *state1;
|
||||
|
||||
#define A S0.x
|
||||
#define B S0.y
|
||||
#define C S0.z
|
||||
#define D S0.w
|
||||
#define E S1.x
|
||||
#define F S1.y
|
||||
#define G S1.z
|
||||
#define H S1.w
|
||||
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[0]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[1]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[2]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[3]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[4]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[5]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[6]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[7]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[8]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[9]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[10]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[11]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[12]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[13]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[14]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[15]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[16]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[17]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[18]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[19]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[20]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[21]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[22]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[23]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[24]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[25]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[26]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[27]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[28]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[29]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[30]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[31]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[32]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[33]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[34]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[35]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[36]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[37]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[38]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[39]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[40]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[41]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[42]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[43]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[44]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[45]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[46]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[47]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[48]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[49]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[50]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[51]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[52]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[53]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[54]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[55]);
|
||||
RND(A,B,C,D,E,F,G,H, fixedW[56]);
|
||||
RND(H,A,B,C,D,E,F,G, fixedW[57]);
|
||||
RND(G,H,A,B,C,D,E,F, fixedW[58]);
|
||||
RND(F,G,H,A,B,C,D,E, fixedW[59]);
|
||||
RND(E,F,G,H,A,B,C,D, fixedW[60]);
|
||||
RND(D,E,F,G,H,A,B,C, fixedW[61]);
|
||||
RND(C,D,E,F,G,H,A,B, fixedW[62]);
|
||||
RND(B,C,D,E,F,G,H,A, fixedW[63]);
|
||||
|
||||
#undef A
|
||||
#undef B
|
||||
#undef C
|
||||
#undef D
|
||||
#undef E
|
||||
#undef F
|
||||
#undef G
|
||||
#undef H
|
||||
*state0 += S0;
|
||||
*state1 += S1;
|
||||
}
|
||||
|
||||
void shittify(uint4 B[8])
|
||||
{
|
||||
uint4 tmp[4];
|
||||
tmp[0] = (uint4)(B[1].x,B[2].y,B[3].z,B[0].w);
|
||||
tmp[1] = (uint4)(B[2].x,B[3].y,B[0].z,B[1].w);
|
||||
tmp[2] = (uint4)(B[3].x,B[0].y,B[1].z,B[2].w);
|
||||
tmp[3] = (uint4)(B[0].x,B[1].y,B[2].z,B[3].w);
|
||||
|
||||
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i] = EndianSwap(tmp[i]);
|
||||
|
||||
tmp[0] = (uint4)(B[5].x,B[6].y,B[7].z,B[4].w);
|
||||
tmp[1] = (uint4)(B[6].x,B[7].y,B[4].z,B[5].w);
|
||||
tmp[2] = (uint4)(B[7].x,B[4].y,B[5].z,B[6].w);
|
||||
tmp[3] = (uint4)(B[4].x,B[5].y,B[6].z,B[7].w);
|
||||
|
||||
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i+4] = EndianSwap(tmp[i]);
|
||||
}
|
||||
|
||||
void unshittify(uint4 B[8])
|
||||
{
|
||||
uint4 tmp[4];
|
||||
tmp[0] = (uint4)(B[3].x,B[2].y,B[1].z,B[0].w);
|
||||
tmp[1] = (uint4)(B[0].x,B[3].y,B[2].z,B[1].w);
|
||||
tmp[2] = (uint4)(B[1].x,B[0].y,B[3].z,B[2].w);
|
||||
tmp[3] = (uint4)(B[2].x,B[1].y,B[0].z,B[3].w);
|
||||
|
||||
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i] = EndianSwap(tmp[i]);
|
||||
|
||||
tmp[0] = (uint4)(B[7].x,B[6].y,B[5].z,B[4].w);
|
||||
tmp[1] = (uint4)(B[4].x,B[7].y,B[6].z,B[5].w);
|
||||
tmp[2] = (uint4)(B[5].x,B[4].y,B[7].z,B[6].w);
|
||||
tmp[3] = (uint4)(B[6].x,B[5].y,B[4].z,B[7].w);
|
||||
|
||||
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i+4] = EndianSwap(tmp[i]);
|
||||
}
|
||||
|
||||
void salsa(uint4 B[8])
|
||||
{
|
||||
uint4 w[4];
|
||||
|
||||
|
||||
for(uint i=0; i<4; ++i)
|
||||
w[i] = (B[i]^=B[i+4]);
|
||||
|
||||
|
||||
for(uint i=0; i<4; ++i)
|
||||
{
|
||||
w[0] ^= rotl(w[3] +w[2] , 7U);
|
||||
w[1] ^= rotl(w[0] +w[3] , 9U);
|
||||
w[2] ^= rotl(w[1] +w[0] ,13U);
|
||||
w[3] ^= rotl(w[2] +w[1] ,18U);
|
||||
w[2] ^= rotl(w[3].wxyz+w[0].zwxy, 7U);
|
||||
w[1] ^= rotl(w[2].wxyz+w[3].zwxy, 9U);
|
||||
w[0] ^= rotl(w[1].wxyz+w[2].zwxy,13U);
|
||||
w[3] ^= rotl(w[0].wxyz+w[1].zwxy,18U);
|
||||
}
|
||||
|
||||
|
||||
for(uint i=0; i<4; ++i)
|
||||
w[i] = (B[i+4]^=(B[i]+=w[i]));
|
||||
|
||||
|
||||
for(uint i=0; i<4; ++i)
|
||||
{
|
||||
w[0] ^= rotl(w[3] +w[2] , 7U);
|
||||
w[1] ^= rotl(w[0] +w[3] , 9U);
|
||||
w[2] ^= rotl(w[1] +w[0] ,13U);
|
||||
w[3] ^= rotl(w[2] +w[1] ,18U);
|
||||
w[2] ^= rotl(w[3].wxyz+w[0].zwxy, 7U);
|
||||
w[1] ^= rotl(w[2].wxyz+w[3].zwxy, 9U);
|
||||
w[0] ^= rotl(w[1].wxyz+w[2].zwxy,13U);
|
||||
w[3] ^= rotl(w[0].wxyz+w[1].zwxy,18U);
|
||||
}
|
||||
|
||||
|
||||
for(uint i=0; i<4; ++i)
|
||||
B[i+4] += w[i];
|
||||
}
|
||||
|
||||
#define Coord(x,y,z) x+y*(x ## SIZE)+z*(y ## SIZE)*(x ## SIZE)
|
||||
#define CO Coord(z,x,y)
|
||||
|
||||
void scrypt_core(uint4 X[8], __global uint4*restrict lookup)
|
||||
{
|
||||
shittify(X);
|
||||
const uint zSIZE = 8;
|
||||
const uint ySIZE = (1024/LOOKUP_GAP+(1024%LOOKUP_GAP>0));
|
||||
const uint xSIZE = CONCURRENT_THREADS;
|
||||
uint x = get_global_id(0)%xSIZE;
|
||||
|
||||
for(uint y=0; y<1024/LOOKUP_GAP; ++y)
|
||||
{
|
||||
|
||||
for(uint z=0; z<zSIZE; ++z)
|
||||
lookup[CO] = X[z];
|
||||
for(uint i=0; i<LOOKUP_GAP; ++i)
|
||||
salsa(X);
|
||||
}
|
||||
for (uint i=0; i<1024; ++i)
|
||||
{
|
||||
uint j = X[7].x & K[85];
|
||||
uint y = (j/LOOKUP_GAP);
|
||||
|
||||
if (j&1)
|
||||
{
|
||||
uint4 V[8];
|
||||
for(uint z=0; z<zSIZE; ++z)
|
||||
V[z] = lookup[CO];
|
||||
|
||||
salsa(V);
|
||||
|
||||
for(uint z=0; z<zSIZE; ++z)
|
||||
X[z] ^= V[z];
|
||||
} else {
|
||||
for(uint z=0; z<zSIZE; ++z)
|
||||
X[z] ^= lookup[CO];
|
||||
}
|
||||
|
||||
salsa(X);
|
||||
}
|
||||
unshittify(X);
|
||||
}
|
||||
|
||||
#define FOUND (0xFF)
|
||||
#define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
|
||||
|
||||
__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
|
||||
__kernel void search(__global const uint4 * restrict input,
|
||||
volatile __global uint*restrict output, __global uint4*restrict padcache,
|
||||
const uint4 midstate0, const uint4 midstate16, const uint target)
|
||||
{
|
||||
uint gid = get_global_id(0);
|
||||
uint4 X[8];
|
||||
uint4 tstate0, tstate1, ostate0, ostate1, tmp0, tmp1;
|
||||
uint4 data = (uint4)(input[4].x,input[4].y,input[4].z,gid);
|
||||
uint4 pad0 = midstate0, pad1 = midstate16;
|
||||
|
||||
SHA256(&pad0,&pad1, data, (uint4)(K[84],0,0,0), (uint4)(0,0,0,0), (uint4)(0,0,0, K[86]));
|
||||
SHA256_fresh(&ostate0,&ostate1, pad0^ K[82], pad1^ K[82], K[82], K[82]);
|
||||
SHA256_fresh(&tstate0,&tstate1, pad0^ K[83], pad1^ K[83], K[83], K[83]);
|
||||
|
||||
tmp0 = tstate0;
|
||||
tmp1 = tstate1;
|
||||
SHA256(&tstate0, &tstate1, input[0],input[1],input[2],input[3]);
|
||||
|
||||
|
||||
for (uint i=0; i<4; i++)
|
||||
{
|
||||
pad0 = tstate0;
|
||||
pad1 = tstate1;
|
||||
X[i<<1 ] = ostate0;
|
||||
X[(i<<1)+1] = ostate1;
|
||||
|
||||
SHA256(&pad0,&pad1, data, (uint4)(i+1,K[84],0,0), (uint4)(0,0,0,0), (uint4)(0,0,0, K[87]));
|
||||
SHA256(X+(i<<1),X+(i<<1)+1, pad0, pad1, (uint4)(K[84], 0U, 0U, 0U), (uint4)(0U, 0U, 0U, K[88]));
|
||||
}
|
||||
scrypt_core(X,padcache);
|
||||
SHA256(&tmp0,&tmp1, X[0], X[1], X[2], X[3]);
|
||||
SHA256(&tmp0,&tmp1, X[4], X[5], X[6], X[7]);
|
||||
SHA256_fixed(&tmp0,&tmp1);
|
||||
SHA256(&ostate0,&ostate1, tmp0, tmp1, (uint4)(K[84], 0U, 0U, 0U), (uint4)(0U, 0U, 0U, K[88]));
|
||||
|
||||
bool result = (EndianSwap(ostate1.w) <= target);
|
||||
if (result)
|
||||
SETFOUND(gid);
|
||||
}
|
||||
5
miner.h
5
miner.h
@ -372,7 +372,10 @@ enum dev_enable {
|
||||
|
||||
enum cl_kernels {
|
||||
KL_NONE,
|
||||
KL_SCRYPT,
|
||||
KL_ALEXKARNEW,
|
||||
KL_ALEXKAROLD,
|
||||
KL_CKOLIVAS,
|
||||
KL_ZUIKKIS,
|
||||
};
|
||||
|
||||
enum dev_reason {
|
||||
|
||||
34
ocl.c
34
ocl.c
@ -400,8 +400,8 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
|
||||
char numbuf[16];
|
||||
|
||||
if (cgpu->kernel == KL_NONE) {
|
||||
applog(LOG_INFO, "Selecting scrypt kernel");
|
||||
clState->chosen_kernel = KL_SCRYPT;
|
||||
applog(LOG_INFO, "Selecting kernel ckolivas");
|
||||
clState->chosen_kernel = KL_CKOLIVAS;
|
||||
cgpu->kernel = clState->chosen_kernel;
|
||||
} else {
|
||||
clState->chosen_kernel = cgpu->kernel;
|
||||
@ -415,12 +415,32 @@ _clState *initCl(unsigned int gpu, char *name, size_t nameSize)
|
||||
else if (preferred_vwidth > 2)
|
||||
preferred_vwidth = 2;
|
||||
|
||||
/* All available kernels only support vector 1 */
|
||||
cgpu->vwidth = 1;
|
||||
|
||||
switch (clState->chosen_kernel) {
|
||||
case KL_SCRYPT:
|
||||
strcpy(filename, SCRYPT_KERNNAME".cl");
|
||||
strcpy(binaryfilename, SCRYPT_KERNNAME);
|
||||
/* Scrypt only supports vector 1 */
|
||||
cgpu->vwidth = 1;
|
||||
case KL_ALEXKARNEW:
|
||||
applog(LOG_WARNING, "Kernel alexkarnew is experimental.");
|
||||
strcpy(filename, ALEXKARNEW_KERNNAME".cl");
|
||||
strcpy(binaryfilename, ALEXKARNEW_KERNNAME);
|
||||
break;
|
||||
case KL_ALEXKAROLD:
|
||||
applog(LOG_WARNING, "Kernel alexkarold is experimental.");
|
||||
strcpy(filename, ALEXKAROLD_KERNNAME".cl");
|
||||
strcpy(binaryfilename, ALEXKAROLD_KERNNAME);
|
||||
break;
|
||||
case KL_CKOLIVAS:
|
||||
strcpy(filename, CKOLIVAS_KERNNAME".cl");
|
||||
strcpy(binaryfilename, CKOLIVAS_KERNNAME);
|
||||
break;
|
||||
case KL_ZUIKKIS:
|
||||
applog(LOG_WARNING, "Kernel zuikkis is experimental.");
|
||||
strcpy(filename, ZUIKKIS_KERNNAME".cl");
|
||||
strcpy(binaryfilename, ZUIKKIS_KERNNAME);
|
||||
/* Kernel only supports lookup-gap 2 */
|
||||
cgpu->lookup_gap = 2;
|
||||
/* Kernel only supports worksize 256 */
|
||||
cgpu->work_size = 256;
|
||||
break;
|
||||
case KL_NONE: /* Shouldn't happen */
|
||||
break;
|
||||
|
||||
13
sgminer.c
13
sgminer.c
@ -4209,8 +4209,17 @@ void write_config(FILE *fcfg)
|
||||
switch (gpus[i].kernel) {
|
||||
case KL_NONE: // Shouldn't happen
|
||||
break;
|
||||
case KL_SCRYPT:
|
||||
fprintf(fcfg, "scrypt");
|
||||
case KL_ALEXKARNEW:
|
||||
fprintf(fcfg, ALEXKARNEW_KERNNAME);
|
||||
break;
|
||||
case KL_ALEXKAROLD:
|
||||
fprintf(fcfg, ALEXKAROLD_KERNNAME);
|
||||
break;
|
||||
case KL_CKOLIVAS:
|
||||
fprintf(fcfg, CKOLIVAS_KERNNAME);
|
||||
break;
|
||||
case KL_ZUIKKIS:
|
||||
fprintf(fcfg, ZUIKKIS_KERNNAME);
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user