GpuKang.cpp

// This file is a part of RCKangaroo software
// (c) 2024, RetiredCoder (RC)
// License: GPLv3, see "LICENSE.TXT" file
// https://github.com/RetiredC


#include <iostream>
#include "cuda_runtime.h"
#include "cuda.h"

#include "GpuKang.h"
#include "tagtrace.h"
#include "p140ops.h"
#include "turbo_cubin.h"

cudaError_t cuSetGpuParams(TKparams Kparams, u64* _jmp2_table);
extern "C" cudaError_t cuUploadTagTrace(const u32* t1, const u32* t2, const u32* t3, u32 enabled);
void CallGpuKernelGen(TKparams Kparams);
void CallGpuKernelABC(TKparams Kparams);
void CallGpuKernelABC_Stream(TKparams Kparams, cudaStream_t stream);
void CallGpuKernelA_Stream(TKparams Kparams, cudaStream_t stream);
void CallGpuKernelBC_Stream(TKparams Kparams, cudaStream_t stream);
void AddPointsToList(u32* data, int cnt, u64 ops_cnt);
extern bool gGenMode; //tames generation mode
extern u32 gTotalErrors;

RCGpuKang::RCGpuKang()
{
	StopFlag = false;
	PauseFlag = false;
	Failed = false;
	TurboAsmWanted = false;
	TurboAsmReady = false;
	sm_inv_cnt = 0;
	Is5xxx = false;
	ComputeMajor = 0;
	ComputeMinor = 0;
	memset(&TurboStatus, 0, sizeof(TurboStatus));
	IsOldGpu = true;
	OverrideBlockSize = 0;
	OverrideGroupCnt = 0;
	RestoredFromCheckpoint = false;
	SavedKangBuf = NULL;
	SavedL1S2 = NULL;
	RndPnts = NULL;
	DPs_out = NULL;
	DPs_out_B = NULL;
	DualLaneOn = true;
	DualLaneReady = false;
	StreamHop = NULL;
	StreamCopy = NULL;
	StreamTrioOn = true;
	StreamTrioReady = false;
	StreamDpPack = NULL;
	EvHopDone = NULL;
	EvPackDone = NULL;
	L2PinOn = true;
	WarpDPOn = true;
	GroupSweepOn = false;
	GroupSweepLastMs = 0;
	GroupSweepSuggest = 0;
	SobolOn = false;
	GravityOn = false;
	HerdSkewOn = false;
	TameFrac = 1.0 / 3.0;
	memset(&GravCluster, 0, sizeof(GravCluster));
	HerdSkew_Init(&SkewState);
	KangCnt = 0;
	memset(dbg, 0, sizeof(dbg));
	memset(SpeedStats, 0, sizeof(SpeedStats));
	cur_stats_ind = 0;
	memset(&Kparams, 0, sizeof(Kparams));
}

void RCGpuKang::SetDualLane(bool on)
{
	DualLaneOn = on;
}

void RCGpuKang::SetStreamTrio(bool on) { StreamTrioOn = on; }
void RCGpuKang::SetL2Pin(bool on) { L2PinOn = on; }
void RCGpuKang::SetWarpDP(bool on) { WarpDPOn = on; }
void RCGpuKang::SetGroupSweep(bool on) { GroupSweepOn = on; }
void RCGpuKang::SetSobol(bool on) { SobolOn = on; }
void RCGpuKang::SetGravity(bool on) { GravityOn = on; }
void RCGpuKang::SetHerdSkew(bool on) { HerdSkewOn = on; }
void RCGpuKang::SetTameFrac(double f)
{
	if (f < 0.15) f = 0.15;
	if (f > 0.55) f = 0.55;
	TameFrac = f;
}
void RCGpuKang::ObserveGravityNearMiss(const u8* x12, int bump)
{
	Gravity_ObserveNearMiss(&GravCluster, x12, bump);
}
void RCGpuKang::ObserveHerdSkew(bool candidate)
{
	if (HerdSkewOn)
	{
		HerdSkew_Observe(&SkewState, candidate);
		TameFrac = HerdSkew_TameFrac(&SkewState);
	}
}

int RCGpuKang::CalcKangCnt()
{
	Kparams.BlockCnt = mpCnt;
	Kparams.BlockSize = OverrideBlockSize ? OverrideBlockSize : (IsOldGpu ? 512 : 256);
	Kparams.GroupCnt = OverrideGroupCnt ? OverrideGroupCnt : (IsOldGpu ? 64 : 24);
	return Kparams.BlockSize * Kparams.GroupCnt * Kparams.BlockCnt;
}

