/*
 * Copyright 2025 Neuraville Inc.
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 */

/**
 * CUDA Kernel: Neural Dynamics with Fire Candidate List (FCL)
 * 
 * Processes only FCL candidate neurons (sparse processing).
 * Updates membrane potentials, checks thresholds, generates fired neurons.
 * 
 * This is a PORT of the WGSL shader: neural_dynamics_fcl.wgsl
 * 
 * Performance characteristics:
 * - Processes ~100K FCL candidates in ~0.1-0.2ms on H100
 * - Scales with FCL size, not genome size (sparse advantage!)
 * - Atomic bitpacking for fired neuron mask
 */

extern "C" {

/**
 * Neural dynamics kernel (FCL-aware)
 * 
 * @param fcl_potentials         Input: Accumulated synaptic potentials
 * @param membrane_potentials    Input/Output: Current membrane potentials
 * @param thresholds            Input: Firing thresholds
 * @param leak_coefficients     Input: Leak/decay coefficients
 * @param resting_potentials    Input: Resting potential values
 * @param excitabilities        Input: Excitability factors (0.0-1.0)
 * @param refractory_countdowns Input/Output: Refractory period counters
 * @param fcl_fired_mask        Output: Bitpacked mask of fired neurons
 * @param neuron_count          Total number of neurons
 * @param burst_count           Current burst counter (for PCG hash)
 */
__global__ void neural_dynamics_fcl(
    const int* fcl_potentials,
    float* membrane_potentials,
    const float* thresholds,
    const float* leak_coefficients,
    const float* resting_potentials,
    const float* excitabilities,
    unsigned short* refractory_countdowns,
    unsigned int* fcl_fired_mask,
    unsigned int neuron_count,
    unsigned long long burst_count
) {
    // Thread indexing
    unsigned int neuron_id = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (neuron_id >= neuron_count) {
        return;
    }
    
    // Skip if no FCL potential (sparse processing optimization)
    int fcl_potential_i32 = fcl_potentials[neuron_id];
    if (fcl_potential_i32 == 0) {
        return;  // No synaptic input this burst
    }
    
    // Convert atomic i32 potential to f32
    float fcl_potential = (float)fcl_potential_i32 / 1000.0f;
    
    // Load neuron state
    float v = membrane_potentials[neuron_id];
    float threshold = thresholds[neuron_id];
    float leak = leak_coefficients[neuron_id];
    float v_rest = resting_potentials[neuron_id];
    float excitability = excitabilities[neuron_id];
    unsigned short refractory = refractory_countdowns[neuron_id];
    
    // Update refractory period (host finish_burst_refractory_period also ticks all neurons; do not
    // decrement in-kernel to avoid double-counting when the host path is wired.)
    if (refractory > 0) {
        return;  // Option B: no synaptic add, no leak, membrane unchanged
    }
    
    // Leak toward resting potential
    v += (v_rest - v) * leak;
    
    // Add synaptic input from FCL
    v += fcl_potential;
    
    // Check for firing with excitability
    bool will_fire = false;
    
    if (excitability < 0.999f) {
        // Probabilistic firing based on excitability
        // PCG hash for pseudo-random number generation
        unsigned long long state = (unsigned long long)neuron_id + burst_count * 1000000ull;
        state = state * 6364136223846793005ull + 1442695040888963407ull;
        float random = (float)((state >> 33) & 0xFFFFFFFFull) / 4294967296.0f;
        
        will_fire = (v >= threshold) && (random < excitability);
    } else {
        // Deterministic firing
        will_fire = (v >= threshold);
    }
    
    // Handle firing
    if (will_fire) {
        // Reset membrane potential
        v = v_rest;
        
        // Set refractory period (hardcoded to 3 for now, should be parameter)
        refractory_countdowns[neuron_id] = 3;
        
        // Atomic bitpacking into fired mask
        unsigned int word_idx = neuron_id / 32;
        unsigned int bit_idx = neuron_id % 32;
        atomicOr(&fcl_fired_mask[word_idx], 1u << bit_idx);
    }
    
    // Store updated membrane potential
    membrane_potentials[neuron_id] = v;
}

} // extern "C"

/*
 * CUDA Kernel Launch Configuration:
 * 
 * Recommended grid/block sizes:
 * - Block size: 256 threads
 * - Grid size: (neuron_count + 255) / 256 blocks
 * 
 * Note: Even though this is FCL-aware, we launch for all neurons but
 * threads return early if fcl_potentials[neuron_id] == 0 (sparse optimization).
 * 
 * Alternative: Launch only for FCL candidates (requires building index list)
 * - Can save 50-90% of threads for very sparse activity
 * - Trade-off: Extra kernel launch overhead to build candidate list
 * 
 * Multi-GPU considerations:
 * - Each GPU processes its own neuron shard independently
 * - No inter-GPU communication needed (embarrassingly parallel)
 * - Gather fired neurons from all GPUs at the end
 */

