Skip to content

Commit

Permalink
converting to monero v7
Browse files Browse the repository at this point in the history
divide nonce

revise to config file

Updated readme

Updated contacts

bitbucket links -> github

delete to manual ping

skein runtime memory corruption checking

support to force cl_khr_int64_base_atomics

change to stratum protocol (difficulty)

add to config file in cmake

source refactoring (prehash, diff zero count)

change to parse nonce logic

change to difficulty length

delete unnecessary parsing

apply modified target

apply modified format, delete legacy code

set variant as 1

apply job id for each block

refactoring protocol define

chage to icon & remove donate

change to ico

Update README.md
  • Loading branch information
Joowon Yun authored and Joon_ committed Jun 19, 2018
1 parent 4cc922f commit c3a49c4
Show file tree
Hide file tree
Showing 24 changed files with 358 additions and 209 deletions.
1 change: 1 addition & 0 deletions .gitignore
Original file line number Diff line number Diff line change
@@ -1,3 +1,4 @@
/build
/CMakeLists.txt.user
/build2
/.vscode
67 changes: 29 additions & 38 deletions README.md
Original file line number Diff line number Diff line change
@@ -1,50 +1,38 @@
# XMRig AMD
# HYCminer AMD

:warning: **You must update miners to version 2.5 before April 6 due [Monero PoW change](https://getmonero.org/2018/02/11/PoW-change-and-key-reuse.html).**
HYCminer is a high performance Hycon (HYC) OpenCL miner, forked from [XMRig](https://github.com/xmrig/xmrig) release v2.6.1.

[![Github All Releases](https://img.shields.io/github/downloads/xmrig/xmrig-amd/total.svg)](https://github.com/xmrig/xmrig-amd/releases)
[![GitHub release](https://img.shields.io/github/release/xmrig/xmrig-amd/all.svg)](https://github.com/xmrig/xmrig-amd/releases)
[![GitHub Release Date](https://img.shields.io/github/release-date-pre/xmrig/xmrig-amd.svg)](https://github.com/xmrig/xmrig-amd/releases)
[![GitHub license](https://img.shields.io/github/license/xmrig/xmrig-amd.svg)](https://github.com/xmrig/xmrig-amd/blob/master/LICENSE)
[![GitHub stars](https://img.shields.io/github/stars/xmrig/xmrig-amd.svg)](https://github.com/xmrig/xmrig-amd/stargazers)
[![GitHub forks](https://img.shields.io/github/forks/xmrig/xmrig-amd.svg)](https://github.com/xmrig/xmrig-amd/network)

XMRig is high performance Monero (XMR) OpenCL miner, with the official full Windows support.

GPU mining part based on [Wolf9466](https://github.com/OhGodAPet) and [psychocrypt](https://github.com/psychocrypt) code.

* This is the AMD (OpenCL) GPU mining version, there is also a [CPU version](https://github.com/xmrig/xmrig) and [NVIDIA GPU version](https://github.com/xmrig/xmrig-nvidia).
* [Roadmap](https://github.com/xmrig/xmrig/issues/106) for next releases.

:warning: Suggested values for GPU auto configuration can be not optimal or not working, you may need tweak your threads options. Please fell free open an [issue](https://github.com/xmrig/xmrig-amd/issues) if auto configuration suggest wrong values.

<img src="https://i.imgur.com/TFncsi7.png" width="696" >
* This is the **AMD GPU** mining version, there is also a [CPU version](https://github.com/team-hycon/xmrig) and [NVIDIA GPU version](https://github.com/team-hycon/xmrig-nvidia).
* We plan on releasing our roadmap for next releases soon.

#### Table of contents
* [Features](#features)
* [Download](#download)
* [Usage](#usage)
* [Build](https://github.com/xmrig/xmrig-amd/wiki/Build)
* [Donations](#donations)
* [Release checksums](#release-checksums)
* [Algorithm variations](#algorithm-variations)
* [Build](https://github.com/team-hycon/xmrig-amd/wiki/Build)
* [Common Issues](#common-issues)
* [Other information](#other-information)
* [Contacts](#contacts)

## Features
* Currently running Cryptonight v7
* High performance.
* Official Windows support.
* Support for backup (failover) mining server.
* CryptoNight-Lite support for AEON.
* Automatic GPU configuration.
* GPU health monitoring (clocks, power, temperature, fan speed)
* Nicehash support.
* It's open source software.

## Download
* Binary releases: https://github.com/xmrig/xmrig-amd/releases
* Git tree: https://github.com/xmrig/xmrig-amd.git
* Clone with `git clone https://github.com/xmrig/xmrig-amd.git` :hammer: [Build instructions](https://github.com/xmrig/xmrig-amd/wiki/Build).
* Binary releases: TBA
* Git tree: https://github.com/team-hycon/xmrig-amd.git
* Clone with `git clone https://github.com/team-hycon/xmrig-amd.git` :hammer: [Build instructions](https://github.com/team-hycon/xmrig-amd/wiki/Build).

## Usage
Use [config.xmrig.com](https://config.xmrig.com/amd) to generate, edit or share configurations.
See the example `config.json` to configure the miner.

### Command line options
```
Expand Down Expand Up @@ -75,19 +63,22 @@ Use [config.xmrig.com](https://config.xmrig.com/amd) to generate, edit or share
-V, --version output version information and exit
```

## Donations
Default donation 5% (5 minutes in 100 minutes) can be reduced to 1% via command line option `--donate-level`.
Also you can use configuration via config file, default **config.json**. You can load multiple config files and combine it with command line options.

* XMR: `48edfHu7V9Z84YzzMa6fUueoELZ9ZRXq9VetWzYGzKt52XU5xvqgzYnDK9URnRoJMk1j8nLwEVsaSWJ4fhdUyZijBGUicoD`
* BTC: `1P7ujsXeX7GxQwHNnJsRMgAdNkFZmNVqJT`
## Other information
* No HTTP support, only stratum protocol support.
* No TLS support.
* Default donation 5% (5 minutes in 100 minutes) can be reduced to 1% via command line option `--donate-level`.

## Release checksums
### SHA-256
```
3763040fa33980f3cd0c56467044a179e8f7f10193503cbee100a64030ba6469 xmrig-amd-2.7.0-beta-win32.zip/xmrig-amd.exe
01834ceb8b451a0094acdb6cca684e8b92791182a7c3a6deb38a5e37e5236f64 xmrig-amd-2.7.0-beta-win64.zip/xmrig-amd.exe
```
## SHA256 for Windows Binary files
ad7871564a0417d8b0d608bc34e61d7ab71dffcb61fabff1d3efb79173223826: certUtil -hashfile hycon-amd-win-0.0.1-without-mhttpd/xmrig-amd.exe sha256
0dd192e71b6df71a912ab3d614b19de240ebfc7be2c03963d4b268e5599154f3: certUtil -hashfile hycon-amd-win-0.0.1/xmrig-nvidia.exe sha256

## Contacts
* [email protected]
* [reddit](https://www.reddit.com/user/XMRig/)
* [email protected]
* [site](https://www.hycon.io)
* [reddit](https://www.reddit.com/r/HYCON)
* [facebook](https://www.facebook.com/teamHycon)
* [instagram](https://www.instagram.com/teamhycon)
* [medium](https://www.medium.com/@teamhycon)
* [twitter](https://www.twitter.com/teamhycon)
38 changes: 38 additions & 0 deletions config.json
Original file line number Diff line number Diff line change
@@ -0,0 +1,38 @@
{
"algo": "cryptonight",
"background": false,
"colors": true,
"donate-level": 1,
"log-file": null,
"print-time": 60,
"retries": 5,
"retry-pause": 5,
"syslog": false,
"threads": [
{
"index": 0,
"threads": 8,
"intensity": 400,
"worksize" : 8,
"blocks": 12,
"bfactor": 0,
"bsleep": 0,
"affine_to_cpu": false
}
],
"pools": [
{
"url": "127.0.0.1:9081",
"user": "userid",
"pass": "pw",
"keepalive": true,
"nicehash": false,
"variant": 1
}
],
"api": {
"port": 0,
"access-token": null,
"worker-id": null
}
}
Binary file modified res/app.ico
Binary file not shown.
3 changes: 2 additions & 1 deletion src/amd/GpuContext.h
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,7 @@
* Copyright 2016 Jay D Dee <[email protected]>
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
* Copyright 2016-2018 XMRig <https://github.com/xmrig>, <[email protected]>
* Copyright 2018 Team-Hycon <https://github.com/Team-Hycon>
*
* 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
Expand Down Expand Up @@ -97,7 +98,7 @@ struct GpuContext
int computeUnits;
std::string name;

uint32_t Nonce;
uint64_t Nonce;
};


Expand Down
20 changes: 12 additions & 8 deletions src/amd/OclGPU.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,7 @@
* Copyright 2017-2018 XMR-Stak <https://github.com/fireice-uk>, <https://github.com/psychocrypt>
* Copyright 2018 Lee Clagett <https://github.com/vtnerd>
* Copyright 2016-2018 XMRig <https://github.com/xmrig>, <[email protected]>
* Copyright 2018 Team-Hycon <https://github.com/Team-Hycon>
*
* 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
Expand Down Expand Up @@ -40,6 +41,7 @@
#include "core/Config.h"
#include "crypto/CryptoNight_constants.h"
#include "cryptonight.h"
#include "common/net/Protocol.h"


constexpr const char *kSetKernelArgErr = "Error %s when calling clSetKernelArg for kernel %d, argument %d.";
Expand Down Expand Up @@ -113,7 +115,7 @@ size_t InitOpenCLGpu(int index, cl_context opencl_ctx, GpuContext* ctx, const ch
return OCL_ERR_API;
}

ctx->InputBuffer = clCreateBuffer(opencl_ctx, CL_MEM_READ_ONLY, 88, NULL, &ret);
ctx->InputBuffer = clCreateBuffer(opencl_ctx, CL_MEM_READ_ONLY, LEN::BLOB, NULL, &ret);
if (ret != CL_SUCCESS) {
LOG_ERR("Error %s when calling clCreateBuffer to create input buffer.", err_to_str(ret));
return OCL_ERR_API;
Expand Down Expand Up @@ -436,20 +438,17 @@ size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, xmrig::Config *config)
return OCL_ERR_SUCCESS;
}

size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, uint32_t variant)
size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, uint32_t variant, uint32_t moneroNonce)
{
cl_int ret;

if (input_len > 84) {
if (input_len > LEN::BLOB) {
return OCL_ERR_BAD_PARAMS;
}

input[input_len] = 0x01;
memset(input + input_len + 1, 0, 88 - input_len - 1);

size_t numThreads = ctx->rawIntensity;

if ((ret = clEnqueueWriteBuffer(ctx->CommandQueues, ctx->InputBuffer, CL_TRUE, 0, 88, input, 0, NULL, NULL)) != CL_SUCCESS) {
if ((ret = clEnqueueWriteBuffer(ctx->CommandQueues, ctx->InputBuffer, CL_TRUE, 0, LEN::BLOB, input, 0, NULL, NULL)) != CL_SUCCESS) {
LOG_ERR("Error %s when calling clEnqueueWriteBuffer to fill input buffer.", err_to_str(ret));
return OCL_ERR_API;
}
Expand Down Expand Up @@ -497,6 +496,11 @@ size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t tar
LOG_ERR(kSetKernelArgErr, err_to_str(ret), 1 + cn_kernel_offset, 4);
return OCL_ERR_API;
}

if ((ret = clSetKernelArg(ctx->Kernels[1 + cn_kernel_offset], 4, sizeof(cl_uint), &moneroNonce)) != CL_SUCCESS) {
LOG_ERR(kSetKernelArgErr, err_to_str(ret), 1 + cn_kernel_offset, 4);
return OCL_ERR_API;
}
}

// CN3 Kernel
Expand Down Expand Up @@ -666,4 +670,4 @@ size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, uint32_t variant)
ctx->Nonce += (uint32_t) g_intensity;

return OCL_ERR_SUCCESS;
}
}
3 changes: 2 additions & 1 deletion src/amd/OclGPU.h
Original file line number Diff line number Diff line change
Expand Up @@ -31,6 +31,7 @@


#include "amd/GpuContext.h"
#include "common/net/Protocol.h"


namespace xmrig {
Expand All @@ -44,7 +45,7 @@ int getAMDPlatformIdx(xmrig::Config *config);
std::vector<GpuContext> getAMDDevices(int index, xmrig::Config *config);

size_t InitOpenCL(GpuContext* ctx, size_t num_gpus, xmrig::Config *config);
size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, uint32_t variant);
size_t XMRSetJob(GpuContext* ctx, uint8_t* input, size_t input_len, uint64_t target, uint32_t variant, uint32_t moneroNonce);
size_t XMRRunJob(GpuContext* ctx, cl_uint* HashOutput, uint32_t variant);

#endif /* __OCLGPU_H__ */
51 changes: 27 additions & 24 deletions src/amd/opencl/cryptonight.cl
Original file line number Diff line number Diff line change
@@ -1,5 +1,7 @@
R"===(
/*
/* Team-Hycon
* Copyright 2018 Team-Hycon <https://github.com/Team-Hycon>
*
* 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
Expand Down Expand Up @@ -484,15 +486,10 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul

((ulong8 *)State)[0] = vload8(0, input);
State[8] = input[8];
State[9] = input[9];
State[10] = input[10];

((uint *)State)[9] &= 0x00FFFFFFU;
((uint *)State)[9] |= ((get_global_id(0)) & 0xFF) << 24;
((uint *)State)[10] &= 0xFF000000U;
((uint *)State)[10] |= ((get_global_id(0) >> 8));
((uint *)State)[17] = get_global_id(0);

for(int i = 11; i < 25; ++i) State[i] = 0x00UL;
for(int i = 9; i < 25; ++i) State[i] = 0x00UL;
State[9] = 0x01;

// Last bit of padding
State[16] = 0x8000000000000000UL;
Expand Down Expand Up @@ -566,15 +563,16 @@ __kernel void cn0(__global ulong *input, __global uint4 *Scratchpad, __global ul

#define VARIANT1_2(p) ((uint2 *)&(p))[0] ^= tweak1_2_0

#define VARIANT1_INIT() \
#define VARIANT1_INIT(p) \
tweak1_2 = as_uint2(input[4]); \
tweak1_2.s0 >>= 24; \
tweak1_2.s0 |= tweak1_2.s1 << 8; \
tweak1_2.s1 = get_global_id(0); \
tweak1_2 ^= as_uint2(states[24])
tweak1_2.s1 = p; \
tweak1_2 ^= as_uint2(states[24]);


__attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
__kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulong Threads, uint variant, __global ulong *input)
__kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulong Threads, uint variant, __global ulong *input, uint moneroNonce)
{
ulong a[2], b[2];
__local uint AES0[256], AES1[256], AES2[256], AES3[256];
Expand Down Expand Up @@ -614,7 +612,8 @@ __kernel void cn1_monero(__global uint4 *Scratchpad, __global ulong *states, ulo
b[1] = states[3] ^ states[7];

b_x = ((uint4 *)b)[0];
VARIANT1_INIT();

VARIANT1_INIT(moneroNonce);
}

mem_fence(CLK_LOCAL_MEM_FENCE);
Expand Down Expand Up @@ -939,6 +938,9 @@ __kernel void Skein(__global ulong *states, __global uint *BranchBuf, __global u
// The tweak for the output transform is Type = Output with the Final bit set
// T[0] for the output is 8, and I don't know why - should be message size...
ulong t[3] = { 0x00UL, 0x7000000000000000UL, 0x00UL };
t[0] = 0x00UL;
t[1] = 0x7000000000000000UL;
t[2] = 0x00UL;
ulong8 p, m;

for(uint i = 0; i < 4; ++i)
Expand All @@ -960,20 +962,22 @@ __kernel void Skein(__global ulong *states, __global uint *BranchBuf, __global u
t[1] = 0xFF00000000000000UL;
t[2] = t[0] ^ t[1];

p = (ulong8)(0);
m = (ulong8)(0);
const ulong h8 = h.s0 ^ h.s1 ^ h.s2 ^ h.s3 ^ h.s4 ^ h.s5 ^ h.s6 ^ h.s7 ^ SKEIN_KS_PARITY;

p = Skein512Block(p, h, h8, t);
p = Skein512Block(m, h, h8, t);

//vstore8(p, 0, output);

// Note that comparison is equivalent to subtraction - we can't just compare 8 32-bit values
// and expect an accurate result for target > 32-bit without implementing carries
if(p.s3 <= Target)

if(p.s3 < Target)
{
ulong outIdx = atomic_inc(output + 0xFF);
if(outIdx < 0xFF)
output[outIdx] = BranchBuf[idx] + get_global_offset(0);
if(outIdx < 0xFF) {
output[outIdx] = BranchBuf[idx] + get_global_offset(0);
}
}
}
mem_fence(CLK_GLOBAL_MEM_FENCE);
Expand Down Expand Up @@ -1045,7 +1049,8 @@ __kernel void JH(__global ulong *states, __global uint *BranchBuf, __global uint

// Note that comparison is equivalent to subtraction - we can't just compare 8 32-bit values
// and expect an accurate result for target > 32-bit without implementing carries
if(h7l <= Target)

if(h7l < Target)
{
ulong outIdx = atomic_inc(output + 0xFF);
if(outIdx < 0xFF)
Expand Down Expand Up @@ -1122,9 +1127,7 @@ __kernel void Blake(__global ulong *states, __global uint *BranchBuf, __global u

// Note that comparison is equivalent to subtraction - we can't just compare 8 32-bit values
// and expect an accurate result for target > 32-bit without implementing carries
uint2 t = (uint2)(h[6],h[7]);
if( as_ulong(t) <= Target)
{
if(as_ulong((uint2)(h[6],h[7])) < Target) {
ulong outIdx = atomic_inc(output + 0xFF);
if(outIdx < 0xFF)
output[outIdx] = BranchBuf[idx] + get_global_offset(0);
Expand Down Expand Up @@ -1184,7 +1187,7 @@ __kernel void Groestl(__global ulong *states, __global uint *BranchBuf, __global

// Note that comparison is equivalent to subtraction - we can't just compare 8 32-bit values
// and expect an accurate result for target > 32-bit without implementing carries
if(State[7] <= Target)
if(State[7] < Target)
{
ulong outIdx = atomic_inc(output + 0xFF);
if(outIdx < 0xFF)
Expand Down
3 changes: 2 additions & 1 deletion src/api/NetworkState.h
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,7 @@
* Copyright 2014-2016 Wolf9466 <https://github.com/OhGodAPet>
* Copyright 2016 Jay D Dee <[email protected]>
* Copyright 2016-2017 XMRig <[email protected]>
* Copyright 2018 Team-Hycon <https://github.com/Team-Hycon>
*
*
* This program is free software: you can redistribute it and/or modify
Expand Down Expand Up @@ -46,7 +47,7 @@ class NetworkState

char pool[256];
std::array<uint64_t, 10> topDiff { { } };
uint32_t diff;
uint64_t diff;
uint64_t accepted;
uint64_t failures;
uint64_t rejected;
Expand Down
Loading

0 comments on commit c3a49c4

Please sign in to comment.