Skip to content
Snippets Groups Projects
Commit 89749c32 authored by psychocrypt's avatar psychocrypt
Browse files

add aeon support to backend nvidia

- add template parameter to kernel to support aeon and xmr
- update auto suggestion
parent 2346c8be
No related branches found
No related tags found
No related merge requests found
......@@ -32,6 +32,7 @@
#include "xmrstak/jconf.hpp"
#include "xmrstak/misc/environment.hpp"
#include "xmrstak/backend/cpu/hwlocMemory.hpp"
#include "../cryptonight.hpp"
#include <assert.h>
#include <cmath>
......@@ -208,7 +209,7 @@ void minethd::work_main()
uint64_t iCount = 0;
cryptonight_ctx* cpu_ctx;
cpu_ctx = cpu::minethd::minethd_alloc_ctx();
cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/);
cn_hash_fun hash_fun = cpu::minethd::func_selector(::jconf::inst()->HaveHardwareAes(), true /*bNoPrefetch*/, ::jconf::inst()->IsCurrencyXMR());
uint32_t iNonce;
globalStates::inst().iConsumeCnt++;
......@@ -218,6 +219,9 @@ void minethd::work_main()
printer::inst()->print_msg(L0, "Setup failed for GPU %d. Exitting.\n", (int)iThreadNo);
std::exit(0);
}
bool useXMR = ::jconf::inst()->GetCurrency().compare("xmr") == 0;
bool useAEON = ::jconf::inst()->GetCurrency().compare("aeon") == 0;
while (bQuit == 0)
{
......@@ -256,7 +260,18 @@ void minethd::work_main()
uint32_t foundCount;
cryptonight_extra_cpu_prepare(&ctx, iNonce);
cryptonight_core_cpu_hash(&ctx);
#ifndef CONF_NO_XMR
if(useXMR)
{
cryptonight_core_cpu_hash<XMR_ITER, XMR_MASK, 19>(&ctx);
}
#endif
#ifndef CONF_NO_AEON
if(useAEON)
{
cryptonight_core_cpu_hash<XMR_ITER, XMR_MASK, 18>(&ctx);
}
#endif
cryptonight_extra_cpu_final(&ctx, iNonce, oWork.iTarget, &foundCount, foundNonce);
for(size_t i = 0; i < foundCount; i++)
......
......@@ -41,8 +41,9 @@ int cuda_get_deviceinfo(nvid_ctx *ctx);
int cryptonight_extra_cpu_init(nvid_ctx *ctx);
void cryptonight_extra_cpu_set_data( nvid_ctx* ctx, const void *data, uint32_t len);
void cryptonight_extra_cpu_prepare(nvid_ctx* ctx, uint32_t startNonce);
void cryptonight_core_cpu_hash(nvid_ctx* ctx);
void cryptonight_extra_cpu_final(nvid_ctx* ctx, uint32_t startNonce, uint64_t target, uint32_t* rescount, uint32_t *resnonce);
}
template<size_t ITERATIONS, size_t THREAD_SHIFT, uint32_t MASK>
void cryptonight_core_cpu_hash(nvid_ctx* ctx);
......@@ -92,6 +92,7 @@ __device__ __forceinline__ void storeGlobal32( T* addr, T const & val )
asm volatile( "st.global.cg.u32 [%0], %1;" : : "l"( addr ), "r"( val ) );
}
template<size_t ITERATIONS, size_t THREAD_SHIFT>
__global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int partidx, uint32_t * __restrict__ long_state, uint32_t * __restrict__ ctx_state, uint32_t * __restrict__ ctx_key1 )
{
__shared__ uint32_t sharedMemory[1024];
......@@ -102,7 +103,7 @@ __global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int parti
const int thread = ( blockDim.x * blockIdx.x + threadIdx.x ) >> 3;
const int sub = ( threadIdx.x & 7 ) << 2;
const int batchsize = 0x80000 >> bfactor;
const int batchsize = ITERATIONS >> bfactor;
const int start = partidx * batchsize;
const int end = start + batchsize;
......@@ -121,13 +122,13 @@ __global__ void cryptonight_core_gpu_phase1( int threads, int bfactor, int parti
else
{
// load previous text data
MEMCPY8( text, &long_state[( (uint64_t) thread << 19 ) + sub + start - 32], 2 );
MEMCPY8( text, &long_state[( (uint64_t) thread << THREAD_SHIFT ) + sub + start - 32], 2 );
}
__syncthreads( );
for ( int i = start; i < end; i += 32 )
{
cn_aes_pseudo_round_mut( sharedMemory, text, key );
MEMCPY8(&long_state[((uint64_t) thread << 19) + (sub + i)], text, 2);
MEMCPY8(&long_state[((uint64_t) thread << THREAD_SHIFT) + (sub + i)], text, 2);
}
}
......@@ -167,6 +168,7 @@ __forceinline__ __device__ uint32_t shuffle(volatile uint32_t* ptr,const uint32_
#ifdef XMR_STAK_THREADS
__launch_bounds__( XMR_STAK_THREADS * 4 )
#endif
template<size_t ITERATIONS, size_t THREAD_SHIFT, uint32_t MASK>
__global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int partidx, uint32_t * d_long_state, uint32_t * d_ctx_a, uint32_t * d_ctx_b )
{
__shared__ uint32_t sharedMemory[1024];
......@@ -190,10 +192,10 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
int i, k;
uint32_t j;
const int batchsize = ITER >> ( 2 + bfactor );
const int batchsize = (ITERATIONS * 2) >> ( 2 + bfactor );
const int start = partidx * batchsize;
const int end = start + batchsize;
uint32_t * long_state = &d_long_state[(IndexType) thread << 19];
uint32_t * long_state = &d_long_state[(IndexType) thread << THREAD_SHIFT];
uint32_t * ctx_a = d_ctx_a + thread * 4;
uint32_t * ctx_b = d_ctx_b + thread * 4;
uint32_t a, d[2];
......@@ -207,7 +209,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
#pragma unroll 2
for ( int x = 0; x < 2; ++x )
{
j = ( ( shuffle(sPtr,sub, a, 0) & 0x1FFFF0 ) >> 2 ) + sub;
j = ( ( shuffle(sPtr,sub, a, 0) & MASK ) >> 2 ) + sub;
const uint32_t x_0 = loadGlobal32<uint32_t>( long_state + j );
const uint32_t x_1 = shuffle(sPtr,sub, x_0, sub + 1);
......@@ -225,8 +227,8 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
//long_state[j] = d[0] ^ d[1];
storeGlobal32( long_state + j, d[0] ^ d[1] );
//MUL_SUM_XOR_DST(c, a, &long_state[((uint32_t *)c)[0] & 0x1FFFF0]);
j = ( ( *t1 & 0x1FFFF0 ) >> 2 ) + sub;
//MUL_SUM_XOR_DST(c, a, &long_state[((uint32_t *)c)[0] & MASK]);
j = ( ( *t1 & MASK ) >> 2 ) + sub;
uint32_t yy[2];
*( (uint64_t*) yy ) = loadGlobal64<uint64_t>( ( (uint64_t *) long_state )+( j >> 1 ) );
......@@ -255,6 +257,7 @@ __global__ void cryptonight_core_gpu_phase2( int threads, int bfactor, int parti
}
}
template<size_t ITERATIONS, size_t THREAD_SHIFT>
__global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int partidx, const uint32_t * __restrict__ long_state, uint32_t * __restrict__ d_ctx_state, uint32_t * __restrict__ d_ctx_key2 )
{
__shared__ uint32_t sharedMemory[1024];
......@@ -265,7 +268,7 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti
int thread = ( blockDim.x * blockIdx.x + threadIdx.x ) >> 3;
int sub = ( threadIdx.x & 7 ) << 2;
const int batchsize = 0x80000 >> bfactor;
const int batchsize = ITERATIONS >> bfactor;
const int start = partidx * batchsize;
const int end = start + batchsize;
......@@ -281,7 +284,7 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti
{
#pragma unroll
for ( int j = 0; j < 4; ++j )
text[j] ^= long_state[((IndexType) thread << 19) + (sub + i + j)];
text[j] ^= long_state[((IndexType) thread << THREAD_SHIFT) + (sub + i + j)];
cn_aes_pseudo_round_mut( sharedMemory, text, key );
}
......@@ -289,7 +292,8 @@ __global__ void cryptonight_core_gpu_phase3( int threads, int bfactor, int parti
MEMCPY8( d_ctx_state + thread * 50 + sub + 16, text, 2 );
}
extern "C" void cryptonight_core_cpu_hash(nvid_ctx* ctx)
template<size_t ITERATIONS, size_t THREAD_SHIFT, uint32_t MASK>
void cryptonight_core_cpu_hash(nvid_ctx* ctx)
{
dim3 grid( ctx->device_blocks );
dim3 block( ctx->device_threads );
......@@ -311,7 +315,7 @@ extern "C" void cryptonight_core_cpu_hash(nvid_ctx* ctx)
for ( int i = 0; i < partcountOneThree; i++ )
{
CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase1<<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads,
CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase1<ITERATIONS,THREAD_SHIFT><<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads,
bfactorOneThree, i,
ctx->d_long_state, ctx->d_ctx_state, ctx->d_ctx_key1 ));
......@@ -321,7 +325,7 @@ extern "C" void cryptonight_core_cpu_hash(nvid_ctx* ctx)
for ( int i = 0; i < partcount; i++ )
{
CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase2<<<
CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase2<ITERATIONS,THREAD_SHIFT,MASK><<<
grid,
block4,
block4.x * sizeof(uint32_t) * static_cast< int >( ctx->device_arch[0] < 3 )
......@@ -339,7 +343,7 @@ extern "C" void cryptonight_core_cpu_hash(nvid_ctx* ctx)
for ( int i = 0; i < partcountOneThree; i++ )
{
CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase3<<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads,
CUDA_CHECK_KERNEL(ctx->device_id, cryptonight_core_gpu_phase3<ITERATIONS,THREAD_SHIFT><<< grid, block8 >>>( ctx->device_blocks*ctx->device_threads,
bfactorOneThree, i,
ctx->d_long_state,
ctx->d_ctx_state, ctx->d_ctx_key2 ));
......
......@@ -5,6 +5,7 @@
#include <cuda_runtime.h>
#include <device_functions.hpp>
#include <algorithm>
#include "../../../jconf.hpp"
#ifdef __CUDACC__
__constant__
......@@ -188,8 +189,18 @@ extern "C" int cryptonight_extra_cpu_init(nvid_ctx* ctx)
cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
size_t hashMemSize;
if(::jconf::inst()->IsCurrencyXMR())
{
hashMemSize = XMR_MEMORY;
}
else
{
hashMemSize = AEON_MEMORY;
}
size_t wsize = ctx->device_blocks * ctx->device_threads;
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_long_state, (size_t)MEMORY * wsize));
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_long_state, hashMemSize * wsize));
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_state, 50 * sizeof(uint32_t) * wsize));
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_key1, 40 * sizeof(uint32_t) * wsize));
CUDA_CHECK(ctx->device_id, cudaMalloc(&ctx->d_ctx_key2, 40 * sizeof(uint32_t) * wsize));
......@@ -343,13 +354,23 @@ extern "C" int cuda_get_deviceinfo(nvid_ctx* ctx)
ctx->total_device_memory = totalMemory;
ctx->free_device_memory = freeMemory;
size_t hashMemSize;
if(::jconf::inst()->IsCurrencyXMR())
{
hashMemSize = XMR_MEMORY;
}
else
{
hashMemSize = AEON_MEMORY;
}
// keep 64MiB memory free (value is randomly chosen)
// 200byte are meta data memory (result nonce, ...)
size_t availableMem = freeMemory - (64u * 1024 * 1024) - 200u;
size_t limitedMemory = std::min(availableMem, maxMemUsage);
// up to 920bytes extra memory is used per thread for some kernel (lmem/local memory)
// 680bytes are extra meta data memory per hash
size_t perThread = size_t(MEMORY) + 740u + 680u;
size_t perThread = hashMemSize + 740u + 680u;
size_t max_intensity = limitedMemory / perThread;
ctx->device_threads = max_intensity / ctx->device_blocks;
// use only odd number of threads
......
#pragma once
#include "../../cryptonight.hpp"
#ifdef __INTELLISENSE__
#define __CUDA_ARCH__ 520
/* avoid red underlining */
......@@ -18,8 +20,6 @@ struct uint3 blockDim;
#define __shfl(a,b,c) 1
#endif
#define MEMORY (1 << 21) // 2 MiB / 2097152 B
#define ITER (1 << 20) // 1048576
#define AES_BLOCK_SIZE 16
#define AES_KEY_SIZE 32
#define INIT_SIZE_BLK 8
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment