harmoniis-wallet 0.1.121

//! GPU mining backend using wgpu compute shaders.
//!
//! Supports range-based mining over the fixed 1M nonce space using dynamic
//! dispatch sizing and adapter capability limits.
//!
//! The shader outputs only (best_difficulty, nonce_id).  The host re-computes
//! the full hash from the winning nonce to guarantee correctness — the same
//! approach used by the CUDA backend.

use async_trait::async_trait;
use wgpu::util::DeviceExt;

use super::nonce_table::NonceTable;
use super::sha256::{leading_zero_bits_words, state_words_to_bytes, Sha256Midstate};
use super::{CancelFlag, MinerBackend, MiningChunkResult, MiningResult, NONCE_SPACE_SIZE};

/// Default workgroup size (must match `@workgroup_size` in shader).
const WORKGROUP_SIZE: u32 = 256;

/// Input buffer words:
/// [0..8] = midstate words
/// [8] = difficulty
/// [9] = prefix_len
/// [10] = nonce_offset
/// [11] = nonce_count
const INPUT_WORDS: usize = 12;

/// Result buffer words:
/// [0] = best difficulty found (0 = no valid solution)
/// [1] = flat nonce id of the winner
/// [2] = reserved
const RESULT_WORDS: usize = 3;
const RESULT_BUFFER_SIZE: u64 = (RESULT_WORDS * 4) as u64;

/// Create a wgpu Instance for the given backends.
pub fn create_instance(backends: wgpu::Backends) -> wgpu::Instance {
    wgpu::Instance::new(wgpu::InstanceDescriptor {
        backends,
        flags: wgpu::InstanceFlags::default(),
        memory_budget_thresholds: wgpu::MemoryBudgetThresholds::default(),
        backend_options: wgpu::BackendOptions::default(),
        display: None,
    })
}

/// Platform-native compute backend. One backend per platform, no mixing.
/// Linux = Vulkan, Windows = DX12, macOS = Metal.
pub fn platform_backend() -> wgpu::Backends {
    #[cfg(target_arch = "wasm32")]
    {
        wgpu::Backends::BROWSER_WEBGPU
    }
    #[cfg(not(target_arch = "wasm32"))]
    {
        if cfg!(target_os = "windows") {
            wgpu::Backends::DX12
        } else if cfg!(target_os = "macos") {
            wgpu::Backends::METAL
        } else {
            wgpu::Backends::VULKAN
        }
    }
}

/// Enumerate Vulkan adapters on Windows to detect the true physical GPU
/// topology.  DX12/DXGI has two known bugs with identical GPUs:
///
///   1. `get_adapter_pci_info()` in wgpu-hal matches SetupDi by vendor+device
///      and returns the FIRST hit — so all identical cards share one PCI bus ID.
///   2. `EnumAdapters1` can return the same adapter multiple times or collapse
///      identical headless cards into one "linked adapter".
///
/// Vulkan's `VK_EXT_pci_bus_info` provides a unique PCI bus address per PCIe
/// slot and `vkEnumeratePhysicalDevices` sees all GPUs regardless of monitors.
///
/// Returns `None` if Vulkan is unavailable (no driver) — caller falls back to
/// DX12 as-is, no regression.
#[cfg(all(target_os = "windows", not(target_arch = "wasm32")))]
pub async fn enumerate_vulkan_gpus() -> Option<Vec<wgpu::Adapter>> {
    let instance = create_instance(wgpu::Backends::VULKAN);
    let adapters: Vec<wgpu::Adapter> = instance
        .enumerate_adapters(wgpu::Backends::VULKAN)
        .await
        .into_iter()
        .filter(|a| a.get_info().device_type != wgpu::DeviceType::Cpu)
        .collect();
    if adapters.is_empty() {
        None
    } else {
        Some(adapters)
    }
}

/// Identity of a physical GPU adapter.
#[derive(Debug, Clone, PartialEq, Eq, Hash)]
pub struct AdapterIdentity {
    pub name: String,
    pub vendor: u32,
    pub device: u32,
    pub backend: String,
    /// PCI bus address (e.g. "0000:01:00.0"). Empty on Metal/non-PCI.
    pub pci_bus: String,
}

impl AdapterIdentity {
    /// Extract identity from a wgpu AdapterInfo.
    pub fn from_info(info: &wgpu::AdapterInfo) -> Self {
        Self {
            name: info.name.clone(),
            vendor: info.vendor,
            device: info.device,
            backend: format!("{:?}", info.backend).to_lowercase(),
            pci_bus: info.device_pci_bus_id.trim().to_string(),
        }
    }