//executes in main thread
bool RCGpuKang::Prepare(EcPoint _PntToSolve, int _Range, int _DP, EcJMP* _EcJumps1, EcJMP* _EcJumps2, EcJMP* _EcJumps3)
{
	PntToSolve = _PntToSolve;
	Range = _Range;
	DP = _DP;
	EcJumps1 = _EcJumps1;
	EcJumps2 = _EcJumps2;
	EcJumps3 = _EcJumps3;
	StopFlag = false;
	PauseFlag = false;
	Failed = false;
	u64 total_mem = 0;
	memset(dbg, 0, sizeof(dbg));
	memset(SpeedStats, 0, sizeof(SpeedStats));
	cur_stats_ind = 0;

	cudaError_t err;
	err = cudaSetDevice(CudaIndex);
	if (err != cudaSuccess)
		return false;

	// Report / load v4 turbo cubin policy for this GPU
	TurboAsmReady = false;
	if (TurboSelectCubin(ComputeMajor, ComputeMinor, &TurboStatus))
	{
		TurboAsmWanted = true;
		if (TurboLoadCubin(CubinCaller, &TurboStatus))
		{
			// Full Asm execute still needs upstream TKparams dual-path (see ROADMAP).
			// Cubin is loaded so Ada/Blackwell users have the module resident; CUDA path
			// remains the active hop kernel until dual-path lands.
			TurboAsmReady = false;
			printf("GPU %d: turbo ASM cubin LOADED (%s) — execute dual-path pending; CUDA StreamTrio active\r\n",
				CudaIndex, TurboStatus.path);
		}
		else
		{
			printf("GPU %d: turbo ASM cubin select OK but load failed: %s\r\n",
				CudaIndex, TurboStatus.detail);
		}
	}
	else
	{
		printf("GPU %d: %s\r\n", CudaIndex, TurboStatus.detail);
		printf("GPU %d: turbo kernel path ACTIVE (sm_%d%d CUDA fatbin + StreamTrio/DualLane, no host-shim)\r\n",
			CudaIndex, ComputeMajor, ComputeMinor);
	}

	Kparams.BlockCnt = mpCnt;
	Kparams.BlockSize = OverrideBlockSize ? OverrideBlockSize : (IsOldGpu ? BLOCK_SIZE_OLD_GPU : BLOCK_SIZE_NEW_GPU);
	Kparams.GroupCnt = OverrideGroupCnt ? OverrideGroupCnt : (IsOldGpu ? PNT_GROUP_OLD_GPU : PNT_GROUP_NEW_GPU);
	KangCnt = Kparams.BlockSize * Kparams.GroupCnt * Kparams.BlockCnt;
	Kparams.KangCnt = KangCnt;
	Kparams.DP = DP;
	Kparams.KernelA_LDS_Size = 64 * JMP_CNT + 16 * Kparams.BlockSize;
	Kparams.KernelB_LDS_Size = 64 * JMP_CNT;
	Kparams.KernelC_LDS_Size = 96 * JMP_CNT;
	Kparams.IsGenMode = gGenMode;
	Kparams.WarpDP = WarpDPOn ? 1u : 0u;
	Kparams.L2Pin = L2PinOn ? 1u : 0u;
	Kparams.TagTrace = TagTrace_IsEnabled() ? 1u : 0u;

//allocate gpu mem
	u64 size;
	if (!IsOldGpu)
	{
		//L2	
		int L2size = Kparams.KangCnt * (3 * 32);
		total_mem += L2size;
		err = cudaMalloc((void**)&Kparams.L2, L2size);
		if (err != cudaSuccess)
		{
			printf("GPU %d, Allocate L2 memory failed: %s\n", CudaIndex, cudaGetErrorString(err));
			return false;
		}
		size = L2size;
		if (size > (u64)persistingL2CacheMaxSize)
			size = persistingL2CacheMaxSize;
		err = cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, size);
		cudaStreamAttrValue stream_attribute;                                                   
		stream_attribute.accessPolicyWindow.base_ptr = Kparams.L2;
		stream_attribute.accessPolicyWindow.num_bytes = size;										
		stream_attribute.accessPolicyWindow.hitRatio = 1.0;                                     
		stream_attribute.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;             
		stream_attribute.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;  	
		err = cudaStreamSetAttribute(NULL, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);
		if (err != cudaSuccess)
		{
			printf("GPU %d, cudaStreamSetAttribute failed: %s\n", CudaIndex, cudaGetErrorString(err));
			return false;
		}
	}
	size = MAX_DP_CNT * GPU_DP_SIZE + 16;
	total_mem += size;
	err = cudaMalloc((void**)&Kparams.DPs_out, size);
	if (err != cudaSuccess)
	{
		printf("GPU %d Allocate GpuOut memory failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}

	size = KangCnt * 96;
	total_mem += size;
	err = cudaMalloc((void**)&Kparams.Kangs, size);
	if (err != cudaSuccess)
	{
		printf("GPU %d Allocate pKangs memory failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}

	total_mem += JMP_CNT * 96;
	err = cudaMalloc((void**)&Kparams.Jumps1, JMP_CNT * 96);
	if (err != cudaSuccess)
	{
		printf("GPU %d Allocate Jumps1 memory failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}

	total_mem += JMP_CNT * 96;
	err = cudaMalloc((void**)&Kparams.Jumps2, JMP_CNT * 96);
	if (err != cudaSuccess)
	{
		printf("GPU %d Allocate Jumps1 memory failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}

	total_mem += JMP_CNT * 96;
	err = cudaMalloc((void**)&Kparams.Jumps3, JMP_CNT * 96);
	if (err != cudaSuccess)
	{
		printf("GPU %d Allocate Jumps3 memory failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}

	size = 2 * (u64)KangCnt * STEP_CNT;
	total_mem += size;
	err = cudaMalloc((void**)&Kparams.JumpsList, size);
	if (err != cudaSuccess)
	{
		printf("GPU %d Allocate JumpsList memory failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}

	size = (u64)KangCnt * (16 * DPTABLE_MAX_CNT + sizeof(u32));
	total_mem += size;
	err = cudaMalloc((void**)&Kparams.DPTable, size);
	if (err != cudaSuccess)
	{
		printf("GPU %d Allocate DPTable memory failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}

	size = mpCnt * Kparams.BlockSize * sizeof(u64);
	total_mem += size;
	err = cudaMalloc((void**)&Kparams.L1S2, size);
	if (err != cudaSuccess)
	{
		printf("GPU %d Allocate L1S2 memory failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}

	size = (u64)KangCnt * MD_LEN * (2 * 32);
	total_mem += size;
	err = cudaMalloc((void**)&Kparams.LastPnts, size);
	if (err != cudaSuccess)
	{
		printf("GPU %d Allocate LastPnts memory failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}

	size = (u64)KangCnt * MD_LEN * sizeof(u64);
	total_mem += size;
	err = cudaMalloc((void**)&Kparams.LoopTable, size);
	if (err != cudaSuccess)
	{
		printf("GPU %d Allocate LastPnts memory failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}

	total_mem += 1024;
	err = cudaMalloc((void**)&Kparams.dbg_buf, 1024);
	if (err != cudaSuccess)
	{
		printf("GPU %d Allocate dbg_buf memory failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}

	size = sizeof(u32) * KangCnt + 8;
	total_mem += size;
	err = cudaMalloc((void**)&Kparams.LoopedKangs, size);
	if (err != cudaSuccess)
	{
		printf("GPU %d Allocate LoopedKangs memory failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}

	DPs_out = (u32*)malloc(MAX_DP_CNT * GPU_DP_SIZE);
	DPs_out_B = (u32*)malloc(MAX_DP_CNT * GPU_DP_SIZE);

	// Streams: prefer StreamTrio (3), else DualLane (2), else classic
	DualLaneReady = false;
	StreamTrioReady = false;
	if (StreamTrioOn)
	{
		cudaError_t e1 = cudaStreamCreate(&StreamHop);
		cudaError_t e2 = cudaStreamCreate(&StreamDpPack);
		cudaError_t e3 = cudaStreamCreate(&StreamCopy);
		cudaError_t e4 = cudaEventCreateWithFlags(&EvHopDone, cudaEventDisableTiming);
		cudaError_t e5 = cudaEventCreateWithFlags(&EvPackDone, cudaEventDisableTiming);
		if (e1 == cudaSuccess && e2 == cudaSuccess && e3 == cudaSuccess &&
			e4 == cudaSuccess && e5 == cudaSuccess && DPs_out_B)
		{
			StreamTrioReady = true;
			DualLaneReady = true; // trio subsumes dual
			printf("GPU %d: StreamTrio ON (unified ABC hop stream + host-copy overlap)\r\n", CudaIndex);
		}
		else
		{
			if (e1 == cudaSuccess) cudaStreamDestroy(StreamHop);
			if (e2 == cudaSuccess) cudaStreamDestroy(StreamDpPack);
			if (e3 == cudaSuccess) cudaStreamDestroy(StreamCopy);
			if (e4 == cudaSuccess) cudaEventDestroy(EvHopDone);
			if (e5 == cudaSuccess) cudaEventDestroy(EvPackDone);
			StreamHop = StreamCopy = StreamDpPack = NULL;
			EvHopDone = EvPackDone = NULL;
			printf("GPU %d: StreamTrio unavailable — trying DualLane\r\n", CudaIndex);
		}
	}
	if (!StreamTrioReady && DualLaneOn)
	{
		cudaError_t e1 = cudaStreamCreate(&StreamHop);
		cudaError_t e2 = cudaStreamCreate(&StreamCopy);
		if (e1 == cudaSuccess && e2 == cudaSuccess && DPs_out_B)
		{
			DualLaneReady = true;
			printf("GPU %d: DualLane ON (hop/copy streams + double-buffer DP download)\r\n", CudaIndex);
		}
		else
		{
			if (e1 == cudaSuccess) cudaStreamDestroy(StreamHop);
			if (e2 == cudaSuccess) cudaStreamDestroy(StreamCopy);
			StreamHop = StreamCopy = NULL;
			printf("GPU %d: DualLane unavailable (%s) — classic path\r\n",
				CudaIndex, cudaGetErrorString(e1 != cudaSuccess ? e1 : e2));
		}
	}

//jmp1
	u64* buf = (u64*)malloc(JMP_CNT * 96);
	for (int i = 0; i < JMP_CNT; i++)
	{
		memcpy(buf + i * 12, EcJumps1[i].p.x.data, 32);
		memcpy(buf + i * 12 + 4, EcJumps1[i].p.y.data, 32);
		memcpy(buf + i * 12 + 8, EcJumps1[i].dist.data, 32);
	}
	err = cudaMemcpy(Kparams.Jumps1, buf, JMP_CNT * 96, cudaMemcpyHostToDevice);
	if (err != cudaSuccess)
	{
		printf("GPU %d, cudaMemcpy Jumps1 failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}
	free(buf);
//jmp2
	buf = (u64*)malloc(JMP_CNT * 96);
	u64* jmp2_table = (u64*)malloc(JMP_CNT * 64);
	for (int i = 0; i < JMP_CNT; i++)
	{
		memcpy(buf + i * 12, EcJumps2[i].p.x.data, 32);
		memcpy(jmp2_table + i * 8, EcJumps2[i].p.x.data, 32);
		memcpy(buf + i * 12 + 4, EcJumps2[i].p.y.data, 32);
		memcpy(jmp2_table + i * 8 + 4, EcJumps2[i].p.y.data, 32);
		memcpy(buf + i * 12 + 8, EcJumps2[i].dist.data, 32);
	}
	err = cudaMemcpy(Kparams.Jumps2, buf, JMP_CNT * 96, cudaMemcpyHostToDevice);
	if (err != cudaSuccess)
	{
		printf("GPU %d, cudaMemcpy Jumps2 failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}
	free(buf);

	err = cuSetGpuParams(Kparams, jmp2_table);
	if (err != cudaSuccess)
	{
		free(jmp2_table);
		printf("GPU %d, cuSetGpuParams failed: %s!\r\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}
	free(jmp2_table);

	// TagTrace: upload jump tags into __constant__ (full GPU implant)
	if (TagTrace_IsEnabled())
	{
		if (!TagTrace_UploadToGpu())
			printf("GPU %d: TagTrace constant upload failed — host filter remains\r\n", CudaIndex);
		else
			printf("GPU %d: JumpDNA/TagTrace tables in device memory (Jumps1/2/3 + tag constants)\r\n", CudaIndex);
	}

	// L2Pin: also pin jump table + a small DP-mask proxy window into persisting L2 (Ampere+)
	if (L2PinOn && persistingL2CacheMaxSize > 0)
	{
		u64 pin_bytes = JMP_CNT * 96;
		if (pin_bytes > (u64)persistingL2CacheMaxSize / 2)
			pin_bytes = (u64)persistingL2CacheMaxSize / 2;
		cudaStreamAttrValue attr;
		attr.accessPolicyWindow.base_ptr = Kparams.Jumps1;
		attr.accessPolicyWindow.num_bytes = pin_bytes;
		attr.accessPolicyWindow.hitRatio = 1.0;
		attr.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;
		attr.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;
		cudaStream_t pin_stream = StreamHop ? StreamHop : NULL;
		err = cudaStreamSetAttribute(pin_stream, cudaStreamAttributeAccessPolicyWindow, &attr);
		if (err == cudaSuccess)
			printf("GPU %d: L2Pin ON (jump tables + DP window in persisting L2, %llu KB)\r\n",
				CudaIndex, pin_bytes / 1024);
		else
			printf("GPU %d: L2Pin requested but attribute failed (%s) — continuing\r\n",
				CudaIndex, cudaGetErrorString(err));
	}
	if (WarpDPOn)
		printf("GPU %d: WarpDP ON (warp ballot coalesce before DPs_out atomics)\r\n", CudaIndex);
	if (GroupSweepOn)
		printf("GPU %d: GroupSweep ON (live PNT_GROUP advice; compile-time group is fatbin-fixed)\r\n", CudaIndex);
//jmp3
	buf = (u64*)malloc(JMP_CNT * 96);
	for (int i = 0; i < JMP_CNT; i++)
	{
		memcpy(buf + i * 12, EcJumps3[i].p.x.data, 32);
		memcpy(buf + i * 12 + 4, EcJumps3[i].p.y.data, 32);
		memcpy(buf + i * 12 + 8, EcJumps3[i].dist.data, 32);
	}
	err = cudaMemcpy(Kparams.Jumps3, buf, JMP_CNT * 96, cudaMemcpyHostToDevice);
	if (err != cudaSuccess)
	{
		printf("GPU %d, cudaMemcpy Jumps3 failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}
	free(buf);

	size_t free_b = 0, total_b = 0;
	if (cudaMemGetInfo(&free_b, &total_b) == cudaSuccess)
	{
		double free_gb = (double)free_b / (1024.0 * 1024.0 * 1024.0);
		if (free_gb < 0.75)
		{
			printf("GPU %d: WARNING only %.2f GB VRAM free after alloc — "
				"close other CUDA apps (keyhunt/games) or hops may stall at 0 Mk/s.\r\n",
				CudaIndex, free_gb);
		}
	}

	printf("GPU %d: allocated %llu MB, %d kangaroos. OldGpuMode: %s (block=%u group=%u)\r\n",
		CudaIndex, total_mem / (1024 * 1024), KangCnt, IsOldGpu ? "Yes" : "No",
		Kparams.BlockSize, Kparams.GroupCnt);
	return true;
}

void RCGpuKang::Release()
{
	if (StreamTrioReady || DualLaneReady)
	{
		if (StreamHop) cudaStreamDestroy(StreamHop);
		if (StreamDpPack) cudaStreamDestroy(StreamDpPack);
		if (StreamCopy) cudaStreamDestroy(StreamCopy);
		if (EvHopDone) cudaEventDestroy(EvHopDone);
		if (EvPackDone) cudaEventDestroy(EvPackDone);
		StreamHop = StreamCopy = StreamDpPack = NULL;
		EvHopDone = EvPackDone = NULL;
		StreamTrioReady = false;
		DualLaneReady = false;
	}
	if (RndPnts)
	{
		free(RndPnts);
		RndPnts = NULL;
	}
	if (DPs_out)
	{
		free(DPs_out);
		DPs_out = NULL;
	}
	if (DPs_out_B)
	{
		free(DPs_out_B);
		DPs_out_B = NULL;
	}
	if (SavedKangBuf)
	{
		free(SavedKangBuf);
		SavedKangBuf = NULL;
	}
	if (SavedL1S2)
	{
		free(SavedL1S2);
		SavedL1S2 = NULL;
	}
	cudaFree(Kparams.LoopedKangs);
	cudaFree(Kparams.dbg_buf);
	cudaFree(Kparams.LoopTable);
	cudaFree(Kparams.LastPnts);
	cudaFree(Kparams.L1S2);
	cudaFree(Kparams.DPTable);
	cudaFree(Kparams.JumpsList);
	cudaFree(Kparams.Jumps3);
	cudaFree(Kparams.Jumps2);
	cudaFree(Kparams.Jumps1);
	cudaFree(Kparams.Kangs);
	cudaFree(Kparams.DPs_out);
	if (!IsOldGpu)
		cudaFree(Kparams.L2);
}

void RCGpuKang::Stop()
{
	StopFlag = true;
}

void RCGpuKang::RequestPause()
{
	PauseFlag = true;
}

void RCGpuKang::ClearPause()
{
	PauseFlag = false;
}

bool RCGpuKang::IsPaused() const
{
	return PauseFlag;
}

void RCGpuKang::SetRestoredState(u8* kang_buf, u32* l1s2_buf)
{
	if (SavedKangBuf)
		free(SavedKangBuf);
	if (SavedL1S2)
		free(SavedL1S2);
	SavedKangBuf = kang_buf;
	SavedL1S2 = l1s2_buf;
	RestoredFromCheckpoint = (kang_buf != NULL);
}

void RCGpuKang::GenerateRndDistances()
{
	if (HerdSkewOn)
		TameFrac = HerdSkew_TameFrac(&SkewState);
	int tame_cnt = (int)(KangCnt * TameFrac);
	if (tame_cnt < 1) tame_cnt = 1;
	if (tame_cnt > KangCnt - 2) tame_cnt = KangCnt - 2;
	int wild1_end = tame_cnt + (KangCnt - tame_cnt) / 2;

	for (int i = 0; i < KangCnt; i++)
	{
		EcInt d;
		if (i < tame_cnt)
			d.RndBits(Range - 4); //TAME kangs
		else
		{
			if (SobolOn)
				Sobol_FillWildDist(&d, Range, (u32)i, (u64)CudaIndex * 0x100000001b3ULL);
			else if (GravityOn && GravCluster.hits > 0)
				Gravity_FillWildDist(&d, Range, &GravCluster, (u64)i * 0x9e3779b97f4a7c15ULL);
			else
			{
				d.RndBits(Range - 1);
				d.data[0] &= 0xFFFFFFFFFFFFFFFEULL; //must be even
			}
		}
		memcpy(RndPnts[i].priv, d.data, 24);
	}
	if (HerdSkewOn || SobolOn || GravityOn)
		printf("GPU %d: herd starts tame_frac=%.2f sobol=%d gravity=%d (hits=%d)\r\n",
			CudaIndex, TameFrac, (int)SobolOn, (int)GravityOn, GravCluster.hits);
}

bool RCGpuKang::Start()
{
	if (Failed)
		return false;

	cudaError_t err;
	err = cudaSetDevice(CudaIndex);
	if (err != cudaSuccess)
		return false;

	HalfRange.Set(1);
	HalfRange.ShiftLeft(Range - 1);
	PntHalfRange = ec.MultiplyG(HalfRange);
	NegPntHalfRange = PntHalfRange;
	NegPntHalfRange.y.NegModP();

	PntA = ec.AddPoints(PntToSolve, NegPntHalfRange);
	PntB = PntA;
	PntB.y.NegModP();

	RndPnts = (TPointPriv*)malloc(KangCnt * 96);

	if (RestoredFromCheckpoint && SavedKangBuf)
	{
		err = cudaMemcpy(Kparams.Kangs, SavedKangBuf, KangCnt * 96, cudaMemcpyHostToDevice);
		if (err != cudaSuccess)
		{
			printf("GPU %d, restore Kangs failed: %s\n", CudaIndex, cudaGetErrorString(err));
			return false;
		}
		if (SavedL1S2)
		{
			err = cudaMemcpy(Kparams.L1S2, SavedL1S2, mpCnt * Kparams.BlockSize * sizeof(u64), cudaMemcpyHostToDevice);
			if (err != cudaSuccess)
				return false;
		}
		else
			cudaMemset(Kparams.L1S2, 0, mpCnt * Kparams.BlockSize * 8);
		cudaMemset(Kparams.dbg_buf, 0, 1024);
		cudaMemset(Kparams.LoopTable, 0, KangCnt * MD_LEN * sizeof(u64));
		printf("GPU %d: restored kangaroo state from checkpoint\r\n", CudaIndex);
		return true;
	}

	GenerateRndDistances();

	int tame_cnt = (int)(KangCnt * TameFrac);
	if (tame_cnt < 1) tame_cnt = 1;
	if (tame_cnt > KangCnt - 2) tame_cnt = KangCnt - 2;
	int wild1_end = tame_cnt + (KangCnt - tame_cnt) / 2;

	u8 buf_PntA[64], buf_PntB[64];
	PntA.SaveToBuffer64(buf_PntA);
	PntB.SaveToBuffer64(buf_PntB);
	for (int i = 0; i < KangCnt; i++)
	{
		if (i < tame_cnt)
			memset(RndPnts[i].x, 0, 64);
		else if (i < wild1_end)
			memcpy(RndPnts[i].x, buf_PntA, 64);
		else
			memcpy(RndPnts[i].x, buf_PntB, 64);
	}
	err = cudaMemcpy(Kparams.Kangs, RndPnts, KangCnt * 96, cudaMemcpyHostToDevice);
	if (err != cudaSuccess)
	{
		printf("GPU %d, cudaMemcpy failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}
	printf("GPU %d: KernelGen start points (range %d, %d kangs)...\r\n", CudaIndex, Range, KangCnt);
	fflush(stdout);
	u64 t_gen0 = GetTickCount64();
	CallGpuKernelGen(Kparams);
	err = cudaDeviceSynchronize();
	if (err != cudaSuccess)
	{
		printf("GPU %d, KernelGen failed: %s\n", CudaIndex, cudaGetErrorString(err));
		return false;
	}
	printf("GPU %d: KernelGen done in %llu ms — hopping\r\n",
		CudaIndex, (unsigned long long)(GetTickCount64() - t_gen0));
	fflush(stdout);

	err = cudaMemset(Kparams.L1S2, 0, mpCnt * Kparams.BlockSize * 8);
	if (err != cudaSuccess)
		return false;
	cudaMemset(Kparams.dbg_buf, 0, 1024);
	cudaMemset(Kparams.LoopTable, 0, KangCnt * MD_LEN * sizeof(u64));
	return true;
}

bool RCGpuKang::SaveKangState(FILE* fp)
{
	cudaSetDevice(CudaIndex);
	u32 kang_bytes = KangCnt * 96;
	u32 l1_bytes = mpCnt * Kparams.BlockSize * (u32)sizeof(u64);
	u8* buf = (u8*)malloc(kang_bytes);
	u32* l1 = (u32*)malloc(l1_bytes);
	if (!buf || !l1)
	{
		free(buf);
		free(l1);
		return false;
	}
	cudaError_t err = cudaMemcpy(buf, Kparams.Kangs, kang_bytes, cudaMemcpyDeviceToHost);
	if (err != cudaSuccess)
	{
		free(buf);
		free(l1);
		return false;
	}
	err = cudaMemcpy(l1, Kparams.L1S2, l1_bytes, cudaMemcpyDeviceToHost);
	if (err != cudaSuccess)
	{
		free(buf);
		free(l1);
		return false;
	}
	bool ok = true;
	ok = ok && fwrite(&CudaIndex, 1, 4, fp) == 4;
	ok = ok && fwrite(&KangCnt, 1, 4, fp) == 4;
	ok = ok && fwrite(&Kparams.BlockSize, 1, 4, fp) == 4;
	ok = ok && fwrite(&Kparams.GroupCnt, 1, 4, fp) == 4;
	ok = ok && fwrite(&mpCnt, 1, 4, fp) == 4;
	u8 old = IsOldGpu ? 1 : 0;
	ok = ok && fwrite(&old, 1, 1, fp) == 1;
	ok = ok && fwrite(buf, 1, kang_bytes, fp) == kang_bytes;
	ok = ok && fwrite(l1, 1, l1_bytes, fp) == l1_bytes;
	free(buf);
	free(l1);
	return ok;
}

bool RCGpuKang::LoadKangState(FILE* fp)
{
	int cuda_index = 0, kang_cnt = 0, bsz = 0, gcnt = 0, mpc = 0;
	u8 old = 0;
	if (fread(&cuda_index, 1, 4, fp) != 4) return false;
	if (fread(&kang_cnt, 1, 4, fp) != 4) return false;
	if (fread(&bsz, 1, 4, fp) != 4) return false;
	if (fread(&gcnt, 1, 4, fp) != 4) return false;
	if (fread(&mpc, 1, 4, fp) != 4) return false;
	if (fread(&old, 1, 1, fp) != 1) return false;

	if (cuda_index != CudaIndex)
	{
		printf("GPU checkpoint index mismatch (%d vs %d)\r\n", cuda_index, CudaIndex);
		return false;
	}

	u32 kang_bytes = kang_cnt * 96;
	u32 l1_bytes = mpc * bsz * (u32)sizeof(u64);
	u8* buf = (u8*)malloc(kang_bytes);
	u32* l1 = (u32*)malloc(l1_bytes);
	if (!buf || !l1)
	{
		free(buf);
		free(l1);
		return false;
	}
	if (fread(buf, 1, kang_bytes, fp) != kang_bytes ||
		fread(l1, 1, l1_bytes, fp) != l1_bytes)
	{
		free(buf);
		free(l1);
		return false;
	}

	OverrideBlockSize = bsz;
	OverrideGroupCnt = gcnt;
	IsOldGpu = old != 0;
	mpCnt = mpc;
	SetRestoredState(buf, l1);
	return true;
}

#ifdef DEBUG_MODE
int RCGpuKang::Dbg_CheckKangs()
{
	int kang_size = mpCnt * Kparams.BlockSize * Kparams.GroupCnt * 96;
	u64* kangs = (u64*)malloc(kang_size);
	cudaError_t err = cudaMemcpy(kangs, Kparams.Kangs, kang_size, cudaMemcpyDeviceToHost);
	int res = 0;
	for (int i = 0; i < KangCnt; i++)
	{
		EcPoint Pnt, p;
		Pnt.LoadFromBuffer64((u8*)&kangs[i * 12 + 0]);
		EcInt dist;
		dist.Set(0);
		memcpy(dist.data, &kangs[i * 12 + 8], 24);
		bool neg = false;
		if (dist.data[2] >> 63)
		{
			neg = true;
			memset(((u8*)dist.data) + 24, 0xFF, 16);
			dist.Neg();
		}
		p = ec.MultiplyG_Fast(dist);
		if (neg)
			p.y.NegModP();
		if (i < KangCnt / 3)
			p = p;
		else
			if (i < 2 * KangCnt / 3)
				p = ec.AddPoints(PntA, p);
			else
				p = ec.AddPoints(PntB, p);
		if (!p.IsEqual(Pnt))
			res++;
	}
	free(kangs);
	return res;
}
#endif

void RCGpuKang::ExecuteClassic()
{
	cudaError_t err;
#ifdef DEBUG_MODE
	u64 iter = 1;
#endif
	while (!StopFlag)
	{
		if (PauseFlag)
		{
			Sleep(20);
			continue;
		}

		u64 t1 = GetTickCount64();
		cudaMemset(Kparams.DPs_out, 0, 4);
		cudaMemset(Kparams.DPTable, 0, KangCnt * sizeof(u32));
		cudaMemset(Kparams.LoopedKangs, 0, 8);
		CallGpuKernelABC(Kparams);
		int cnt;
		err = cudaMemcpy(&cnt, Kparams.DPs_out, 4, cudaMemcpyDeviceToHost);
		if (err != cudaSuccess)
		{
			printf("GPU %d, CallGpuKernel failed: %s\r\n", CudaIndex, cudaGetErrorString(err));
			gTotalErrors++;
			break;
		}

		if (cnt >= MAX_DP_CNT)
		{
			cnt = MAX_DP_CNT;
			printf("GPU %d, gpu DP buffer overflow, some points lost, increase DP value!\r\n", CudaIndex);
		}
		u64 pnt_cnt = (u64)KangCnt * STEP_CNT;

		if (cnt)
		{
			err = cudaMemcpy(DPs_out, Kparams.DPs_out + 4, cnt * GPU_DP_SIZE, cudaMemcpyDeviceToHost);
			if (err != cudaSuccess)
			{
				gTotalErrors++;
				break;
			}
			if (TagTrace_IsEnabled())
				cnt = TagTrace_FilterDpBatch(DPs_out, cnt);
			if (cnt)
				AddPointsToList(DPs_out, cnt, (u64)KangCnt * STEP_CNT);
		}

		cudaMemcpy(dbg, Kparams.dbg_buf, 1024, cudaMemcpyDeviceToHost);
		if (TagTrace_IsEnabled() && dbg[200])
		{
			TagTrace_AddGpuSkipped(dbg[200]);
			cudaMemset(Kparams.dbg_buf + 200, 0, 4);
		}

		u32 lcnt;
		cudaMemcpy(&lcnt, Kparams.LoopedKangs, 4, cudaMemcpyDeviceToHost);

		u64 t2 = GetTickCount64();
		u64 tm = t2 - t1;
		if (!tm)
			tm = 1;
		int cur_speed = (int)(pnt_cnt / (tm * 1000));

		SpeedStats[cur_stats_ind] = cur_speed;
		cur_stats_ind = (cur_stats_ind + 1) % STATS_WND_SIZE;

#ifdef DEBUG_MODE
		if ((iter % 300) == 0)
		{
			int corr_cnt = Dbg_CheckKangs();
			if (corr_cnt)
			{
				printf("DBG: GPU %d, KANGS CORRUPTED: %d\r\n", CudaIndex, corr_cnt);
				gTotalErrors++;
			}
			else
				printf("DBG: GPU %d, ALL KANGS OK!\r\n", CudaIndex);
		}
		iter++;
#endif
	}
}

void RCGpuKang::ExecuteDualLane()
{
	// Pipeline: KernelA/B/C on StreamHop; DP download on StreamCopy (double host buffer).
	// Host insert of batch N-1 overlaps hop kernel N (PCIe/host work under GPU hops).
	cudaError_t err;
	int pending_cnt = 0;
	u32* pending_buf = NULL;
	bool have_pending = false;
	int flip = 0;
	u32* host_bufs[2] = { DPs_out, DPs_out_B };

	while (!StopFlag)
	{
		if (PauseFlag)
		{
			Sleep(20);
			continue;
		}

		u64 t1 = GetTickCount64();

		err = cudaMemsetAsync(Kparams.DPs_out, 0, 4, StreamHop);
		if (err != cudaSuccess) { gTotalErrors++; break; }
		err = cudaMemsetAsync(Kparams.DPTable, 0, KangCnt * sizeof(u32), StreamHop);
		if (err != cudaSuccess) { gTotalErrors++; break; }
		err = cudaMemsetAsync(Kparams.LoopedKangs, 0, 8, StreamHop);
		if (err != cudaSuccess) { gTotalErrors++; break; }

		CallGpuKernelABC_Stream(Kparams, StreamHop);

		// OVERLAP: insert previous DPs on host while hop kernels run
		if (have_pending && pending_cnt > 0 && pending_buf)
		{
			int pc = pending_cnt;
			if (TagTrace_IsEnabled())
				pc = TagTrace_FilterDpBatch(pending_buf, pc);
			if (pc)
				AddPointsToList(pending_buf, pc, (u64)KangCnt * STEP_CNT);
			have_pending = false;
		}

		err = cudaStreamSynchronize(StreamHop);
		if (err != cudaSuccess)
		{
			printf("GPU %d, DualLane hop failed: %s — falling back\r\n", CudaIndex, cudaGetErrorString(err));
			gTotalErrors++;
			break;
		}

		int cnt = 0;
		err = cudaMemcpyAsync(&cnt, Kparams.DPs_out, 4, cudaMemcpyDeviceToHost, StreamCopy);
		if (err != cudaSuccess) { gTotalErrors++; break; }
		err = cudaStreamSynchronize(StreamCopy);
		if (err != cudaSuccess) { gTotalErrors++; break; }

		if (cnt >= MAX_DP_CNT)
		{
			cnt = MAX_DP_CNT;
			printf("GPU %d, gpu DP buffer overflow, some points lost, increase DP value!\r\n", CudaIndex);
		}

		u32* dest = host_bufs[flip];
		if (cnt)
		{
			err = cudaMemcpyAsync(dest, Kparams.DPs_out + 4, cnt * GPU_DP_SIZE,
				cudaMemcpyDeviceToHost, StreamCopy);
			if (err != cudaSuccess) { gTotalErrors++; break; }
			err = cudaStreamSynchronize(StreamCopy);
			if (err != cudaSuccess) { gTotalErrors++; break; }
			pending_buf = dest;
			pending_cnt = cnt;
			have_pending = true;
			flip ^= 1;
		}

		cudaMemcpyAsync(dbg, Kparams.dbg_buf, 1024, cudaMemcpyDeviceToHost, StreamCopy);
		cudaStreamSynchronize(StreamCopy);
		if (TagTrace_IsEnabled() && dbg[200])
		{
			TagTrace_AddGpuSkipped(dbg[200]);
			cudaMemsetAsync(Kparams.dbg_buf + 200, 0, 4, StreamCopy);
		}

		u64 pnt_cnt = (u64)KangCnt * STEP_CNT;
		u64 t2 = GetTickCount64();
		u64 tm = t2 - t1;
		if (!tm) tm = 1;
		int cur_speed = (int)(pnt_cnt / (tm * 1000));
		SpeedStats[cur_stats_ind] = cur_speed;
		cur_stats_ind = (cur_stats_ind + 1) % STATS_WND_SIZE;
	}

	if (have_pending && pending_cnt > 0 && pending_buf)
	{
		int pc = pending_cnt;
		if (TagTrace_IsEnabled())
			pc = TagTrace_FilterDpBatch(pending_buf, pc);
		if (pc)
			AddPointsToList(pending_buf, pc, (u64)KangCnt * STEP_CNT);
	}
}

void RCGpuKang::MaybeGroupSweep()
{
	if (!GroupSweepOn)
		return;
	u64 now = GetTickCount64();
	if (GroupSweepLastMs && now - GroupSweepLastMs < 30000)
		return;
	GroupSweepLastMs = now;
	// PNT_GROUP_CNT is fatbin-compile-time; we advise only (honest).
	int cur = (int)Kparams.GroupCnt;
	int suggest = cur;
	// Simple occupancy heuristic from speed trend
	int sp = GetStatsSpeed();
	if (sp > 0 && SpeedStats[(cur_stats_ind + STATS_WND_SIZE - 1) % STATS_WND_SIZE] > 0)
	{
		int prev = SpeedStats[(cur_stats_ind + STATS_WND_SIZE - 1) % STATS_WND_SIZE];
		if (sp < prev * 85 / 100 && cur > 16)
			suggest = cur - 8;
		else if (sp > prev * 105 / 100 && cur < 64)
			suggest = cur + 8;
	}
	if (suggest != GroupSweepSuggest)
	{
		GroupSweepSuggest = suggest;
		printf("GPU %d GroupSweep: current group=%d  suggested=%d (restart to apply; kernel group is compile-fixed)\r\n",
			CudaIndex, cur, suggest);
	}
}

void RCGpuKang::ExecuteStreamTrio()
{
	// Unified ABC on StreamHop (see CallGpuKernelA_Stream); host insert overlaps next hop.
	cudaError_t err;
	int pending_cnt = 0;
	u32* pending_buf = NULL;
	bool have_pending = false;
	int flip = 0;
	u32* host_bufs[2] = { DPs_out, DPs_out_B };
	bool logged_first = false;

	while (!StopFlag)
	{
		if (PauseFlag)
		{
			Sleep(20);
			continue;
		}

		u64 t1 = GetTickCount64();
		MaybeGroupSweep();

		err = cudaMemsetAsync(Kparams.DPs_out, 0, 4, StreamHop);
		if (err != cudaSuccess) { gTotalErrors++; break; }
		err = cudaMemsetAsync(Kparams.DPTable, 0, KangCnt * sizeof(u32), StreamHop);
		if (err != cudaSuccess) { gTotalErrors++; break; }
		err = cudaMemsetAsync(Kparams.LoopedKangs, 0, 8, StreamHop);
		if (err != cudaSuccess) { gTotalErrors++; break; }

		CallGpuKernelA_Stream(Kparams, StreamHop);
		err = cudaEventRecord(EvHopDone, StreamHop);
		if (err != cudaSuccess) { gTotalErrors++; break; }

		// Host insert of previous batch while hop runs
		if (have_pending && pending_cnt > 0 && pending_buf)
		{
			int pc = pending_cnt;
			if (TagTrace_IsEnabled())
				pc = TagTrace_FilterDpBatch(pending_buf, pc);
			if (pc)
				AddPointsToList(pending_buf, pc, (u64)KangCnt * STEP_CNT);
			have_pending = false;
		}

		err = cudaStreamWaitEvent(StreamCopy, EvHopDone, 0);
		if (err != cudaSuccess) { gTotalErrors++; break; }

		int cnt = 0;
		err = cudaMemcpyAsync(&cnt, Kparams.DPs_out, 4, cudaMemcpyDeviceToHost, StreamCopy);
		if (err != cudaSuccess) { gTotalErrors++; break; }
		err = cudaStreamSynchronize(StreamCopy);
		if (err != cudaSuccess)
		{
			printf("GPU %d, StreamTrio failed: %s — falling back\r\n", CudaIndex, cudaGetErrorString(err));
			fflush(stdout);
			gTotalErrors++;
			break;
		}
		if (!logged_first)
		{
			printf("GPU %d: first hop batch OK (DPs_this_batch=%d)\r\n", CudaIndex, cnt);
			fflush(stdout);
			logged_first = true;
		}

		if (cnt >= MAX_DP_CNT)
		{
			cnt = MAX_DP_CNT;
			printf("GPU %d, gpu DP buffer overflow, some points lost, increase DP value!\r\n", CudaIndex);
		}

		u32* dest = host_bufs[flip];
		if (cnt)
		{
			err = cudaMemcpyAsync(dest, Kparams.DPs_out + 4, cnt * GPU_DP_SIZE,
				cudaMemcpyDeviceToHost, StreamCopy);
			if (err != cudaSuccess) { gTotalErrors++; break; }
			err = cudaStreamSynchronize(StreamCopy);
			if (err != cudaSuccess) { gTotalErrors++; break; }
			pending_buf = dest;
			pending_cnt = cnt;
			have_pending = true;
			flip ^= 1;
		}

		cudaMemcpyAsync(dbg, Kparams.dbg_buf, 1024, cudaMemcpyDeviceToHost, StreamCopy);
		cudaStreamSynchronize(StreamCopy);
		if (TagTrace_IsEnabled() && dbg[200])
		{
			TagTrace_AddGpuSkipped(dbg[200]);
			cudaMemsetAsync(Kparams.dbg_buf + 200, 0, 4, StreamCopy);
		}

		u64 pnt_cnt = (u64)KangCnt * STEP_CNT;
		u64 t2 = GetTickCount64();
		u64 tm = t2 - t1;
		if (!tm) tm = 1;
		int cur_speed = (int)(pnt_cnt / (tm * 1000));
		SpeedStats[cur_stats_ind] = cur_speed;
		cur_stats_ind = (cur_stats_ind + 1) % STATS_WND_SIZE;
	}

	if (have_pending && pending_cnt > 0 && pending_buf)
	{
		int pc = pending_cnt;
		if (TagTrace_IsEnabled())
			pc = TagTrace_FilterDpBatch(pending_buf, pc);
		if (pc)
			AddPointsToList(pending_buf, pc, (u64)KangCnt * STEP_CNT);
	}
}

//executes in separate thread
void RCGpuKang::Execute()
{
	cudaSetDevice(CudaIndex);

	if (!Start())
	{
		gTotalErrors++;
		return;
	}

	if (StreamTrioReady)
		ExecuteStreamTrio();
	else if (DualLaneReady)
		ExecuteDualLane();
	else
		ExecuteClassic();

	Release();
}

int RCGpuKang::GetStatsSpeed()
{
	// Average only samples written so far — zero-filled window diluted FILL to 0 Mk/s.
	int n = 0;
	long long sum = 0;
	for (int i = 0; i < STATS_WND_SIZE; i++)
	{
		if (SpeedStats[i] > 0)
		{
			sum += SpeedStats[i];
			n++;
		}
	}
	if (!n)
		return 0;
	return (int)(sum / n);
}