You can not select more than 25 topics
Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
61 lines
1.4 KiB
61 lines
1.4 KiB
/*
|
|
* Genoil's CUDA mining kernel for Ethereum
|
|
* based on Tim Hughes' opencl kernel.
|
|
* thanks to sp_, trpuvot, djm34, cbuchner for things i took from ccminer.
|
|
*/
|
|
|
|
#include "ethash_cuda_miner_kernel.h"
|
|
#include "ethash_cuda_miner_kernel_globals.h"
|
|
#include "cuda_helper.h"
|
|
|
|
#include "dagger_shuffled.cuh"
|
|
|
|
__global__ void
|
|
__launch_bounds__(896, 1)
|
|
ethash_search(
|
|
volatile uint32_t* g_output,
|
|
uint64_t start_nonce
|
|
)
|
|
{
|
|
uint32_t const gid = blockIdx.x * blockDim.x + threadIdx.x;
|
|
uint64_t hash = compute_hash_shuffle(start_nonce + gid);
|
|
if (cuda_swab64(hash) > d_target) return;
|
|
uint32_t index = atomicInc(const_cast<uint32_t*>(g_output), SEARCH_RESULT_BUFFER_SIZE - 1) + 1;
|
|
g_output[index] = gid;
|
|
__threadfence_system();
|
|
}
|
|
|
|
void run_ethash_search(
|
|
uint32_t blocks,
|
|
uint32_t threads,
|
|
cudaStream_t stream,
|
|
volatile uint32_t* g_output,
|
|
uint64_t start_nonce
|
|
)
|
|
{
|
|
ethash_search <<<blocks, threads, 0, stream >>>(g_output, start_nonce);
|
|
CUDA_SAFE_CALL(cudaGetLastError());
|
|
}
|
|
|
|
void set_constants(
|
|
hash128_t* _dag,
|
|
uint32_t _dag_size
|
|
)
|
|
{
|
|
CUDA_SAFE_CALL(cudaMemcpyToSymbol(d_dag, &_dag, sizeof(hash128_t *)));
|
|
CUDA_SAFE_CALL(cudaMemcpyToSymbol(d_dag_size, &_dag_size, sizeof(uint32_t)));
|
|
}
|
|
|
|
void set_header(
|
|
hash32_t _header
|
|
)
|
|
{
|
|
CUDA_SAFE_CALL(cudaMemcpyToSymbol(d_header, &_header, sizeof(hash32_t)));
|
|
}
|
|
|
|
void set_target(
|
|
uint64_t _target
|
|
)
|
|
{
|
|
CUDA_SAFE_CALL(cudaMemcpyToSymbol(d_target, &_target, sizeof(uint64_t)));
|
|
}
|
|
|