    /// Check if this identity matches a wgpu AdapterInfo.
    /// Uses PCI bus ID on Vulkan (unique per PCIe slot).
    /// Falls back to vendor+device+name on Metal/DX12.
    pub fn matches(&self, info: &wgpu::AdapterInfo) -> bool {
        let backend_ok = format!("{:?}", info.backend).to_lowercase() == self.backend;
        if !self.pci_bus.is_empty() {
            info.device_pci_bus_id.trim() == self.pci_bus && backend_ok
        } else {
            info.vendor == self.vendor
                && info.device == self.device
                && info.name == self.name
                && backend_ok
        }
    }
}

/// Run a GPU pipeline probe for an adapter identified by vendor+device+backend.
/// Native-only: spawns subprocess to test GPU driver stability.
#[cfg(not(target_arch = "wasm32"))]
pub async fn probe_adapter(identity: &AdapterIdentity) -> anyhow::Result<()> {
    let backend = platform_backend();
    let instance = create_instance(backend);
    let adapters = instance.enumerate_adapters(backend).await;
    let adapter = adapters
        .into_iter()
        .find(|a| identity.matches(&a.get_info()))
        .ok_or_else(|| {
            anyhow::anyhow!(
                "no adapter matches {} ({})",
                identity.name,
                identity.backend,
            )
        })?;
    let probe_info = adapter.get_info();
    // Probe runs in a subprocess — keep quiet.
    let (device, _queue) = adapter
        .request_device(&wgpu::DeviceDescriptor {
            label: Some("probe"),
            required_features: wgpu::Features::empty(),
            required_limits: wgpu::Limits::downlevel_defaults(),
            experimental_features: wgpu::ExperimentalFeatures::default(),
            memory_hints: wgpu::MemoryHints::default(),
            trace: wgpu::Trace::default(),
        })
        .await
        .map_err(|e| anyhow::anyhow!("device request failed: {e}"))?;

    let shader = if probe_info.backend == wgpu::Backend::Vulkan {
        let spirv_bytes: &[u8] = include_bytes!("shader/sha256_mine_opt.spv");
        let spirv_words: Vec<u32> = spirv_bytes
            .chunks_exact(4)
            .map(|c| u32::from_le_bytes([c[0], c[1], c[2], c[3]]))
            .collect();
        unsafe {
            device.create_shader_module_passthrough(wgpu::ShaderModuleDescriptorPassthrough {
                label: Some("probe_shader_spirv"),
                spirv: Some(std::borrow::Cow::Owned(spirv_words)),
                num_workgroups: (0, 0, 0),
                dxil: None,
                metallib: None,
                msl: None,
                hlsl: None,
                glsl: None,
                wgsl: None,
            })
        }
    } else {
        device.create_shader_module(wgpu::ShaderModuleDescriptor {
            label: Some("probe_shader"),
            source: wgpu::ShaderSource::Wgsl(include_str!("shader/sha256_mine.wgsl").into()),
        })
    };

    let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
        label: None,
        entries: &[
            wgpu::BindGroupLayoutEntry {
                binding: 0,
                visibility: wgpu::ShaderStages::COMPUTE,
                ty: wgpu::BindingType::Buffer {
                    ty: wgpu::BufferBindingType::Storage { read_only: true },
                    has_dynamic_offset: false,
                    min_binding_size: None,
                },
                count: None,
            },
            wgpu::BindGroupLayoutEntry {
                binding: 1,
                visibility: wgpu::ShaderStages::COMPUTE,
                ty: wgpu::BindingType::Buffer {
                    ty: wgpu::BufferBindingType::Storage { read_only: true },
                    has_dynamic_offset: false,
                    min_binding_size: None,
                },
                count: None,
            },
            wgpu::BindGroupLayoutEntry {
                binding: 2,
                visibility: wgpu::ShaderStages::COMPUTE,
                ty: wgpu::BindingType::Buffer {
                    ty: wgpu::BufferBindingType::Storage { read_only: false },
                    has_dynamic_offset: false,
                    min_binding_size: None,
                },
                count: None,
            },
        ],
    });
    let pl = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
        label: None,
        bind_group_layouts: &[Some(&bgl)],
        immediate_size: 0,
    });
    // This is the call that can segfault on buggy AMD Vulkan drivers.
    let _pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
        label: Some("probe_pipeline"),
        layout: Some(&pl),
        module: &shader,
        entry_point: Some("main"),
        compilation_options: Default::default(),
        cache: None,
    });
    // Success — parent process will use this adapter.
    Ok(())
}

/// Probe an adapter by spawning the current binary with `gpu-probe`.
/// Native-only: uses std::process::Command.
#[cfg(not(target_arch = "wasm32"))]
#[allow(dead_code)]
pub(crate) fn subprocess_probe(identity: &AdapterIdentity) -> bool {
    let exe = match std::env::current_exe() {
        Ok(p) => p,
        Err(e) => {
            eprintln!(
                "GPU probe: cannot find exe for {} ({}) — {e}",
                identity.backend, identity.vendor,
            );
            return false;
        }
    };
    let mut cmd = std::process::Command::new(exe);
    cmd.arg("gpu-probe")
        .arg("--vendor")
        .arg(identity.vendor.to_string())
        .arg("--device")
        .arg(identity.device.to_string())
        .arg("--backend")
        .arg(&identity.backend);
    if !identity.pci_bus.is_empty() {
        cmd.arg("--pci-bus").arg(&identity.pci_bus);
    }
    let status = cmd
        .stdout(std::process::Stdio::null())
        .stderr(std::process::Stdio::null())
        .status();
    let ok = match status {
        Ok(s) => s.success(),
        Err(e) => {
            eprintln!(
                "GPU probe: failed to spawn for {} ({}) — {e}",
                identity.backend, identity.vendor,
            );
            false
        }
    };
    if !ok {
        eprintln!(
            "GPU probe: adapter {} (vendor={:#x}, device={:#x}) failed — skipping",
            identity.backend, identity.vendor, identity.device,
        );
    }
    ok
}

/// Maximum work units batched per GPU in a single submit (matches CUDA).
const MAX_BATCH: usize = 64;

/// One slot = one concurrent dispatch with its own input/result/staging/bind_group.
struct BatchSlot {
    input_buffer: wgpu::Buffer,
    result_buffer: wgpu::Buffer,
    staging_buffer: wgpu::Buffer,
    bind_group: wgpu::BindGroup,
}

pub struct GpuMiner {
    device: wgpu::Device,
    queue: wgpu::Queue,
    pipeline: wgpu::ComputePipeline,
    slots: Vec<BatchSlot>,
    nonce_words: Vec<u32>,
    adapter_name: String,
    adapter_backend: wgpu::Backend,
    max_dispatch_nonces: u32,
}

impl GpuMiner {
    /// Try to initialize the default high-performance adapter.
    pub async fn try_new() -> Option<Self> {
        let compute_backends = platform_backend();
        let instance = create_instance(compute_backends);

        // Fast path: ask wgpu for a high-performance adapter.
        let preferred = instance
            .request_adapter(&wgpu::RequestAdapterOptions {
                power_preference: wgpu::PowerPreference::HighPerformance,
                compatible_surface: None,
                force_fallback_adapter: false,
            })
            .await;
        if let Ok(adapter) = preferred {
            let info = adapter.get_info();
            eprintln!(
                "GPU: preferred adapter: {} ({:?}, {:?})",
                info.name, info.backend, info.device_type
            );
            if let Some(miner) = Self::try_from_adapter(adapter).await {
                return Some(miner);
            }
        }

        // Fallback: scan all adapters and pick the first one we can open.
        let adapters = instance.enumerate_adapters(wgpu::Backends::all()).await;
        if adapters.is_empty() {
            eprintln!("GPU: no adapters visible to wgpu (enumerate_adapters returned 0)");
            return None;
        }
        eprintln!("GPU: scanning {} adapters for fallback...", adapters.len());
        for adapter in adapters {
            let info = adapter.get_info();
            eprintln!("GPU: trying {} ({:?})", info.name, info.backend);
            if let Some(miner) = Self::try_from_adapter(adapter).await {
                return Some(miner);
            }
        }

        eprintln!("GPU: no compatible adapter could be initialized");
        None
    }

    /// Try to initialize from a specific adapter.
    pub async fn try_from_adapter(adapter: wgpu::Adapter) -> Option<Self> {
        let info = adapter.get_info();
        if info.device_type == wgpu::DeviceType::Cpu {
            eprintln!("GPU: skipping CPU adapter: {}", info.name);
            return None;
        }

        let adapter_name = info.name.clone();
        eprintln!(
            "GPU: initializing {} ({:?}, vendor={:#x}, device={:#x})",
            adapter_name, info.backend, info.vendor, info.device,
        );

        let req_default = adapter
            .request_device(&wgpu::DeviceDescriptor {
                label: Some("webminer"),
                required_features: wgpu::Features::empty(),
                required_limits: wgpu::Limits::default(),
                experimental_features: wgpu::ExperimentalFeatures::default(),
                memory_hints: wgpu::MemoryHints::default(),
                trace: wgpu::Trace::default(),
            })
            .await;
        let (device, queue) = match req_default {
            Ok(ok) => ok,
            Err(err_default) => {
                eprintln!(
                    "GPU adapter '{}' failed default limits ({}), retrying with downlevel limits",
                    adapter_name, err_default
                );
                match adapter
                    .request_device(&wgpu::DeviceDescriptor {
                        label: Some("webminer-downlevel"),
                        required_features: wgpu::Features::empty(),
                        required_limits: wgpu::Limits::downlevel_defaults(),
                        experimental_features: wgpu::ExperimentalFeatures::default(),
                        memory_hints: wgpu::MemoryHints::default(),
                        trace: wgpu::Trace::default(),
                    })
                    .await
                {
                    Ok(ok) => ok,
                    Err(e) => {
                        eprintln!(
                            "GPU adapter '{}' failed downlevel limits too: {}",
                            adapter_name, e
                        );
                        return None;
                    }
                }
            }
        };

        // Native Vulkan: load pre-compiled SPIR-V (bypasses naga).
        // All other backends (Metal, DX12, WebGPU): use WGSL.
        #[cfg(all(not(target_arch = "wasm32"), feature = "gpu"))]
        let shader_module = if info.backend == wgpu::Backend::Vulkan {
            let spirv_bytes: &[u8] = include_bytes!("shader/sha256_mine_opt.spv");
            let spirv_words: Vec<u32> = spirv_bytes
                .chunks_exact(4)
                .map(|c| u32::from_le_bytes([c[0], c[1], c[2], c[3]]))
                .collect();
            eprintln!(
                "GPU: using optimized SPIR-V shader (unrolled, {} words)",
                spirv_words.len()
            );
            unsafe {
                device.create_shader_module_passthrough(wgpu::ShaderModuleDescriptorPassthrough {
                    label: Some("sha256_mine_spirv"),
                    spirv: Some(std::borrow::Cow::Owned(spirv_words)),
                    num_workgroups: (0, 0, 0),
                    dxil: None,
                    metallib: None,
                    msl: None,
                    hlsl: None,
                    glsl: None,
                    wgsl: None,
                })
            }
        } else {
            device.create_shader_module(wgpu::ShaderModuleDescriptor {
                label: Some("sha256_mine"),
                source: wgpu::ShaderSource::Wgsl(include_str!("shader/sha256_mine.wgsl").into()),
            })
        };
        #[cfg(any(target_arch = "wasm32", not(feature = "gpu")))]
        let shader_module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
            label: Some("sha256_mine"),
            source: wgpu::ShaderSource::Wgsl(include_str!("shader/sha256_mine.wgsl").into()),
        });

        let bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
            label: Some("miner_bind_group_layout"),
            entries: &[
                // binding 0: nonce_table (read-only storage)
                wgpu::BindGroupLayoutEntry {
                    binding: 0,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::Buffer {
                        ty: wgpu::BufferBindingType::Storage { read_only: true },
                        has_dynamic_offset: false,
                        min_binding_size: None,
                    },
                    count: None,
                },
                // binding 1: input (midstate + run params, read-only storage)
                wgpu::BindGroupLayoutEntry {
                    binding: 1,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::Buffer {
                        ty: wgpu::BufferBindingType::Storage { read_only: true },
                        has_dynamic_offset: false,
                        min_binding_size: None,
                    },
                    count: None,
                },
                // binding 2: output (result buffer, read-write storage)
                wgpu::BindGroupLayoutEntry {
                    binding: 2,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::Buffer {
                        ty: wgpu::BufferBindingType::Storage { read_only: false },
                        has_dynamic_offset: false,
                        min_binding_size: None,
                    },
                    count: None,
                },
            ],
        });

        let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
            label: Some("miner_pipeline_layout"),
            bind_group_layouts: &[Some(&bind_group_layout)],
            immediate_size: 0,
        });

        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
            label: Some("sha256_mine_pipeline"),
            layout: Some(&pipeline_layout),
            module: &shader_module,
            entry_point: Some("main"),
            compilation_options: Default::default(),
            cache: None,
        });

        let nonce_table = NonceTable::new();
        let nonce_words = nonce_table.as_u32_slice();
        let nonce_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
            label: Some("nonce_table"),
            contents: bytemuck::cast_slice(&nonce_words),
            usage: wgpu::BufferUsages::STORAGE,
        });

        // Pre-allocate MAX_BATCH slots (input + result + staging + bind_group).
        // This enables mine_batch: encode N dispatches in one command buffer,
        // submit once, sync once — matching CUDA's batched sync-once pattern.
        let mut slots = Vec::with_capacity(MAX_BATCH);
        for i in 0..MAX_BATCH {
            let input_buffer = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some(&format!("input_{i}")),
                size: (INPUT_WORDS * 4) as u64,
                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            let result_buffer = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some(&format!("result_{i}")),
                size: RESULT_BUFFER_SIZE,
                usage: wgpu::BufferUsages::STORAGE
                    | wgpu::BufferUsages::COPY_SRC
                    | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            let staging_buffer = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some(&format!("staging_{i}")),
                size: RESULT_BUFFER_SIZE,
                usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
                label: Some(&format!("bind_{i}")),
                layout: &bind_group_layout,
                entries: &[
                    wgpu::BindGroupEntry {
                        binding: 0,
                        resource: nonce_buffer.as_entire_binding(),
                    },
                    wgpu::BindGroupEntry {
                        binding: 1,
                        resource: input_buffer.as_entire_binding(),
                    },
                    wgpu::BindGroupEntry {
                        binding: 2,
                        resource: result_buffer.as_entire_binding(),
                    },
                ],
            });
            slots.push(BatchSlot {
                input_buffer,
                result_buffer,
                staging_buffer,
                bind_group,
            });
        }

        let limits = device.limits();
        let max_dispatch_nonces = limits
            .max_compute_workgroups_per_dimension
            .max(1)
            .saturating_mul(WORKGROUP_SIZE)
            .max(WORKGROUP_SIZE);

        Some(GpuMiner {
            device,
            queue,
            pipeline,
            slots,
            nonce_words,
            adapter_name,
            adapter_backend: info.backend,
            max_dispatch_nonces,
        })
    }

    pub fn adapter_name(&self) -> &str {
        &self.adapter_name
    }

    pub fn max_dispatch_nonces(&self) -> u32 {
        self.max_dispatch_nonces
    }

    /// Mine a batch of midstates with ONE submit + ONE sync (matches CUDA
    /// mine_batch). All dispatches are encoded into a single command buffer.
    pub async fn mine_batch(
        &self,
        midstates: &[Sha256Midstate],
        difficulty: u32,
    ) -> anyhow::Result<Vec<MiningChunkResult>> {
        let batch_size = midstates.len().min(self.slots.len());
        if batch_size == 0 {
            return Ok(Vec::new());
        }

        // Phase 1: write inputs + clear results for all slots.
        for (i, midstate) in midstates[..batch_size].iter().enumerate() {
            let slot = &self.slots[i];
            let mut input_data = [0u32; INPUT_WORDS];
            input_data[..8].copy_from_slice(midstate.state_words());
            input_data[8] = difficulty;
            input_data[9] = midstate.prefix_len as u32;
            input_data[10] = 0; // nonce_offset
            input_data[11] = NONCE_SPACE_SIZE; // nonce_count
            self.queue
                .write_buffer(&slot.input_buffer, 0, bytemuck::cast_slice(&input_data));
            self.queue
                .write_buffer(&slot.result_buffer, 0, &[0u8; RESULT_WORDS * 4]);
        }

        // Phase 2: encode ALL dispatches + copies in ONE command buffer.
        let mut encoder = self
            .device
            .create_command_encoder(&wgpu::CommandEncoderDescriptor {
                label: Some("batch_encoder"),
            });

        let num_workgroups = NONCE_SPACE_SIZE.div_ceil(WORKGROUP_SIZE);
        for i in 0..batch_size {
            let slot = &self.slots[i];
            {
                let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
                    label: None,
                    timestamp_writes: None,
                });
                pass.set_pipeline(&self.pipeline);
                pass.set_bind_group(0, &slot.bind_group, &[]);
                pass.dispatch_workgroups(num_workgroups, 1, 1);
            }
            encoder.copy_buffer_to_buffer(
                &slot.result_buffer,
                0,
                &slot.staging_buffer,
                0,
                RESULT_BUFFER_SIZE,
            );
        }

        // Phase 3: ONE submit, ONE sync.
        #[allow(unused_variables)]
        let submission = self.queue.submit(std::iter::once(encoder.finish()));

        // Map all staging buffers for reading (all issued in parallel).
        {
            let mut receivers = Vec::with_capacity(batch_size);
            for i in 0..batch_size {
                #[cfg(not(target_arch = "wasm32"))]
                let (tx, rx) = tokio::sync::oneshot::channel();
                #[cfg(target_arch = "wasm32")]
                let (tx, rx) =
                    futures_channel::oneshot::channel::<Result<(), wgpu::BufferAsyncError>>();
                self.slots[i]
                    .staging_buffer
                    .slice(0..RESULT_BUFFER_SIZE)
                    .map_async(wgpu::MapMode::Read, move |result| {
                        let _ = tx.send(result);
                    });
                receivers.push(rx);
            }
            #[cfg(not(target_arch = "wasm32"))]
            let _ = self.device.poll(wgpu::PollType::Wait {
                submission_index: Some(submission),
                timeout: None,
            });
            for rx in receivers {
                rx.await.map_err(|_| anyhow::anyhow!("map cancelled"))??;
            }
        }

        // Phase 4: read all results.
        #[cfg(not(target_arch = "wasm32"))]
        let started = std::time::Instant::now();
        let mut results = Vec::with_capacity(batch_size);
        for i in 0..batch_size {
            let data = self.slots[i].staging_buffer.slice(..).get_mapped_range();
            let words: &[u32] = bytemuck::cast_slice(&data[..RESULT_BUFFER_SIZE as usize]);
            let best_zeros = words[0];
            let nonce_id = words[1];
            drop(data);
            self.slots[i].staging_buffer.unmap();

            let result = self.verify_result(&midstates[i], difficulty, best_zeros, nonce_id);
            results.push(MiningChunkResult {
                result,
                attempted: NONCE_SPACE_SIZE as u64,
                #[cfg(not(target_arch = "wasm32"))]
                elapsed: started.elapsed(),
                #[cfg(target_arch = "wasm32")]
                elapsed: std::time::Duration::ZERO,
            });
        }
        Ok(results)
    }

    fn verify_result(
        &self,
        midstate: &Sha256Midstate,
        difficulty: u32,
        best_zeros: u32,
        nonce_id: u32,
    ) -> Option<MiningResult> {
        if best_zeros < difficulty || nonce_id >= NONCE_SPACE_SIZE {
            return None;
        }
        let n1 = (nonce_id / 1000) as usize;
        let n2 = (nonce_id % 1000) as usize;
        let state_words =
            midstate.finalize_words_from_nonce_u32(self.nonce_words[n1], self.nonce_words[n2]);
        let achieved = leading_zero_bits_words(&state_words);
        if achieved < difficulty {
            return None;
        }
        Some(MiningResult {
            nonce1_idx: n1 as u16,
            nonce2_idx: n2 as u16,
            hash: state_words_to_bytes(&state_words),
            difficulty_achieved: achieved,
        })
    }

    async fn dispatch_range(
        &self,
        midstate: &Sha256Midstate,
        difficulty: u32,
        nonce_offset: u32,
        nonce_count: u32,
    ) -> anyhow::Result<Option<MiningResult>> {
        let slot = &self.slots[0];
        let mut input_data = [0u32; INPUT_WORDS];
        input_data[..8].copy_from_slice(midstate.state_words());
        input_data[8] = difficulty;
        input_data[9] = midstate.prefix_len as u32;
        input_data[10] = nonce_offset;
        input_data[11] = nonce_count;
        self.queue
            .write_buffer(&slot.input_buffer, 0, bytemuck::cast_slice(&input_data));
        self.queue
            .write_buffer(&slot.result_buffer, 0, &[0u8; RESULT_WORDS * 4]);

        let mut encoder = self
            .device
            .create_command_encoder(&wgpu::CommandEncoderDescriptor {
                label: Some("miner_encoder"),
            });
        {
            let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
                label: Some("sha256_mine"),
                timestamp_writes: None,
            });
            pass.set_pipeline(&self.pipeline);
            pass.set_bind_group(0, &slot.bind_group, &[]);
            pass.dispatch_workgroups(nonce_count.div_ceil(WORKGROUP_SIZE), 1, 1);
        }
        encoder.copy_buffer_to_buffer(
            &slot.result_buffer,
            0,
            &slot.staging_buffer,
            0,
            RESULT_BUFFER_SIZE,
        );
        #[allow(unused_variables)]
        let submission = self.queue.submit(std::iter::once(encoder.finish()));

        {
            #[cfg(not(target_arch = "wasm32"))]
            let (tx, rx) = tokio::sync::oneshot::channel();
            #[cfg(target_arch = "wasm32")]
            let (tx, rx) =
                futures_channel::oneshot::channel::<Result<(), wgpu::BufferAsyncError>>();
            slot.staging_buffer.slice(0..RESULT_BUFFER_SIZE).map_async(
                wgpu::MapMode::Read,
                move |result| {
                    let _ = tx.send(result);
                },
            );
            #[cfg(not(target_arch = "wasm32"))]
            let _ = self.device.poll(wgpu::PollType::Wait {
                submission_index: Some(submission),
                timeout: None,
            });
            rx.await.map_err(|_| anyhow::anyhow!("map cancelled"))??;
        }
        let buffer_slice = slot.staging_buffer.slice(..);

        let data = buffer_slice.get_mapped_range();
        let words: &[u32] = bytemuck::cast_slice(&data[..RESULT_BUFFER_SIZE as usize]);
        let best_zeros = words[0];
        let nonce_id = words[1];
        drop(data);
        slot.staging_buffer.unmap();

        Ok(self.verify_result(midstate, difficulty, best_zeros, nonce_id))
    }
}

#[cfg_attr(target_arch = "wasm32", async_trait(?Send))]
#[cfg_attr(not(target_arch = "wasm32"), async_trait)]
impl MinerBackend for GpuMiner {
    fn name(&self) -> &str {
        &self.adapter_name
    }

    fn startup_summary(&self) -> Vec<String> {
        vec![
            format!("gpu_name={}", self.adapter_name),
            format!("gpu_backend={:?}", self.adapter_backend),
            format!("workgroup_size={}", WORKGROUP_SIZE),
            format!("max_dispatch_nonces={}", self.max_dispatch_nonces),
        ]
    }

    async fn benchmark(&self) -> anyhow::Result<f64> {
        let nonce_table = NonceTable::new();
        let midstate = Sha256Midstate::from_prefix(&[0u8; 64]);

        // Warm up GPU pipeline/driver state.
        let _ = self
            .mine_range(&midstate, &nonce_table, 256, 0, NONCE_SPACE_SIZE, None)
            .await?;

        let mut samples = Vec::with_capacity(8);
        for _ in 0..8 {
            let chunk = self
                .mine_range(&midstate, &nonce_table, 256, 0, NONCE_SPACE_SIZE, None)
                .await?;
            let secs = chunk.elapsed.as_secs_f64();
            if secs > 0.0 {
                samples.push(chunk.attempted as f64 / secs);
            }
        }

        if samples.is_empty() {
            return Ok(0.0);
        }
        samples.sort_by(|a, b| a.partial_cmp(b).unwrap_or(std::cmp::Ordering::Equal));
        Ok(samples[samples.len() / 2])
    }

    fn max_batch_hint(&self) -> u32 {
        NONCE_SPACE_SIZE
    }

    async fn mine_range(
        &self,
        midstate: &Sha256Midstate,
        _nonce_table: &NonceTable,
        difficulty: u32,
        start_nonce: u32,
        nonce_count: u32,
        _cancel: Option<CancelFlag>,
    ) -> anyhow::Result<MiningChunkResult> {
        let range_start = start_nonce.min(NONCE_SPACE_SIZE);
        let range_end = range_start
            .saturating_add(nonce_count)
            .min(NONCE_SPACE_SIZE);
        if range_start >= range_end {
            return Ok(MiningChunkResult::empty());
        }

        #[cfg(not(target_arch = "wasm32"))]
        let started = std::time::Instant::now();
        // Single large dispatch for GPU stability (especially AMD consumer GPUs).
        let result = self
            .dispatch_range(midstate, difficulty, range_start, range_end - range_start)
            .await?;

        Ok(MiningChunkResult {
            result,
            attempted: (range_end - range_start) as u64,
            #[cfg(not(target_arch = "wasm32"))]
            elapsed: started.elapsed(),
            #[cfg(target_arch = "wasm32")]
            elapsed: std::time::Duration::ZERO,
        })
    }
}

#[cfg(test)]
mod tests {
    use super::*;

    #[test]
    fn platform_backend_returns_one() {
        let b = platform_backend();
        // Should be exactly one backend flag.
        assert!(
            b == wgpu::Backends::VULKAN || b == wgpu::Backends::DX12 || b == wgpu::Backends::METAL
        );
    }

    #[test]
    fn identity_from_info_captures_name() {
        // AdapterIdentity should carry all fields from AdapterInfo.
        let id = AdapterIdentity {
            name: "AMD Radeon RX 6800".into(),
            vendor: 4098,
            device: 0x73BF,
            backend: "vulkan".into(),
            pci_bus: "0000:01:00.0".into(),
        };
        assert_eq!(id.name, "AMD Radeon RX 6800");
        assert_eq!(id.pci_bus, "0000:01:00.0");
    }
}

// ── WASM: complete mining batch with non-blocking claim ─────────

#[cfg(target_arch = "wasm32")]
use webylib::wallet::Wallet as WebcashWallet;
#[cfg(target_arch = "wasm32")]
use webylib::{Amount, SecretWebcash};

/// Result of one GPU mining batch.
#[cfg(target_arch = "wasm32")]
#[derive(serde::Serialize)]
pub struct GpuMineBatchResult {
    pub found: bool,
    pub state: String,
    pub attempted: u64,
    pub difficulty: u32,
    pub mining_amount: String,
    #[serde(skip_serializing_if = "Option::is_none")]
    pub preimage_b64: Option<String>,
    #[serde(skip_serializing_if = "Option::is_none")]
    pub hash_hex: Option<String>,
    #[serde(skip_serializing_if = "Option::is_none")]
    pub difficulty_achieved: Option<u32>,
}

