on-GPU DAG generation

Author: Genoil

CMakeLists.txt

@@ -2,7 +2,7 @@
 cmake_minimum_required(VERSION 2.8.12)
 
 set(PROJECT_VERSION "0.9.41")
-set(GENOIL_VERSION "1.0.8")
+set(GENOIL_VERSION "1.1")
 if (${CMAKE_VERSION} VERSION_GREATER 3.0)
 	cmake_policy(SET CMP0042 OLD) 	# fix MACOSX_RPATH
 	cmake_policy(SET CMP0048 NEW) 	# allow VERSION argument in project()

ethminer/MinerAux.h

@@ -362,6 +362,7 @@ class MinerCLI
 		}
 		else if (arg == "--current-block" && i + 1 < argc)
 			m_currentBlock = stol(argv[++i]);
+		/*
 		else if ((arg == "-R" || arg == "--dag-dir") && i + 1 < argc)
 		{
 			strcpy(s_dagDir, argv[++i]);
@@ -395,6 +396,7 @@ class MinerCLI
 				BOOST_THROW_EXCEPTION(BadArgument());
 			}
 		}
+		*/
 		else if ((arg == "-w" || arg == "--check-pow") && i + 4 < argc)
 		{
 			string m;
@@ -494,13 +496,15 @@ class MinerCLI
 
 	void execute()
 	{
+		/*
 		EthashAux::setDAGDirName(s_dagDir);
 		EthashAux::setDAGEraseMode(m_eraseMode);
 		EthashAux::eraseDAGs();
 		if (m_eraseMode == DAGEraseMode::All)
 		{
 			m_eraseMode = DAGEraseMode::None;
 		}
+		*/
 
 		if (m_shouldListDevices)
 		{
@@ -596,7 +600,7 @@ class MinerCLI
 			<< "	-FS, --failover-stratum <host:port>  Failover stratum server at host:port" << endl
 			<< "    -O, --userpass <username.workername:password> Stratum login credentials" << endl
 			<< "    -FO, --failover-userpass <username.workername:password> Failover stratum login credentials (optional, will use normal credentials when omitted)" << endl
-			<< "    --work-timeout <n> reconnect/failover after n seconds of working on the same (stratum) job. Defaults to 60. Don't set lower than max. avg. block time" << endl
+			<< "    --work-timeout <n> reconnect/failover after n seconds of working on the same (stratum) job. Defaults to 180. Don't set lower than max. avg. block time" << endl
 #endif
 #if ETH_JSONRPC || ETH_STRATUM || !ETH_TRUE
 			<< "    --farm-recheck <n>  Leave n ms between checks for changed work (default: 500). When using stratum, use a high value (i.e. 2000) to get more stable hashrate output" << endl
@@ -615,14 +619,14 @@ class MinerCLI
 #if ETH_JSONRPC || !ETH_TRUE
 			<< "    --phone-home <on/off>  When benchmarking, publish results (default: off)" << endl
 #endif
-			<< "DAG file management:" << endl
-			<< "    -D,--create-dag <number>  Create the DAG in preparation for mining on given block and exit." << endl
-			<< "    -R <s>, --dag-dir <s> Store/Load DAG files in/from the specified directory. Useful for running multiple instances with different configurations." << endl
-			<< "    -E <mode>, --erase-dags <mode> Erase unneeded DAG files. Default is 'none'. Possible values are:" << endl
-			<< "        none  - don't erase DAG files (default)" << endl
-			<< "        old   - erase all DAG files older than current epoch" << endl
-			<< "		bench - like old, but keep epoch 0 for benchmarking" << endl
-			<< "        all   - erase all DAG files. After deleting all files, setting changes to none." << endl
+//			<< "DAG file management:" << endl
+//			<< "    -D,--create-dag <number>  Create the DAG in preparation for mining on given block and exit." << endl
+//			<< "    -R <s>, --dag-dir <s> Store/Load DAG files in/from the specified directory. Useful for running multiple instances with different configurations." << endl
+//			<< "    -E <mode>, --erase-dags <mode> Erase unneeded DAG files. Default is 'none'. Possible values are:" << endl
+//			<< "        none  - don't erase DAG files (default)" << endl
+//			<< "        old   - erase all DAG files older than current epoch" << endl
+//			<< "		bench - like old, but keep epoch 0 for benchmarking" << endl
+//			<< "        all   - erase all DAG files. After deleting all files, setting changes to none." << endl
 			<< "Mining configuration:" << endl
 			<< "    -C,--cpu  When mining, use the CPU." << endl
 			<< "    -G,--opencl  When mining use the GPU via OpenCL." << endl
@@ -691,7 +695,7 @@ class MinerCLI
 		cout << "Benchmarking on platform: " << platformInfo << endl;
 
 		cout << "Preparing DAG for block #" << m_benchmarkBlock << endl;
-		genesis.prep();
+		//genesis.prep();
 
 		genesis.setDifficulty(u256(1) << 63);
 		f.setWork(genesis);
@@ -774,7 +778,7 @@ class MinerCLI
 		cout << "Running mining simulation on platform: " << platformInfo << endl;
 
 		cout << "Preparing DAG for block #" << m_benchmarkBlock << endl;
-		genesis.prep();
+		//genesis.prep();
 
 		genesis.setDifficulty(u256(1) << difficulty);
 		f.setWork(genesis);
@@ -911,6 +915,7 @@ class MinerCLI
 					Json::Value v = prpc->eth_getWork();
 					h256 hh(v[0].asString());
 					h256 newSeedHash(v[1].asString());
+					/*
 					if (current.seedHash != newSeedHash)
 					{
 						minelog << "Grabbing DAG for" << newSeedHash;
@@ -923,6 +928,7 @@ class MinerCLI
 					{
 						EthashAux::computeFull(sha3(newSeedHash), true);
 					}
+					*/
 					if (hh != current.headerHash)
 					{
 						x_current.lock();
@@ -1104,7 +1110,7 @@ class MinerCLI
 
 	/// Benchmarking params
 	bool m_phoneHome = false;
-	unsigned m_benchmarkWarmup = 3;
+	unsigned m_benchmarkWarmup = 15;
 	unsigned m_benchmarkTrial = 3;
 	unsigned m_benchmarkTrials = 5;
 	unsigned m_benchmarkBlock = 0;
@@ -1119,7 +1125,7 @@ class MinerCLI
 	unsigned m_farmRecheckPeriod = 500;
 	unsigned m_defaultStratumFarmRecheckPeriod = 2000;
 	bool m_farmRecheckSet = false;
-	int m_worktimeout = 90;
+	int m_worktimeout = 180;
 	bool m_precompute = true;
 
 #if ETH_STRATUM || !ETH_TRUE

libethash-cl/ethash_cl_miner.cpp

@@ -324,9 +324,11 @@ void ethash_cl_miner::finish()
 		m_queue.finish();
 }
 
+
 bool ethash_cl_miner::init(
-	uint8_t const* _dag,
-	uint64_t _dagSize,
+	ethash_light_t _light, 
+	uint8_t const* _lightData, 
+	uint64_t _lightSize,
 	unsigned _platformId,
 	unsigned _deviceId
 )
@@ -399,12 +401,17 @@ bool ethash_cl_miner::init(
 		if (m_globalWorkSize % s_workgroupSize != 0)
 			m_globalWorkSize = ((m_globalWorkSize / s_workgroupSize) + 1) * s_workgroupSize;
 
+		uint64_t dagSize = ethash_get_datasize(_light->block_number);
+		uint32_t dagSize128 = (unsigned)(dagSize / ETHASH_MIX_BYTES);
+		uint32_t lightSize64 = (unsigned)(_lightSize / sizeof(node));
+
 		// patch source code
 		// note: ETHASH_CL_MINER_KERNEL is simply ethash_cl_miner_kernel.cl compiled
 		// into a byte array by bin2h.cmake. There is no need to load the file by hand in runtime
 		string code(ETHASH_CL_MINER_KERNEL, ETHASH_CL_MINER_KERNEL + ETHASH_CL_MINER_KERNEL_SIZE);
 		addDefinition(code, "GROUP_SIZE", s_workgroupSize);
-		addDefinition(code, "DAG_SIZE", (unsigned)(_dagSize / ETHASH_MIX_BYTES));
+		addDefinition(code, "DAG_SIZE", dagSize128);
+		addDefinition(code, "LIGHT_SIZE", lightSize64);
 		addDefinition(code, "ACCESSES", ETHASH_ACCESSES);
 		addDefinition(code, "MAX_OUTPUTS", c_maxSearchResults);
 		addDefinition(code, "PLATFORM", platformId);
@@ -430,16 +437,19 @@ bool ethash_cl_miner::init(
 		// create buffer for dag
 		try
 		{
-			ETHCL_LOG("Creating one big buffer for the DAG");
-			m_dag = cl::Buffer(m_context, CL_MEM_READ_ONLY, _dagSize);
-			ETHCL_LOG("Loading single big chunk kernels");
+			ETHCL_LOG("Creating cache buffer");
+			m_light = cl::Buffer(m_context, CL_MEM_READ_ONLY, _lightSize);
+			ETHCL_LOG("Creating DAG buffer");
+			m_dag = cl::Buffer(m_context, CL_MEM_READ_ONLY, dagSize);
+			ETHCL_LOG("Loading kernels");
 			m_searchKernel = cl::Kernel(program, "ethash_search");
-			ETHCL_LOG("Mapping one big chunk.");
-			m_queue.enqueueWriteBuffer(m_dag, CL_TRUE, 0, _dagSize, _dag);
+			m_dagKernel = cl::Kernel(program, "ethash_calculate_dag_item");
+			ETHCL_LOG("Writing cache buffer");
+			m_queue.enqueueWriteBuffer(m_light, CL_TRUE, 0, _lightSize, _lightData);
 		}
 		catch (cl::Error const& err)
 		{
-			ETHCL_LOG("Allocating/mapping single buffer failed with: " << err.what() << "(" << err.err() << "). GPU can't allocate the DAG in a single chunk. Bailing.");
+			ETHCL_LOG("Allocating/mapping DAG buffer failed with: " << err.what() << "(" << err.err() << "). GPU can't allocate the DAG in a single chunk. Bailing.");
 			return false;
 		}
 		// create buffer for header
@@ -456,6 +466,28 @@ bool ethash_cl_miner::init(
 			ETHCL_LOG("Creating mining buffer " << i);
 			m_searchBuffer[i] = cl::Buffer(m_context, CL_MEM_WRITE_ONLY, (c_maxSearchResults + 1) * sizeof(uint32_t));
 		}
+
+		ETHCL_LOG("Generating DAG data");
+
+		uint32_t const work = (uint32_t)(dagSize / sizeof(node));
+		//while (work < blocks * threads) blocks /= 2;
+
+		uint32_t fullRuns = work / m_globalWorkSize;
+		uint32_t const restWork = work % m_globalWorkSize;
+		if (restWork > 0) fullRuns++;
+
+		m_dagKernel.setArg(1, m_light);
+		m_dagKernel.setArg(2, m_dag);
+		m_dagKernel.setArg(3, ~0u);
+
+		for (uint32_t i = 0; i < fullRuns; i++)
+		{
+			m_dagKernel.setArg(0, i * m_globalWorkSize);
+			m_queue.enqueueNDRangeKernel(m_dagKernel, cl::NullRange, m_globalWorkSize, s_workgroupSize);
+			m_queue.finish();
+			printf("%.0f%%\n", 100.0f * (float)i / (float)fullRuns);
+		}
+
 	}
 	catch (cl::Error const& err)
 	{

libethash-cl/ethash_cl_miner.h

@@ -51,12 +51,13 @@ class ethash_cl_miner
 		uint64_t _currentBlock
 	);
 
-	bool init(
-		uint8_t const* _dag,
-		uint64_t _dagSize,
-		unsigned _platformId = 0,
-		unsigned _deviceId = 0
-	);
+	bool ethash_cl_miner::init(
+		ethash_light_t _light,
+		uint8_t const* _lightData,
+		uint64_t _lightSize,
+		unsigned _platformId,
+		unsigned _deviceId
+		);
 	void finish();
 	void search(uint8_t const* _header, uint64_t _target, search_hook& _hook);
 
@@ -74,7 +75,9 @@ class ethash_cl_miner
 	cl::Context m_context;
 	cl::CommandQueue m_queue;
 	cl::Kernel m_searchKernel;
+	cl::Kernel m_dagKernel;
 	cl::Buffer m_dag;
+	cl::Buffer m_light;
 	cl::Buffer m_header;
 	cl::Buffer m_searchBuffer[c_bufferCount];
 	unsigned m_globalWorkSize;

libethash-cl/ethash_cl_miner_kernel.cl

@@ -1,11 +1,12 @@
 #define OPENCL_PLATFORM_UNKNOWN 0
 #define OPENCL_PLATFORM_NVIDIA  1
-#define OPENCL_PLATFORM_AMD			2
+#define OPENCL_PLATFORM_AMD		2
 
+#define ETHASH_DATASET_PARENTS 256
+#define NODE_WORDS (64/4)
 
 #define THREADS_PER_HASH (128 / 16)
 #define HASHES_PER_LOOP (GROUP_SIZE / THREADS_PER_HASH)
-
 #define FNV_PRIME	0x01000193
 
 __constant uint2 const Keccak_f1600_RC[24] = {
@@ -176,8 +177,6 @@ static void keccak_f1600_round(uint2* a, uint r)
 
 static void keccak_f1600_no_absorb(uint2* a, uint out_size, uint isolate)
 {
-
-
 	// Originally I unrolled the first and last rounds to interface
 	// better with surrounding code, however I haven't done this
 	// without causing the AMD compiler to blow up the VGPR usage.
@@ -227,6 +226,18 @@ typedef struct
 	ulong ulongs[32 / sizeof(ulong)];
 } hash32_t;
 
+typedef union {
+	uint	 words[64 / sizeof(uint)];
+	uint2	 uint2s[64 / sizeof(uint2)];
+	uint4	 uint4s[64 / sizeof(uint4)];
+} hash64_t;
+
+typedef union {
+	uint	 words[200 / sizeof(uint)];
+	uint2	 uint2s[200 / sizeof(uint2)];
+	uint4	 uint4s[200 / sizeof(uint4)];
+} hash200_t;
+
 typedef struct
 {
 	uint4 uint4s[128 / sizeof(uint4)];
@@ -334,3 +345,35 @@ __kernel void ethash_search(
 		g_output[slot] = gid;
 	}
 }
+
+static void SHA3_512(uint2* s, uint isolate)
+{
+	for (uint i = 8; i != 25; ++i)
+	{
+		s[i] = (uint2){ 0, 0 };
+	}
+	s[8].x = 0x00000001;
+	s[8].y = 0x80000000;
+	keccak_f1600_no_absorb(s, 8, isolate);
+}
+
+__kernel void ethash_calculate_dag_item(uint start, __global hash64_t const* g_light, __global hash64_t * g_dag, uint isolate)
+{
+	uint const node_index = start + get_global_id(0);
+	if (node_index > DAG_SIZE * 2) return;
+
+	hash200_t dag_node;
+	copy(dag_node.uint4s, g_light[node_index % LIGHT_SIZE].uint4s, 4);
+	dag_node.words[0] ^= node_index;
+	SHA3_512(dag_node.uint2s, isolate);
+
+	for (uint i = 0; i != ETHASH_DATASET_PARENTS; ++i) {
+		uint parent_index = fnv(node_index ^ i, dag_node.words[i % NODE_WORDS]) % LIGHT_SIZE;
+
+		for (uint w = 0; w != 4; ++w) {
+			dag_node.uint4s[w] = fnv4(dag_node.uint4s[w], g_light[parent_index].uint4s[w]);
+		}
+	}
+	SHA3_512(dag_node.uint2s, isolate);
+	copy(g_dag[node_index].uint4s, dag_node.uint4s, 4);
+}

libethash-cuda/dagger_shared.cuh

@@ -1,9 +1,5 @@
 #include "ethash_cuda_miner_kernel_globals.h"
 #include "ethash_cuda_miner_kernel.h"
-#include "keccak_u64.cuh"
-#include "fnv.cuh"
-
-#define copy(dst, src, count) for (int i = 0; i != count; ++i) { (dst)[i] = (src)[i]; }
 
 typedef union {
 	uint4	 uint4s[4];

libethash-cuda/dagger_shuffled.cuh

@@ -1,7 +1,5 @@
 #include "ethash_cuda_miner_kernel_globals.h"
 #include "ethash_cuda_miner_kernel.h"
-#include "keccak.cuh"
-#include "fnv.cuh"
 
 #define HASHES_PER_LOOP (GROUP_SIZE / THREADS_PER_HASH)
 

libethash-cuda/ethash_cuda_miner.cpp

@@ -199,7 +199,7 @@ void ethash_cuda_miner::finish()
 	CUDA_SAFE_CALL(cudaDeviceReset());
 }
 
-bool ethash_cuda_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned _deviceId)
+bool ethash_cuda_miner::init(ethash_light_t _light, uint8_t const* _lightData, uint64_t _lightSize, unsigned _deviceId)
 {
 	try
 	{
@@ -224,25 +224,38 @@ bool ethash_cuda_miner::init(uint8_t const* _dag, uint64_t _dagSize, unsigned _d
 		m_search_buf = new volatile uint32_t *[s_numStreams];
 		m_streams = new cudaStream_t[s_numStreams];
 
-		uint32_t dagSize128 = (unsigned)(_dagSize / ETHASH_MIX_BYTES);
+		uint64_t dagSize = ethash_get_datasize(_light->block_number);
+		uint32_t dagSize128   = (unsigned)(dagSize / ETHASH_MIX_BYTES);
+		uint32_t lightSize64 = (unsigned)(_lightSize / sizeof(node));
+
+		// create buffer for cache
+		hash64_t * light;
+		CUDA_SAFE_CALL(cudaMalloc(reinterpret_cast<void**>(&light), _lightSize));
+		// copy dag to CPU.
+		CUDA_SAFE_CALL(cudaMemcpy(reinterpret_cast<void*>(light), _lightData, _lightSize, cudaMemcpyHostToDevice));
 
 		// create buffer for dag
 		hash128_t * dag;
-		CUDA_SAFE_CALL(cudaMalloc(reinterpret_cast<void**>(&dag), _dagSize));
+		CUDA_SAFE_CALL(cudaMalloc(reinterpret_cast<void**>(&dag), dagSize));
 		// copy dag to CPU.
-		CUDA_SAFE_CALL(cudaMemcpy(reinterpret_cast<void*>(dag), _dag, _dagSize, cudaMemcpyHostToDevice));
+		//CUDA_SAFE_CALL(cudaMemcpy(reinterpret_cast<void*>(dag), _dag, _dagSize, cudaMemcpyHostToDevice));
 
+		
 		// create mining buffers
 		for (unsigned i = 0; i != s_numStreams; ++i)
 		{
 			CUDA_SAFE_CALL(cudaMallocHost(&m_search_buf[i], SEARCH_RESULT_BUFFER_SIZE * sizeof(uint32_t)));
 			CUDA_SAFE_CALL(cudaStreamCreate(&m_streams[i]));
 		}
-		set_constants(dag, dagSize128);
+		set_constants(dag, dagSize128, light, lightSize64);
 		memset(&m_current_header, 0, sizeof(hash32_t));
 		m_current_target = 0;
 		m_current_nonce = 0;
 		m_current_index = 0;
+
+		cout << "Generating DAG..." << endl;
+		ethash_generate_dag(dagSize, s_gridSize, s_blockSize, m_streams[0]);
+
 		return true;
 	}
 	catch (runtime_error)