mirror of
https://github.com/xmrig/xmrig.git
synced 2024-09-30 06:21:17 +00:00
Merge branch 'dev'
This commit is contained in:
XMRig
commit
f9c65f3bbf
35 changed files with 2114 additions and 6 deletions
|
|
@ -1,4 +1,5 @@
|
||||||
# v5.10.0
|
# v5.10.0
|
||||||
|
- [#1602](https://github.com/xmrig/xmrig/pull/1602) Added AMD GPUs support for AstroBWT algorithm.
|
||||||
- [#1590](https://github.com/xmrig/xmrig/pull/1590) MSR mod automatically deactivated after switching from RandomX algorithms.
|
- [#1590](https://github.com/xmrig/xmrig/pull/1590) MSR mod automatically deactivated after switching from RandomX algorithms.
|
||||||
- [#1592](https://github.com/xmrig/xmrig/pull/1592) Added AVX2 optimized code for AstroBWT algorithm.
|
- [#1592](https://github.com/xmrig/xmrig/pull/1592) Added AVX2 optimized code for AstroBWT algorithm.
|
||||||
- Added new config option `astrobwt-avx2` in `cpu` object and command line option `--astrobwt-avx2`.
|
- Added new config option `astrobwt-avx2` in `cpu` object and command line option `--astrobwt-avx2`.
|
||||||
|
|
|
||||||
|
|
@ -76,6 +76,15 @@ function rx()
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
function astrobwt()
|
||||||
|
{
|
||||||
|
const astrobwt = opencl_minify(addIncludes('astrobwt.cl', [ 'BWT.cl', 'salsa20.cl', 'sha3.cl' ]));
|
||||||
|
|
||||||
|
// fs.writeFileSync('astrobwt_gen.cl', astrobwt);
|
||||||
|
fs.writeFileSync('astrobwt_cl.h', text2h(astrobwt, 'xmrig', 'astrobwt_cl'));
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
process.chdir(path.resolve('src/backend/opencl/cl/cn'));
|
process.chdir(path.resolve('src/backend/opencl/cl/cn'));
|
||||||
|
|
||||||
cn();
|
cn();
|
||||||
|
|
@ -86,3 +95,8 @@ process.chdir(cwd);
|
||||||
process.chdir(path.resolve('src/backend/opencl/cl/rx'));
|
process.chdir(path.resolve('src/backend/opencl/cl/rx'));
|
||||||
|
|
||||||
rx();
|
rx();
|
||||||
|
|
||||||
|
process.chdir(cwd);
|
||||||
|
process.chdir(path.resolve('src/backend/opencl/cl/astrobwt'));
|
||||||
|
|
||||||
|
astrobwt();
|
||||||
|
|
|
||||||
|
|
@ -36,6 +36,7 @@
|
||||||
#include "backend/opencl/OclLaunchData.h"
|
#include "backend/opencl/OclLaunchData.h"
|
||||||
#include "backend/opencl/OclWorker.h"
|
#include "backend/opencl/OclWorker.h"
|
||||||
#include "backend/opencl/runners/tools/OclSharedState.h"
|
#include "backend/opencl/runners/tools/OclSharedState.h"
|
||||||
|
#include "backend/opencl/runners/OclAstroBWTRunner.h"
|
||||||
#include "backend/opencl/wrappers/OclContext.h"
|
#include "backend/opencl/wrappers/OclContext.h"
|
||||||
#include "backend/opencl/wrappers/OclLib.h"
|
#include "backend/opencl/wrappers/OclLib.h"
|
||||||
#include "base/io/log/Log.h"
|
#include "base/io/log/Log.h"
|
||||||
|
|
@ -202,6 +203,14 @@ public:
|
||||||
|
|
||||||
Log::print(WHITE_BOLD("| # | GPU | BUS ID | I | W | SI | MC | U | MEM | NAME"));
|
Log::print(WHITE_BOLD("| # | GPU | BUS ID | I | W | SI | MC | U | MEM | NAME"));
|
||||||
|
|
||||||
|
size_t algo_l3 = algo.l3();
|
||||||
|
|
||||||
|
# ifdef XMRIG_ALGO_ASTROBWT
|
||||||
|
if (algo.family() == Algorithm::ASTROBWT) {
|
||||||
|
algo_l3 = OclAstroBWTRunner::BWT_DATA_STRIDE * 17 + 324;
|
||||||
|
}
|
||||||
|
# endif
|
||||||
|
|
||||||
size_t i = 0;
|
size_t i = 0;
|
||||||
for (const auto &data : threads) {
|
for (const auto &data : threads) {
|
||||||
Log::print("|" CYAN_BOLD("%3zu") " |" CYAN_BOLD("%4u") " |" YELLOW(" %7s") " |" CYAN_BOLD("%5u") " |" CYAN_BOLD("%3u") " |"
|
Log::print("|" CYAN_BOLD("%3zu") " |" CYAN_BOLD("%4u") " |" YELLOW(" %7s") " |" CYAN_BOLD("%5u") " |" CYAN_BOLD("%3u") " |"
|
||||||
|
|
@ -214,7 +223,7 @@ public:
|
||||||
data.thread.stridedIndex(),
|
data.thread.stridedIndex(),
|
||||||
data.thread.stridedIndex() == 2 ? std::to_string(data.thread.memChunk()).c_str() : "-",
|
data.thread.stridedIndex() == 2 ? std::to_string(data.thread.memChunk()).c_str() : "-",
|
||||||
data.thread.unrollFactor(),
|
data.thread.unrollFactor(),
|
||||||
data.thread.intensity() * algo.l3() / oneMiB,
|
data.thread.intensity() * algo_l3 / oneMiB,
|
||||||
data.device.printableName().data()
|
data.device.printableName().data()
|
||||||
);
|
);
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -220,6 +220,7 @@ void xmrig::OclConfig::generate()
|
||||||
count += xmrig::generate<Algorithm::CN_HEAVY>(m_threads, devices);
|
count += xmrig::generate<Algorithm::CN_HEAVY>(m_threads, devices);
|
||||||
count += xmrig::generate<Algorithm::CN_PICO>(m_threads, devices);
|
count += xmrig::generate<Algorithm::CN_PICO>(m_threads, devices);
|
||||||
count += xmrig::generate<Algorithm::RANDOM_X>(m_threads, devices);
|
count += xmrig::generate<Algorithm::RANDOM_X>(m_threads, devices);
|
||||||
|
count += xmrig::generate<Algorithm::ASTROBWT>(m_threads, devices);
|
||||||
|
|
||||||
m_shouldSave = count > 0;
|
m_shouldSave = count > 0;
|
||||||
}
|
}
|
||||||
|
|
|
||||||
|
|
@ -130,6 +130,15 @@ size_t inline generate<Algorithm::RANDOM_X>(Threads<OclThreads> &threads, const
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
|
|
||||||
|
#ifdef XMRIG_ALGO_ASTROBWT
|
||||||
|
template<>
|
||||||
|
size_t inline generate<Algorithm::ASTROBWT>(Threads<OclThreads>& threads, const std::vector<OclDevice>& devices)
|
||||||
|
{
|
||||||
|
return generate("astrobwt", threads, Algorithm::ASTROBWT_DERO, devices);
|
||||||
|
}
|
||||||
|
#endif
|
||||||
|
|
||||||
|
|
||||||
static inline std::vector<OclDevice> filterDevices(const std::vector<OclDevice> &devices, const std::vector<uint32_t> &hints)
|
static inline std::vector<OclDevice> filterDevices(const std::vector<OclDevice> &devices, const std::vector<uint32_t> &hints)
|
||||||
{
|
{
|
||||||
std::vector<OclDevice> out;
|
std::vector<OclDevice> out;
|
||||||
|
|
|
||||||
|
|
@ -123,7 +123,9 @@ rapidjson::Value xmrig::OclThread::toJSON(rapidjson::Document &doc) const
|
||||||
|
|
||||||
out.AddMember(StringRef(kIndex), index(), allocator);
|
out.AddMember(StringRef(kIndex), index(), allocator);
|
||||||
out.AddMember(StringRef(kIntensity), intensity(), allocator);
|
out.AddMember(StringRef(kIntensity), intensity(), allocator);
|
||||||
out.AddMember(StringRef(kWorksize), worksize(), allocator);
|
if (!m_fields.test(ASTROBWT_FIELDS)) {
|
||||||
|
out.AddMember(StringRef(kWorksize), worksize(), allocator);
|
||||||
|
}
|
||||||
|
|
||||||
if (m_fields.test(STRIDED_INDEX_FIELD)) {
|
if (m_fields.test(STRIDED_INDEX_FIELD)) {
|
||||||
Value si(kArrayType);
|
Value si(kArrayType);
|
||||||
|
|
@ -149,7 +151,7 @@ rapidjson::Value xmrig::OclThread::toJSON(rapidjson::Document &doc) const
|
||||||
out.AddMember(StringRef(kDatasetHost), isDatasetHost(), allocator);
|
out.AddMember(StringRef(kDatasetHost), isDatasetHost(), allocator);
|
||||||
# endif
|
# endif
|
||||||
}
|
}
|
||||||
else {
|
else if (!m_fields.test(ASTROBWT_FIELDS)) {
|
||||||
out.AddMember(StringRef(kUnroll), unrollFactor(), allocator);
|
out.AddMember(StringRef(kUnroll), unrollFactor(), allocator);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -81,6 +81,20 @@ public:
|
||||||
}
|
}
|
||||||
# endif
|
# endif
|
||||||
|
|
||||||
|
# ifdef XMRIG_ALGO_ASTROBWT
|
||||||
|
OclThread(uint32_t index, uint32_t intensity, uint32_t threads) :
|
||||||
|
m_fields(4),
|
||||||
|
m_threads(threads, -1),
|
||||||
|
m_index(index),
|
||||||
|
m_memChunk(0),
|
||||||
|
m_stridedIndex(0),
|
||||||
|
m_unrollFactor(1),
|
||||||
|
m_worksize(1)
|
||||||
|
{
|
||||||
|
setIntensity(intensity);
|
||||||
|
}
|
||||||
|
# endif
|
||||||
|
|
||||||
OclThread(const rapidjson::Value &value);
|
OclThread(const rapidjson::Value &value);
|
||||||
|
|
||||||
inline bool isAsm() const { return m_gcnAsm; }
|
inline bool isAsm() const { return m_gcnAsm; }
|
||||||
|
|
@ -105,6 +119,7 @@ private:
|
||||||
enum Fields {
|
enum Fields {
|
||||||
STRIDED_INDEX_FIELD,
|
STRIDED_INDEX_FIELD,
|
||||||
RANDOMX_FIELDS,
|
RANDOMX_FIELDS,
|
||||||
|
ASTROBWT_FIELDS,
|
||||||
FIELD_MAX
|
FIELD_MAX
|
||||||
};
|
};
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -41,6 +41,10 @@
|
||||||
# include "backend/opencl/runners/OclRxVmRunner.h"
|
# include "backend/opencl/runners/OclRxVmRunner.h"
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
|
#ifdef XMRIG_ALGO_ASTROBWT
|
||||||
|
# include "backend/opencl/runners/OclAstroBWTRunner.h"
|
||||||
|
#endif
|
||||||
|
|
||||||
#ifdef XMRIG_ALGO_CN_GPU
|
#ifdef XMRIG_ALGO_CN_GPU
|
||||||
# include "backend/opencl/runners/OclRyoRunner.h"
|
# include "backend/opencl/runners/OclRyoRunner.h"
|
||||||
#endif
|
#endif
|
||||||
|
|
@ -96,6 +100,12 @@ xmrig::OclWorker::OclWorker(size_t id, const OclLaunchData &data) :
|
||||||
# endif
|
# endif
|
||||||
break;
|
break;
|
||||||
|
|
||||||
|
case Algorithm::ASTROBWT:
|
||||||
|
# ifdef XMRIG_ALGO_ASTROBWT
|
||||||
|
m_runner = new OclAstroBWTRunner(id, data);
|
||||||
|
# endif
|
||||||
|
break;
|
||||||
|
|
||||||
default:
|
default:
|
||||||
# ifdef XMRIG_ALGO_CN_GPU
|
# ifdef XMRIG_ALGO_CN_GPU
|
||||||
if (m_algorithm == Algorithm::CN_GPU) {
|
if (m_algorithm == Algorithm::CN_GPU) {
|
||||||
|
|
@ -148,6 +158,8 @@ void xmrig::OclWorker::start()
|
||||||
{
|
{
|
||||||
cl_uint results[0x100];
|
cl_uint results[0x100];
|
||||||
|
|
||||||
|
const uint32_t runnerRoundSize = m_runner->roundSize();
|
||||||
|
|
||||||
while (Nonce::sequence(Nonce::OPENCL) > 0) {
|
while (Nonce::sequence(Nonce::OPENCL) > 0) {
|
||||||
if (!isReady()) {
|
if (!isReady()) {
|
||||||
m_sharedData.setResumeCounter(0);
|
m_sharedData.setResumeCounter(0);
|
||||||
|
|
@ -186,7 +198,7 @@ void xmrig::OclWorker::start()
|
||||||
JobResults::submit(m_job.currentJob(), results, results[0xFF]);
|
JobResults::submit(m_job.currentJob(), results, results[0xFF]);
|
||||||
}
|
}
|
||||||
|
|
||||||
if (!m_job.nextRound(roundSize(m_intensity), m_intensity)) {
|
if (!m_job.nextRound(roundSize(runnerRoundSize), runnerRoundSize)) {
|
||||||
JobResults::done(m_job.currentJob());
|
JobResults::done(m_job.currentJob());
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
@ -228,7 +240,7 @@ void xmrig::OclWorker::storeStats(uint64_t t)
|
||||||
return;
|
return;
|
||||||
}
|
}
|
||||||
|
|
||||||
m_count += m_intensity;
|
m_count += m_runner->processedHashes();
|
||||||
|
|
||||||
m_sharedData.setRunTime(Chrono::steadyMSecs() - t);
|
m_sharedData.setRunTime(Chrono::steadyMSecs() - t);
|
||||||
|
|
||||||
|
|
|
||||||
|
|
@ -36,6 +36,10 @@
|
||||||
# include "backend/opencl/cl/rx/randomx_cl.h"
|
# include "backend/opencl/cl/rx/randomx_cl.h"
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
|
#ifdef XMRIG_ALGO_ASTROBWT
|
||||||
|
# include "backend/opencl/cl/astrobwt/astrobwt_cl.h"
|
||||||
|
#endif
|
||||||
|
|
||||||
|
|
||||||
const char *xmrig::OclSource::get(const Algorithm &algorithm)
|
const char *xmrig::OclSource::get(const Algorithm &algorithm)
|
||||||
{
|
{
|
||||||
|
|
@ -45,6 +49,12 @@ const char *xmrig::OclSource::get(const Algorithm &algorithm)
|
||||||
}
|
}
|
||||||
# endif
|
# endif
|
||||||
|
|
||||||
|
# ifdef XMRIG_ALGO_ASTROBWT
|
||||||
|
if (algorithm.family() == Algorithm::ASTROBWT) {
|
||||||
|
return astrobwt_cl;
|
||||||
|
}
|
||||||
|
# endif
|
||||||
|
|
||||||
# ifdef XMRIG_ALGO_CN_GPU
|
# ifdef XMRIG_ALGO_CN_GPU
|
||||||
if (algorithm == Algorithm::CN_GPU) {
|
if (algorithm == Algorithm::CN_GPU) {
|
||||||
return cryptonight_gpu_cl;
|
return cryptonight_gpu_cl;
|
||||||
|
|
|
||||||
199
src/backend/opencl/cl/astrobwt/BWT.cl
Normal file
199
src/backend/opencl/cl/astrobwt/BWT.cl
Normal file
|
|
@ -0,0 +1,199 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#define STAGE1_SIZE 147253
|
||||||
|
|
||||||
|
#define COUNTING_SORT_BITS 11
|
||||||
|
#define COUNTING_SORT_SIZE (1 << COUNTING_SORT_BITS)
|
||||||
|
#define FINAL_SORT_BATCH_SIZE COUNTING_SORT_SIZE
|
||||||
|
#define FINAL_SORT_OVERLAP_SIZE 32
|
||||||
|
|
||||||
|
__attribute__((reqd_work_group_size(BWT_GROUP_SIZE, 1, 1)))
|
||||||
|
__kernel void BWT(__global uint8_t* datas, __global uint32_t* data_sizes, uint32_t data_stride, __global uint64_t* indices, __global uint64_t* tmp_indices)
|
||||||
|
{
|
||||||
|
const uint32_t tid = get_local_id(0);
|
||||||
|
const uint32_t gid = get_group_id(0);
|
||||||
|
|
||||||
|
__local int counters[COUNTING_SORT_SIZE][2];
|
||||||
|
|
||||||
|
for (uint32_t i = tid; i < COUNTING_SORT_SIZE * 2; i += BWT_GROUP_SIZE)
|
||||||
|
((__local int*)counters)[i] = 0;
|
||||||
|
|
||||||
|
const uint64_t data_offset = (uint64_t)(gid) * data_stride;
|
||||||
|
|
||||||
|
__global uint8_t* input = datas + data_offset + 128;
|
||||||
|
const uint32_t N = data_sizes[gid] + 1;
|
||||||
|
|
||||||
|
__global uint64_t* p = (__global uint64_t*)(input);
|
||||||
|
volatile __local uint8_t* counters_atomic = (volatile __local uint8_t*)(counters);
|
||||||
|
|
||||||
|
indices += data_offset;
|
||||||
|
tmp_indices += data_offset;
|
||||||
|
|
||||||
|
for (uint32_t i = tid; i < N; i += BWT_GROUP_SIZE)
|
||||||
|
{
|
||||||
|
const uint32_t index = i >> 3;
|
||||||
|
const uint32_t bit_offset = (i & 7) << 3;
|
||||||
|
|
||||||
|
const uint64_t a = p[index];
|
||||||
|
uint64_t b = p[index + 1];
|
||||||
|
if (bit_offset == 0)
|
||||||
|
b = 0;
|
||||||
|
|
||||||
|
uint64_t value = (a >> bit_offset) | (b << (64 - bit_offset));
|
||||||
|
|
||||||
|
uint2 tmp;
|
||||||
|
const uchar4 mask = (uchar4)(3, 2, 1, 0);
|
||||||
|
|
||||||
|
tmp.x = as_uint(shuffle(as_uchar4(as_uint2(value).y), mask));
|
||||||
|
tmp.y = as_uint(shuffle(as_uchar4(as_uint2(value).x), mask));
|
||||||
|
|
||||||
|
value = as_ulong(tmp);
|
||||||
|
|
||||||
|
indices[i] = (value & ((uint64_t)(-1) << 21)) | i;
|
||||||
|
atomic_add((volatile __local int*)(counters_atomic + (((value >> (64 - COUNTING_SORT_BITS * 2)) & (COUNTING_SORT_SIZE - 1)) << 3)), 1);
|
||||||
|
atomic_add((volatile __local int*)(counters_atomic + ((value >> (64 - COUNTING_SORT_BITS)) << 3) + 4), 1);
|
||||||
|
}
|
||||||
|
|
||||||
|
if (tid == 0)
|
||||||
|
{
|
||||||
|
int t0 = counters[0][0];
|
||||||
|
int t1 = counters[0][1];
|
||||||
|
counters[0][0] = t0 - 1;
|
||||||
|
counters[0][1] = t1 - 1;
|
||||||
|
for (uint32_t i = 1; i < COUNTING_SORT_SIZE; ++i)
|
||||||
|
{
|
||||||
|
t0 += counters[i][0];
|
||||||
|
t1 += counters[i][1];
|
||||||
|
counters[i][0] = t0 - 1;
|
||||||
|
counters[i][1] = t1 - 1;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
for (int i = tid; i < N; i += BWT_GROUP_SIZE)
|
||||||
|
{
|
||||||
|
const uint64_t data = indices[i];
|
||||||
|
const int k = atomic_sub((volatile __local int*)(counters_atomic + (((data >> (64 - COUNTING_SORT_BITS * 2)) & (COUNTING_SORT_SIZE - 1)) << 3)), 1);
|
||||||
|
tmp_indices[k] = data;
|
||||||
|
}
|
||||||
|
|
||||||
|
for (int i = N - 1 - tid; i >= 0; i -= BWT_GROUP_SIZE)
|
||||||
|
{
|
||||||
|
const uint64_t data = tmp_indices[i];
|
||||||
|
const int k = atomic_sub((volatile __local int*)(counters_atomic + ((data >> (64 - COUNTING_SORT_BITS)) << 3) + 4), 1);
|
||||||
|
indices[k] = data;
|
||||||
|
}
|
||||||
|
|
||||||
|
__local uint64_t* buf = (__local uint64_t*)(counters);
|
||||||
|
for (uint32_t i = 0; i < N; i += FINAL_SORT_BATCH_SIZE - FINAL_SORT_OVERLAP_SIZE)
|
||||||
|
{
|
||||||
|
const uint32_t len = (N - i < FINAL_SORT_BATCH_SIZE) ? (N - i) : FINAL_SORT_BATCH_SIZE;
|
||||||
|
|
||||||
|
for (uint32_t j = tid; j < len; j += BWT_GROUP_SIZE)
|
||||||
|
buf[j] = indices[i + j];
|
||||||
|
|
||||||
|
if (tid == 0)
|
||||||
|
{
|
||||||
|
uint64_t prev_t = buf[0];
|
||||||
|
for (int i = 1; i < len; ++i)
|
||||||
|
{
|
||||||
|
uint64_t t = buf[i];
|
||||||
|
if (t < prev_t)
|
||||||
|
{
|
||||||
|
const uint64_t t2 = prev_t;
|
||||||
|
int j = i - 1;
|
||||||
|
do
|
||||||
|
{
|
||||||
|
buf[j + 1] = prev_t;
|
||||||
|
--j;
|
||||||
|
if (j < 0)
|
||||||
|
break;
|
||||||
|
prev_t = buf[j];
|
||||||
|
} while (t < prev_t);
|
||||||
|
buf[j + 1] = t;
|
||||||
|
t = t2;
|
||||||
|
}
|
||||||
|
prev_t = t;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
for (uint32_t j = tid; j < len; j += BWT_GROUP_SIZE)
|
||||||
|
indices[i + j] = buf[j];
|
||||||
|
}
|
||||||
|
|
||||||
|
--input;
|
||||||
|
__global uint8_t* output = (__global uint8_t*)(tmp_indices);
|
||||||
|
for (int i = tid; i <= N; i += BWT_GROUP_SIZE)
|
||||||
|
output[i] = input[indices[i] & ((1 << 21) - 1)];
|
||||||
|
}
|
||||||
|
|
||||||
|
__kernel void filter(uint32_t nonce, uint32_t bwt_max_size, __global const uint32_t* hashes, __global uint32_t* filtered_hashes)
|
||||||
|
{
|
||||||
|
const uint32_t global_id = get_global_id(0);
|
||||||
|
|
||||||
|
__global const uint32_t* hash = hashes + global_id * (32 / sizeof(uint32_t));
|
||||||
|
const uint32_t stage2_size = STAGE1_SIZE + (*hash & 0xfffff);
|
||||||
|
|
||||||
|
if (stage2_size < bwt_max_size)
|
||||||
|
{
|
||||||
|
const int index = atomic_add((volatile __global int*)(filtered_hashes), 1) * (36 / sizeof(uint32_t)) + 1;
|
||||||
|
|
||||||
|
filtered_hashes[index] = nonce + global_id;
|
||||||
|
|
||||||
|
#pragma unroll(8)
|
||||||
|
for (uint32_t i = 0; i < 8; ++i)
|
||||||
|
filtered_hashes[index + i + 1] = hash[i];
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
__kernel void prepare_batch2(__global uint32_t* hashes, __global uint32_t* filtered_hashes, __global uint32_t* data_sizes)
|
||||||
|
{
|
||||||
|
const uint32_t global_id = get_global_id(0);
|
||||||
|
const uint32_t N = filtered_hashes[0] - get_global_size(0);
|
||||||
|
|
||||||
|
if (global_id == 0)
|
||||||
|
filtered_hashes[0] = N;
|
||||||
|
|
||||||
|
__global uint32_t* hash = hashes + global_id * 8;
|
||||||
|
__global uint32_t* filtered_hash = filtered_hashes + (global_id + N) * 9 + 1;
|
||||||
|
|
||||||
|
const uint32_t stage2_size = STAGE1_SIZE + (filtered_hash[1] & 0xfffff);
|
||||||
|
data_sizes[global_id] = stage2_size;
|
||||||
|
|
||||||
|
#pragma unroll(8)
|
||||||
|
for (uint32_t i = 0; i < 8; ++i)
|
||||||
|
hash[i] = filtered_hash[i + 1];
|
||||||
|
}
|
||||||
|
|
||||||
|
__kernel void find_shares(__global const uint64_t* hashes, __global const uint32_t* filtered_hashes, uint64_t target, __global uint32_t* shares)
|
||||||
|
{
|
||||||
|
const uint32_t global_index = get_global_id(0);
|
||||||
|
|
||||||
|
if (hashes[global_index * 4 + 3] < target)
|
||||||
|
{
|
||||||
|
const uint32_t idx = atomic_inc(shares + 0xFF);
|
||||||
|
if (idx < 0xFF)
|
||||||
|
shares[idx] = filtered_hashes[(filtered_hashes[0] + global_index) * 9 + 1];
|
||||||
|
}
|
||||||
|
}
|
||||||
35
src/backend/opencl/cl/astrobwt/astrobwt.cl
Normal file
35
src/backend/opencl/cl/astrobwt/astrobwt.cl
Normal file
|
|
@ -0,0 +1,35 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
typedef uchar uint8_t;
|
||||||
|
typedef ushort uint16_t;
|
||||||
|
typedef uint uint32_t;
|
||||||
|
typedef ulong uint64_t;
|
||||||
|
|
||||||
|
typedef int int32_t;
|
||||||
|
typedef long int64_t;
|
||||||
|
|
||||||
|
#include "BWT.cl"
|
||||||
|
#include "salsa20.cl"
|
||||||
|
#include "sha3.cl"
|
||||||
395
src/backend/opencl/cl/astrobwt/astrobwt_cl.h
Normal file
395
src/backend/opencl/cl/astrobwt/astrobwt_cl.h
Normal file
|
|
@ -0,0 +1,395 @@
|
||||||
|
#pragma once
|
||||||
|
|
||||||
|
namespace xmrig {
|
||||||
|
|
||||||
|
static char astrobwt_cl[12378] = {
|
||||||
|
0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x75,0x63,0x68,0x61,0x72,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x3b,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x75,
|
||||||
|
0x73,0x68,0x6f,0x72,0x74,0x20,0x75,0x69,0x6e,0x74,0x31,0x36,0x5f,0x74,0x3b,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x75,0x69,0x6e,0x74,0x20,0x75,0x69,0x6e,
|
||||||
|
0x74,0x33,0x32,0x5f,0x74,0x3b,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x75,0x6c,0x6f,0x6e,0x67,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x3b,0x0a,0x74,
|
||||||
|
0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x69,0x6e,0x74,0x20,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x3b,0x0a,0x74,0x79,0x70,0x65,0x64,0x65,0x66,0x20,0x6c,0x6f,0x6e,0x67,
|
||||||
|
0x20,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x3b,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x53,0x54,0x41,0x47,0x45,0x31,0x5f,0x53,0x49,0x5a,0x45,0x20,0x31,0x34,
|
||||||
|
0x37,0x32,0x35,0x33,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x49,0x54,0x53,0x20,
|
||||||
|
0x31,0x31,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x20,0x28,0x31,
|
||||||
|
0x20,0x3c,0x3c,0x20,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x49,0x54,0x53,0x29,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,
|
||||||
|
0x46,0x49,0x4e,0x41,0x4c,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x41,0x54,0x43,0x48,0x5f,0x53,0x49,0x5a,0x45,0x20,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,
|
||||||
|
0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x46,0x49,0x4e,0x41,0x4c,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x4f,0x56,0x45,0x52,
|
||||||
|
0x4c,0x41,0x50,0x5f,0x53,0x49,0x5a,0x45,0x20,0x33,0x32,0x0a,0x5f,0x5f,0x61,0x74,0x74,0x72,0x69,0x62,0x75,0x74,0x65,0x5f,0x5f,0x28,0x28,0x72,0x65,0x71,0x64,0x5f,
|
||||||
|
0x77,0x6f,0x72,0x6b,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,0x73,0x69,0x7a,0x65,0x28,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x2c,0x31,
|
||||||
|
0x2c,0x31,0x29,0x29,0x29,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,0x42,0x57,0x54,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,
|
||||||
|
0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x64,0x61,0x74,0x61,0x73,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,
|
||||||
|
0x74,0x2a,0x20,0x64,0x61,0x74,0x61,0x5f,0x73,0x69,0x7a,0x65,0x73,0x2c,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x64,0x61,0x74,0x61,0x5f,0x73,0x74,0x72,0x69,
|
||||||
|
0x64,0x65,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x2c,0x5f,0x5f,
|
||||||
|
0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x74,0x6d,0x70,0x5f,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x29,0x0a,0x7b,0x0a,
|
||||||
|
0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x74,0x69,0x64,0x3d,0x67,0x65,0x74,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x5f,0x69,0x64,0x28,
|
||||||
|
0x30,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x69,0x64,0x3d,0x67,0x65,0x74,0x5f,0x67,0x72,0x6f,0x75,0x70,
|
||||||
|
0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x69,0x6e,0x74,0x20,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5b,0x43,0x4f,0x55,
|
||||||
|
0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x5d,0x5b,0x32,0x5d,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,
|
||||||
|
0x5f,0x74,0x20,0x69,0x3d,0x74,0x69,0x64,0x3b,0x20,0x69,0x3c,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x2a,0x32,
|
||||||
|
0x3b,0x20,0x69,0x2b,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x28,0x28,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x69,
|
||||||
|
0x6e,0x74,0x2a,0x29,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x29,0x5b,0x69,0x5d,0x3d,0x30,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,
|
||||||
|
0x5f,0x74,0x20,0x64,0x61,0x74,0x61,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x3d,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x67,0x69,0x64,0x29,0x2a,0x64,
|
||||||
|
0x61,0x74,0x61,0x5f,0x73,0x74,0x72,0x69,0x64,0x65,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x69,0x6e,
|
||||||
|
0x70,0x75,0x74,0x3d,0x64,0x61,0x74,0x61,0x73,0x2b,0x64,0x61,0x74,0x61,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x2b,0x31,0x32,0x38,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,
|
||||||
|
0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x4e,0x3d,0x64,0x61,0x74,0x61,0x5f,0x73,0x69,0x7a,0x65,0x73,0x5b,0x67,0x69,0x64,0x5d,0x2b,0x31,0x3b,0x0a,0x5f,
|
||||||
|
0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x70,0x3d,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,
|
||||||
|
0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x29,0x28,0x69,0x6e,0x70,0x75,0x74,0x29,0x3b,0x0a,0x76,0x6f,0x6c,0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x6c,0x6f,0x63,0x61,
|
||||||
|
0x6c,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5f,0x61,0x74,0x6f,0x6d,0x69,0x63,0x3d,0x28,0x76,0x6f,0x6c,0x61,
|
||||||
|
0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x29,0x28,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x29,
|
||||||
|
0x3b,0x0a,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x2b,0x3d,0x64,0x61,0x74,0x61,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x3b,0x0a,0x74,0x6d,0x70,0x5f,0x69,0x6e,0x64,0x69,
|
||||||
|
0x63,0x65,0x73,0x2b,0x3d,0x64,0x61,0x74,0x61,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,
|
||||||
|
0x69,0x3d,0x74,0x69,0x64,0x3b,0x20,0x69,0x3c,0x4e,0x3b,0x20,0x69,0x2b,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x7b,
|
||||||
|
0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x6e,0x64,0x65,0x78,0x3d,0x69,0x3e,0x3e,0x33,0x3b,0x0a,0x63,0x6f,0x6e,0x73,
|
||||||
|
0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x62,0x69,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x3d,0x28,0x69,0x26,0x37,0x29,0x3c,0x3c,0x33,0x3b,0x0a,
|
||||||
|
0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x61,0x3d,0x70,0x5b,0x69,0x6e,0x64,0x65,0x78,0x5d,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x36,
|
||||||
|
0x34,0x5f,0x74,0x20,0x62,0x3d,0x70,0x5b,0x69,0x6e,0x64,0x65,0x78,0x2b,0x31,0x5d,0x3b,0x0a,0x69,0x66,0x28,0x62,0x69,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x3d,
|
||||||
|
0x3d,0x30,0x29,0x0a,0x62,0x3d,0x30,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x76,0x61,0x6c,0x75,0x65,0x3d,0x28,0x61,0x3e,0x3e,0x62,0x69,0x74,0x5f,
|
||||||
|
0x6f,0x66,0x66,0x73,0x65,0x74,0x29,0x7c,0x28,0x62,0x3c,0x3c,0x28,0x36,0x34,0x2d,0x62,0x69,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x29,0x29,0x3b,0x0a,0x75,0x69,
|
||||||
|
0x6e,0x74,0x32,0x20,0x74,0x6d,0x70,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x63,0x68,0x61,0x72,0x34,0x20,0x6d,0x61,0x73,0x6b,0x3d,0x28,0x75,0x63,0x68,0x61,
|
||||||
|
0x72,0x34,0x29,0x28,0x33,0x2c,0x32,0x2c,0x31,0x2c,0x30,0x29,0x3b,0x0a,0x74,0x6d,0x70,0x2e,0x78,0x3d,0x61,0x73,0x5f,0x75,0x69,0x6e,0x74,0x28,0x73,0x68,0x75,0x66,
|
||||||
|
0x66,0x6c,0x65,0x28,0x61,0x73,0x5f,0x75,0x63,0x68,0x61,0x72,0x34,0x28,0x61,0x73,0x5f,0x75,0x69,0x6e,0x74,0x32,0x28,0x76,0x61,0x6c,0x75,0x65,0x29,0x2e,0x79,0x29,
|
||||||
|
0x2c,0x6d,0x61,0x73,0x6b,0x29,0x29,0x3b,0x0a,0x74,0x6d,0x70,0x2e,0x79,0x3d,0x61,0x73,0x5f,0x75,0x69,0x6e,0x74,0x28,0x73,0x68,0x75,0x66,0x66,0x6c,0x65,0x28,0x61,
|
||||||
|
0x73,0x5f,0x75,0x63,0x68,0x61,0x72,0x34,0x28,0x61,0x73,0x5f,0x75,0x69,0x6e,0x74,0x32,0x28,0x76,0x61,0x6c,0x75,0x65,0x29,0x2e,0x78,0x29,0x2c,0x6d,0x61,0x73,0x6b,
|
||||||
|
0x29,0x29,0x3b,0x0a,0x76,0x61,0x6c,0x75,0x65,0x3d,0x61,0x73,0x5f,0x75,0x6c,0x6f,0x6e,0x67,0x28,0x74,0x6d,0x70,0x29,0x3b,0x0a,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,
|
||||||
|
0x5b,0x69,0x5d,0x3d,0x28,0x76,0x61,0x6c,0x75,0x65,0x26,0x28,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x2d,0x31,0x29,0x3c,0x3c,0x32,0x31,0x29,0x29,
|
||||||
|
0x7c,0x69,0x3b,0x0a,0x61,0x74,0x6f,0x6d,0x69,0x63,0x5f,0x61,0x64,0x64,0x28,0x28,0x76,0x6f,0x6c,0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,
|
||||||
|
0x20,0x69,0x6e,0x74,0x2a,0x29,0x28,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5f,0x61,0x74,0x6f,0x6d,0x69,0x63,0x2b,0x28,0x28,0x28,0x76,0x61,0x6c,0x75,0x65,0x3e,
|
||||||
|
0x3e,0x28,0x36,0x34,0x2d,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x49,0x54,0x53,0x2a,0x32,0x29,0x29,0x26,0x28,0x43,0x4f,0x55,
|
||||||
|
0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x2d,0x31,0x29,0x29,0x3c,0x3c,0x33,0x29,0x29,0x2c,0x31,0x29,0x3b,0x0a,0x61,0x74,0x6f,
|
||||||
|
0x6d,0x69,0x63,0x5f,0x61,0x64,0x64,0x28,0x28,0x76,0x6f,0x6c,0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x69,0x6e,0x74,0x2a,0x29,0x28,
|
||||||
|
0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5f,0x61,0x74,0x6f,0x6d,0x69,0x63,0x2b,0x28,0x28,0x76,0x61,0x6c,0x75,0x65,0x3e,0x3e,0x28,0x36,0x34,0x2d,0x43,0x4f,0x55,
|
||||||
|
0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x49,0x54,0x53,0x29,0x29,0x3c,0x3c,0x33,0x29,0x2b,0x34,0x29,0x2c,0x31,0x29,0x3b,0x0a,0x7d,0x0a,0x69,
|
||||||
|
0x66,0x28,0x74,0x69,0x64,0x3d,0x3d,0x30,0x29,0x0a,0x7b,0x0a,0x69,0x6e,0x74,0x20,0x74,0x30,0x3d,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5b,0x30,0x5d,0x5b,0x30,
|
||||||
|
0x5d,0x3b,0x0a,0x69,0x6e,0x74,0x20,0x74,0x31,0x3d,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5b,0x30,0x5d,0x5b,0x31,0x5d,0x3b,0x0a,0x63,0x6f,0x75,0x6e,0x74,0x65,
|
||||||
|
0x72,0x73,0x5b,0x30,0x5d,0x5b,0x30,0x5d,0x3d,0x74,0x30,0x2d,0x31,0x3b,0x0a,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5b,0x30,0x5d,0x5b,0x31,0x5d,0x3d,0x74,0x31,
|
||||||
|
0x2d,0x31,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x31,0x3b,0x20,0x69,0x3c,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,
|
||||||
|
0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x3b,0x20,0x2b,0x2b,0x69,0x29,0x0a,0x7b,0x0a,0x74,0x30,0x2b,0x3d,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,
|
||||||
|
0x5b,0x69,0x5d,0x5b,0x30,0x5d,0x3b,0x0a,0x74,0x31,0x2b,0x3d,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5b,0x69,0x5d,0x5b,0x31,0x5d,0x3b,0x0a,0x63,0x6f,0x75,0x6e,
|
||||||
|
0x74,0x65,0x72,0x73,0x5b,0x69,0x5d,0x5b,0x30,0x5d,0x3d,0x74,0x30,0x2d,0x31,0x3b,0x0a,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5b,0x69,0x5d,0x5b,0x31,0x5d,0x3d,
|
||||||
|
0x74,0x31,0x2d,0x31,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x74,0x69,0x64,0x3b,0x20,0x69,0x3c,0x4e,0x3b,0x20,0x69,
|
||||||
|
0x2b,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,
|
||||||
|
0x5f,0x74,0x20,0x64,0x61,0x74,0x61,0x3d,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x5b,0x69,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x6b,0x3d,
|
||||||
|
0x61,0x74,0x6f,0x6d,0x69,0x63,0x5f,0x73,0x75,0x62,0x28,0x28,0x76,0x6f,0x6c,0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x69,0x6e,0x74,
|
||||||
|
0x2a,0x29,0x28,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5f,0x61,0x74,0x6f,0x6d,0x69,0x63,0x2b,0x28,0x28,0x28,0x64,0x61,0x74,0x61,0x3e,0x3e,0x28,0x36,0x34,0x2d,
|
||||||
|
0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x49,0x54,0x53,0x2a,0x32,0x29,0x29,0x26,0x28,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,
|
||||||
|
0x5f,0x53,0x4f,0x52,0x54,0x5f,0x53,0x49,0x5a,0x45,0x2d,0x31,0x29,0x29,0x3c,0x3c,0x33,0x29,0x29,0x2c,0x31,0x29,0x3b,0x0a,0x74,0x6d,0x70,0x5f,0x69,0x6e,0x64,0x69,
|
||||||
|
0x63,0x65,0x73,0x5b,0x6b,0x5d,0x3d,0x64,0x61,0x74,0x61,0x3b,0x0a,0x7d,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x4e,0x2d,0x31,0x2d,0x74,0x69,
|
||||||
|
0x64,0x3b,0x20,0x69,0x3e,0x3d,0x30,0x3b,0x20,0x69,0x2d,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x7b,0x0a,0x63,0x6f,
|
||||||
|
0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x64,0x61,0x74,0x61,0x3d,0x74,0x6d,0x70,0x5f,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x5b,0x69,0x5d,
|
||||||
|
0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x6b,0x3d,0x61,0x74,0x6f,0x6d,0x69,0x63,0x5f,0x73,0x75,0x62,0x28,0x28,0x76,0x6f,0x6c,0x61,0x74,0x69,
|
||||||
|
0x6c,0x65,0x20,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x69,0x6e,0x74,0x2a,0x29,0x28,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x5f,0x61,0x74,0x6f,0x6d,0x69,0x63,
|
||||||
|
0x2b,0x28,0x28,0x64,0x61,0x74,0x61,0x3e,0x3e,0x28,0x36,0x34,0x2d,0x43,0x4f,0x55,0x4e,0x54,0x49,0x4e,0x47,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x49,0x54,0x53,0x29,
|
||||||
|
0x29,0x3c,0x3c,0x33,0x29,0x2b,0x34,0x29,0x2c,0x31,0x29,0x3b,0x0a,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x5b,0x6b,0x5d,0x3d,0x64,0x61,0x74,0x61,0x3b,0x0a,0x7d,0x0a,
|
||||||
|
0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x62,0x75,0x66,0x3d,0x28,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,
|
||||||
|
0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x29,0x28,0x63,0x6f,0x75,0x6e,0x74,0x65,0x72,0x73,0x29,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,
|
||||||
|
0x5f,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x4e,0x3b,0x20,0x69,0x2b,0x3d,0x46,0x49,0x4e,0x41,0x4c,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x42,0x41,0x54,0x43,0x48,
|
||||||
|
0x5f,0x53,0x49,0x5a,0x45,0x2d,0x46,0x49,0x4e,0x41,0x4c,0x5f,0x53,0x4f,0x52,0x54,0x5f,0x4f,0x56,0x45,0x52,0x4c,0x41,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x7b,
|
||||||
|
0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6c,0x65,0x6e,0x3d,0x28,0x4e,0x2d,0x69,0x3c,0x46,0x49,0x4e,0x41,0x4c,0x5f,0x53,
|
||||||
|
0x4f,0x52,0x54,0x5f,0x42,0x41,0x54,0x43,0x48,0x5f,0x53,0x49,0x5a,0x45,0x29,0x3f,0x28,0x4e,0x2d,0x69,0x29,0x3a,0x46,0x49,0x4e,0x41,0x4c,0x5f,0x53,0x4f,0x52,0x54,
|
||||||
|
0x5f,0x42,0x41,0x54,0x43,0x48,0x5f,0x53,0x49,0x5a,0x45,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x3d,0x74,0x69,0x64,
|
||||||
|
0x3b,0x20,0x6a,0x3c,0x6c,0x65,0x6e,0x3b,0x20,0x6a,0x2b,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x62,0x75,0x66,0x5b,
|
||||||
|
0x6a,0x5d,0x3d,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x5b,0x69,0x2b,0x6a,0x5d,0x3b,0x0a,0x69,0x66,0x28,0x74,0x69,0x64,0x3d,0x3d,0x30,0x29,0x0a,0x7b,0x0a,0x75,0x69,
|
||||||
|
0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x70,0x72,0x65,0x76,0x5f,0x74,0x3d,0x62,0x75,0x66,0x5b,0x30,0x5d,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,
|
||||||
|
0x3d,0x31,0x3b,0x20,0x69,0x3c,0x6c,0x65,0x6e,0x3b,0x20,0x2b,0x2b,0x69,0x29,0x0a,0x7b,0x0a,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x74,0x3d,0x62,0x75,0x66,
|
||||||
|
0x5b,0x69,0x5d,0x3b,0x0a,0x69,0x66,0x28,0x74,0x3c,0x70,0x72,0x65,0x76,0x5f,0x74,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,
|
||||||
|
0x5f,0x74,0x20,0x74,0x32,0x3d,0x70,0x72,0x65,0x76,0x5f,0x74,0x3b,0x0a,0x69,0x6e,0x74,0x20,0x6a,0x3d,0x69,0x2d,0x31,0x3b,0x0a,0x64,0x6f,0x0a,0x7b,0x0a,0x62,0x75,
|
||||||
|
0x66,0x5b,0x6a,0x2b,0x31,0x5d,0x3d,0x70,0x72,0x65,0x76,0x5f,0x74,0x3b,0x0a,0x2d,0x2d,0x6a,0x3b,0x0a,0x69,0x66,0x28,0x6a,0x3c,0x30,0x29,0x0a,0x62,0x72,0x65,0x61,
|
||||||
|
0x6b,0x3b,0x0a,0x70,0x72,0x65,0x76,0x5f,0x74,0x3d,0x62,0x75,0x66,0x5b,0x6a,0x5d,0x3b,0x0a,0x7d,0x20,0x77,0x68,0x69,0x6c,0x65,0x20,0x28,0x74,0x3c,0x70,0x72,0x65,
|
||||||
|
0x76,0x5f,0x74,0x29,0x3b,0x0a,0x62,0x75,0x66,0x5b,0x6a,0x2b,0x31,0x5d,0x3d,0x74,0x3b,0x0a,0x74,0x3d,0x74,0x32,0x3b,0x0a,0x7d,0x0a,0x70,0x72,0x65,0x76,0x5f,0x74,
|
||||||
|
0x3d,0x74,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x3d,0x74,0x69,0x64,0x3b,0x20,0x6a,0x3c,0x6c,
|
||||||
|
0x65,0x6e,0x3b,0x20,0x6a,0x2b,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x5b,0x69,
|
||||||
|
0x2b,0x6a,0x5d,0x3d,0x62,0x75,0x66,0x5b,0x6a,0x5d,0x3b,0x0a,0x7d,0x0a,0x2d,0x2d,0x69,0x6e,0x70,0x75,0x74,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,
|
||||||
|
0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x3d,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,
|
||||||
|
0x74,0x2a,0x29,0x28,0x74,0x6d,0x70,0x5f,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x29,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x74,0x69,0x64,
|
||||||
|
0x3b,0x20,0x69,0x3c,0x3d,0x4e,0x3b,0x20,0x69,0x2b,0x3d,0x42,0x57,0x54,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x6f,0x75,0x74,0x70,0x75,
|
||||||
|
0x74,0x5b,0x69,0x5d,0x3d,0x69,0x6e,0x70,0x75,0x74,0x5b,0x69,0x6e,0x64,0x69,0x63,0x65,0x73,0x5b,0x69,0x5d,0x26,0x28,0x28,0x31,0x3c,0x3c,0x32,0x31,0x29,0x2d,0x31,
|
||||||
|
0x29,0x5d,0x3b,0x0a,0x7d,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,0x66,0x69,0x6c,0x74,0x65,0x72,0x28,0x75,0x69,0x6e,0x74,0x33,
|
||||||
|
0x32,0x5f,0x74,0x20,0x6e,0x6f,0x6e,0x63,0x65,0x2c,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x62,0x77,0x74,0x5f,0x6d,0x61,0x78,0x5f,0x73,0x69,0x7a,0x65,0x2c,
|
||||||
|
0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x68,0x61,0x73,0x68,0x65,0x73,0x2c,
|
||||||
|
0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,
|
||||||
|
0x65,0x73,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x3d,0x67,
|
||||||
|
0x65,0x74,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,
|
||||||
|
0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x68,0x61,0x73,0x68,0x3d,0x68,0x61,0x73,0x68,0x65,0x73,0x2b,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x2a,
|
||||||
|
0x28,0x33,0x32,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,
|
||||||
|
0x74,0x33,0x32,0x5f,0x74,0x20,0x73,0x74,0x61,0x67,0x65,0x32,0x5f,0x73,0x69,0x7a,0x65,0x3d,0x53,0x54,0x41,0x47,0x45,0x31,0x5f,0x53,0x49,0x5a,0x45,0x2b,0x28,0x2a,
|
||||||
|
0x68,0x61,0x73,0x68,0x26,0x30,0x78,0x66,0x66,0x66,0x66,0x66,0x29,0x3b,0x0a,0x69,0x66,0x28,0x73,0x74,0x61,0x67,0x65,0x32,0x5f,0x73,0x69,0x7a,0x65,0x3c,0x62,0x77,
|
||||||
|
0x74,0x5f,0x6d,0x61,0x78,0x5f,0x73,0x69,0x7a,0x65,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x69,0x6e,0x64,0x65,0x78,0x3d,0x61,0x74,
|
||||||
|
0x6f,0x6d,0x69,0x63,0x5f,0x61,0x64,0x64,0x28,0x28,0x76,0x6f,0x6c,0x61,0x74,0x69,0x6c,0x65,0x20,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x69,0x6e,0x74,0x2a,
|
||||||
|
0x29,0x28,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x29,0x2c,0x31,0x29,0x2a,0x28,0x33,0x36,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,
|
||||||
|
0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x29,0x2b,0x31,0x3b,0x0a,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x69,
|
||||||
|
0x6e,0x64,0x65,0x78,0x5d,0x3d,0x6e,0x6f,0x6e,0x63,0x65,0x2b,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,
|
||||||
|
0x6e,0x72,0x6f,0x6c,0x6c,0x28,0x38,0x29,0x0a,0x66,0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x38,0x3b,
|
||||||
|
0x20,0x2b,0x2b,0x69,0x29,0x0a,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x69,0x6e,0x64,0x65,0x78,0x2b,0x69,0x2b,0x31,0x5d,
|
||||||
|
0x3d,0x68,0x61,0x73,0x68,0x5b,0x69,0x5d,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,0x70,0x72,0x65,0x70,
|
||||||
|
0x61,0x72,0x65,0x5f,0x62,0x61,0x74,0x63,0x68,0x32,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x68,0x61,
|
||||||
|
0x73,0x68,0x65,0x73,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,
|
||||||
|
0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x64,0x61,0x74,0x61,0x5f,
|
||||||
|
0x73,0x69,0x7a,0x65,0x73,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,
|
||||||
|
0x64,0x3d,0x67,0x65,0x74,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,
|
||||||
|
0x5f,0x74,0x20,0x4e,0x3d,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x30,0x5d,0x2d,0x67,0x65,0x74,0x5f,0x67,0x6c,0x6f,0x62,
|
||||||
|
0x61,0x6c,0x5f,0x73,0x69,0x7a,0x65,0x28,0x30,0x29,0x3b,0x0a,0x69,0x66,0x28,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x3d,0x3d,0x30,0x29,0x0a,0x66,0x69,0x6c,
|
||||||
|
0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x30,0x5d,0x3d,0x4e,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,
|
||||||
|
0x33,0x32,0x5f,0x74,0x2a,0x20,0x68,0x61,0x73,0x68,0x3d,0x68,0x61,0x73,0x68,0x65,0x73,0x2b,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x2a,0x38,0x3b,0x0a,0x5f,
|
||||||
|
0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x3d,
|
||||||
|
0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x2b,0x28,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x2b,0x4e,0x29,0x2a,0x39,0x2b,
|
||||||
|
0x31,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x73,0x74,0x61,0x67,0x65,0x32,0x5f,0x73,0x69,0x7a,0x65,0x3d,0x53,0x54,
|
||||||
|
0x41,0x47,0x45,0x31,0x5f,0x53,0x49,0x5a,0x45,0x2b,0x28,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x5b,0x31,0x5d,0x26,0x30,0x78,0x66,0x66,
|
||||||
|
0x66,0x66,0x66,0x29,0x3b,0x0a,0x64,0x61,0x74,0x61,0x5f,0x73,0x69,0x7a,0x65,0x73,0x5b,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x5d,0x3d,0x73,0x74,0x61,0x67,
|
||||||
|
0x65,0x32,0x5f,0x73,0x69,0x7a,0x65,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x28,0x38,0x29,0x0a,0x66,0x6f,0x72,0x20,0x28,
|
||||||
|
0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x38,0x3b,0x20,0x2b,0x2b,0x69,0x29,0x0a,0x68,0x61,0x73,0x68,0x5b,0x69,0x5d,0x3d,
|
||||||
|
0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x5b,0x69,0x2b,0x31,0x5d,0x3b,0x0a,0x7d,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,
|
||||||
|
0x6f,0x69,0x64,0x20,0x66,0x69,0x6e,0x64,0x5f,0x73,0x68,0x61,0x72,0x65,0x73,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,
|
||||||
|
0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x68,0x61,0x73,0x68,0x65,0x73,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,
|
||||||
|
0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x2c,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,
|
||||||
|
0x74,0x20,0x74,0x61,0x72,0x67,0x65,0x74,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x73,0x68,0x61,0x72,
|
||||||
|
0x65,0x73,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x6e,0x64,0x65,
|
||||||
|
0x78,0x3d,0x67,0x65,0x74,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x69,0x66,0x28,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x67,0x6c,
|
||||||
|
0x6f,0x62,0x61,0x6c,0x5f,0x69,0x6e,0x64,0x65,0x78,0x2a,0x34,0x2b,0x33,0x5d,0x3c,0x74,0x61,0x72,0x67,0x65,0x74,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,
|
||||||
|
0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x64,0x78,0x3d,0x61,0x74,0x6f,0x6d,0x69,0x63,0x5f,0x69,0x6e,0x63,0x28,0x73,0x68,0x61,0x72,0x65,0x73,0x2b,0x30,
|
||||||
|
0x78,0x46,0x46,0x29,0x3b,0x0a,0x69,0x66,0x28,0x69,0x64,0x78,0x3c,0x30,0x78,0x46,0x46,0x29,0x0a,0x73,0x68,0x61,0x72,0x65,0x73,0x5b,0x69,0x64,0x78,0x5d,0x3d,0x66,
|
||||||
|
0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x28,0x66,0x69,0x6c,0x74,0x65,0x72,0x65,0x64,0x5f,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,
|
||||||
|
0x30,0x5d,0x2b,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x5f,0x69,0x6e,0x64,0x65,0x78,0x29,0x2a,0x39,0x2b,0x31,0x5d,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x23,0x64,0x65,0x66,0x69,
|
||||||
|
0x6e,0x65,0x20,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x76,0x2c,0x63,0x29,0x20,0x28,0x72,0x6f,0x74,0x61,0x74,0x65,0x28,0x76,0x2c,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,
|
||||||
|
0x5f,0x74,0x29,0x63,0x29,0x29,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x58,0x4f,0x52,0x28,0x76,0x2c,0x77,0x29,0x20,0x28,0x28,0x76,0x29,0x20,0x5e,0x20,0x28,
|
||||||
|
0x77,0x29,0x29,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x50,0x4c,0x55,0x53,0x28,0x76,0x2c,0x77,0x29,0x20,0x28,0x28,0x76,0x29,0x20,0x2b,0x20,0x28,0x77,0x29,
|
||||||
|
0x29,0x0a,0x5f,0x5f,0x61,0x74,0x74,0x72,0x69,0x62,0x75,0x74,0x65,0x5f,0x5f,0x28,0x28,0x72,0x65,0x71,0x64,0x5f,0x77,0x6f,0x72,0x6b,0x5f,0x67,0x72,0x6f,0x75,0x70,
|
||||||
|
0x5f,0x73,0x69,0x7a,0x65,0x28,0x53,0x41,0x4c,0x53,0x41,0x32,0x30,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x2c,0x31,0x2c,0x31,0x29,0x29,0x29,0x0a,
|
||||||
|
0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,0x53,0x61,0x6c,0x73,0x61,0x32,0x30,0x5f,0x58,0x4f,0x52,0x4b,0x65,0x79,0x53,0x74,0x72,0x65,
|
||||||
|
0x61,0x6d,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x6b,0x65,0x79,0x73,
|
||||||
|
0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x73,0x2c,0x5f,0x5f,0x67,0x6c,
|
||||||
|
0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x73,0x2c,0x75,0x69,0x6e,0x74,
|
||||||
|
0x33,0x32,0x5f,0x74,0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x74,0x72,0x69,0x64,0x65,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,
|
||||||
|
0x33,0x32,0x5f,0x74,0x20,0x74,0x3d,0x67,0x65,0x74,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,
|
||||||
|
0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x3d,0x67,0x65,0x74,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,
|
||||||
|
0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x3d,0x67,0x2a,0x28,0x28,0x75,0x69,0x6e,0x74,0x36,
|
||||||
|
0x34,0x5f,0x74,0x29,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x74,0x72,0x69,0x64,0x65,0x29,0x2b,0x31,0x32,0x38,0x3b,0x0a,0x7b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,
|
||||||
|
0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x70,0x3d,0x6f,0x75,0x74,0x70,0x75,0x74,0x73,0x2b,0x28,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x6f,
|
||||||
|
0x66,0x66,0x73,0x65,0x74,0x2d,0x31,0x32,0x38,0x29,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x3b,0x0a,0x66,0x6f,0x72,
|
||||||
|
0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x74,0x3b,0x20,0x69,0x3c,0x31,0x32,0x38,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,
|
||||||
|
0x74,0x33,0x32,0x5f,0x74,0x29,0x3b,0x20,0x69,0x2b,0x3d,0x53,0x41,0x4c,0x53,0x41,0x32,0x30,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x29,0x0a,0x70,
|
||||||
|
0x5b,0x69,0x5d,0x3d,0x30,0x3b,0x0a,0x7d,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,
|
||||||
|
0x2a,0x20,0x6b,0x3d,0x6b,0x65,0x79,0x73,0x2b,0x67,0x2a,0x38,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,
|
||||||
|
0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x3d,0x6f,0x75,0x74,0x70,0x75,0x74,0x73,0x2b,0x28,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x2b,0x28,
|
||||||
|
0x74,0x2a,0x36,0x34,0x29,0x29,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,
|
||||||
|
0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x3d,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x73,
|
||||||
|
0x5b,0x67,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x31,0x3d,0x6b,0x5b,0x30,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,
|
||||||
|
0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x32,0x3d,0x6b,0x5b,0x31,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,
|
||||||
|
0x32,0x5f,0x74,0x20,0x6a,0x33,0x3d,0x6b,0x5b,0x32,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x34,0x3d,0x6b,
|
||||||
|
0x5b,0x33,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x31,0x31,0x3d,0x6b,0x5b,0x34,0x5d,0x3b,0x0a,0x63,0x6f,
|
||||||
|
0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x31,0x32,0x3d,0x6b,0x5b,0x35,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,
|
||||||
|
0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x31,0x33,0x3d,0x6b,0x5b,0x36,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,
|
||||||
|
0x31,0x34,0x3d,0x6b,0x5b,0x37,0x5d,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x30,0x3d,0x30,0x78,0x36,0x31,0x37,
|
||||||
|
0x30,0x37,0x38,0x36,0x35,0x55,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x35,0x3d,0x30,0x78,0x33,0x33,0x32,0x30,
|
||||||
|
0x36,0x34,0x36,0x45,0x55,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x31,0x30,0x3d,0x30,0x78,0x37,0x39,0x36,0x32,
|
||||||
|
0x32,0x44,0x33,0x32,0x55,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x31,0x35,0x3d,0x30,0x78,0x36,0x42,0x32,0x30,
|
||||||
|
0x36,0x35,0x37,0x34,0x55,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x36,0x3d,0x30,0x3b,0x0a,0x63,0x6f,0x6e,0x73,
|
||||||
|
0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x37,0x3d,0x30,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,
|
||||||
|
0x6a,0x38,0x3d,0x30,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x39,0x3d,0x30,0x3b,0x0a,0x66,0x6f,0x72,0x20,0x28,
|
||||||
|
0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x74,0x2a,0x36,0x34,0x3b,0x20,0x69,0x3c,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x3b,0x20,
|
||||||
|
0x69,0x2b,0x3d,0x53,0x41,0x4c,0x53,0x41,0x32,0x30,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,0x53,0x49,0x5a,0x45,0x2a,0x36,0x34,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,
|
||||||
|
0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x38,0x5f,0x31,0x3d,0x6a,0x38,0x2b,0x28,0x69,0x2f,0x36,0x34,0x29,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,
|
||||||
|
0x32,0x5f,0x74,0x20,0x78,0x30,0x3d,0x6a,0x30,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x31,0x3d,0x6a,0x31,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,
|
||||||
|
0x32,0x5f,0x74,0x20,0x78,0x32,0x3d,0x6a,0x32,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x33,0x3d,0x6a,0x33,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,
|
||||||
|
0x32,0x5f,0x74,0x20,0x78,0x34,0x3d,0x6a,0x34,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x35,0x3d,0x6a,0x35,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,
|
||||||
|
0x32,0x5f,0x74,0x20,0x78,0x36,0x3d,0x6a,0x36,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x37,0x3d,0x6a,0x37,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,
|
||||||
|
0x32,0x5f,0x74,0x20,0x78,0x38,0x3d,0x6a,0x38,0x5f,0x31,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x39,0x3d,0x6a,0x39,0x3b,0x0a,0x75,0x69,0x6e,
|
||||||
|
0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x31,0x30,0x3d,0x6a,0x31,0x30,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x31,0x31,0x3d,0x6a,0x31,0x31,0x3b,
|
||||||
|
0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x31,0x32,0x3d,0x6a,0x31,0x32,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x31,0x33,0x3d,
|
||||||
|
0x6a,0x31,0x33,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x78,0x31,0x34,0x3d,0x6a,0x31,0x34,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,
|
||||||
|
0x78,0x31,0x35,0x3d,0x6a,0x31,0x35,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x28,0x35,0x29,0x0a,0x66,0x6f,0x72,0x20,0x28,
|
||||||
|
0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6a,0x3d,0x30,0x3b,0x20,0x6a,0x3c,0x31,0x30,0x3b,0x20,0x2b,0x2b,0x6a,0x29,0x0a,0x7b,0x0a,0x78,0x34,0x3d,0x58,0x4f,
|
||||||
|
0x52,0x28,0x20,0x78,0x34,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x30,0x2c,0x78,0x31,0x32,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,
|
||||||
|
0x78,0x38,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x38,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x34,0x2c,0x78,0x30,0x29,0x2c,0x39,
|
||||||
|
0x29,0x29,0x3b,0x0a,0x78,0x31,0x32,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x32,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x38,0x2c,
|
||||||
|
0x78,0x34,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x30,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x30,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,
|
||||||
|
0x28,0x78,0x31,0x32,0x2c,0x78,0x38,0x29,0x2c,0x31,0x38,0x29,0x29,0x3b,0x0a,0x78,0x39,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x39,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,
|
||||||
|
0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x35,0x2c,0x78,0x31,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,0x78,0x31,0x33,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x33,0x2c,0x52,
|
||||||
|
0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x39,0x2c,0x78,0x35,0x29,0x2c,0x39,0x29,0x29,0x3b,0x0a,0x78,0x31,0x3d,0x58,0x4f,0x52,0x28,0x20,
|
||||||
|
0x78,0x31,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x33,0x2c,0x78,0x39,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x35,0x3d,
|
||||||
|
0x58,0x4f,0x52,0x28,0x20,0x78,0x35,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x31,0x2c,0x78,0x31,0x33,0x29,0x2c,0x31,0x38,0x29,
|
||||||
|
0x29,0x3b,0x0a,0x78,0x31,0x34,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x34,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x30,0x2c,0x78,
|
||||||
|
0x36,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,0x78,0x32,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x32,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,
|
||||||
|
0x31,0x34,0x2c,0x78,0x31,0x30,0x29,0x2c,0x39,0x29,0x29,0x3b,0x0a,0x78,0x36,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x36,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,
|
||||||
|
0x4c,0x55,0x53,0x28,0x20,0x78,0x32,0x2c,0x78,0x31,0x34,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x31,0x30,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x30,0x2c,0x52,
|
||||||
|
0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x36,0x2c,0x78,0x32,0x29,0x2c,0x31,0x38,0x29,0x29,0x3b,0x0a,0x78,0x33,0x3d,0x58,0x4f,0x52,0x28,
|
||||||
|
0x20,0x78,0x33,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x35,0x2c,0x78,0x31,0x31,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,0x78,0x37,
|
||||||
|
0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x37,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x33,0x2c,0x78,0x31,0x35,0x29,0x2c,0x39,0x29,
|
||||||
|
0x29,0x3b,0x0a,0x78,0x31,0x31,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x31,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x37,0x2c,0x78,
|
||||||
|
0x33,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x31,0x35,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x35,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,
|
||||||
|
0x28,0x78,0x31,0x31,0x2c,0x78,0x37,0x29,0x2c,0x31,0x38,0x29,0x29,0x3b,0x0a,0x78,0x31,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x31,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,
|
||||||
|
0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x30,0x2c,0x78,0x33,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,0x78,0x32,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x32,0x2c,0x52,0x4f,
|
||||||
|
0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x31,0x2c,0x78,0x30,0x29,0x2c,0x39,0x29,0x29,0x3b,0x0a,0x78,0x33,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,
|
||||||
|
0x33,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x32,0x2c,0x78,0x31,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x30,0x3d,0x58,
|
||||||
|
0x4f,0x52,0x28,0x20,0x78,0x30,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x33,0x2c,0x78,0x32,0x29,0x2c,0x31,0x38,0x29,0x29,0x3b,
|
||||||
|
0x0a,0x78,0x36,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x36,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x35,0x2c,0x78,0x34,0x29,0x2c,
|
||||||
|
0x37,0x29,0x29,0x3b,0x0a,0x78,0x37,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x37,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x36,0x2c,
|
||||||
|
0x78,0x35,0x29,0x2c,0x39,0x29,0x29,0x3b,0x0a,0x78,0x34,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x34,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,
|
||||||
|
0x20,0x78,0x37,0x2c,0x78,0x36,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x35,0x3d,0x58,0x4f,0x52,0x28,0x20,0x78,0x35,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,
|
||||||
|
0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x34,0x2c,0x78,0x37,0x29,0x2c,0x31,0x38,0x29,0x29,0x3b,0x0a,0x78,0x31,0x31,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x31,0x2c,0x52,
|
||||||
|
0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x30,0x2c,0x78,0x39,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,0x78,0x38,0x3d,0x58,0x4f,0x52,0x28,0x20,
|
||||||
|
0x78,0x38,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x31,0x2c,0x78,0x31,0x30,0x29,0x2c,0x39,0x29,0x29,0x3b,0x0a,0x78,0x39,0x3d,
|
||||||
|
0x58,0x4f,0x52,0x28,0x20,0x78,0x39,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x38,0x2c,0x78,0x31,0x31,0x29,0x2c,0x31,0x33,0x29,
|
||||||
|
0x29,0x3b,0x0a,0x78,0x31,0x30,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x30,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x20,0x78,0x39,0x2c,0x78,
|
||||||
|
0x38,0x29,0x2c,0x31,0x38,0x29,0x29,0x3b,0x0a,0x78,0x31,0x32,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x32,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,
|
||||||
|
0x28,0x78,0x31,0x35,0x2c,0x78,0x31,0x34,0x29,0x2c,0x37,0x29,0x29,0x3b,0x0a,0x78,0x31,0x33,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x33,0x2c,0x52,0x4f,0x54,0x41,0x54,
|
||||||
|
0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x32,0x2c,0x78,0x31,0x35,0x29,0x2c,0x39,0x29,0x29,0x3b,0x0a,0x78,0x31,0x34,0x3d,0x58,0x4f,0x52,0x28,0x78,0x31,0x34,
|
||||||
|
0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x33,0x2c,0x78,0x31,0x32,0x29,0x2c,0x31,0x33,0x29,0x29,0x3b,0x0a,0x78,0x31,0x35,0x3d,
|
||||||
|
0x58,0x4f,0x52,0x28,0x78,0x31,0x35,0x2c,0x52,0x4f,0x54,0x41,0x54,0x45,0x28,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x34,0x2c,0x78,0x31,0x33,0x29,0x2c,0x31,0x38,0x29,
|
||||||
|
0x29,0x3b,0x0a,0x7d,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x30,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x30,0x2c,0x6a,0x30,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,
|
||||||
|
0x75,0x74,0x5b,0x31,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x2c,0x6a,0x31,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x32,0x5d,0x3d,0x50,0x4c,0x55,
|
||||||
|
0x53,0x28,0x78,0x32,0x2c,0x6a,0x32,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x33,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x33,0x2c,0x6a,0x33,0x29,0x3b,
|
||||||
|
0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x34,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x34,0x2c,0x6a,0x34,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x35,
|
||||||
|
0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x35,0x2c,0x6a,0x35,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x36,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x36,
|
||||||
|
0x2c,0x6a,0x36,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x37,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x37,0x2c,0x6a,0x37,0x29,0x3b,0x0a,0x6f,0x75,0x74,
|
||||||
|
0x70,0x75,0x74,0x5b,0x38,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x38,0x2c,0x6a,0x38,0x5f,0x31,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x39,0x5d,0x3d,
|
||||||
|
0x50,0x4c,0x55,0x53,0x28,0x78,0x39,0x2c,0x6a,0x39,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x31,0x30,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x30,
|
||||||
|
0x2c,0x6a,0x31,0x30,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x31,0x31,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x31,0x2c,0x6a,0x31,0x31,0x29,0x3b,
|
||||||
|
0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x31,0x32,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x32,0x2c,0x6a,0x31,0x32,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,
|
||||||
|
0x74,0x5b,0x31,0x33,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x33,0x2c,0x6a,0x31,0x33,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x31,0x34,0x5d,0x3d,
|
||||||
|
0x50,0x4c,0x55,0x53,0x28,0x78,0x31,0x34,0x2c,0x6a,0x31,0x34,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x5b,0x31,0x35,0x5d,0x3d,0x50,0x4c,0x55,0x53,0x28,0x78,
|
||||||
|
0x31,0x35,0x2c,0x6a,0x31,0x35,0x29,0x3b,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,0x2b,0x3d,0x28,0x53,0x41,0x4c,0x53,0x41,0x32,0x30,0x5f,0x47,0x52,0x4f,0x55,0x50,0x5f,
|
||||||
|
0x53,0x49,0x5a,0x45,0x2a,0x36,0x34,0x29,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x3b,0x0a,0x7d,0x0a,0x62,0x61,0x72,
|
||||||
|
0x72,0x69,0x65,0x72,0x28,0x43,0x4c,0x4b,0x5f,0x47,0x4c,0x4f,0x42,0x41,0x4c,0x5f,0x4d,0x45,0x4d,0x5f,0x46,0x45,0x4e,0x43,0x45,0x29,0x3b,0x0a,0x69,0x66,0x28,0x74,
|
||||||
|
0x3c,0x31,0x36,0x29,0x0a,0x7b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x70,0x3d,0x6f,0x75,0x74,0x70,
|
||||||
|
0x75,0x74,0x73,0x2b,0x28,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x2b,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x2b,0x33,
|
||||||
|
0x29,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x3b,0x0a,0x70,0x5b,0x74,0x5d,0x3d,0x30,0x3b,0x0a,0x7d,0x0a,0x69,0x66,
|
||||||
|
0x28,0x28,0x74,0x3d,0x3d,0x30,0x29,0x26,0x26,0x28,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x26,0x33,0x29,0x29,0x0a,0x6f,0x75,0x74,0x70,0x75,0x74,
|
||||||
|
0x73,0x5b,0x28,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,0x2b,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x29,0x2f,0x73,0x69,
|
||||||
|
0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x29,0x5d,0x20,0x26,0x3d,0x20,0x30,0x78,0x46,0x46,0x46,0x46,0x46,0x46,0x46,0x46,0x55,0x3e,0x3e,
|
||||||
|
0x28,0x28,0x34,0x2d,0x28,0x6f,0x75,0x74,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x26,0x33,0x29,0x29,0x3c,0x3c,0x33,0x29,0x3b,0x0a,0x7d,0x0a,0x23,0x64,0x65,0x66,
|
||||||
|
0x69,0x6e,0x65,0x20,0x52,0x4f,0x55,0x4e,0x44,0x53,0x20,0x32,0x34,0x20,0x0a,0x23,0x64,0x65,0x66,0x69,0x6e,0x65,0x20,0x52,0x36,0x34,0x28,0x61,0x2c,0x62,0x2c,0x63,
|
||||||
|
0x29,0x20,0x28,0x28,0x28,0x61,0x29,0x20,0x3c,0x3c,0x20,0x62,0x29,0x20,0x7c,0x20,0x28,0x28,0x61,0x29,0x20,0x3e,0x3e,0x20,0x63,0x29,0x29,0x0a,0x5f,0x5f,0x63,0x6f,
|
||||||
|
0x6e,0x73,0x74,0x61,0x6e,0x74,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x72,0x63,0x5b,0x32,0x5d,0x5b,0x52,0x4f,0x55,0x4e,
|
||||||
|
0x44,0x53,0x5d,0x3d,0x7b,0x0a,0x7b,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x31,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,
|
||||||
|
0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x38,0x32,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,
|
||||||
|
0x30,0x38,0x41,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,
|
||||||
|
0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x38,0x42,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x30,
|
||||||
|
0x30,0x30,0x31,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x38,0x30,0x38,0x31,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,
|
||||||
|
0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x39,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,
|
||||||
|
0x30,0x38,0x41,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x38,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,
|
||||||
|
0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x38,0x30,0x30,0x39,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x30,
|
||||||
|
0x30,0x30,0x41,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x38,0x30,0x38,0x42,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,
|
||||||
|
0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x42,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,
|
||||||
|
0x30,0x38,0x39,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x33,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,
|
||||||
|
0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x32,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,
|
||||||
|
0x30,0x38,0x30,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x41,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,
|
||||||
|
0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x41,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x38,
|
||||||
|
0x30,0x38,0x31,0x55,0x4c,0x2c,0x0a,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x38,0x30,0x55,0x4c,0x2c,0x30,0x78,0x30,0x30,
|
||||||
|
0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x31,0x55,0x4c,0x2c,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x38,0x30,0x30,0x30,0x38,
|
||||||
|
0x30,0x30,0x38,0x55,0x4c,0x7d,0x2c,0x0a,0x7b,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,
|
||||||
|
0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x0a,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,
|
||||||
|
0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x0a,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,
|
||||||
|
0x55,0x4c,0x2c,0x30,0x55,0x4c,0x2c,0x30,0x55,0x4c,0x7d,0x0a,0x7d,0x3b,0x0a,0x5f,0x5f,0x63,0x6f,0x6e,0x73,0x74,0x61,0x6e,0x74,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,
|
||||||
|
0x69,0x6e,0x74,0x20,0x72,0x6f,0x5b,0x32,0x35,0x5d,0x5b,0x32,0x5d,0x3d,0x7b,0x0a,0x7b,0x20,0x30,0x2c,0x36,0x34,0x7d,0x2c,0x7b,0x34,0x34,0x2c,0x32,0x30,0x7d,0x2c,
|
||||||
|
0x7b,0x34,0x33,0x2c,0x32,0x31,0x7d,0x2c,0x7b,0x32,0x31,0x2c,0x34,0x33,0x7d,0x2c,0x7b,0x31,0x34,0x2c,0x35,0x30,0x7d,0x2c,0x0a,0x7b,0x20,0x31,0x2c,0x36,0x33,0x7d,
|
||||||
|
0x2c,0x7b,0x20,0x36,0x2c,0x35,0x38,0x7d,0x2c,0x7b,0x32,0x35,0x2c,0x33,0x39,0x7d,0x2c,0x7b,0x20,0x38,0x2c,0x35,0x36,0x7d,0x2c,0x7b,0x31,0x38,0x2c,0x34,0x36,0x7d,
|
||||||
|
0x2c,0x0a,0x7b,0x36,0x32,0x2c,0x32,0x7d,0x2c,0x7b,0x35,0x35,0x2c,0x39,0x7d,0x2c,0x7b,0x33,0x39,0x2c,0x32,0x35,0x7d,0x2c,0x7b,0x34,0x31,0x2c,0x32,0x33,0x7d,0x2c,
|
||||||
|
0x7b,0x20,0x32,0x2c,0x36,0x32,0x7d,0x2c,0x0a,0x7b,0x32,0x38,0x2c,0x33,0x36,0x7d,0x2c,0x7b,0x32,0x30,0x2c,0x34,0x34,0x7d,0x2c,0x7b,0x20,0x33,0x2c,0x36,0x31,0x7d,
|
||||||
|
0x2c,0x7b,0x34,0x35,0x2c,0x31,0x39,0x7d,0x2c,0x7b,0x36,0x31,0x2c,0x33,0x7d,0x2c,0x0a,0x7b,0x32,0x37,0x2c,0x33,0x37,0x7d,0x2c,0x7b,0x33,0x36,0x2c,0x32,0x38,0x7d,
|
||||||
|
0x2c,0x7b,0x31,0x30,0x2c,0x35,0x34,0x7d,0x2c,0x7b,0x31,0x35,0x2c,0x34,0x39,0x7d,0x2c,0x7b,0x35,0x36,0x2c,0x38,0x7d,0x0a,0x7d,0x3b,0x0a,0x5f,0x5f,0x63,0x6f,0x6e,
|
||||||
|
0x73,0x74,0x61,0x6e,0x74,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x61,0x5b,0x32,0x35,0x5d,0x3d,0x7b,0x0a,0x30,0x2c,0x36,0x2c,0x31,0x32,0x2c,0x31,
|
||||||
|
0x38,0x2c,0x32,0x34,0x2c,0x0a,0x31,0x2c,0x37,0x2c,0x31,0x33,0x2c,0x31,0x39,0x2c,0x32,0x30,0x2c,0x0a,0x32,0x2c,0x38,0x2c,0x31,0x34,0x2c,0x31,0x35,0x2c,0x32,0x31,
|
||||||
|
0x2c,0x0a,0x33,0x2c,0x39,0x2c,0x31,0x30,0x2c,0x31,0x36,0x2c,0x32,0x32,0x2c,0x0a,0x34,0x2c,0x35,0x2c,0x31,0x31,0x2c,0x31,0x37,0x2c,0x32,0x33,0x0a,0x7d,0x3b,0x0a,
|
||||||
|
0x5f,0x5f,0x63,0x6f,0x6e,0x73,0x74,0x61,0x6e,0x74,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x62,0x5b,0x32,0x35,0x5d,0x3d,0x7b,0x0a,0x30,0x2c,0x31,
|
||||||
|
0x2c,0x32,0x2c,0x33,0x2c,0x34,0x2c,0x0a,0x31,0x2c,0x32,0x2c,0x33,0x2c,0x34,0x2c,0x30,0x2c,0x0a,0x32,0x2c,0x33,0x2c,0x34,0x2c,0x30,0x2c,0x31,0x2c,0x0a,0x33,0x2c,
|
||||||
|
0x34,0x2c,0x30,0x2c,0x31,0x2c,0x32,0x2c,0x0a,0x34,0x2c,0x30,0x2c,0x31,0x2c,0x32,0x2c,0x33,0x0a,0x7d,0x3b,0x0a,0x5f,0x5f,0x63,0x6f,0x6e,0x73,0x74,0x61,0x6e,0x74,
|
||||||
|
0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x63,0x5b,0x32,0x35,0x5d,0x5b,0x33,0x5d,0x3d,0x7b,0x0a,0x7b,0x20,0x30,0x2c,0x31,0x2c,0x32,0x7d,0x2c,0x7b,
|
||||||
|
0x20,0x31,0x2c,0x32,0x2c,0x33,0x7d,0x2c,0x7b,0x20,0x32,0x2c,0x33,0x2c,0x34,0x7d,0x2c,0x7b,0x20,0x33,0x2c,0x34,0x2c,0x30,0x7d,0x2c,0x7b,0x20,0x34,0x2c,0x30,0x2c,
|
||||||
|
0x31,0x7d,0x2c,0x0a,0x7b,0x20,0x35,0x2c,0x36,0x2c,0x37,0x7d,0x2c,0x7b,0x20,0x36,0x2c,0x37,0x2c,0x38,0x7d,0x2c,0x7b,0x20,0x37,0x2c,0x38,0x2c,0x39,0x7d,0x2c,0x7b,
|
||||||
|
0x20,0x38,0x2c,0x39,0x2c,0x35,0x7d,0x2c,0x7b,0x20,0x39,0x2c,0x35,0x2c,0x36,0x7d,0x2c,0x0a,0x7b,0x31,0x30,0x2c,0x31,0x31,0x2c,0x31,0x32,0x7d,0x2c,0x7b,0x31,0x31,
|
||||||
|
0x2c,0x31,0x32,0x2c,0x31,0x33,0x7d,0x2c,0x7b,0x31,0x32,0x2c,0x31,0x33,0x2c,0x31,0x34,0x7d,0x2c,0x7b,0x31,0x33,0x2c,0x31,0x34,0x2c,0x31,0x30,0x7d,0x2c,0x7b,0x31,
|
||||||
|
0x34,0x2c,0x31,0x30,0x2c,0x31,0x31,0x7d,0x2c,0x0a,0x7b,0x31,0x35,0x2c,0x31,0x36,0x2c,0x31,0x37,0x7d,0x2c,0x7b,0x31,0x36,0x2c,0x31,0x37,0x2c,0x31,0x38,0x7d,0x2c,
|
||||||
|
0x7b,0x31,0x37,0x2c,0x31,0x38,0x2c,0x31,0x39,0x7d,0x2c,0x7b,0x31,0x38,0x2c,0x31,0x39,0x2c,0x31,0x35,0x7d,0x2c,0x7b,0x31,0x39,0x2c,0x31,0x35,0x2c,0x31,0x36,0x7d,
|
||||||
|
0x2c,0x0a,0x7b,0x32,0x30,0x2c,0x32,0x31,0x2c,0x32,0x32,0x7d,0x2c,0x7b,0x32,0x31,0x2c,0x32,0x32,0x2c,0x32,0x33,0x7d,0x2c,0x7b,0x32,0x32,0x2c,0x32,0x33,0x2c,0x32,
|
||||||
|
0x34,0x7d,0x2c,0x7b,0x32,0x33,0x2c,0x32,0x34,0x2c,0x32,0x30,0x7d,0x2c,0x7b,0x32,0x34,0x2c,0x32,0x30,0x2c,0x32,0x31,0x7d,0x0a,0x7d,0x3b,0x0a,0x5f,0x5f,0x63,0x6f,
|
||||||
|
0x6e,0x73,0x74,0x61,0x6e,0x74,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x69,0x6e,0x74,0x20,0x64,0x5b,0x32,0x35,0x5d,0x3d,0x7b,0x0a,0x30,0x2c,0x31,0x2c,0x32,0x2c,0x33,
|
||||||
|
0x2c,0x34,0x2c,0x0a,0x31,0x30,0x2c,0x31,0x31,0x2c,0x31,0x32,0x2c,0x31,0x33,0x2c,0x31,0x34,0x2c,0x0a,0x32,0x30,0x2c,0x32,0x31,0x2c,0x32,0x32,0x2c,0x32,0x33,0x2c,
|
||||||
|
0x32,0x34,0x2c,0x0a,0x35,0x2c,0x36,0x2c,0x37,0x2c,0x38,0x2c,0x39,0x2c,0x0a,0x31,0x35,0x2c,0x31,0x36,0x2c,0x31,0x37,0x2c,0x31,0x38,0x2c,0x31,0x39,0x0a,0x7d,0x3b,
|
||||||
|
0x0a,0x5f,0x5f,0x61,0x74,0x74,0x72,0x69,0x62,0x75,0x74,0x65,0x5f,0x5f,0x28,0x28,0x72,0x65,0x71,0x64,0x5f,0x77,0x6f,0x72,0x6b,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,
|
||||||
|
0x73,0x69,0x7a,0x65,0x28,0x33,0x32,0x2c,0x31,0x2c,0x31,0x29,0x29,0x29,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,0x73,0x68,0x61,
|
||||||
|
0x33,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x69,0x6e,0x70,0x75,0x74,0x73,
|
||||||
|
0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x69,0x6e,0x70,0x75,0x74,0x5f,
|
||||||
|
0x73,0x69,0x7a,0x65,0x73,0x2c,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x74,0x72,0x69,0x64,0x65,0x2c,0x5f,0x5f,0x67,0x6c,
|
||||||
|
0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x68,0x61,0x73,0x68,0x65,0x73,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,
|
||||||
|
0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x74,0x3d,0x67,0x65,0x74,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,
|
||||||
|
0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x3d,0x67,0x65,0x74,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x69,0x66,0x28,
|
||||||
|
0x74,0x3e,0x3d,0x32,0x35,0x29,0x0a,0x72,0x65,0x74,0x75,0x72,0x6e,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x73,0x3d,
|
||||||
|
0x74,0x20,0x25,0x20,0x35,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x69,0x6e,0x70,0x75,0x74,0x5f,0x6f,0x66,0x66,0x73,
|
||||||
|
0x65,0x74,0x3d,0x28,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x74,0x72,0x69,0x64,0x65,0x29,0x2a,0x67,0x3b,0x0a,0x5f,
|
||||||
|
0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x69,0x6e,0x70,0x75,0x74,0x3d,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,
|
||||||
|
0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x29,0x28,0x69,0x6e,0x70,0x75,0x74,0x73,0x2b,0x69,0x6e,0x70,0x75,0x74,0x5f,0x6f,0x66,0x66,0x73,0x65,0x74,
|
||||||
|
0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x3d,0x69,0x6e,0x70,
|
||||||
|
0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x73,0x5b,0x67,0x5d,0x2b,0x31,0x3b,0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,
|
||||||
|
0x41,0x5b,0x32,0x35,0x5d,0x3b,0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x43,0x5b,0x32,0x35,0x5d,0x3b,0x0a,0x5f,
|
||||||
|
0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x44,0x5b,0x32,0x35,0x5d,0x3b,0x0a,0x41,0x5b,0x74,0x5d,0x3d,0x30,0x3b,0x0a,0x63,
|
||||||
|
0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x77,0x6f,0x72,0x64,0x73,0x3d,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x2f,0x73,
|
||||||
|
0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,
|
||||||
|
0x74,0x61,0x69,0x6c,0x5f,0x73,0x69,0x7a,0x65,0x3d,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x20,0x25,0x20,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,
|
||||||
|
0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,0x3d,0x30,0x3b,0x0a,0x66,
|
||||||
|
0x6f,0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x77,0x6f,0x72,0x64,0x73,0x3b,0x20,0x2b,0x2b,0x69,0x2c,0x2b,
|
||||||
|
0x2b,0x69,0x6e,0x70,0x75,0x74,0x29,0x0a,0x7b,0x0a,0x41,0x5b,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,0x5d,0x20,0x5e,0x3d,0x20,0x2a,0x69,0x6e,0x70,0x75,0x74,
|
||||||
|
0x3b,0x0a,0x2b,0x2b,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,0x3b,0x0a,0x69,0x66,0x28,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,0x3d,0x3d,0x31,0x37,0x29,
|
||||||
|
0x0a,0x7b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x28,0x52,0x4f,0x55,0x4e,0x44,0x53,0x29,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,
|
||||||
|
0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x52,0x4f,0x55,0x4e,0x44,0x53,0x3b,0x20,0x2b,0x2b,0x69,0x29,0x0a,0x7b,0x0a,0x43,0x5b,0x74,0x5d,0x3d,0x41,0x5b,
|
||||||
|
0x73,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x35,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x31,0x30,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x31,0x35,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x32,0x30,
|
||||||
|
0x5d,0x3b,0x0a,0x44,0x5b,0x74,0x5d,0x3d,0x43,0x5b,0x62,0x5b,0x32,0x30,0x2b,0x73,0x5d,0x5d,0x5e,0x52,0x36,0x34,0x28,0x43,0x5b,0x62,0x5b,0x35,0x2b,0x73,0x5d,0x5d,
|
||||||
|
0x2c,0x31,0x2c,0x36,0x33,0x29,0x3b,0x0a,0x43,0x5b,0x74,0x5d,0x3d,0x52,0x36,0x34,0x28,0x41,0x5b,0x61,0x5b,0x74,0x5d,0x5d,0x5e,0x44,0x5b,0x62,0x5b,0x74,0x5d,0x5d,
|
||||||
|
0x2c,0x72,0x6f,0x5b,0x74,0x5d,0x5b,0x30,0x5d,0x2c,0x72,0x6f,0x5b,0x74,0x5d,0x5b,0x31,0x5d,0x29,0x3b,0x0a,0x41,0x5b,0x64,0x5b,0x74,0x5d,0x5d,0x3d,0x43,0x5b,0x63,
|
||||||
|
0x5b,0x74,0x5d,0x5b,0x30,0x5d,0x5d,0x5e,0x28,0x28,0x7e,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x31,0x5d,0x5d,0x29,0x26,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x32,0x5d,
|
||||||
|
0x5d,0x29,0x3b,0x0a,0x41,0x5b,0x74,0x5d,0x20,0x5e,0x3d,0x20,0x72,0x63,0x5b,0x28,0x74,0x3d,0x3d,0x30,0x29,0x3f,0x30,0x3a,0x31,0x5d,0x5b,0x69,0x5d,0x3b,0x20,0x0a,
|
||||||
|
0x7d,0x0a,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,0x3d,0x30,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x74,0x61,0x69,0x6c,
|
||||||
|
0x3d,0x30,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x70,0x3d,0x28,0x5f,
|
||||||
|
0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x29,0x69,0x6e,0x70,0x75,0x74,0x3b,0x0a,0x66,0x6f,
|
||||||
|
0x72,0x20,0x28,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x74,0x61,0x69,0x6c,0x5f,0x73,0x69,0x7a,0x65,0x3b,0x20,0x2b,0x2b,
|
||||||
|
0x69,0x29,0x0a,0x7b,0x0a,0x74,0x61,0x69,0x6c,0x7c,0x3d,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x70,0x5b,0x69,0x5d,0x29,0x3c,0x3c,0x28,0x69,0x2a,
|
||||||
|
0x38,0x29,0x3b,0x0a,0x7d,0x0a,0x41,0x5b,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,0x5d,0x20,0x5e,0x3d,0x20,0x74,0x61,0x69,0x6c,0x5e,0x28,0x28,0x75,0x69,0x6e,
|
||||||
|
0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x28,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x30,0x78,0x30,0x32,0x7c,0x28,0x31,0x3c,0x3c,0x32,0x29,0x29,0x29,
|
||||||
|
0x3c,0x3c,0x28,0x74,0x61,0x69,0x6c,0x5f,0x73,0x69,0x7a,0x65,0x2a,0x38,0x29,0x29,0x29,0x3b,0x0a,0x41,0x5b,0x31,0x36,0x5d,0x20,0x5e,0x3d,0x20,0x30,0x78,0x38,0x30,
|
||||||
|
0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x55,0x4c,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,
|
||||||
|
0x28,0x31,0x29,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,0x69,0x3c,0x52,0x4f,0x55,0x4e,0x44,0x53,0x3b,0x20,0x2b,0x2b,0x69,0x29,
|
||||||
|
0x0a,0x7b,0x0a,0x43,0x5b,0x74,0x5d,0x3d,0x41,0x5b,0x73,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x35,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x31,0x30,0x5d,0x5e,0x41,0x5b,0x73,0x2b,
|
||||||
|
0x31,0x35,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x32,0x30,0x5d,0x3b,0x0a,0x44,0x5b,0x74,0x5d,0x3d,0x43,0x5b,0x62,0x5b,0x32,0x30,0x2b,0x73,0x5d,0x5d,0x5e,0x52,0x36,0x34,
|
||||||
|
0x28,0x43,0x5b,0x62,0x5b,0x35,0x2b,0x73,0x5d,0x5d,0x2c,0x31,0x2c,0x36,0x33,0x29,0x3b,0x0a,0x43,0x5b,0x74,0x5d,0x3d,0x52,0x36,0x34,0x28,0x41,0x5b,0x61,0x5b,0x74,
|
||||||
|
0x5d,0x5d,0x5e,0x44,0x5b,0x62,0x5b,0x74,0x5d,0x5d,0x2c,0x72,0x6f,0x5b,0x74,0x5d,0x5b,0x30,0x5d,0x2c,0x72,0x6f,0x5b,0x74,0x5d,0x5b,0x31,0x5d,0x29,0x3b,0x0a,0x41,
|
||||||
|
0x5b,0x64,0x5b,0x74,0x5d,0x5d,0x3d,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x30,0x5d,0x5d,0x5e,0x28,0x28,0x7e,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x31,0x5d,0x5d,0x29,
|
||||||
|
0x26,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x32,0x5d,0x5d,0x29,0x3b,0x0a,0x41,0x5b,0x74,0x5d,0x20,0x5e,0x3d,0x20,0x72,0x63,0x5b,0x28,0x74,0x3d,0x3d,0x30,0x29,0x3f,
|
||||||
|
0x30,0x3a,0x31,0x5d,0x5b,0x69,0x5d,0x3b,0x20,0x0a,0x7d,0x0a,0x69,0x66,0x28,0x74,0x3c,0x34,0x29,0x0a,0x7b,0x0a,0x68,0x61,0x73,0x68,0x65,0x73,0x2b,0x3d,0x67,0x2a,
|
||||||
|
0x28,0x33,0x32,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x29,0x3b,0x0a,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x74,0x5d,
|
||||||
|
0x3d,0x41,0x5b,0x74,0x5d,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x5f,0x5f,0x61,0x74,0x74,0x72,0x69,0x62,0x75,0x74,0x65,0x5f,0x5f,0x28,0x28,0x72,0x65,0x71,0x64,0x5f,0x77,
|
||||||
|
0x6f,0x72,0x6b,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,0x73,0x69,0x7a,0x65,0x28,0x33,0x32,0x2c,0x31,0x2c,0x31,0x29,0x29,0x29,0x0a,0x5f,0x5f,0x6b,0x65,0x72,0x6e,0x65,
|
||||||
|
0x6c,0x20,0x76,0x6f,0x69,0x64,0x20,0x73,0x68,0x61,0x33,0x5f,0x69,0x6e,0x69,0x74,0x69,0x61,0x6c,0x28,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x63,0x6f,0x6e,
|
||||||
|
0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x38,0x5f,0x74,0x2a,0x20,0x69,0x6e,0x70,0x75,0x74,0x5f,0x64,0x61,0x74,0x61,0x2c,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,
|
||||||
|
0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x2c,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x6e,0x6f,0x6e,0x63,0x65,0x2c,0x5f,0x5f,0x67,0x6c,0x6f,0x62,
|
||||||
|
0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x68,0x61,0x73,0x68,0x65,0x73,0x29,0x0a,0x7b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,
|
||||||
|
0x74,0x33,0x32,0x5f,0x74,0x20,0x74,0x3d,0x67,0x65,0x74,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,
|
||||||
|
0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x67,0x3d,0x67,0x65,0x74,0x5f,0x67,0x72,0x6f,0x75,0x70,0x5f,0x69,0x64,0x28,0x30,0x29,0x3b,0x0a,0x69,0x66,0x28,0x74,0x3e,
|
||||||
|
0x3d,0x32,0x35,0x29,0x0a,0x72,0x65,0x74,0x75,0x72,0x6e,0x3b,0x0a,0x63,0x6f,0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x73,0x3d,0x74,0x20,
|
||||||
|
0x25,0x20,0x35,0x3b,0x0a,0x5f,0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x20,0x69,0x6e,0x70,0x75,0x74,0x3d,0x28,0x5f,
|
||||||
|
0x5f,0x67,0x6c,0x6f,0x62,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x2a,0x29,0x28,0x69,0x6e,0x70,0x75,0x74,0x5f,0x64,0x61,0x74,0x61,0x29,0x3b,0x0a,
|
||||||
|
0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x41,0x5b,0x32,0x35,0x5d,0x3b,0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,
|
||||||
|
0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x20,0x43,0x5b,0x32,0x35,0x5d,0x3b,0x0a,0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,
|
||||||
|
0x20,0x44,0x5b,0x32,0x35,0x5d,0x3b,0x0a,0x41,0x5b,0x74,0x5d,0x3d,0x28,0x74,0x3c,0x31,0x36,0x29,0x3f,0x69,0x6e,0x70,0x75,0x74,0x5b,0x74,0x5d,0x3a,0x30,0x3b,0x0a,
|
||||||
|
0x5f,0x5f,0x6c,0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x20,0x6e,0x6f,0x6e,0x63,0x65,0x5f,0x70,0x6f,0x73,0x3d,0x28,0x5f,0x5f,0x6c,
|
||||||
|
0x6f,0x63,0x61,0x6c,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x2a,0x29,0x28,0x41,0x29,0x2b,0x39,0x3b,0x0a,0x6e,0x6f,0x6e,0x63,0x65,0x2b,0x3d,0x67,0x3b,0x0a,
|
||||||
|
0x6e,0x6f,0x6e,0x63,0x65,0x5f,0x70,0x6f,0x73,0x5b,0x30,0x5d,0x3d,0x28,0x6e,0x6f,0x6e,0x63,0x65,0x5f,0x70,0x6f,0x73,0x5b,0x30,0x5d,0x26,0x30,0x78,0x46,0x46,0x46,
|
||||||
|
0x46,0x46,0x46,0x55,0x29,0x7c,0x28,0x28,0x6e,0x6f,0x6e,0x63,0x65,0x26,0x30,0x78,0x46,0x46,0x29,0x3c,0x3c,0x32,0x34,0x29,0x3b,0x0a,0x6e,0x6f,0x6e,0x63,0x65,0x5f,
|
||||||
|
0x70,0x6f,0x73,0x5b,0x31,0x5d,0x3d,0x28,0x6e,0x6f,0x6e,0x63,0x65,0x5f,0x70,0x6f,0x73,0x5b,0x31,0x5d,0x26,0x30,0x78,0x46,0x46,0x30,0x30,0x30,0x30,0x30,0x30,0x55,
|
||||||
|
0x29,0x7c,0x28,0x6e,0x6f,0x6e,0x63,0x65,0x3e,0x3e,0x38,0x29,0x3b,0x0a,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,0x78,
|
||||||
|
0x3d,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,0x65,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x3b,0x0a,0x63,0x6f,
|
||||||
|
0x6e,0x73,0x74,0x20,0x75,0x69,0x6e,0x74,0x33,0x32,0x5f,0x74,0x20,0x74,0x61,0x69,0x6c,0x5f,0x73,0x69,0x7a,0x65,0x3d,0x69,0x6e,0x70,0x75,0x74,0x5f,0x73,0x69,0x7a,
|
||||||
|
0x65,0x20,0x25,0x20,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x3b,0x0a,0x41,0x5b,0x77,0x6f,0x72,0x64,0x49,0x6e,0x64,0x65,
|
||||||
|
0x78,0x5d,0x20,0x5e,0x3d,0x20,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x28,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,0x74,0x29,0x28,0x30,0x78,0x30,
|
||||||
|
0x32,0x7c,0x28,0x31,0x3c,0x3c,0x32,0x29,0x29,0x29,0x3c,0x3c,0x28,0x74,0x61,0x69,0x6c,0x5f,0x73,0x69,0x7a,0x65,0x2a,0x38,0x29,0x29,0x3b,0x0a,0x41,0x5b,0x31,0x36,
|
||||||
|
0x5d,0x20,0x5e,0x3d,0x20,0x30,0x78,0x38,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x30,0x55,0x4c,0x3b,0x0a,0x23,0x70,0x72,0x61,0x67,
|
||||||
|
0x6d,0x61,0x20,0x75,0x6e,0x72,0x6f,0x6c,0x6c,0x28,0x52,0x4f,0x55,0x4e,0x44,0x53,0x29,0x0a,0x66,0x6f,0x72,0x20,0x28,0x69,0x6e,0x74,0x20,0x69,0x3d,0x30,0x3b,0x20,
|
||||||
|
0x69,0x3c,0x52,0x4f,0x55,0x4e,0x44,0x53,0x3b,0x20,0x2b,0x2b,0x69,0x29,0x0a,0x7b,0x0a,0x43,0x5b,0x74,0x5d,0x3d,0x41,0x5b,0x73,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x35,
|
||||||
|
0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x31,0x30,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x31,0x35,0x5d,0x5e,0x41,0x5b,0x73,0x2b,0x32,0x30,0x5d,0x3b,0x0a,0x44,0x5b,0x74,0x5d,0x3d,
|
||||||
|
0x43,0x5b,0x62,0x5b,0x32,0x30,0x2b,0x73,0x5d,0x5d,0x5e,0x52,0x36,0x34,0x28,0x43,0x5b,0x62,0x5b,0x35,0x2b,0x73,0x5d,0x5d,0x2c,0x31,0x2c,0x36,0x33,0x29,0x3b,0x0a,
|
||||||
|
0x43,0x5b,0x74,0x5d,0x3d,0x52,0x36,0x34,0x28,0x41,0x5b,0x61,0x5b,0x74,0x5d,0x5d,0x5e,0x44,0x5b,0x62,0x5b,0x74,0x5d,0x5d,0x2c,0x72,0x6f,0x5b,0x74,0x5d,0x5b,0x30,
|
||||||
|
0x5d,0x2c,0x72,0x6f,0x5b,0x74,0x5d,0x5b,0x31,0x5d,0x29,0x3b,0x0a,0x41,0x5b,0x64,0x5b,0x74,0x5d,0x5d,0x3d,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x30,0x5d,0x5d,0x5e,
|
||||||
|
0x28,0x28,0x7e,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x31,0x5d,0x5d,0x29,0x26,0x43,0x5b,0x63,0x5b,0x74,0x5d,0x5b,0x32,0x5d,0x5d,0x29,0x3b,0x0a,0x41,0x5b,0x74,0x5d,
|
||||||
|
0x20,0x5e,0x3d,0x20,0x72,0x63,0x5b,0x28,0x74,0x3d,0x3d,0x30,0x29,0x3f,0x30,0x3a,0x31,0x5d,0x5b,0x69,0x5d,0x3b,0x20,0x0a,0x7d,0x0a,0x69,0x66,0x28,0x74,0x3c,0x34,
|
||||||
|
0x29,0x0a,0x7b,0x0a,0x68,0x61,0x73,0x68,0x65,0x73,0x2b,0x3d,0x67,0x2a,0x28,0x33,0x32,0x2f,0x73,0x69,0x7a,0x65,0x6f,0x66,0x28,0x75,0x69,0x6e,0x74,0x36,0x34,0x5f,
|
||||||
|
0x74,0x29,0x29,0x3b,0x0a,0x68,0x61,0x73,0x68,0x65,0x73,0x5b,0x74,0x5d,0x3d,0x41,0x5b,0x74,0x5d,0x3b,0x0a,0x7d,0x0a,0x7d,0x0a,0x00
|
||||||
|
};
|
||||||
|
|
||||||
|
} // namespace xmrig
|
||||||
153
src/backend/opencl/cl/astrobwt/salsa20.cl
Normal file
153
src/backend/opencl/cl/astrobwt/salsa20.cl
Normal file
|
|
@ -0,0 +1,153 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#define ROTATE(v,c) (rotate(v,(uint32_t)c))
|
||||||
|
#define XOR(v,w) ((v) ^ (w))
|
||||||
|
#define PLUS(v,w) ((v) + (w))
|
||||||
|
|
||||||
|
__attribute__((reqd_work_group_size(SALSA20_GROUP_SIZE, 1, 1)))
|
||||||
|
__kernel void Salsa20_XORKeyStream(__global const uint32_t* keys, __global uint32_t* outputs, __global uint32_t* output_sizes, uint32_t output_stride)
|
||||||
|
{
|
||||||
|
const uint32_t t = get_local_id(0);
|
||||||
|
const uint32_t g = get_group_id(0);
|
||||||
|
|
||||||
|
// Put zeroes in the beginning
|
||||||
|
const uint64_t output_offset = g * ((uint64_t)output_stride) + 128;
|
||||||
|
{
|
||||||
|
__global uint32_t* p = outputs + (output_offset - 128) / sizeof(uint32_t);
|
||||||
|
for (uint32_t i = t; i < 128 / sizeof(uint32_t); i += SALSA20_GROUP_SIZE)
|
||||||
|
p[i] = 0;
|
||||||
|
}
|
||||||
|
|
||||||
|
__global const uint32_t* k = keys + g * 8;
|
||||||
|
__global uint32_t* output = outputs + (output_offset + (t * 64)) / sizeof(uint32_t);
|
||||||
|
const uint32_t output_size = output_sizes[g];
|
||||||
|
|
||||||
|
const uint32_t j1 = k[0];
|
||||||
|
const uint32_t j2 = k[1];
|
||||||
|
const uint32_t j3 = k[2];
|
||||||
|
const uint32_t j4 = k[3];
|
||||||
|
const uint32_t j11 = k[4];
|
||||||
|
const uint32_t j12 = k[5];
|
||||||
|
const uint32_t j13 = k[6];
|
||||||
|
const uint32_t j14 = k[7];
|
||||||
|
const uint32_t j0 = 0x61707865U;
|
||||||
|
const uint32_t j5 = 0x3320646EU;
|
||||||
|
const uint32_t j10 = 0x79622D32U;
|
||||||
|
const uint32_t j15 = 0x6B206574U;
|
||||||
|
const uint32_t j6 = 0;
|
||||||
|
const uint32_t j7 = 0;
|
||||||
|
const uint32_t j8 = 0;
|
||||||
|
const uint32_t j9 = 0;
|
||||||
|
|
||||||
|
for (uint32_t i = t * 64; i < output_size; i += SALSA20_GROUP_SIZE * 64)
|
||||||
|
{
|
||||||
|
const uint32_t j8_1 = j8 + (i / 64);
|
||||||
|
|
||||||
|
uint32_t x0 = j0;
|
||||||
|
uint32_t x1 = j1;
|
||||||
|
uint32_t x2 = j2;
|
||||||
|
uint32_t x3 = j3;
|
||||||
|
uint32_t x4 = j4;
|
||||||
|
uint32_t x5 = j5;
|
||||||
|
uint32_t x6 = j6;
|
||||||
|
uint32_t x7 = j7;
|
||||||
|
uint32_t x8 = j8_1;
|
||||||
|
uint32_t x9 = j9;
|
||||||
|
uint32_t x10 = j10;
|
||||||
|
uint32_t x11 = j11;
|
||||||
|
uint32_t x12 = j12;
|
||||||
|
uint32_t x13 = j13;
|
||||||
|
uint32_t x14 = j14;
|
||||||
|
uint32_t x15 = j15;
|
||||||
|
|
||||||
|
#pragma unroll(5)
|
||||||
|
for (uint32_t j = 0; j < 10; ++j)
|
||||||
|
{
|
||||||
|
x4 = XOR( x4,ROTATE(PLUS( x0,x12), 7));
|
||||||
|
x8 = XOR( x8,ROTATE(PLUS( x4, x0), 9));
|
||||||
|
x12 = XOR(x12,ROTATE(PLUS( x8, x4),13));
|
||||||
|
x0 = XOR( x0,ROTATE(PLUS(x12, x8),18));
|
||||||
|
x9 = XOR( x9,ROTATE(PLUS( x5, x1), 7));
|
||||||
|
x13 = XOR(x13,ROTATE(PLUS( x9, x5), 9));
|
||||||
|
x1 = XOR( x1,ROTATE(PLUS(x13, x9),13));
|
||||||
|
x5 = XOR( x5,ROTATE(PLUS( x1,x13),18));
|
||||||
|
x14 = XOR(x14,ROTATE(PLUS(x10, x6), 7));
|
||||||
|
x2 = XOR( x2,ROTATE(PLUS(x14,x10), 9));
|
||||||
|
x6 = XOR( x6,ROTATE(PLUS( x2,x14),13));
|
||||||
|
x10 = XOR(x10,ROTATE(PLUS( x6, x2),18));
|
||||||
|
x3 = XOR( x3,ROTATE(PLUS(x15,x11), 7));
|
||||||
|
x7 = XOR( x7,ROTATE(PLUS( x3,x15), 9));
|
||||||
|
x11 = XOR(x11,ROTATE(PLUS( x7, x3),13));
|
||||||
|
x15 = XOR(x15,ROTATE(PLUS(x11, x7),18));
|
||||||
|
x1 = XOR( x1,ROTATE(PLUS( x0, x3), 7));
|
||||||
|
x2 = XOR( x2,ROTATE(PLUS( x1, x0), 9));
|
||||||
|
x3 = XOR( x3,ROTATE(PLUS( x2, x1),13));
|
||||||
|
x0 = XOR( x0,ROTATE(PLUS( x3, x2),18));
|
||||||
|
x6 = XOR( x6,ROTATE(PLUS( x5, x4), 7));
|
||||||
|
x7 = XOR( x7,ROTATE(PLUS( x6, x5), 9));
|
||||||
|
x4 = XOR( x4,ROTATE(PLUS( x7, x6),13));
|
||||||
|
x5 = XOR( x5,ROTATE(PLUS( x4, x7),18));
|
||||||
|
x11 = XOR(x11,ROTATE(PLUS(x10, x9), 7));
|
||||||
|
x8 = XOR( x8,ROTATE(PLUS(x11,x10), 9));
|
||||||
|
x9 = XOR( x9,ROTATE(PLUS( x8,x11),13));
|
||||||
|
x10 = XOR(x10,ROTATE(PLUS( x9, x8),18));
|
||||||
|
x12 = XOR(x12,ROTATE(PLUS(x15,x14), 7));
|
||||||
|
x13 = XOR(x13,ROTATE(PLUS(x12,x15), 9));
|
||||||
|
x14 = XOR(x14,ROTATE(PLUS(x13,x12),13));
|
||||||
|
x15 = XOR(x15,ROTATE(PLUS(x14,x13),18));
|
||||||
|
}
|
||||||
|
|
||||||
|
output[0] = PLUS(x0, j0);
|
||||||
|
output[1] = PLUS(x1, j1);
|
||||||
|
output[2] = PLUS(x2, j2);
|
||||||
|
output[3] = PLUS(x3, j3);
|
||||||
|
output[4] = PLUS(x4, j4);
|
||||||
|
output[5] = PLUS(x5, j5);
|
||||||
|
output[6] = PLUS(x6, j6);
|
||||||
|
output[7] = PLUS(x7, j7);
|
||||||
|
output[8] = PLUS(x8, j8_1);
|
||||||
|
output[9] = PLUS(x9, j9);
|
||||||
|
output[10] = PLUS(x10,j10);
|
||||||
|
output[11] = PLUS(x11,j11);
|
||||||
|
output[12] = PLUS(x12,j12);
|
||||||
|
output[13] = PLUS(x13,j13);
|
||||||
|
output[14] = PLUS(x14,j14);
|
||||||
|
output[15] = PLUS(x15,j15);
|
||||||
|
|
||||||
|
output += (SALSA20_GROUP_SIZE * 64) / sizeof(uint32_t);
|
||||||
|
}
|
||||||
|
|
||||||
|
barrier(CLK_GLOBAL_MEM_FENCE);
|
||||||
|
|
||||||
|
// Put zeroes after output's end
|
||||||
|
if (t < 16)
|
||||||
|
{
|
||||||
|
__global uint32_t* p = outputs + (output_offset + output_size + 3) / sizeof(uint32_t);
|
||||||
|
p[t] = 0;
|
||||||
|
}
|
||||||
|
|
||||||
|
if ((t == 0) && (output_size & 3))
|
||||||
|
outputs[(output_offset + output_size) / sizeof(uint32_t)] &= 0xFFFFFFFFU >> ((4 - (output_size & 3)) << 3);
|
||||||
|
}
|
||||||
198
src/backend/opencl/cl/astrobwt/sha3.cl
Normal file
198
src/backend/opencl/cl/astrobwt/sha3.cl
Normal file
|
|
@ -0,0 +1,198 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#define ROUNDS 24
|
||||||
|
#define R64(a,b,c) (((a) << b) | ((a) >> c))
|
||||||
|
|
||||||
|
__constant const uint64_t rc[2][ROUNDS] = {
|
||||||
|
{0x0000000000000001UL, 0x0000000000008082UL, 0x800000000000808AUL,
|
||||||
|
0x8000000080008000UL, 0x000000000000808BUL, 0x0000000080000001UL,
|
||||||
|
0x8000000080008081UL, 0x8000000000008009UL, 0x000000000000008AUL,
|
||||||
|
0x0000000000000088UL, 0x0000000080008009UL, 0x000000008000000AUL,
|
||||||
|
0x000000008000808BUL, 0x800000000000008BUL, 0x8000000000008089UL,
|
||||||
|
0x8000000000008003UL, 0x8000000000008002UL, 0x8000000000000080UL,
|
||||||
|
0x000000000000800AUL, 0x800000008000000AUL, 0x8000000080008081UL,
|
||||||
|
0x8000000000008080UL, 0x0000000080000001UL, 0x8000000080008008UL},
|
||||||
|
{0UL, 0UL, 0UL, 0UL, 0UL, 0UL, 0UL, 0UL,
|
||||||
|
0UL, 0UL, 0UL, 0UL, 0UL, 0UL, 0UL, 0UL,
|
||||||
|
0UL, 0UL, 0UL, 0UL, 0UL, 0UL, 0UL, 0UL}
|
||||||
|
};
|
||||||
|
|
||||||
|
|
||||||
|
__constant const int ro[25][2] = {
|
||||||
|
{ 0,64}, {44,20}, {43,21}, {21,43}, {14,50},
|
||||||
|
{ 1,63}, { 6,58}, {25,39}, { 8,56}, {18,46},
|
||||||
|
{62, 2}, {55, 9}, {39,25}, {41,23}, { 2,62},
|
||||||
|
{28,36}, {20,44}, { 3,61}, {45,19}, {61, 3},
|
||||||
|
{27,37}, {36,28}, {10,54}, {15,49}, {56, 8}
|
||||||
|
};
|
||||||
|
|
||||||
|
__constant const int a[25] = {
|
||||||
|
0, 6, 12, 18, 24,
|
||||||
|
1, 7, 13, 19, 20,
|
||||||
|
2, 8, 14, 15, 21,
|
||||||
|
3, 9, 10, 16, 22,
|
||||||
|
4, 5, 11, 17, 23
|
||||||
|
};
|
||||||
|
|
||||||
|
__constant const int b[25] = {
|
||||||
|
0, 1, 2, 3, 4,
|
||||||
|
1, 2, 3, 4, 0,
|
||||||
|
2, 3, 4, 0, 1,
|
||||||
|
3, 4, 0, 1, 2,
|
||||||
|
4, 0, 1, 2, 3
|
||||||
|
};
|
||||||
|
|
||||||
|
__constant const int c[25][3] = {
|
||||||
|
{ 0, 1, 2}, { 1, 2, 3}, { 2, 3, 4}, { 3, 4, 0}, { 4, 0, 1},
|
||||||
|
{ 5, 6, 7}, { 6, 7, 8}, { 7, 8, 9}, { 8, 9, 5}, { 9, 5, 6},
|
||||||
|
{10,11,12}, {11,12,13}, {12,13,14}, {13,14,10}, {14,10,11},
|
||||||
|
{15,16,17}, {16,17,18}, {17,18,19}, {18,19,15}, {19,15,16},
|
||||||
|
{20,21,22}, {21,22,23}, {22,23,24}, {23,24,20}, {24,20,21}
|
||||||
|
};
|
||||||
|
|
||||||
|
__constant const int d[25] = {
|
||||||
|
0, 1, 2, 3, 4,
|
||||||
|
10, 11, 12, 13, 14,
|
||||||
|
20, 21, 22, 23, 24,
|
||||||
|
5, 6, 7, 8, 9,
|
||||||
|
15, 16, 17, 18, 19
|
||||||
|
};
|
||||||
|
|
||||||
|
__attribute__((reqd_work_group_size(32, 1, 1)))
|
||||||
|
__kernel void sha3(__global const uint8_t* inputs, __global const uint32_t* input_sizes, uint32_t input_stride, __global uint64_t* hashes)
|
||||||
|
{
|
||||||
|
const uint32_t t = get_local_id(0);
|
||||||
|
const uint32_t g = get_group_id(0);
|
||||||
|
|
||||||
|
if (t >= 25)
|
||||||
|
return;
|
||||||
|
|
||||||
|
const uint32_t s = t % 5;
|
||||||
|
|
||||||
|
const uint64_t input_offset = ((uint64_t)input_stride) * g;
|
||||||
|
__global uint64_t* input = (__global uint64_t*)(inputs + input_offset);
|
||||||
|
const uint32_t input_size = input_sizes[g] + 1;
|
||||||
|
|
||||||
|
__local uint64_t A[25];
|
||||||
|
__local uint64_t C[25];
|
||||||
|
__local uint64_t D[25];
|
||||||
|
|
||||||
|
A[t] = 0;
|
||||||
|
|
||||||
|
const uint32_t words = input_size / sizeof(uint64_t);
|
||||||
|
const uint32_t tail_size = input_size % sizeof(uint64_t);
|
||||||
|
|
||||||
|
uint32_t wordIndex = 0;
|
||||||
|
for (uint32_t i = 0; i < words; ++i, ++input)
|
||||||
|
{
|
||||||
|
A[wordIndex] ^= *input;
|
||||||
|
++wordIndex;
|
||||||
|
if (wordIndex == 17)
|
||||||
|
{
|
||||||
|
#pragma unroll(ROUNDS)
|
||||||
|
for (int i = 0; i < ROUNDS; ++i)
|
||||||
|
{
|
||||||
|
C[t] = A[s] ^ A[s+5] ^ A[s+10] ^ A[s+15] ^ A[s+20];
|
||||||
|
D[t] = C[b[20+s]] ^ R64(C[b[5+s]], 1, 63);
|
||||||
|
C[t] = R64(A[a[t]] ^ D[b[t]], ro[t][0], ro[t][1]);
|
||||||
|
A[d[t]] = C[c[t][0]] ^ ((~C[c[t][1]]) & C[c[t][2]]);
|
||||||
|
A[t] ^= rc[(t == 0) ? 0 : 1][i];
|
||||||
|
}
|
||||||
|
wordIndex = 0;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
uint64_t tail = 0;
|
||||||
|
__global const uint8_t* p = (__global const uint8_t*)input;
|
||||||
|
for (uint32_t i = 0; i < tail_size; ++i)
|
||||||
|
{
|
||||||
|
tail |= (uint64_t)(p[i]) << (i * 8);
|
||||||
|
}
|
||||||
|
A[wordIndex] ^= tail ^ ((uint64_t)(((uint64_t)(0x02 | (1 << 2))) << (tail_size * 8)));
|
||||||
|
A[16] ^= 0x8000000000000000UL;
|
||||||
|
|
||||||
|
#pragma unroll(1)
|
||||||
|
for (int i = 0; i < ROUNDS; ++i)
|
||||||
|
{
|
||||||
|
C[t] = A[s] ^ A[s+5] ^ A[s+10] ^ A[s+15] ^ A[s+20];
|
||||||
|
D[t] = C[b[20+s]] ^ R64(C[b[5+s]], 1, 63);
|
||||||
|
C[t] = R64(A[a[t]] ^ D[b[t]], ro[t][0], ro[t][1]);
|
||||||
|
A[d[t]] = C[c[t][0]] ^ ((~C[c[t][1]]) & C[c[t][2]]);
|
||||||
|
A[t] ^= rc[(t == 0) ? 0 : 1][i];
|
||||||
|
}
|
||||||
|
|
||||||
|
if (t < 4)
|
||||||
|
{
|
||||||
|
hashes += g * (32 / sizeof(uint64_t));
|
||||||
|
hashes[t] = A[t];
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
__attribute__((reqd_work_group_size(32, 1, 1)))
|
||||||
|
__kernel void sha3_initial(__global const uint8_t* input_data, uint32_t input_size, uint32_t nonce, __global uint64_t* hashes)
|
||||||
|
{
|
||||||
|
const uint32_t t = get_local_id(0);
|
||||||
|
const uint32_t g = get_group_id(0);
|
||||||
|
|
||||||
|
if (t >= 25)
|
||||||
|
return;
|
||||||
|
|
||||||
|
const uint32_t s = t % 5;
|
||||||
|
|
||||||
|
__global uint64_t* input = (__global uint64_t*)(input_data);
|
||||||
|
|
||||||
|
__local uint64_t A[25];
|
||||||
|
__local uint64_t C[25];
|
||||||
|
__local uint64_t D[25];
|
||||||
|
|
||||||
|
A[t] = (t < 16) ? input[t] : 0;
|
||||||
|
|
||||||
|
__local uint32_t* nonce_pos = (__local uint32_t*)(A) + 9;
|
||||||
|
nonce += g;
|
||||||
|
nonce_pos[0] = (nonce_pos[0] & 0xFFFFFFU) | ((nonce & 0xFF) << 24);
|
||||||
|
nonce_pos[1] = (nonce_pos[1] & 0xFF000000U) | (nonce >> 8);
|
||||||
|
|
||||||
|
uint32_t wordIndex = input_size / sizeof(uint64_t);
|
||||||
|
const uint32_t tail_size = input_size % sizeof(uint64_t);
|
||||||
|
|
||||||
|
A[wordIndex] ^= (uint64_t)(((uint64_t)(0x02 | (1 << 2))) << (tail_size * 8));
|
||||||
|
A[16] ^= 0x8000000000000000UL;
|
||||||
|
|
||||||
|
#pragma unroll(ROUNDS)
|
||||||
|
for (int i = 0; i < ROUNDS; ++i)
|
||||||
|
{
|
||||||
|
C[t] = A[s] ^ A[s+5] ^ A[s+10] ^ A[s+15] ^ A[s+20];
|
||||||
|
D[t] = C[b[20+s]] ^ R64(C[b[5+s]], 1, 63);
|
||||||
|
C[t] = R64(A[a[t]] ^ D[b[t]], ro[t][0], ro[t][1]);
|
||||||
|
A[d[t]] = C[c[t][0]] ^ ((~C[c[t][1]]) & C[c[t][2]]);
|
||||||
|
A[t] ^= rc[(t == 0) ? 0 : 1][i];
|
||||||
|
}
|
||||||
|
|
||||||
|
if (t < 4)
|
||||||
|
{
|
||||||
|
hashes += g * (32 / sizeof(uint64_t));
|
||||||
|
hashes[t] = A[t];
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,64 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018 Lee Clagett <https://github.com/vtnerd>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/OclThreads.h"
|
||||||
|
#include "backend/opencl/wrappers/OclDevice.h"
|
||||||
|
#include "base/crypto/Algorithm.h"
|
||||||
|
#include "crypto/randomx/randomx.h"
|
||||||
|
#include "crypto/rx/RxAlgo.h"
|
||||||
|
|
||||||
|
|
||||||
|
namespace xmrig {
|
||||||
|
|
||||||
|
|
||||||
|
bool ocl_generic_astrobwt_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads)
|
||||||
|
{
|
||||||
|
if (algorithm.family() != Algorithm::ASTROBWT) {
|
||||||
|
return false;
|
||||||
|
}
|
||||||
|
|
||||||
|
const size_t mem = device.globalMemSize();
|
||||||
|
|
||||||
|
uint32_t per_thread_mem = 10 << 20;
|
||||||
|
uint32_t intensity = static_cast<uint32_t>((mem - (128 << 20)) / per_thread_mem / 2);
|
||||||
|
|
||||||
|
intensity &= ~63U;
|
||||||
|
|
||||||
|
if (!intensity) {
|
||||||
|
return false;
|
||||||
|
}
|
||||||
|
|
||||||
|
if (intensity > 256) {
|
||||||
|
intensity = 256;
|
||||||
|
}
|
||||||
|
|
||||||
|
threads.add(OclThread(device.index(), intensity, 2));
|
||||||
|
|
||||||
|
return true;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
} // namespace xmrig
|
||||||
|
|
@ -59,6 +59,8 @@ public:
|
||||||
virtual const OclLaunchData &data() const = 0;
|
virtual const OclLaunchData &data() const = 0;
|
||||||
virtual size_t intensity() const = 0;
|
virtual size_t intensity() const = 0;
|
||||||
virtual size_t threadId() const = 0;
|
virtual size_t threadId() const = 0;
|
||||||
|
virtual uint32_t roundSize() const = 0;
|
||||||
|
virtual uint32_t processedHashes() const = 0;
|
||||||
virtual uint32_t deviceIndex() const = 0;
|
virtual uint32_t deviceIndex() const = 0;
|
||||||
virtual void build() = 0;
|
virtual void build() = 0;
|
||||||
virtual void init() = 0;
|
virtual void init() = 0;
|
||||||
|
|
|
||||||
|
|
@ -0,0 +1,42 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.h"
|
||||||
|
#include "backend/opencl/wrappers/OclLib.h"
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_FilterKernel::enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size)
|
||||||
|
{
|
||||||
|
enqueueNDRange(queue, 1, nullptr, &threads, &workgroup_size);
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_FilterKernel::setArgs(uint32_t nonce, uint32_t bwt_max_size, cl_mem hashes, cl_mem filtered_hashes)
|
||||||
|
{
|
||||||
|
setArg(0, sizeof(uint32_t), &nonce);
|
||||||
|
setArg(1, sizeof(uint32_t), &bwt_max_size);
|
||||||
|
setArg(2, sizeof(cl_mem), &hashes);
|
||||||
|
setArg(3, sizeof(cl_mem), &filtered_hashes);
|
||||||
|
}
|
||||||
48
src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.h
Normal file
48
src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.h
Normal file
|
|
@ -0,0 +1,48 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#ifndef XMRIG_ASTROBWT_FILTERKERNEL_H
|
||||||
|
#define XMRIG_ASTROBWT_FILTERKERNEL_H
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/wrappers/OclKernel.h"
|
||||||
|
|
||||||
|
|
||||||
|
namespace xmrig {
|
||||||
|
|
||||||
|
|
||||||
|
class AstroBWT_FilterKernel : public OclKernel
|
||||||
|
{
|
||||||
|
public:
|
||||||
|
inline AstroBWT_FilterKernel(cl_program program) : OclKernel(program, "filter") {}
|
||||||
|
|
||||||
|
void enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size);
|
||||||
|
void setArgs(uint32_t nonce, uint32_t bwt_max_size, cl_mem hashes, cl_mem filtered_hashes);
|
||||||
|
};
|
||||||
|
|
||||||
|
|
||||||
|
} // namespace xmrig
|
||||||
|
|
||||||
|
|
||||||
|
#endif /* XMRIG_ASTROBWT_FILTERKERNEL_H */
|
||||||
|
|
@ -0,0 +1,47 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2019 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2019 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.h"
|
||||||
|
#include "backend/opencl/wrappers/OclLib.h"
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_FindSharesKernel::enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size)
|
||||||
|
{
|
||||||
|
enqueueNDRange(queue, 1, nullptr, &threads, &workgroup_size);
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_FindSharesKernel::setArgs(cl_mem hashes, cl_mem filtered_hashes, cl_mem shares)
|
||||||
|
{
|
||||||
|
setArg(0, sizeof(cl_mem), &hashes);
|
||||||
|
setArg(1, sizeof(cl_mem), &filtered_hashes);
|
||||||
|
setArg(3, sizeof(cl_mem), &shares);
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_FindSharesKernel::setTarget(uint64_t target)
|
||||||
|
{
|
||||||
|
setArg(2, sizeof(uint64_t), &target);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,49 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2019 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2019 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#ifndef XMRIG_ASTROBWTFINDSHARESKERNEL_H
|
||||||
|
#define XMRIG_ASTROBWTFINDSHARESKERNEL_H
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/wrappers/OclKernel.h"
|
||||||
|
|
||||||
|
|
||||||
|
namespace xmrig {
|
||||||
|
|
||||||
|
|
||||||
|
class AstroBWT_FindSharesKernel : public OclKernel
|
||||||
|
{
|
||||||
|
public:
|
||||||
|
inline AstroBWT_FindSharesKernel(cl_program program) : OclKernel(program, "find_shares") {}
|
||||||
|
|
||||||
|
void enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size);
|
||||||
|
void setArgs(cl_mem hashes, cl_mem filtered_hashes, cl_mem shares);
|
||||||
|
void setTarget(uint64_t target);
|
||||||
|
};
|
||||||
|
|
||||||
|
|
||||||
|
} // namespace xmrig
|
||||||
|
|
||||||
|
|
||||||
|
#endif /* XMRIG_ASTROBWTFINDSHARESKERNEL_H */
|
||||||
44
src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.cpp
Normal file
44
src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.cpp
Normal file
|
|
@ -0,0 +1,44 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.h"
|
||||||
|
#include "backend/opencl/wrappers/OclLib.h"
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_MainKernel::enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size)
|
||||||
|
{
|
||||||
|
const size_t gthreads = threads * workgroup_size;
|
||||||
|
enqueueNDRange(queue, 1, nullptr, >hreads, &workgroup_size);
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_MainKernel::setArgs(cl_mem inputs, cl_mem input_sizes, uint32_t input_stride, cl_mem indices, cl_mem tmp_indices)
|
||||||
|
{
|
||||||
|
setArg(0, sizeof(cl_mem), &inputs);
|
||||||
|
setArg(1, sizeof(cl_mem), &input_sizes);
|
||||||
|
setArg(2, sizeof(uint32_t), &input_stride);
|
||||||
|
setArg(3, sizeof(cl_mem), &indices);
|
||||||
|
setArg(4, sizeof(cl_mem), &tmp_indices);
|
||||||
|
}
|
||||||
48
src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.h
Normal file
48
src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.h
Normal file
|
|
@ -0,0 +1,48 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#ifndef XMRIG_ASTROBWT_MAINKERNEL_H
|
||||||
|
#define XMRIG_ASTROBWT_MAINKERNEL_H
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/wrappers/OclKernel.h"
|
||||||
|
|
||||||
|
|
||||||
|
namespace xmrig {
|
||||||
|
|
||||||
|
|
||||||
|
class AstroBWT_MainKernel : public OclKernel
|
||||||
|
{
|
||||||
|
public:
|
||||||
|
inline AstroBWT_MainKernel(cl_program program) : OclKernel(program, "BWT") {}
|
||||||
|
|
||||||
|
void enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size);
|
||||||
|
void setArgs(cl_mem inputs, cl_mem input_sizes, uint32_t input_stride, cl_mem indices, cl_mem tmp_indices);
|
||||||
|
};
|
||||||
|
|
||||||
|
|
||||||
|
} // namespace xmrig
|
||||||
|
|
||||||
|
|
||||||
|
#endif /* XMRIG_ASTROBWT_MAINKERNEL_H */
|
||||||
|
|
@ -0,0 +1,41 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.h"
|
||||||
|
#include "backend/opencl/wrappers/OclLib.h"
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_PrepareBatch2Kernel::enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size)
|
||||||
|
{
|
||||||
|
enqueueNDRange(queue, 1, nullptr, &threads, &workgroup_size);
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_PrepareBatch2Kernel::setArgs(cl_mem hashes, cl_mem filtered_hashes, cl_mem bwt_data_sizes)
|
||||||
|
{
|
||||||
|
setArg(0, sizeof(cl_mem), &hashes);
|
||||||
|
setArg(1, sizeof(cl_mem), &filtered_hashes);
|
||||||
|
setArg(2, sizeof(cl_mem), &bwt_data_sizes);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,48 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#ifndef XMRIG_ASTROBWT_PREPAREBATCH2KERNEL_H
|
||||||
|
#define XMRIG_ASTROBWT_PREPAREBATCH2KERNEL_H
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/wrappers/OclKernel.h"
|
||||||
|
|
||||||
|
|
||||||
|
namespace xmrig {
|
||||||
|
|
||||||
|
|
||||||
|
class AstroBWT_PrepareBatch2Kernel : public OclKernel
|
||||||
|
{
|
||||||
|
public:
|
||||||
|
inline AstroBWT_PrepareBatch2Kernel(cl_program program) : OclKernel(program, "prepare_batch2") {}
|
||||||
|
|
||||||
|
void enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size);
|
||||||
|
void setArgs(cl_mem hashes, cl_mem filtered_hashes, cl_mem bwt_data_sizes);
|
||||||
|
};
|
||||||
|
|
||||||
|
|
||||||
|
} // namespace xmrig
|
||||||
|
|
||||||
|
|
||||||
|
#endif /* XMRIG_ASTROBWT_PREPAREBATCH2KERNEL_H */
|
||||||
|
|
@ -0,0 +1,44 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.h"
|
||||||
|
#include "backend/opencl/wrappers/OclLib.h"
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_SHA3InitialKernel::enqueue(cl_command_queue queue, size_t threads)
|
||||||
|
{
|
||||||
|
const size_t workgroup_size = 32;
|
||||||
|
const size_t gthreads = threads * workgroup_size;
|
||||||
|
enqueueNDRange(queue, 1, nullptr, >hreads, &workgroup_size);
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_SHA3InitialKernel::setArgs(cl_mem input, uint32_t input_size, uint32_t nonce, cl_mem output_salsa20_keys)
|
||||||
|
{
|
||||||
|
setArg(0, sizeof(cl_mem), &input);
|
||||||
|
setArg(1, sizeof(uint32_t), &input_size);
|
||||||
|
setArg(2, sizeof(uint32_t), &nonce);
|
||||||
|
setArg(3, sizeof(cl_mem), &output_salsa20_keys);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,48 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#ifndef XMRIG_ASTROBWT_SHA3INITIALKERNEL_H
|
||||||
|
#define XMRIG_ASTROBWT_SHA3INITIALKERNEL_H
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/wrappers/OclKernel.h"
|
||||||
|
|
||||||
|
|
||||||
|
namespace xmrig {
|
||||||
|
|
||||||
|
|
||||||
|
class AstroBWT_SHA3InitialKernel : public OclKernel
|
||||||
|
{
|
||||||
|
public:
|
||||||
|
inline AstroBWT_SHA3InitialKernel(cl_program program) : OclKernel(program, "sha3_initial") {}
|
||||||
|
|
||||||
|
void enqueue(cl_command_queue queue, size_t threads);
|
||||||
|
void setArgs(cl_mem input, uint32_t input_size, uint32_t nonce, cl_mem output_salsa20_keys);
|
||||||
|
};
|
||||||
|
|
||||||
|
|
||||||
|
} // namespace xmrig
|
||||||
|
|
||||||
|
|
||||||
|
#endif /* XMRIG_ASTROBWT_SHA3INITIALKERNEL_H */
|
||||||
44
src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.cpp
Normal file
44
src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.cpp
Normal file
|
|
@ -0,0 +1,44 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.h"
|
||||||
|
#include "backend/opencl/wrappers/OclLib.h"
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_SHA3Kernel::enqueue(cl_command_queue queue, size_t threads)
|
||||||
|
{
|
||||||
|
const size_t workgroup_size = 32;
|
||||||
|
const size_t gthreads = threads * workgroup_size;
|
||||||
|
enqueueNDRange(queue, 1, nullptr, >hreads, &workgroup_size);
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_SHA3Kernel::setArgs(cl_mem inputs, cl_mem input_sizes, uint32_t input_stride, cl_mem output_salsa20_keys)
|
||||||
|
{
|
||||||
|
setArg(0, sizeof(cl_mem), &inputs);
|
||||||
|
setArg(1, sizeof(cl_mem), &input_sizes);
|
||||||
|
setArg(2, sizeof(uint32_t), &input_stride);
|
||||||
|
setArg(3, sizeof(cl_mem), &output_salsa20_keys);
|
||||||
|
}
|
||||||
48
src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.h
Normal file
48
src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.h
Normal file
|
|
@ -0,0 +1,48 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#ifndef XMRIG_ASTROBWT_SHA3KERNEL_H
|
||||||
|
#define XMRIG_ASTROBWT_SHA3KERNEL_H
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/wrappers/OclKernel.h"
|
||||||
|
|
||||||
|
|
||||||
|
namespace xmrig {
|
||||||
|
|
||||||
|
|
||||||
|
class AstroBWT_SHA3Kernel : public OclKernel
|
||||||
|
{
|
||||||
|
public:
|
||||||
|
inline AstroBWT_SHA3Kernel(cl_program program) : OclKernel(program, "sha3") {}
|
||||||
|
|
||||||
|
void enqueue(cl_command_queue queue, size_t threads);
|
||||||
|
void setArgs(cl_mem inputs, cl_mem input_sizes, uint32_t input_stride, cl_mem output_salsa20_keys);
|
||||||
|
};
|
||||||
|
|
||||||
|
|
||||||
|
} // namespace xmrig
|
||||||
|
|
||||||
|
|
||||||
|
#endif /* XMRIG_ASTROBWT_SHA3KERNEL_H */
|
||||||
|
|
@ -0,0 +1,43 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.h"
|
||||||
|
#include "backend/opencl/wrappers/OclLib.h"
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_Salsa20Kernel::enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size)
|
||||||
|
{
|
||||||
|
const size_t gthreads = threads * workgroup_size;
|
||||||
|
enqueueNDRange(queue, 1, nullptr, >hreads, &workgroup_size);
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::AstroBWT_Salsa20Kernel::setArgs(cl_mem salsa20_keys, cl_mem outputs, cl_mem output_sizes, uint32_t output_stride)
|
||||||
|
{
|
||||||
|
setArg(0, sizeof(cl_mem), &salsa20_keys);
|
||||||
|
setArg(1, sizeof(cl_mem), &outputs);
|
||||||
|
setArg(2, sizeof(cl_mem), &output_sizes);
|
||||||
|
setArg(3, sizeof(uint32_t), &output_stride);
|
||||||
|
}
|
||||||
48
src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.h
Normal file
48
src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.h
Normal file
|
|
@ -0,0 +1,48 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#ifndef XMRIG_ASTROBWT_SALSA20KERNEL_H
|
||||||
|
#define XMRIG_ASTROBWT_SALSA20KERNEL_H
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/wrappers/OclKernel.h"
|
||||||
|
|
||||||
|
|
||||||
|
namespace xmrig {
|
||||||
|
|
||||||
|
|
||||||
|
class AstroBWT_Salsa20Kernel : public OclKernel
|
||||||
|
{
|
||||||
|
public:
|
||||||
|
inline AstroBWT_Salsa20Kernel(cl_program program) : OclKernel(program, "Salsa20_XORKeyStream") {}
|
||||||
|
|
||||||
|
void enqueue(cl_command_queue queue, size_t threads, size_t workgroup_size);
|
||||||
|
void setArgs(cl_mem salsa20_keys, cl_mem outputs, cl_mem output_sizes, uint32_t output_stride);
|
||||||
|
};
|
||||||
|
|
||||||
|
|
||||||
|
} // namespace xmrig
|
||||||
|
|
||||||
|
|
||||||
|
#endif /* XMRIG_ASTROBWT_SALSA20KERNEL_H */
|
||||||
|
|
@ -100,6 +100,31 @@ if (WITH_OPENCL)
|
||||||
)
|
)
|
||||||
endif()
|
endif()
|
||||||
|
|
||||||
|
if (WITH_ASTROBWT)
|
||||||
|
list(APPEND HEADERS_BACKEND_OPENCL
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.h
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.h
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.h
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.h
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.h
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.h
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.h
|
||||||
|
src/backend/opencl/runners/OclAstroBWTRunner.h
|
||||||
|
)
|
||||||
|
|
||||||
|
list(APPEND SOURCES_BACKEND_OPENCL
|
||||||
|
src/backend/opencl/generators/ocl_generic_astrobwt_generator.cpp
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.cpp
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.cpp
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.cpp
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.cpp
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.cpp
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.cpp
|
||||||
|
src/backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.cpp
|
||||||
|
src/backend/opencl/runners/OclAstroBWTRunner.cpp
|
||||||
|
)
|
||||||
|
endif()
|
||||||
|
|
||||||
if (WITH_CN_GPU AND CMAKE_SIZEOF_VOID_P EQUAL 8)
|
if (WITH_CN_GPU AND CMAKE_SIZEOF_VOID_P EQUAL 8)
|
||||||
list(APPEND HEADERS_BACKEND_OPENCL
|
list(APPEND HEADERS_BACKEND_OPENCL
|
||||||
src/backend/opencl/kernels/Cn00RyoKernel.h
|
src/backend/opencl/kernels/Cn00RyoKernel.h
|
||||||
|
|
|
||||||
218
src/backend/opencl/runners/OclAstroBWTRunner.cpp
Normal file
218
src/backend/opencl/runners/OclAstroBWTRunner.cpp
Normal file
|
|
@ -0,0 +1,218 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/runners/OclAstroBWTRunner.h"
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_FilterKernel.h"
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_FindSharesKernel.h"
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_MainKernel.h"
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_PrepareBatch2Kernel.h"
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_Salsa20Kernel.h"
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_SHA3InitialKernel.h"
|
||||||
|
#include "backend/opencl/kernels/astrobwt/AstroBWT_SHA3Kernel.h"
|
||||||
|
#include "backend/opencl/OclLaunchData.h"
|
||||||
|
#include "backend/opencl/wrappers/OclLib.h"
|
||||||
|
#include "base/io/log/Log.h"
|
||||||
|
#include "base/net/stratum/Job.h"
|
||||||
|
|
||||||
|
|
||||||
|
namespace xmrig {
|
||||||
|
|
||||||
|
|
||||||
|
constexpr int STAGE1_SIZE = 147253;
|
||||||
|
constexpr uint32_t STAGE1_DATA_STRIDE = (STAGE1_SIZE + 256 + 255) & ~255U;
|
||||||
|
constexpr uint32_t OclAstroBWTRunner::BWT_DATA_STRIDE;
|
||||||
|
|
||||||
|
|
||||||
|
} // namespace xmrig
|
||||||
|
|
||||||
|
|
||||||
|
xmrig::OclAstroBWTRunner::OclAstroBWTRunner(size_t index, const OclLaunchData &data) : OclBaseRunner(index, data)
|
||||||
|
{
|
||||||
|
switch (data.device.type())
|
||||||
|
{
|
||||||
|
case OclDevice::Baffin:
|
||||||
|
case OclDevice::Polaris:
|
||||||
|
case OclDevice::Lexa:
|
||||||
|
case OclDevice::Vega_10:
|
||||||
|
case OclDevice::Vega_20:
|
||||||
|
case OclDevice::Raven:
|
||||||
|
m_workgroup_size = 64;
|
||||||
|
break;
|
||||||
|
|
||||||
|
default:
|
||||||
|
m_workgroup_size = 32;
|
||||||
|
break;
|
||||||
|
}
|
||||||
|
|
||||||
|
m_options += " -DSALSA20_GROUP_SIZE=" + std::to_string(m_workgroup_size);
|
||||||
|
m_options += " -DBWT_GROUP_SIZE=" + std::to_string(m_workgroup_size);
|
||||||
|
|
||||||
|
m_bwt_allocation_size = static_cast<uint64_t>(m_intensity) * BWT_DATA_STRIDE;
|
||||||
|
m_batch_size1 = static_cast<uint32_t>(m_bwt_allocation_size / STAGE1_DATA_STRIDE) & ~255U;
|
||||||
|
|
||||||
|
m_bwt_data_sizes_host = new uint32_t[m_batch_size1];
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
xmrig::OclAstroBWTRunner::~OclAstroBWTRunner()
|
||||||
|
{
|
||||||
|
delete m_sha3_initial_kernel;
|
||||||
|
delete m_sha3_kernel;
|
||||||
|
delete m_salsa20_kernel;
|
||||||
|
delete m_bwt_kernel;
|
||||||
|
delete m_filter_kernel;
|
||||||
|
delete m_prepare_batch2_kernel;
|
||||||
|
delete m_find_shares_kernel;
|
||||||
|
|
||||||
|
OclLib::release(m_salsa20_keys);
|
||||||
|
OclLib::release(m_bwt_data);
|
||||||
|
OclLib::release(m_bwt_data_sizes);
|
||||||
|
OclLib::release(m_indices);
|
||||||
|
OclLib::release(m_tmp_indices);
|
||||||
|
OclLib::release(m_filtered_hashes);
|
||||||
|
|
||||||
|
delete m_bwt_data_sizes_host;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
size_t xmrig::OclAstroBWTRunner::bufferSize() const
|
||||||
|
{
|
||||||
|
return OclBaseRunner::bufferSize() +
|
||||||
|
align(m_batch_size1 * 32) + // m_salsa20_keys
|
||||||
|
align(m_bwt_allocation_size) + // m_bwt_data
|
||||||
|
align(m_batch_size1 * 4) + // m_bwt_data_sizes
|
||||||
|
align(m_bwt_allocation_size * 8) + // m_indices
|
||||||
|
align(m_bwt_allocation_size * 8) + // m_tmp_indices
|
||||||
|
align((m_batch_size1 + m_intensity) * 36 + 4); // m_filtered_hashes
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::OclAstroBWTRunner::run(uint32_t nonce, uint32_t *hashOutput)
|
||||||
|
{
|
||||||
|
m_sha3_initial_kernel->setArg(2, sizeof(nonce), &nonce);
|
||||||
|
m_salsa20_kernel->setArg(3, sizeof(STAGE1_DATA_STRIDE), &STAGE1_DATA_STRIDE);
|
||||||
|
m_bwt_kernel->setArg(2, sizeof(STAGE1_DATA_STRIDE), &STAGE1_DATA_STRIDE);
|
||||||
|
|
||||||
|
const uint32_t t = STAGE1_DATA_STRIDE * 8;
|
||||||
|
m_sha3_kernel->setArg(2, sizeof(t), &t);
|
||||||
|
m_filter_kernel->setArg(0, sizeof(nonce), &nonce);
|
||||||
|
|
||||||
|
const uint32_t zero = 0;
|
||||||
|
enqueueWriteBuffer(m_output, CL_FALSE, sizeof(cl_uint) * 0xFF, sizeof(uint32_t), &zero);
|
||||||
|
|
||||||
|
m_sha3_initial_kernel->enqueue(m_queue, m_batch_size1);
|
||||||
|
|
||||||
|
for (uint32_t i = 0; i < m_batch_size1; ++i)
|
||||||
|
m_bwt_data_sizes_host[i] = STAGE1_SIZE;
|
||||||
|
|
||||||
|
enqueueWriteBuffer(m_bwt_data_sizes, CL_FALSE, 0, m_batch_size1 * sizeof(uint32_t), m_bwt_data_sizes_host);
|
||||||
|
|
||||||
|
m_salsa20_kernel->enqueue(m_queue, m_batch_size1, m_workgroup_size);
|
||||||
|
m_bwt_kernel->enqueue(m_queue, m_batch_size1, m_workgroup_size);
|
||||||
|
m_sha3_kernel->enqueue(m_queue, m_batch_size1);
|
||||||
|
m_filter_kernel->enqueue(m_queue, m_batch_size1, m_workgroup_size);
|
||||||
|
|
||||||
|
uint32_t num_filtered_hashes = 0;
|
||||||
|
enqueueReadBuffer(m_filtered_hashes, CL_TRUE, 0, sizeof(num_filtered_hashes), &num_filtered_hashes);
|
||||||
|
|
||||||
|
m_processedHashes = 0;
|
||||||
|
while (num_filtered_hashes >= m_intensity)
|
||||||
|
{
|
||||||
|
num_filtered_hashes -= m_intensity;
|
||||||
|
m_processedHashes += m_intensity;
|
||||||
|
|
||||||
|
m_salsa20_kernel->setArg(3, sizeof(BWT_DATA_STRIDE), &BWT_DATA_STRIDE);
|
||||||
|
m_bwt_kernel->setArg(2, sizeof(BWT_DATA_STRIDE), &BWT_DATA_STRIDE);
|
||||||
|
|
||||||
|
const uint32_t t = BWT_DATA_STRIDE * 8;
|
||||||
|
m_sha3_kernel->setArg(2, sizeof(t), &t);
|
||||||
|
|
||||||
|
m_prepare_batch2_kernel->enqueue(m_queue, m_intensity, m_workgroup_size);
|
||||||
|
m_salsa20_kernel->enqueue(m_queue, m_intensity, m_workgroup_size);
|
||||||
|
m_bwt_kernel->enqueue(m_queue, m_intensity, m_workgroup_size);
|
||||||
|
m_sha3_kernel->enqueue(m_queue, m_intensity);
|
||||||
|
|
||||||
|
m_find_shares_kernel->enqueue(m_queue, m_intensity, m_workgroup_size);
|
||||||
|
|
||||||
|
finalize(hashOutput);
|
||||||
|
|
||||||
|
OclLib::finish(m_queue);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::OclAstroBWTRunner::set(const Job &job, uint8_t *blob)
|
||||||
|
{
|
||||||
|
if (job.size() > (Job::kMaxBlobSize - 4)) {
|
||||||
|
throw std::length_error("job size too big");
|
||||||
|
}
|
||||||
|
|
||||||
|
if (job.size() < Job::kMaxBlobSize) {
|
||||||
|
memset(blob + job.size(), 0, Job::kMaxBlobSize - job.size());
|
||||||
|
}
|
||||||
|
|
||||||
|
enqueueWriteBuffer(m_input, CL_TRUE, 0, Job::kMaxBlobSize, blob);
|
||||||
|
|
||||||
|
m_sha3_initial_kernel->setArgs(m_input, static_cast<uint32_t>(job.size()), *job.nonce(), m_salsa20_keys);
|
||||||
|
m_salsa20_kernel->setArgs(m_salsa20_keys, m_bwt_data, m_bwt_data_sizes, STAGE1_DATA_STRIDE);
|
||||||
|
m_bwt_kernel->setArgs(m_bwt_data, m_bwt_data_sizes, STAGE1_DATA_STRIDE, m_indices, m_tmp_indices);
|
||||||
|
m_sha3_kernel->setArgs(m_tmp_indices, m_bwt_data_sizes, STAGE1_DATA_STRIDE * 8, m_salsa20_keys);
|
||||||
|
m_filter_kernel->setArgs(*job.nonce(), BWT_DATA_MAX_SIZE, m_salsa20_keys, m_filtered_hashes);
|
||||||
|
m_prepare_batch2_kernel->setArgs(m_salsa20_keys, m_filtered_hashes, m_bwt_data_sizes);
|
||||||
|
m_find_shares_kernel->setArgs(m_salsa20_keys, m_filtered_hashes, m_output);
|
||||||
|
m_find_shares_kernel->setTarget(job.target());
|
||||||
|
|
||||||
|
const uint32_t zero = 0;
|
||||||
|
enqueueWriteBuffer(m_filtered_hashes, CL_TRUE, 0, sizeof(uint32_t), &zero);
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::OclAstroBWTRunner::build()
|
||||||
|
{
|
||||||
|
OclBaseRunner::build();
|
||||||
|
|
||||||
|
m_sha3_initial_kernel = new AstroBWT_SHA3InitialKernel(m_program);
|
||||||
|
m_sha3_kernel = new AstroBWT_SHA3Kernel(m_program);
|
||||||
|
m_salsa20_kernel = new AstroBWT_Salsa20Kernel(m_program);
|
||||||
|
m_bwt_kernel = new AstroBWT_MainKernel(m_program);
|
||||||
|
m_filter_kernel = new AstroBWT_FilterKernel(m_program);
|
||||||
|
m_prepare_batch2_kernel = new AstroBWT_PrepareBatch2Kernel(m_program);
|
||||||
|
m_find_shares_kernel = new AstroBWT_FindSharesKernel(m_program);
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
void xmrig::OclAstroBWTRunner::init()
|
||||||
|
{
|
||||||
|
OclBaseRunner::init();
|
||||||
|
|
||||||
|
const cl_mem_flags f = CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS;
|
||||||
|
|
||||||
|
m_salsa20_keys = createSubBuffer(f, m_batch_size1 * 32);
|
||||||
|
m_bwt_data = createSubBuffer(f, m_bwt_allocation_size);
|
||||||
|
m_bwt_data_sizes = createSubBuffer(CL_MEM_READ_WRITE | CL_MEM_HOST_WRITE_ONLY, m_batch_size1 * 4);
|
||||||
|
m_indices = createSubBuffer(f, m_bwt_allocation_size * 8);
|
||||||
|
m_tmp_indices = createSubBuffer(f, m_bwt_allocation_size * 8);
|
||||||
|
m_filtered_hashes = createSubBuffer(CL_MEM_READ_WRITE, (m_batch_size1 + m_intensity) * 36 + 4);
|
||||||
|
}
|
||||||
95
src/backend/opencl/runners/OclAstroBWTRunner.h
Normal file
95
src/backend/opencl/runners/OclAstroBWTRunner.h
Normal file
|
|
@ -0,0 +1,95 @@
|
||||||
|
/* XMRig
|
||||||
|
* Copyright 2010 Jeff Garzik <jgarzik@pobox.com>
|
||||||
|
* Copyright 2012-2014 pooler <pooler@litecoinpool.org>
|
||||||
|
* Copyright 2014 Lucas Jones <https://github.com/lucasjones>
|
||||||
|
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
|
||||||
|
* Copyright 2016 Jay D Dee <jayddee246@gmail.com>
|
||||||
|
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
|
||||||
|
* Copyright 2018-2020 SChernykh <https://github.com/SChernykh>
|
||||||
|
* Copyright 2016-2020 XMRig <https://github.com/xmrig>, <support@xmrig.com>
|
||||||
|
*
|
||||||
|
* This program is free software: you can redistribute it and/or modify
|
||||||
|
* it under the terms of the GNU General Public License as published by
|
||||||
|
* the Free Software Foundation, either version 3 of the License, or
|
||||||
|
* (at your option) any later version.
|
||||||
|
*
|
||||||
|
* This program is distributed in the hope that it will be useful,
|
||||||
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||||
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||||
|
* GNU General Public License for more details.
|
||||||
|
*
|
||||||
|
* You should have received a copy of the GNU General Public License
|
||||||
|
* along with this program. If not, see <http://www.gnu.org/licenses/>.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#ifndef XMRIG_OCLASTROBWTRUNNER_H
|
||||||
|
#define XMRIG_OCLASTROBWTRUNNER_H
|
||||||
|
|
||||||
|
|
||||||
|
#include "backend/opencl/runners/OclBaseRunner.h"
|
||||||
|
|
||||||
|
namespace xmrig {
|
||||||
|
|
||||||
|
|
||||||
|
class AstroBWT_FilterKernel;
|
||||||
|
class AstroBWT_MainKernel;
|
||||||
|
class AstroBWT_PrepareBatch2Kernel;
|
||||||
|
class AstroBWT_Salsa20Kernel;
|
||||||
|
class AstroBWT_SHA3InitialKernel;
|
||||||
|
class AstroBWT_SHA3Kernel;
|
||||||
|
class AstroBWT_FindSharesKernel;
|
||||||
|
|
||||||
|
|
||||||
|
class OclAstroBWTRunner : public OclBaseRunner
|
||||||
|
{
|
||||||
|
public:
|
||||||
|
static constexpr uint32_t BWT_DATA_MAX_SIZE = 600 * 1024 - 256;
|
||||||
|
static constexpr uint32_t BWT_DATA_STRIDE = (BWT_DATA_MAX_SIZE + 256 + 255) & ~255U;
|
||||||
|
|
||||||
|
XMRIG_DISABLE_COPY_MOVE_DEFAULT(OclAstroBWTRunner)
|
||||||
|
|
||||||
|
OclAstroBWTRunner(size_t index, const OclLaunchData &data);
|
||||||
|
~OclAstroBWTRunner() override;
|
||||||
|
|
||||||
|
inline uint32_t roundSize() const override { return m_batch_size1; }
|
||||||
|
|
||||||
|
// ~0.5% of all hashes are invalid
|
||||||
|
inline uint32_t processedHashes() const override { return static_cast<uint32_t>(m_processedHashes * 0.995); }
|
||||||
|
|
||||||
|
protected:
|
||||||
|
size_t bufferSize() const override;
|
||||||
|
void run(uint32_t nonce, uint32_t *hashOutput) override;
|
||||||
|
void set(const Job &job, uint8_t *blob) override;
|
||||||
|
void build() override;
|
||||||
|
void init() override;
|
||||||
|
|
||||||
|
private:
|
||||||
|
AstroBWT_SHA3InitialKernel* m_sha3_initial_kernel = nullptr;
|
||||||
|
AstroBWT_SHA3Kernel* m_sha3_kernel = nullptr;
|
||||||
|
AstroBWT_Salsa20Kernel* m_salsa20_kernel = nullptr;
|
||||||
|
AstroBWT_MainKernel* m_bwt_kernel = nullptr;
|
||||||
|
AstroBWT_FilterKernel* m_filter_kernel = nullptr;
|
||||||
|
AstroBWT_PrepareBatch2Kernel* m_prepare_batch2_kernel = nullptr;
|
||||||
|
AstroBWT_FindSharesKernel* m_find_shares_kernel = nullptr;
|
||||||
|
|
||||||
|
|
||||||
|
cl_mem m_salsa20_keys = nullptr;
|
||||||
|
cl_mem m_bwt_data = nullptr;
|
||||||
|
cl_mem m_bwt_data_sizes = nullptr;
|
||||||
|
cl_mem m_indices = nullptr;
|
||||||
|
cl_mem m_tmp_indices = nullptr;
|
||||||
|
cl_mem m_filtered_hashes = nullptr;
|
||||||
|
|
||||||
|
uint32_t m_workgroup_size = 0;
|
||||||
|
uint64_t m_bwt_allocation_size = 0;
|
||||||
|
uint64_t m_batch_size1 = 0;
|
||||||
|
uint32_t m_processedHashes = 0;
|
||||||
|
|
||||||
|
uint32_t* m_bwt_data_sizes_host = nullptr;
|
||||||
|
};
|
||||||
|
|
||||||
|
|
||||||
|
} /* namespace xmrig */
|
||||||
|
|
||||||
|
|
||||||
|
#endif // XMRIG_OCLASTROBWTRUNNER_H
|
||||||
|
|
@ -57,6 +57,8 @@ protected:
|
||||||
inline const OclLaunchData &data() const override { return m_data; }
|
inline const OclLaunchData &data() const override { return m_data; }
|
||||||
inline size_t intensity() const override { return m_intensity; }
|
inline size_t intensity() const override { return m_intensity; }
|
||||||
inline size_t threadId() const override { return m_threadId; }
|
inline size_t threadId() const override { return m_threadId; }
|
||||||
|
inline uint32_t roundSize() const override { return m_intensity; }
|
||||||
|
inline uint32_t processedHashes() const override { return m_intensity; }
|
||||||
|
|
||||||
size_t bufferSize() const override;
|
size_t bufferSize() const override;
|
||||||
uint32_t deviceIndex() const override;
|
uint32_t deviceIndex() const override;
|
||||||
|
|
|
||||||
|
|
@ -52,6 +52,10 @@ namespace xmrig {
|
||||||
extern bool ocl_generic_rx_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads);
|
extern bool ocl_generic_rx_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads);
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
|
#ifdef XMRIG_ALGO_ASTROBWT
|
||||||
|
extern bool ocl_generic_astrobwt_generator(const OclDevice& device, const Algorithm& algorithm, OclThreads& threads);
|
||||||
|
#endif
|
||||||
|
|
||||||
#ifdef XMRIG_ALGO_CN_GPU
|
#ifdef XMRIG_ALGO_CN_GPU
|
||||||
extern bool ocl_generic_cn_gpu_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads);
|
extern bool ocl_generic_cn_gpu_generator(const OclDevice &device, const Algorithm &algorithm, OclThreads &threads);
|
||||||
#endif
|
#endif
|
||||||
|
|
@ -64,6 +68,9 @@ static ocl_gen_config_fun generators[] = {
|
||||||
# ifdef XMRIG_ALGO_RANDOMX
|
# ifdef XMRIG_ALGO_RANDOMX
|
||||||
ocl_generic_rx_generator,
|
ocl_generic_rx_generator,
|
||||||
# endif
|
# endif
|
||||||
|
# ifdef XMRIG_ALGO_ASTROBWT
|
||||||
|
ocl_generic_astrobwt_generator,
|
||||||
|
# endif
|
||||||
# ifdef XMRIG_ALGO_CN_GPU
|
# ifdef XMRIG_ALGO_CN_GPU
|
||||||
ocl_generic_cn_gpu_generator,
|
ocl_generic_cn_gpu_generator,
|
||||||
# endif
|
# endif
|
||||||
|
|
|
||||||
Loading…
Reference in a new issue