/// Cached mining target — refreshed every 30 seconds.
#[cfg(target_arch = "wasm32")]
thread_local! {
    static CACHED_TARGET: std::cell::RefCell<Option<(super::protocol::TargetInfo, f64)>> = std::cell::RefCell::new(None);
    static CACHED_NONCE_TABLE: std::cell::RefCell<Option<NonceTable>> = std::cell::RefCell::new(None);
}

/// Run one GPU mining batch. Target is cached (refreshed every 30s).
/// On solution: spawn non-blocking claim (submit + HD insert).
#[cfg(target_arch = "wasm32")]
impl GpuMiner {
    pub async fn mine_and_claim(
        &self,
        wallet: &WebcashWallet,
        network: webylib::server::NetworkMode,
        batch_size: usize,
    ) -> anyhow::Result<GpuMineBatchResult> {
        use super::protocol::MiningProtocol;
        use super::work_unit::WorkUnit;

        // Fetch target with 30-second cache
        let now = js_sys::Date::now();
        let target = CACHED_TARGET.with(|c| {
            let cached = c.borrow();
            if let Some((ref t, ts)) = *cached {
                if now - ts < 30_000.0 {
                    return Some(t.clone());
                }
            }
            None
        });
        let target = match target {
            Some(t) => t,
            None => {
                let t = MiningProtocol::from_network(&network)?.get_target().await?;
                CACHED_TARGET.with(|c| *c.borrow_mut() = Some((t.clone(), now)));
                t
            }
        };
        let difficulty = target.difficulty;
        let mining_amount = target.mining_amount;
        let subsidy_amount = target.subsidy_amount;

        let works: Vec<WorkUnit> = (0..batch_size)
            .map(|_| WorkUnit::new(difficulty, mining_amount, subsidy_amount))
            .collect();
        let midstates: Vec<_> = works.iter().map(|w| w.midstate.clone()).collect();
        let chunks = self.mine_batch(&midstates, difficulty).await?;
        let total_attempted: u64 = chunks.iter().map(|c| c.attempted).sum();

        // Check for solution (nonce table cached)
        CACHED_NONCE_TABLE.with(|c| {
            if c.borrow().is_none() {
                *c.borrow_mut() = Some(NonceTable::new());
            }
        });
        for (chunk, work) in chunks.into_iter().zip(works.into_iter()) {
            if let Some(r) = chunk.result {
                let preimage = CACHED_NONCE_TABLE.with(|c| {
                    let nt = c.borrow();
                    work.preimage_string(nt.as_ref().unwrap(), r.nonce1_idx, r.nonce2_idx)
                });
                let hash_hex = hex::encode(r.hash);
                let keep_str = work.keep_secret.to_string();
                let difficulty_achieved = r.difficulty_achieved;

                // Submit report + insert with HD RECEIVE secret (inline, must complete before return)
                super::super::wallet::webcash::submit_and_claim_mining_solution(
                    wallet, &network, &preimage, &r.hash, &keep_str,
                )
                .await
                .map_err(|e| anyhow::anyhow!("claim failed: {e}"))?;
                let state = wallet
                    .to_json()
                    .map_err(|e| anyhow::anyhow!("to_json: {e}"))?;

                return Ok(GpuMineBatchResult {
                    found: true,
                    state,
                    attempted: total_attempted,
                    difficulty,
                    mining_amount: mining_amount.to_string(),
                    preimage_b64: Some(preimage),
                    hash_hex: Some(hash_hex),
                    difficulty_achieved: Some(difficulty_achieved),
                });
            }
        }

        Ok(GpuMineBatchResult {
            found: false,
            state: wallet
                .to_json()
                .map_err(|e| anyhow::anyhow!("to_json: {e}"))?,
            attempted: total_attempted,
            difficulty,
            mining_amount: mining_amount.to_string(),
            preimage_b64: None,
            hash_hex: None,
            difficulty_achieved: None,
        })
    }
}