app.c

#define UNICODE 1
#define _UNICODE 1
#include <windows.h>

typedef unsigned char u8;
typedef unsigned short u16;
typedef unsigned int u32;
typedef unsigned long long u64;
typedef signed int i32;

/*
 * Zig/LLD's Windows-GNU driver roots this symbol even in a no-CRT image.
 * The program does not use TLS, but defining the standard loader slot keeps
 * the final PE on the normal COFF/x64 link path instead of relying on a
 * hand-built GNU PE import image.
 */
DWORD _tls_index = 0;

#define ARRAY_COUNT(a) ((u32)(sizeof(a) / sizeof((a)[0])))
#define MAX_N 32u
#define MAX_CELLS (MAX_N * MAX_N)
#define MAX_NODES (MAX_CELLS * 8u)
#define MAX_ROUTE MAX_NODES
#define WALL_N 1u
#define WALL_E 2u
#define WALL_S 4u
#define WALL_W 8u
#define INF64 (~0ull)
#define PACK_SHIFT 20u
#define LOW_MASK ((1ull << PACK_SHIFT) - 1ull)
#define LEFT_PANEL 530
#define WM_APP_CUDA (WM_APP + 1)
#define ID_PROFILE 100
#define ID_SEED 101
#define ID_GENERATE 102
#define ID_SOLVE 103
#define ID_DIAGONAL 104
#define ID_MULTIRUN 105
#define ID_STATUS 106
#define ID_RESULT 107
#define ID_OFFICIAL_SERIES 108
#define ID_OFFICIAL_MAZE 109
#define ID_OFFICIAL_LOAD 110
#define ID_PARAM_BASE 200
#define PTX_RESOURCE_ID 101

enum ParamIndex {
    PARAM_STRAIGHT_V,
    PARAM_DIAGONAL_V,
    PARAM_TURN_V,
    PARAM_ACCEL,
    PARAM_BRAKE,
    PARAM_JERK,
    PARAM_MU_X,
    PARAM_MU_Y,
    PARAM_SLIP,
    PARAM_BATTERY,
    PARAM_SAG,
    PARAM_BATTERY_R,
    PARAM_BATTERY_CAPACITY,
    PARAM_MASS,
    PARAM_CHASSIS_LENGTH,
    PARAM_TRACK,
    PARAM_WHEEL_RADIUS,
    PARAM_WHEEL_MASS,
    PARAM_CG_HEIGHT,
    PARAM_MOTOR_R,
    PARAM_MOTOR_KT,
    PARAM_MOTOR_KE,
    PARAM_CURRENT_LIMIT,
    PARAM_CONTROLLER_KP,
    PARAM_RADIUS,
    PARAM_MAX_RUN,
    PARAM_DT_US,
    PARAM_COUNT
};

typedef struct SolverParams {
    u32 n;
    u32 cell_mm;
    u32 straight_v;
    u32 diagonal_v;
    u32 turn_v;
    u32 accel;
    u32 brake;
    u32 jerk;
    u32 mu_milli;
    u32 slip_milli;
    u32 battery_mv;
    u32 sag_mv_s;
    u32 turn_radius_mm;
    u32 max_run;
    u32 mu_y_milli;
    u32 battery_r_mohm;
    u32 battery_capacity_mah;
    u32 mass_g;
    u32 chassis_length_mm;
    u32 track_mm;
    u32 wheel_radius_mm;
    u32 wheel_mass_g;
    u32 cg_height_mm;
    u32 motor_r_mohm;
    u32 motor_kt_mnm_a;
    u32 motor_ke_mv_rad;
    u32 current_limit_ma;
    u32 controller_kp_mv_rad;
    u32 dt_us;
} SolverParams;

typedef struct DynamicsParams {
    float cell_m;
    float straight_v;
    float diagonal_v;
    float turn_v;
    float accel;
    float brake;
    float jerk;
    float mu_x;
    float mu_y;
    float slip_peak;
    float battery_v;
    float battery_sag_v_s;
    float battery_r;
    float battery_capacity_c;
    float mass;
    float chassis_length;
    float track;
    float wheel_radius;
    float wheel_mass;
    float cg_height;
    float motor_r;
    float motor_kt;
    float motor_ke;
    float current_limit;
    float controller_kp;
    float turn_radius;
    float wheel_inertia;
    float yaw_inertia;
    float dt;
    u32 max_steps;
} DynamicsParams;

typedef struct DynamicsResult {
    float time_s;
    float max_slip;
    float final_voltage;
    float max_lateral;
    float final_heading_error;
    float final_soc;
    float final_position_error;
    u32 cone_hits;
    u32 steps;
    u32 valid;
    u32 failure_reason;
    u32 failure_action;
} DynamicsResult;

enum DynamicsFailureReason {
    DYN_FAIL_NONE,
    DYN_FAIL_INVALID_PARAMETERS,
    DYN_FAIL_NONFINITE,
    DYN_FAIL_STATE_LIMIT,
    DYN_FAIL_LATERAL_ENVELOPE,
    DYN_FAIL_SLIP_LIMIT,
    DYN_FAIL_BATTERY_EMPTY,
    DYN_FAIL_TOTAL_STEP_LIMIT,
    DYN_FAIL_SEGMENT_STEP_LIMIT,
    DYN_FAIL_SEGMENT_HEADING,
    DYN_FAIL_FINAL_POSITION,
    DYN_FAIL_FINAL_HEADING,
    DYN_FAIL_COUNT
};

typedef struct GpuCandidateResult {
    float score_time_s;
    float predicted_time_s;
    float max_slip;
    float final_voltage;
    float max_cone_ratio;
    float straight_scale;
    float turn_scale;
    float final_soc;
    u32 cone_hits;
    u32 valid;
} GpuCandidateResult;

_Static_assert(sizeof(DynamicsParams) == 120u, "CUDA DynamicsParams ABI changed");
_Static_assert(__builtin_offsetof(DynamicsParams, battery_v) == 40u,
               "CUDA battery offset changed");
_Static_assert(__builtin_offsetof(DynamicsParams, mass) == 56u,
               "CUDA mass offset changed");
_Static_assert(__builtin_offsetof(DynamicsParams, turn_radius) == 100u,
               "CUDA radius offset changed");
_Static_assert(sizeof(GpuCandidateResult) == 40u, "CUDA result ABI changed");

typedef struct TrajectorySample {
    float x_m;
    float y_m;
    float heading_rad;
    float speed_m_s;
    float lateral_m_s;
    float yaw_rate_rad_s;
    float slip;
    float voltage;
} TrajectorySample;

#define DYNAMICS_CANDIDATES 4096u
#define MAX_TRAJECTORY_SAMPLES 65536u

typedef struct Maze {
    u32 n;
    u32 cell_mm;
    u64 seed;
    u8 walls[MAX_CELLS];
    u32 goal_entrance_cell;
    u32 goal_entrance_dir;
    u8 goal_x0, goal_y0, goal_x1, goal_y1;
    u8 require_single_goal_entrance;
    u8 require_full_connectivity;
    u8 require_peg_walls;
    u8 require_wall_follower_trap;
} Maze;

#include "official_maze_data.h"

typedef int CUresult;
typedef int CUdevice;
typedef void *CUcontext;
typedef void *CUmodule;
typedef void *CUfunction;
typedef u64 CUdeviceptr;

typedef CUresult (WINAPI *PFN_cuInit)(u32);
typedef CUresult (WINAPI *PFN_cuDriverGetVersion)(int *);
typedef CUresult (WINAPI *PFN_cuDeviceGetCount)(int *);
typedef CUresult (WINAPI *PFN_cuDeviceGet)(CUdevice *, int);
typedef CUresult (WINAPI *PFN_cuDeviceGetName)(char *, int, CUdevice);
typedef CUresult (WINAPI *PFN_cuDeviceGetAttribute)(int *, int, CUdevice);
typedef CUresult (WINAPI *PFN_cuDeviceTotalMem)(SIZE_T *, CUdevice);
typedef CUresult (WINAPI *PFN_cuCtxCreate)(CUcontext *, u32, CUdevice);
typedef CUresult (WINAPI *PFN_cuCtxDestroy)(CUcontext);
typedef CUresult (WINAPI *PFN_cuModuleLoadDataEx)(CUmodule *, const void *, u32, void *, void **);
typedef CUresult (WINAPI *PFN_cuModuleUnload)(CUmodule);
typedef CUresult (WINAPI *PFN_cuModuleGetFunction)(CUfunction *, CUmodule, const char *);
typedef CUresult (WINAPI *PFN_cuMemAlloc)(CUdeviceptr *, SIZE_T);
typedef CUresult (WINAPI *PFN_cuMemFree)(CUdeviceptr);
typedef CUresult (WINAPI *PFN_cuMemcpyHtoD)(CUdeviceptr, const void *, SIZE_T);
typedef CUresult (WINAPI *PFN_cuMemcpyDtoH)(void *, CUdeviceptr, SIZE_T);
typedef CUresult (WINAPI *PFN_cuLaunchKernel)(CUfunction, u32, u32, u32,
                                             u32, u32, u32, u32,
                                             void *, void **, void **);
typedef CUresult (WINAPI *PFN_cuCtxSynchronize)(void);
typedef CUresult (WINAPI *PFN_cuGetErrorString)(CUresult, const char **);

typedef struct CudaState {
    HMODULE library;
    CUcontext context;
    CUmodule module;
    CUfunction relax;
    CUfunction dynamics;
    CUfunction fem_clear;
    CUfunction fem_tie;
    CUfunction fem_force;
    CUfunction fem_integrate;
    CUdevice device;
    int driver_version;
    int cc_major;
    int cc_minor;
    int multiprocessors;
    int clock_khz;
    SIZE_T total_memory;
    int ready;
    char gpu_name[128];
    WCHAR error[256];
    PFN_cuInit cuInit;
    PFN_cuDriverGetVersion cuDriverGetVersion;
    PFN_cuDeviceGetCount cuDeviceGetCount;
    PFN_cuDeviceGet cuDeviceGet;
    PFN_cuDeviceGetName cuDeviceGetName;
    PFN_cuDeviceGetAttribute cuDeviceGetAttribute;
    PFN_cuDeviceTotalMem cuDeviceTotalMem;
    PFN_cuCtxCreate cuCtxCreate;
    PFN_cuCtxDestroy cuCtxDestroy;
    PFN_cuModuleLoadDataEx cuModuleLoadDataEx;
    PFN_cuModuleUnload cuModuleUnload;
    PFN_cuModuleGetFunction cuModuleGetFunction;
    PFN_cuMemAlloc cuMemAlloc;
    PFN_cuMemFree cuMemFree;
    PFN_cuMemcpyHtoD cuMemcpyHtoD;
    PFN_cuMemcpyDtoH cuMemcpyDtoH;
    PFN_cuLaunchKernel cuLaunchKernel;
    PFN_cuCtxSynchronize cuCtxSynchronize;
    PFN_cuGetErrorString cuGetErrorString;
} CudaState;

typedef struct AppState {
    HINSTANCE instance;
    HWND window;
    HWND profile;
    HWND seed;
    HWND generate;
    HWND solve;
    HWND diagonal;
    HWND multirun;
    HWND status;
    HWND result;
    HWND official_series;
    HWND official_maze;
    HWND official_load;
    HWND param_edits[PARAM_COUNT];
    HFONT font;
    Maze maze;
    SolverParams params;
    CudaState cuda;
    u32 path_nodes[MAX_ROUTE];
    u8 path_actions[MAX_ROUTE];
    u32 path_count;
    u64 best_time_us;
    u32 solve_iterations;
    int exact_pass;
    int dynamics_pass;
    u64 graph_time_us;
    u32 dynamics_candidate;
    u32 dynamics_candidates;
    u32 dynamics_cpu_replays;
    u32 dynamics_route_attempts;
    u32 dynamics_failure_counts[DYN_FAIL_COUNT];
    float best_straight_scale;
    float best_turn_scale;
    float gpu_predicted_time;
    float gpu_predicted_slip;
    float gpu_max_cone_ratio;
    DynamicsResult dynamics_result;
    TrajectorySample *trajectory;
    u32 trajectory_count;
    u32 trajectory_capacity;
    int seed_dirty;
    int setting_seed_text;
    int official_loaded;
    u32 official_index;
} AppState;

static AppState g_app;
static AppState *volatile g_relocation_anchor = &g_app;
static u8 g_visit[MAX_CELLS];
static u16 g_stack[MAX_CELLS];
static u16 g_queue[MAX_CELLS];
static u8 g_follower_visit[MAX_CELLS * 4u];
static u64 g_dist[MAX_NODES];
static u32 g_trace_nodes[MAX_ROUTE];
static u8 g_trace_actions[MAX_ROUTE];
static WCHAR g_text[16384];
static WCHAR g_small[256];
static u8 g_candidate_tested[DYNAMICS_CANDIDATES];
static float g_segment_distance[MAX_ROUTE];
static float g_segment_speed_limit[MAX_ROUTE];
static float g_boundary_speed[MAX_ROUTE + 1u];

void *memset(void *destination, int value, SIZE_T size) {
    u8 *d = (u8 *)destination;
    for (SIZE_T i = 0; i < size; ++i) d[i] = (u8)value;
    return destination;
}

void *memcpy(void *destination, const void *source, SIZE_T size) {
    u8 *d = (u8 *)destination;
    const u8 *s = (const u8 *)source;
    for (SIZE_T i = 0; i < size; ++i) d[i] = s[i];
    return destination;
}

static SIZE_T wlen(const WCHAR *text) {
    SIZE_T n = 0;
    if (!text) return 0;
    while (text[n]) ++n;
    return n;
}

static void wcopy(WCHAR *destination, SIZE_T capacity, const WCHAR *source) {
    SIZE_T i = 0;
    if (!capacity) return;
    while (source && source[i] && i + 1 < capacity) {
        destination[i] = source[i];
        ++i;
    }
    destination[i] = 0;
}

static void wcat(WCHAR *destination, SIZE_T capacity, const WCHAR *source) {
    SIZE_T at = wlen(destination);
    SIZE_T i = 0;
    if (at >= capacity) return;
    while (source && source[i] && at + 1 < capacity) destination[at++] = source[i++];
    destination[at] = 0;
}

static void wcat_char(WCHAR *destination, SIZE_T capacity, WCHAR value) {
    SIZE_T at = wlen(destination);
    if (at + 1 >= capacity) return;
    destination[at] = value;
    destination[at + 1] = 0;
}

static void wcat_u64(WCHAR *destination, SIZE_T capacity, u64 value) {
    WCHAR digits[32];
    u32 count = 0;
    if (!value) {
        wcat_char(destination, capacity, L'0');
        return;
    }
    while (value && count < ARRAY_COUNT(digits)) {
        digits[count++] = (WCHAR)(L'0' + (value % 10ull));
        value /= 10ull;
    }
    while (count) wcat_char(destination, capacity, digits[--count]);
}

static void wcat_fixed6(WCHAR *destination, SIZE_T capacity, u64 microseconds) {
    u64 seconds = microseconds / 1000000ull;
    u32 fraction = (u32)(microseconds % 1000000ull);
    wcat_u64(destination, capacity, seconds);
    wcat_char(destination, capacity, L'.');
    u32 divisor = 100000u;
    while (divisor) {
        wcat_char(destination, capacity, (WCHAR)(L'0' + (fraction / divisor) % 10u));
        divisor /= 10u;
    }
}

static void wcat_float3(WCHAR *destination, SIZE_T capacity, float value) {
    if (!(value == value)) {
        wcat(destination, capacity, L"NaN");
        return;
    }
    if (value < 0.0f) {
        wcat_char(destination, capacity, L'-');
        value = -value;
    }
    if (value > 4294967.0f) {
        wcat(destination, capacity, L">4294967");
        return;
    }
    u64 milli = (u64)(value * 1000.0f + 0.5f);
    wcat_u64(destination, capacity, milli / 1000ull);
    wcat_char(destination, capacity, L'.');
    u32 fraction = (u32)(milli % 1000ull);
    wcat_char(destination, capacity, (WCHAR)(L'0' + fraction / 100u));
    wcat_char(destination, capacity, (WCHAR)(L'0' + (fraction / 10u) % 10u));
    wcat_char(destination, capacity, (WCHAR)(L'0' + fraction % 10u));
}

static void ascii_to_wide(WCHAR *destination, SIZE_T capacity, const char *source) {
    SIZE_T i = 0;
    if (!capacity) return;
    while (source && source[i] && i + 1 < capacity) {
        unsigned char c = (unsigned char)source[i];
        destination[i] = c < 128 ? (WCHAR)c : L'?';
        ++i;
    }
    destination[i] = 0;
}

static u32 parse_u32_text(const WCHAR *text, u32 fallback, u32 minimum, u32 maximum) {
    u64 value = 0;
    int found = 0;
    for (u32 i = 0; text && text[i]; ++i) {
        if (text[i] >= L'0' && text[i] <= L'9') {
            found = 1;
            value = value * 10ull + (u32)(text[i] - L'0');
            if (value > maximum) return maximum;
        } else if (text[i] == L' ' || text[i] == L'\t') {
            continue;
        } else {
            break;
        }
    }
    if (!found) return fallback;
    if (value < minimum) return minimum;
    return (u32)value;
}

static u32 parse_milli_text(const WCHAR *text, u32 fallback, u32 minimum, u32 maximum) {
    u64 whole = 0;
    u32 fraction = 0;
    u32 fraction_digits = 0;
    int decimal = 0;
    int found = 0;
    for (u32 i = 0; text && text[i]; ++i) {
        WCHAR c = text[i];
        if (c >= L'0' && c <= L'9') {
            found = 1;
            if (!decimal) whole = whole * 10ull + (u32)(c - L'0');
            else if (fraction_digits < 3u) {
                fraction = fraction * 10u + (u32)(c - L'0');
                ++fraction_digits;
            }
        } else if ((c == L'.' || c == L',') && !decimal) {
            decimal = 1;
        } else if (c == L' ' || c == L'\t') {
            continue;
        } else {
            break;
        }
    }
    if (!found) return fallback;
    while (fraction_digits < 3u) {
        fraction *= 10u;
        ++fraction_digits;
    }
    u64 result = whole * 1000ull + fraction;
    if (result < minimum) result = minimum;
    if (result > maximum) result = maximum;
    return (u32)result;
}

static u32 get_edit_u32(HWND edit, u32 fallback, u32 minimum, u32 maximum) {
    WCHAR buffer[64];
    GetWindowTextW(edit, buffer, (int)ARRAY_COUNT(buffer));
    return parse_u32_text(buffer, fallback, minimum, maximum);
}

static u32 get_edit_milli(HWND edit, u32 fallback, u32 minimum, u32 maximum) {
    WCHAR buffer[64];
    GetWindowTextW(edit, buffer, (int)ARRAY_COUNT(buffer));
    return parse_milli_text(buffer, fallback, minimum, maximum);
}

static u64 rng_next(u64 *state) {
    u64 x = *state;
    if (!x) x = 0x9e3779b97f4a7c15ull;
    x ^= x >> 12u;
    x ^= x << 25u;
    x ^= x >> 27u;
    *state = x;
    return x * 2685821657736338717ull;
}

static u32 next_seed_u32(u32 current) {
    u64 state = (u64)current ^ 0xd1b54a32d192ed03ull;
    u64 mixed = rng_next(&state);
    u32 next = (u32)(mixed ^ (mixed >> 32u));
    if (!next || next == current) next ^= 0xa511e9b3u;
    return next ? next : 1u;
}

static void set_seed_text(u32 seed) {
    g_small[0] = 0;
    wcat_u64(g_small, ARRAY_COUNT(g_small), seed);
    g_app.setting_seed_text = 1;
    SetWindowTextW(g_app.seed, g_small);
    g_app.setting_seed_text = 0;
}

static int is_goal_cell(const Maze *maze, u32 cell) {
    u32 x = cell % maze->n;
    u32 y = cell / maze->n;
    return x >= maze->goal_x0 && x <= maze->goal_x1 &&
           y >= maze->goal_y0 && y <= maze->goal_y1;
}

static const i32 g_dx4[4] = {0, 1, 0, -1};
static const i32 g_dy4[4] = {1, 0, -1, 0};
static const u8 g_bit4[4] = {WALL_N, WALL_E, WALL_S, WALL_W};
static const u8 g_opp4[4] = {WALL_S, WALL_W, WALL_N, WALL_E};

static void carve_edge(Maze *maze, u32 cell, u32 direction) {
    i32 x = (i32)(cell % maze->n);
    i32 y = (i32)(cell / maze->n);
    i32 nx = x + g_dx4[direction];
    i32 ny = y + g_dy4[direction];
    if (nx < 0 || ny < 0 || nx >= (i32)maze->n || ny >= (i32)maze->n) return;
    u32 neighbor = (u32)ny * maze->n + (u32)nx;
    maze->walls[cell] &= (u8)~g_bit4[direction];
    maze->walls[neighbor] &= (u8)~g_opp4[direction];
}

static void add_edge(Maze *maze, u32 cell, u32 direction) {
    i32 x = (i32)(cell % maze->n);
    i32 y = (i32)(cell / maze->n);
    i32 nx = x + g_dx4[direction];
    i32 ny = y + g_dy4[direction];
    if (nx < 0 || ny < 0 || nx >= (i32)maze->n || ny >= (i32)maze->n) return;
    u32 neighbor = (u32)ny * maze->n + (u32)nx;
    maze->walls[cell] |= g_bit4[direction];
    maze->walls[neighbor] |= g_opp4[direction];
}

static int is_start_allowed_edge(u32 n, u32 a, u32 b) {
    if (a != 0u && b != 0u) return 1;
    return (a == 0u && b == n) || (b == 0u && a == n);
}

static int maze_is_fully_connected(const Maze *maze) {
    u32 cells = maze->n * maze->n;
    memset(g_visit, 0, cells);
    u32 head = 0u, tail = 0u;
    g_queue[tail++] = 0u;
    g_visit[0] = 1u;
    while (head < tail) {
        u32 cell = g_queue[head++];
        u32 x = cell % maze->n;
        u32 y = cell / maze->n;
        for (u32 direction = 0; direction < 4u; ++direction) {
            if (maze->walls[cell] & g_bit4[direction]) continue;
            i32 nx = (i32)x + g_dx4[direction];
            i32 ny = (i32)y + g_dy4[direction];
            if (nx < 0 || ny < 0 || nx >= (i32)maze->n || ny >= (i32)maze->n)
                return 0;
            u32 next = (u32)ny * maze->n + (u32)nx;
            if (!g_visit[next]) {
                g_visit[next] = 1u;
                g_queue[tail++] = (u16)next;
            }
        }
    }
    return tail == cells;
}

static int interior_post_has_wall(const Maze *maze, u32 px, u32 py) {
    u32 n = maze->n;
    if (!px || !py || px >= n || py >= n) return 1;
    if (maze->walls[(py - 1u) * n + (px - 1u)] & WALL_N) return 1;
    if (maze->walls[(py - 1u) * n + px] & WALL_N) return 1;
    if (maze->walls[(py - 1u) * n + (px - 1u)] & WALL_E) return 1;
    if (maze->walls[py * n + (px - 1u)] & WALL_E) return 1;
    return 0;
}

static int wall_follower_reaches_goal(const Maze *maze, int keep_left) {
    u32 cells = maze->n * maze->n;
    memset(g_follower_visit, 0, cells * 4u);
    u32 cell = 0u;
    u32 heading = 0u;
    static const i32 left_order[4] = {-1, 0, 1, 2};
    static const i32 right_order[4] = {1, 0, -1, 2};
    const i32 *order = keep_left ? left_order : right_order;
    for (u32 step = 0; step <= cells * 4u; ++step) {
        u32 state = cell * 4u + heading;
        if (g_follower_visit[state]) return 0;
        g_follower_visit[state] = 1u;
        u32 direction = heading;
        int found = 0;
        for (u32 choice = 0; choice < 4u; ++choice) {
            direction = (u32)((i32)heading + order[choice] + 4) & 3u;
            if (!(maze->walls[cell] & g_bit4[direction])) { found = 1; break; }
        }
        if (!found) return 0;
        i32 x = (i32)(cell % maze->n) + g_dx4[direction];
        i32 y = (i32)(cell / maze->n) + g_dy4[direction];
        if (x < 0 || y < 0 || x >= (i32)maze->n || y >= (i32)maze->n) return 0;
        cell = (u32)y * maze->n + (u32)x;
        heading = direction;
        if (is_goal_cell(maze, cell)) return 1;
    }
    return 0;
}

static int validate_maze(const Maze *maze) {
    u32 n = maze->n;
    if (n < 2u || n > MAX_N ||
        maze->goal_x0 > maze->goal_x1 || maze->goal_y0 > maze->goal_y1 ||
        maze->goal_x1 >= n || maze->goal_y1 >= n) return 0;
    for (u32 y = 0; y < n; ++y) {
        for (u32 x = 0; x < n; ++x) {
            u32 c = y * n + x;
            u8 w = maze->walls[c];
            if (y == n - 1u && !(w & WALL_N)) return 0;
            if (x == n - 1u && !(w & WALL_E)) return 0;
            if (y == 0u && !(w & WALL_S)) return 0;
            if (x == 0u && !(w & WALL_W)) return 0;
            if (y + 1u < n && ((w & WALL_N) != 0) != ((maze->walls[c + n] & WALL_S) != 0)) return 0;
            if (x + 1u < n && ((w & WALL_E) != 0) != ((maze->walls[c + 1u] & WALL_W) != 0)) return 0;
        }
    }
    if ((maze->walls[0] & (WALL_E | WALL_S | WALL_W)) != (WALL_E | WALL_S | WALL_W)) return 0;
    if (maze->walls[0] & WALL_N) return 0;

    u32 goal_openings = 0;
    for (u32 y = maze->goal_y0; y <= maze->goal_y1; ++y) {
        for (u32 x = maze->goal_x0; x <= maze->goal_x1; ++x) {
            u32 c = y * n + x;
            if (x < maze->goal_x1 && (maze->walls[c] & WALL_E)) return 0;
            if (y < maze->goal_y1 && (maze->walls[c] & WALL_N)) return 0;
            if (x == maze->goal_x0 && !(maze->walls[c] & WALL_W)) ++goal_openings;
            if (x == maze->goal_x1 && !(maze->walls[c] & WALL_E)) ++goal_openings;
            if (y == maze->goal_y0 && !(maze->walls[c] & WALL_S)) ++goal_openings;
            if (y == maze->goal_y1 && !(maze->walls[c] & WALL_N)) ++goal_openings;
        }
    }
    if (!goal_openings || (maze->require_single_goal_entrance && goal_openings != 1u)) return 0;

    if (!maze_is_fully_connected(maze) && maze->require_full_connectivity) return 0;
    int goal_reached = 0;
    for (u32 y = maze->goal_y0; y <= maze->goal_y1; ++y) {
        for (u32 x = maze->goal_x0; x <= maze->goal_x1; ++x) {
            if (g_visit[y * n + x]) goal_reached = 1;
        }
    }
    if (!goal_reached) return 0;
    if (maze->require_peg_walls) {
        for (u32 py = 1u; py < n; ++py) {
            for (u32 px = 1u; px < n; ++px) {
                if (px == n / 2u && py == n / 2u) {
                    if (interior_post_has_wall(maze, px, py)) return 0;
                } else if (!interior_post_has_wall(maze, px, py)) return 0;
            }
        }
    }
    if (maze->require_wall_follower_trap &&
        (wall_follower_reaches_goal(maze, 1) || wall_follower_reaches_goal(maze, 0)))
        return 0;
    return 1;
}

static u32 official_maze_hash(const Maze *maze) {
    u32 value = 2166136261u;
    u8 prefix[7] = {
        (u8)maze->n, (u8)(maze->cell_mm & 255u), (u8)(maze->cell_mm >> 8u),
        maze->goal_x0, maze->goal_y0, maze->goal_x1, maze->goal_y1
    };
    for (u32 i = 0; i < ARRAY_COUNT(prefix); ++i)
        value = (value ^ prefix[i]) * 16777619u;
    for (u32 i = 0; i < maze->n * maze->n; ++i)
        value = (value ^ maze->walls[i]) * 16777619u;
    return value;
}

static int decode_official_maze(Maze *maze, u32 index) {
    if (index >= OFFICIAL_MAZE_COUNT) return 0;
    const OfficialMazeRecord *record = &g_official_mazes[index];
    u32 cells = (u32)record->n * record->n;
    if (record->n < 2u || record->n > MAX_N ||
        record->data_offset > ARRAY_COUNT(g_official_maze_data) ||
        cells > ARRAY_COUNT(g_official_maze_data) - record->data_offset) return 0;
    memset(maze, 0, sizeof(*maze));
    maze->n = record->n;
    maze->cell_mm = record->cell_mm;
    maze->goal_x0 = record->goal_x0;
    maze->goal_y0 = record->goal_y0;
    maze->goal_x1 = record->goal_x1;
    maze->goal_y1 = record->goal_y1;
    maze->require_single_goal_entrance = 0u;
    maze->require_full_connectivity = 0u;
    memcpy(maze->walls, g_official_maze_data + record->data_offset, cells);
    return official_maze_hash(maze) == record->data_hash && validate_maze(maze);
}

static int repair_empty_posts(Maze *maze, u64 *random_state) {
    u32 n = maze->n;
    for (u32 py = 1u; py < n; ++py) {
        for (u32 px = 1u; px < n; ++px) {
            if ((px == n / 2u && py == n / 2u) ||
                interior_post_has_wall(maze, px, py)) continue;
            u32 cells[4] = {
                (py - 1u) * n + (px - 1u),
                (py - 1u) * n + px,
                (py - 1u) * n + (px - 1u),
                py * n + (px - 1u)
            };
            u32 directions[4] = {0u, 0u, 1u, 1u};
            u32 start = (u32)(rng_next(random_state) & 3ull);
            int repaired = 0;
            for (u32 offset = 0; offset < 4u; ++offset) {
                u32 pick = (start + offset) & 3u;
                u32 cell = cells[pick];
                u32 direction = directions[pick];
                i32 nx = (i32)(cell % n) + g_dx4[direction];
                i32 ny = (i32)(cell / n) + g_dy4[direction];
                if (nx < 0 || ny < 0 || nx >= (i32)n || ny >= (i32)n) continue;
                u32 neighbor = (u32)ny * n + (u32)nx;
                if (is_goal_cell(maze, cell) || is_goal_cell(maze, neighbor) ||
                    !is_start_allowed_edge(n, cell, neighbor) ||
                    (maze->walls[cell] & g_bit4[direction])) continue;
                add_edge(maze, cell, direction);
                if (maze_is_fully_connected(maze)) {
                    repaired = 1;
                    break;
                }
                carve_edge(maze, cell, direction);
            }
            if (!repaired) return 0;
        }
    }
    return 1;
}

static int generate_maze_attempt(Maze *maze, u32 n, u32 cell_mm,
                                 u64 display_seed, u64 working_seed) {
    if (n != 16u && n != 32u) return 0;
    maze->n = n;
    maze->cell_mm = cell_mm;
    maze->seed = display_seed;
    maze->goal_x0 = maze->goal_y0 = (u8)(n / 2u - 1u);
    maze->goal_x1 = maze->goal_y1 = (u8)(n / 2u);
    maze->require_single_goal_entrance = 1u;
    maze->require_full_connectivity = 1u;
    maze->require_peg_walls = 1u;
    maze->require_wall_follower_trap = 1u;
    u32 cells = n * n;
    for (u32 i = 0; i < cells; ++i) maze->walls[i] = 0x0fu;
    memset(g_visit, 0, cells);

    u64 random_state = working_seed ? working_seed : 1ull;
    u32 stack_size = 1u;
    g_stack[0] = 0u;
    g_visit[0] = 1u;
    while (stack_size) {
        u32 current = g_stack[stack_size - 1u];
        u32 cx = current % n;
        u32 cy = current / n;
        u32 candidates[4];
        u32 directions[4];
        u32 count = 0;
        for (u32 d = 0; d < 4u; ++d) {
            i32 nx = (i32)cx + g_dx4[d];
            i32 ny = (i32)cy + g_dy4[d];
            if (nx < 0 || ny < 0 || nx >= (i32)n || ny >= (i32)n) continue;
            u32 next = (u32)ny * n + (u32)nx;
            if (is_goal_cell(maze, next) || g_visit[next]) continue;
            if (!is_start_allowed_edge(n, current, next)) continue;
            candidates[count] = next;
            directions[count] = d;
            ++count;
        }
        if (!count) {
            --stack_size;
            continue;
        }
        u32 choice = (u32)(rng_next(&random_state) % count);
        u32 next = candidates[choice];
        carve_edge(maze, current, directions[choice]);
        g_visit[next] = 1u;
        g_stack[stack_size++] = (u16)next;
    }

    for (u32 i = 0; i < cells; ++i) {
        if (!is_goal_cell(maze, i) && !g_visit[i]) return 0;
    }

    u32 g0 = n / 2u - 1u;
    u32 g1 = n / 2u;
    carve_edge(maze, g0 * n + g0, 1u);
    carve_edge(maze, g0 * n + g0, 0u);
    carve_edge(maze, g0 * n + g1, 0u);
    carve_edge(maze, g1 * n + g0, 1u);

    u32 perimeter_cells[8] = {
        g0 * n + g0, g1 * n + g0,
        g0 * n + g1, g1 * n + g1,
        g0 * n + g0, g0 * n + g1,
        g1 * n + g0, g1 * n + g1
    };
    u32 perimeter_dirs[8] = {3u, 3u, 1u, 1u, 2u, 2u, 0u, 0u};
    u32 entrance = (u32)(rng_next(&random_state) & 7ull);
    maze->goal_entrance_cell = perimeter_cells[entrance];
    maze->goal_entrance_dir = perimeter_dirs[entrance];
    carve_edge(maze, maze->goal_entrance_cell, maze->goal_entrance_dir);

    for (u32 y = 0; y < n; ++y) {
        for (u32 x = 0; x < n; ++x) {
            u32 c = y * n + x;
            if (is_goal_cell(maze, c)) continue;
            for (u32 d = 0; d < 2u; ++d) {
                i32 nx = (i32)x + g_dx4[d];
                i32 ny = (i32)y + g_dy4[d];
                if (nx < 0 || ny < 0 || nx >= (i32)n || ny >= (i32)n) continue;
                u32 next = (u32)ny * n + (u32)nx;
                if (is_goal_cell(maze, next)) continue;
                if (!is_start_allowed_edge(n, c, next)) continue;
                if ((maze->walls[c] & g_bit4[d]) && (rng_next(&random_state) % 100ull) < 14ull)
                    carve_edge(maze, c, d);
            }
        }
    }
    if (!repair_empty_posts(maze, &random_state)) return 0;
    return validate_maze(maze);
}

static int generate_maze(Maze *maze, u32 n, u32 cell_mm, u64 seed) {
    u64 base = seed ? seed : 1ull;
    for (u32 attempt = 0; attempt < 512u; ++attempt) {
        u64 working = base ^ (0x9e3779b97f4a7c15ull * (u64)(attempt + 1u));
        if (generate_maze_attempt(maze, n, cell_mm, base, working)) return 1;
    }
    return 0;
}

static u64 isqrt_u64(u64 value) {
    u64 result = 0;
    u64 bit = 1ull << 62u;
    while (bit > value) bit >>= 2u;
    while (bit) {
        if (value >= result + bit) {
            value -= result + bit;
            result = (result >> 1u) + bit;
        } else {
            result >>= 1u;
        }
        bit >>= 2u;
    }
    return result;
}

static void read_solver_params(SolverParams *p) {
    int profile = (int)SendMessageW(g_app.profile, CB_GETCURSEL, 0, 0);
    if (g_app.official_loaded) {
        p->n = g_app.maze.n;
        p->cell_mm = g_app.maze.cell_mm;
    } else {
        p->n = profile == 1 ? 32u : 16u;
        p->cell_mm = profile == 1 ? 90u : 180u;
    }
    p->straight_v = get_edit_u32(g_app.param_edits[PARAM_STRAIGHT_V], 3500u, 100u, 20000u);
    p->diagonal_v = get_edit_u32(g_app.param_edits[PARAM_DIAGONAL_V], 3000u, 100u, 20000u);
    p->turn_v = get_edit_u32(g_app.param_edits[PARAM_TURN_V], 1800u, 50u, 10000u);
    p->accel = get_edit_u32(g_app.param_edits[PARAM_ACCEL], 12000u, 100u, 100000u);
    p->brake = get_edit_u32(g_app.param_edits[PARAM_BRAKE], 14000u, 100u, 100000u);
    p->jerk = get_edit_u32(g_app.param_edits[PARAM_JERK], 120000u, 1000u, 2000000u);
    p->mu_milli = get_edit_milli(g_app.param_edits[PARAM_MU_X], 900u, 100u, 3000u);
    p->mu_y_milli = get_edit_milli(g_app.param_edits[PARAM_MU_Y], 1050u, 100u, 3000u);
    p->slip_milli = get_edit_milli(g_app.param_edits[PARAM_SLIP], 80u, 0u, 2000u);
    p->battery_mv = get_edit_u32(g_app.param_edits[PARAM_BATTERY], 8400u, 3000u, 30000u);
    p->sag_mv_s = get_edit_u32(g_app.param_edits[PARAM_SAG], 90u, 0u, 5000u);
    p->battery_r_mohm = get_edit_u32(g_app.param_edits[PARAM_BATTERY_R], 45u, 0u, 2000u);
    p->battery_capacity_mah = get_edit_u32(g_app.param_edits[PARAM_BATTERY_CAPACITY], 350u, 10u, 20000u);
    p->mass_g = get_edit_u32(g_app.param_edits[PARAM_MASS], 100u, 10u, 2000u);
    p->chassis_length_mm = get_edit_u32(g_app.param_edits[PARAM_CHASSIS_LENGTH], 90u, 20u, 300u);
    p->track_mm = get_edit_u32(g_app.param_edits[PARAM_TRACK], 72u, 10u, 200u);
    p->wheel_radius_mm = get_edit_u32(g_app.param_edits[PARAM_WHEEL_RADIUS], 12u, 2u, 100u);
    p->wheel_mass_g = get_edit_u32(g_app.param_edits[PARAM_WHEEL_MASS], 5u, 1u, 100u);
    p->cg_height_mm = get_edit_u32(g_app.param_edits[PARAM_CG_HEIGHT], 15u, 0u, 100u);
    p->motor_r_mohm = get_edit_u32(g_app.param_edits[PARAM_MOTOR_R], 320u, 10u, 5000u);
    p->motor_kt_mnm_a = get_edit_u32(g_app.param_edits[PARAM_MOTOR_KT], 18u, 1u, 500u);
    p->motor_ke_mv_rad = get_edit_u32(g_app.param_edits[PARAM_MOTOR_KE], 2u, 0u, 100u);
    p->current_limit_ma = get_edit_u32(g_app.param_edits[PARAM_CURRENT_LIMIT], 6000u, 100u, 100000u);
    p->controller_kp_mv_rad = get_edit_u32(g_app.param_edits[PARAM_CONTROLLER_KP], 40u, 1u, 5000u);
    p->turn_radius_mm = get_edit_u32(g_app.param_edits[PARAM_RADIUS], 45u, 5u, 500u);
    p->max_run = get_edit_u32(g_app.param_edits[PARAM_MAX_RUN], 8u, 1u, 8u);
    p->dt_us = get_edit_u32(g_app.param_edits[PARAM_DT_US], 500u, 50u, 5000u);
    if (SendMessageW(g_app.multirun, BM_GETCHECK, 0, 0) != BST_CHECKED) p->max_run = 1u;
    if (SendMessageW(g_app.diagonal, BM_GETCHECK, 0, 0) != BST_CHECKED) p->diagonal_v = 0u;

    u64 lateral_square = ((u64)p->mu_milli * 9810ull * p->turn_radius_mm) / 1000ull;
    u32 lateral_limit = (u32)isqrt_u64(lateral_square);
    if (lateral_limit && p->turn_v > lateral_limit) p->turn_v = lateral_limit;
}

static int step_ok_host(const Maze *maze, i32 x, i32 y, u32 heading) {
    static const i32 dx[8] = {0, 1, 1, 1, 0, -1, -1, -1};
    static const i32 dy[8] = {1, 1, 0, -1, -1, -1, 0, 1};
    i32 nx = x + dx[heading];
    i32 ny = y + dy[heading];
    if (nx < 0 || ny < 0 || nx >= (i32)maze->n || ny >= (i32)maze->n) return 0;
    u32 c = (u32)y * maze->n + (u32)x;
    u8 w = maze->walls[c];
    if (!(heading & 1u)) {
        u8 bit = heading == 0u ? WALL_N : heading == 2u ? WALL_E : heading == 4u ? WALL_S : WALL_W;
        return !(w & bit);
    }
    if (heading == 1u) return !(w & (WALL_N | WALL_E)) && !(maze->walls[c + 1u] & WALL_N) && !(maze->walls[c + maze->n] & WALL_E);
    if (heading == 3u) return !(w & (WALL_S | WALL_E)) && !(maze->walls[c + 1u] & WALL_S) && !(maze->walls[c - maze->n] & WALL_E);
    if (heading == 5u) return !(w & (WALL_S | WALL_W)) && !(maze->walls[c - 1u] & WALL_S) && !(maze->walls[c - maze->n] & WALL_W);
    return !(w & (WALL_N | WALL_W)) && !(maze->walls[c - 1u] & WALL_N) && !(maze->walls[c + maze->n] & WALL_W);
}

static u32 voltage_scaled_host(u32 value, u64 elapsed_us, const SolverParams *p) {
    u64 drop = ((u64)p->sag_mv_s * elapsed_us) / 1000000ull;
    u32 floor_voltage = p->battery_mv * 60u / 100u;
    u32 voltage = drop >= p->battery_mv ? floor_voltage : p->battery_mv - (u32)drop;
    if (voltage < floor_voltage) voltage = floor_voltage;
    u64 scaled = (u64)value * voltage / (p->battery_mv ? p->battery_mv : 1u);
    return scaled ? (u32)scaled : 1u;
}

static u64 forward_cost_host(u32 heading, u32 steps, u64 elapsed_us, const SolverParams *p) {
    u64 distance = (u64)p->cell_mm * steps;
    u32 vmax = (heading & 1u) ? p->diagonal_v : p->straight_v;
    if (heading & 1u) distance = (distance * 1414ull + 500ull) / 1000ull;
    vmax = voltage_scaled_host(vmax, elapsed_us, p);
    u32 traction_accel = p->mu_milli * 9810u / 1000u;
    u32 accel = voltage_scaled_host(p->accel, elapsed_us, p);
    u32 brake = voltage_scaled_host(p->brake, elapsed_us, p);
    if (accel > traction_accel) accel = traction_accel;
    if (brake > traction_accel) brake = traction_accel;
    if (!accel) accel = 1u;
    if (!brake) brake = 1u;
    u64 cruise = (distance * 1000000ull + vmax - 1u) / vmax;
    u64 ramp = (u64)vmax * 500000ull / accel + (u64)vmax * 500000ull / brake;
    u32 jerk = p->jerk ? p->jerk : 1u;
    u64 jerk_penalty = (u64)(accel + brake) * 500000ull / jerk;
    return cruise + ramp + jerk_penalty + 1ull;
}

static u64 turn_cost_host(u32 new_heading, u32 angle_mrad, u64 elapsed_us, const SolverParams *p) {
    u64 distance = angle_mrad == 3142u ? 0u : p->cell_mm;
    if (new_heading & 1u) distance = (distance * 1414ull + 500ull) / 1000ull;
    u32 speed = voltage_scaled_host(p->turn_v, elapsed_us, p);
    u64 travel = (distance * 1000000ull + speed - 1u) / speed;
    u64 arc = ((u64)p->turn_radius_mm * angle_mrad * 1000ull + speed - 1u) / speed;
    arc = (arc * (1000u + p->slip_milli) + 999u) / 1000u;
    u32 jerk = p->jerk ? p->jerk : 1u;
    u64 settle = (u64)p->accel * 250000ull / jerk;
    return travel + arc + settle + 1ull;
}

static int evaluate_action(u32 source, u32 action, u64 source_time,
                           const SolverParams *p, const Maze *maze,
                           u32 *target_out, u64 *cost_out) {
    static const i32 dx[8] = {0, 1, 1, 1, 0, -1, -1, -1};
    static const i32 dy[8] = {1, 1, 0, -1, -1, -1, 0, 1};
    u32 heading = source & 7u;
    u32 cell = source >> 3u;
    i32 x = (i32)(cell % p->n);
    i32 y = (i32)(cell / p->n);
    if (action < 8u) {
        if ((heading & 1u) && !p->diagonal_v) return 0;
        u32 steps = action + 1u;
        for (u32 i = 0; i < steps; ++i) {
            if (!step_ok_host(maze, x, y, heading)) return 0;
            x += dx[heading]; y += dy[heading];
        }
        *target_out = ((((u32)y * p->n) + (u32)x) << 3u) | heading;
        *cost_out = forward_cost_host(heading, steps, source_time, p);
        return 1;
    }
    if (action >= 8u && action <= 13u) {
        static const i32 delta[6] = {-1, 1, -2, 2, -3, 3};
        static const u32 angle[6] = {785u, 785u, 1571u, 1571u, 2356u, 2356u};
        u32 i = action - 8u;
        u32 nh = (u32)((i32)heading + delta[i] + 8) & 7u;
        if ((nh & 1u) && !p->diagonal_v) return 0;
        if (!step_ok_host(maze, x, y, nh)) return 0;
        x += dx[nh]; y += dy[nh];
        *target_out = ((((u32)y * p->n) + (u32)x) << 3u) | nh;
        *cost_out = turn_cost_host(nh, angle[i], source_time, p);
        return 1;
    }
    if (action == 14u) {
        u32 nh = (heading + 4u) & 7u;
        *target_out = (cell << 3u) | nh;
        *cost_out = turn_cost_host(nh, 3142u, source_time, p);
        return 1;
    }
    return 0;
}

typedef struct DynamicsState {
    float x;
    float y;
    float heading;
    float vx;
    float vy;
    float yaw_rate;
    float wheel_l;
    float wheel_r;
    float current_l;
    float current_r;
    float command_v;
    float command_yaw_rate;
    float command_accel;
    float actual_accel_previous;
    float soc;
    float voltage;
    float time;
} DynamicsState;

static float f_abs(float value) { return value < 0.0f ? -value : value; }
static float f_min(float a, float b) { return a < b ? a : b; }
static float f_max(float a, float b) { return a > b ? a : b; }
static float f_clamp(float value, float low, float high) {
    return value < low ? low : value > high ? high : value;
}

static float f_sqrt(float value) {
    if (value <= 0.0f) return 0.0f;
    union { float f; u32 u; } seed;
    seed.f = value;
    seed.u = (seed.u >> 1u) + 0x1fc00000u;
    float estimate = seed.f;
    for (u32 i = 0; i < 5u; ++i)
        estimate = 0.5f * (estimate + value / estimate);
    return estimate;
}

static float wrap_pi(float angle) {
    const float pi = 3.14159265358979323846f;
    const float two_pi = 6.28318530717958647692f;
    while (angle > pi) angle -= two_pi;
    while (angle < -pi) angle += two_pi;
    return angle;
}

static float f_sin(float angle) {
    const float pi = 3.14159265358979323846f;
    const float half_pi = 1.57079632679489661923f;
    float x = wrap_pi(angle);
    if (x > half_pi) x = pi - x;
    else if (x < -half_pi) x = -pi - x;
    float x2 = x * x;
    return x * (1.0f + x2 * (-0.1666666716f + x2 * (0.0083333477f - x2 * 0.0001984090f)));
}

static float f_cos(float angle) {
    return f_sin(angle + 1.57079632679489661923f);
}

static float f_log(float value) {
    if (value <= 0.0f) return -80.0f;
    union { float f; u32 u; } bits;
    bits.f = value;
    i32 exponent = (i32)((bits.u >> 23u) & 255u) - 127;
    bits.u = (bits.u & 0x007fffffu) | 0x3f800000u;
    float mantissa = bits.f;
    float y = (mantissa - 1.0f) / (mantissa + 1.0f);
    float y2 = y * y;
    float term = y;
    float series = term;
    term *= y2; series += term / 3.0f;
    term *= y2; series += term / 5.0f;
    term *= y2; series += term / 7.0f;
    term *= y2; series += term / 9.0f;
    term *= y2; series += term / 11.0f;
    term *= y2; series += term / 13.0f;
    term *= y2; series += term / 15.0f;
    term *= y2; series += term / 17.0f;
    term *= y2; series += term / 19.0f;
    return 2.0f * series + (float)exponent * 0.6931471805599453f;
}

static float f_atan2(float y, float x) {
    const float pi_over_4 = 0.78539816339744830962f;
    const float three_pi_over_4 = 2.35619449019234492885f;
    float abs_y = f_abs(y) + 0.0000001f;
    float ratio;
    float angle;
    if (x >= 0.0f) {
        ratio = (x - abs_y) / (x + abs_y);
        angle = pi_over_4 + (0.1963f * ratio * ratio - 0.9817f) * ratio;
    } else {
        ratio = (x + abs_y) / (abs_y - x);
        angle = three_pi_over_4 + (0.1963f * ratio * ratio - 0.9817f) * ratio;
    }
    return y < 0.0f ? -angle : angle;
}

#include "fem_core.h"

static FemResult g_fem_result;

static void node_reference_position_m(const AppState *app, const DynamicsParams *p,
                                      u32 node, float *x_out, float *y_out) {
    u32 cell = node >> 3u;
    *x_out = ((float)(cell % app->maze.n) + 0.5f) * p->cell_m;
    *y_out = ((float)(cell / app->maze.n) + 0.5f) * p->cell_m;
}

static void hermite_reference(float t,
                              float start_x, float start_y,
                              float tangent_start_x, float tangent_start_y,
                              float end_x, float end_y,
                              float tangent_end_x, float tangent_end_y,
                              float *x, float *y,
                              float *dx, float *dy,
                              float *ddx, float *ddy) {
    float t2 = t * t;
    float t3 = t2 * t;
    float h00 = 2.0f * t3 - 3.0f * t2 + 1.0f;
    float h10 = t3 - 2.0f * t2 + t;
    float h01 = -2.0f * t3 + 3.0f * t2;
    float h11 = t3 - t2;
    float dh00 = 6.0f * t2 - 6.0f * t;
    float dh10 = 3.0f * t2 - 4.0f * t + 1.0f;
    float dh01 = -dh00;
    float dh11 = 3.0f * t2 - 2.0f * t;
    float ddh00 = 12.0f * t - 6.0f;
    float ddh10 = 6.0f * t - 4.0f;
    float ddh01 = -ddh00;
    float ddh11 = 6.0f * t - 2.0f;
    *x = h00 * start_x + h10 * tangent_start_x +
         h01 * end_x + h11 * tangent_end_x;
    *y = h00 * start_y + h10 * tangent_start_y +
         h01 * end_y + h11 * tangent_end_y;
    *dx = dh00 * start_x + dh10 * tangent_start_x +
          dh01 * end_x + dh11 * tangent_end_x;
    *dy = dh00 * start_y + dh10 * tangent_start_y +
          dh01 * end_y + dh11 * tangent_end_y;
    *ddx = ddh00 * start_x + ddh10 * tangent_start_x +
           ddh01 * end_x + ddh11 * tangent_end_x;
    *ddy = ddh00 * start_y + ddh10 * tangent_start_y +
           ddh01 * end_y + ddh11 * tangent_end_y;
}

static void build_dynamics_params(const SolverParams *source, DynamicsParams *target) {
    target->cell_m = (float)source->cell_mm * 0.001f;
    target->straight_v = (float)source->straight_v * 0.001f;
    target->diagonal_v = (float)source->diagonal_v * 0.001f;
    target->turn_v = (float)source->turn_v * 0.001f;
    target->accel = (float)source->accel * 0.001f;
    target->brake = (float)source->brake * 0.001f;
    target->jerk = (float)source->jerk * 0.001f;
    target->mu_x = (float)source->mu_milli * 0.001f;
    target->mu_y = (float)source->mu_y_milli * 0.001f;
    target->slip_peak = f_max((float)source->slip_milli * 0.001f, 0.005f);
    target->battery_v = (float)source->battery_mv * 0.001f;
    target->battery_sag_v_s = (float)source->sag_mv_s * 0.001f;
    target->battery_r = (float)source->battery_r_mohm * 0.001f;
    target->battery_capacity_c = (float)source->battery_capacity_mah * 3.6f;
    target->mass = (float)source->mass_g * 0.001f;
    target->chassis_length = (float)source->chassis_length_mm * 0.001f;
    target->track = (float)source->track_mm * 0.001f;
    target->wheel_radius = (float)source->wheel_radius_mm * 0.001f;
    target->wheel_mass = (float)source->wheel_mass_g * 0.001f;
    target->cg_height = (float)source->cg_height_mm * 0.001f;
    target->motor_r = (float)source->motor_r_mohm * 0.001f;
    target->motor_kt = (float)source->motor_kt_mnm_a * 0.001f;
    target->motor_ke = (float)source->motor_ke_mv_rad * 0.001f;
    target->current_limit = (float)source->current_limit_ma * 0.001f;
    target->controller_kp = (float)source->controller_kp_mv_rad * 0.001f;
    target->turn_radius = (float)source->turn_radius_mm * 0.001f;
    target->wheel_inertia = 0.5f * target->wheel_mass * target->wheel_radius * target->wheel_radius;
    target->yaw_inertia = target->mass *
        (target->chassis_length * target->chassis_length + target->track * target->track) / 12.0f;
    target->dt = (float)source->dt_us * 0.000001f;
    target->max_steps = 300000u;
}

static void dynamics_candidate_scales(u32 candidate, float *straight_scale, float *turn_scale) {
    u32 straight_index = candidate & 63u;
    u32 turn_index = (candidate >> 6u) & 63u;
    *straight_scale = 0.10f + 0.90f * (float)straight_index / 63.0f;
    *turn_scale = 0.10f + 0.90f * (float)turn_index / 63.0f;
}

static u32 build_segment_speed_envelope(const AppState *app, const DynamicsParams *p,
                                        float straight_scale, float turn_scale) {
    const float gravity = 9.81f;
    u32 action_count = app->path_count > 0u ? app->path_count - 1u : 0u;
    if (!action_count || action_count >= MAX_ROUTE) return 0u;

    float curve_friction_limit = f_sqrt(f_max(p->mu_y, 0.05f) * gravity *
                                               f_max(p->turn_radius, 0.005f));
    for (u32 i = 0; i < action_count; ++i) {
        u32 action = app->path_actions[i];
        float source_x, source_y, target_x, target_y;
        node_reference_position_m(app, p, app->path_nodes[i], &source_x, &source_y);
        node_reference_position_m(app, p, app->path_nodes[i + 1u], &target_x, &target_y);
        float dx = target_x - source_x;
        float dy = target_y - source_y;
        float distance = f_sqrt(dx * dx + dy * dy);
        float limit = 0.0f;
        if (action < 8u) {
            u32 heading = app->path_nodes[i] & 7u;
            limit = (heading & 1u ? p->diagonal_v : p->straight_v) * straight_scale;
        } else {
            limit = f_min(p->turn_v * turn_scale, curve_friction_limit * 0.82f);
        }
        float short_primitive_cap = f_sqrt(2.0f * f_min(p->accel, p->brake) *
                                                 f_max(distance, 0.001f) *
                                                 (action < 8u ? 0.24f : 0.16f));
        limit = f_min(limit, short_primitive_cap);
        g_segment_distance[i] = distance;
        g_segment_speed_limit[i] = f_max(limit, 0.025f);
    }

    g_boundary_speed[0] = 0.0f;
    g_boundary_speed[action_count] = 0.0f;
    for (u32 i = 1u; i < action_count; ++i) {
        float limit = f_min(g_segment_speed_limit[i - 1u], g_segment_speed_limit[i]);
        if (app->path_actions[i - 1u] >= 8u || app->path_actions[i] >= 8u) limit = 0.0f;
        g_boundary_speed[i] = limit;
    }

    float envelope_accel = f_max(p->accel * 0.72f, 0.05f);
    float envelope_brake = f_max(p->brake * 0.72f, 0.05f);
    for (u32 pass = 0; pass < 3u; ++pass) {
        for (u32 end = action_count; end > 0u; --end) {
            u32 i = end - 1u;
            float reachable = f_sqrt(g_boundary_speed[i + 1u] * g_boundary_speed[i + 1u] +
                                     2.0f * envelope_brake * g_segment_distance[i]);
            if (g_boundary_speed[i] > reachable) g_boundary_speed[i] = reachable;
        }
        for (u32 i = 0; i < action_count; ++i) {
            float reachable = f_sqrt(g_boundary_speed[i] * g_boundary_speed[i] +
                                     2.0f * envelope_accel * g_segment_distance[i]);
            if (g_boundary_speed[i + 1u] > reachable) g_boundary_speed[i + 1u] = reachable;
        }
    }
    return action_count;
}

static void trajectory_reset(AppState *app) {
    app->trajectory_count = 0;
    if (!app->trajectory) {
        app->trajectory_capacity = MAX_TRAJECTORY_SAMPLES;
        app->trajectory = (TrajectorySample *)VirtualAlloc(
            0, (SIZE_T)app->trajectory_capacity * sizeof(TrajectorySample),
            MEM_COMMIT | MEM_RESERVE, PAGE_READWRITE);
        if (!app->trajectory) app->trajectory_capacity = 0;
    }
}

static void trajectory_append(AppState *app, const DynamicsState *state, float slip) {
    if (!app->trajectory || app->trajectory_count >= app->trajectory_capacity) return;
    TrajectorySample *sample = &app->trajectory[app->trajectory_count++];
    sample->x_m = state->x;
    sample->y_m = state->y;
    sample->heading_rad = state->heading;
    sample->speed_m_s = f_sqrt(state->vx * state->vx + state->vy * state->vy);
    sample->lateral_m_s = state->vy;
    sample->yaw_rate_rad_s = state->yaw_rate;
    sample->slip = slip;
    sample->voltage = state->voltage;
}

static void motor_current_pair(const DynamicsParams *p, float terminal_voltage,
                               float target_l, float target_r,
                               float wheel_l, float wheel_r,
                               float *current_l, float *current_r) {
    float command_l = p->motor_ke * target_l + p->controller_kp * (target_l - wheel_l);
    float command_r = p->motor_ke * target_r + p->controller_kp * (target_r - wheel_r);
    command_l = f_clamp(command_l, -terminal_voltage, terminal_voltage);
    command_r = f_clamp(command_r, -terminal_voltage, terminal_voltage);
    *current_l = f_clamp((command_l - p->motor_ke * wheel_l) / p->motor_r,
                         -p->current_limit, p->current_limit);
    *current_r = f_clamp((command_r - p->motor_ke * wheel_r) / p->motor_r,
                         -p->current_limit, p->current_limit);
}

static int simulate_dynamics_cpu(AppState *app, const DynamicsParams *p,
                                 u32 candidate, int record,
                                 DynamicsResult *result) {
    const float pi = 3.14159265358979323846f;
    const float gravity = 9.81f;
    float straight_scale, turn_scale;
    dynamics_candidate_scales(candidate, &straight_scale, &turn_scale);
    memset(result, 0, sizeof(*result));
    if (app->path_count < 2u || p->mass <= 0.0f || p->track <= 0.0f ||
        p->wheel_radius <= 0.0f || p->motor_r <= 0.0f ||
        p->wheel_inertia <= 0.0f || p->yaw_inertia <= 0.0f || p->dt <= 0.0f) {
        result->failure_reason = DYN_FAIL_INVALID_PARAMETERS;
        return 0;
    }
    u32 action_count = build_segment_speed_envelope(app, p, straight_scale, turn_scale);
    if (!action_count) {
        result->failure_reason = DYN_FAIL_INVALID_PARAMETERS;
        return 0;
    }

    DynamicsState state;
    memset(&state, 0, sizeof(state));
    node_reference_position_m(app, p, app->path_nodes[0], &state.x, &state.y);
    state.heading = 0.5f * pi;
    state.soc = 1.0f;
    state.voltage = p->battery_v;
    if (record) {
        trajectory_reset(app);
        trajectory_append(app, &state, 0.0f);
    }

    u32 total_steps = 0;
    u32 cone_hits = 0;
    float max_slip = 0.0f;
    float max_lateral = 0.0f;
    int valid = 1;

    for (u32 action_index = 0; action_index < action_count && valid; ++action_index) {
        u32 action = app->path_actions[action_index];
        float distance = g_segment_distance[action_index];
        float target_angle = 0.0f;
        float max_speed = g_segment_speed_limit[action_index];
        float exit_speed = g_boundary_speed[action_index + 1u];
        if (action >= 8u && action <= 13u) {
            static const float angle_magnitude[3] = {
                0.7853981633974483f, 1.5707963267948966f, 2.356194490192345f
            };
            target_angle = angle_magnitude[(action - 8u) / 2u] * ((action & 1u) ? -1.0f : 1.0f);
        } else if (action == 14u) {
            target_angle = pi;
            exit_speed = 0.0f;
        }

        float segment_start_x = state.x;
        float segment_start_y = state.y;
        float segment_end_x, segment_end_y;
        node_reference_position_m(app, p, app->path_nodes[action_index + 1u],
                                  &segment_end_x, &segment_end_y);
        float target_reference_heading = 0.5f * pi -
            (float)(app->path_nodes[action_index + 1u] & 7u) * 0.25f * pi;
        float tangent_magnitude = distance * (action >= 8u ? 0.55f : 1.0f);
        float tangent_start_x = f_cos(state.heading) * tangent_magnitude;
        float tangent_start_y = f_sin(state.heading) * tangent_magnitude;
        float tangent_end_x = f_cos(target_reference_heading) * tangent_magnitude;
        float tangent_end_y = f_sin(target_reference_heading) * tangent_magnitude;
        float segment_heading = state.heading;
        float progress = 0.0f;
        float lateral_offset = 0.0f;
        u32 segment_steps = 0;
        int completed = 0;
        while (total_steps < p->max_steps && segment_steps < 40000u) {
            float yaw_progress = wrap_pi(state.heading - segment_heading);
            float remaining = f_max(distance - progress, 0.0f);
            float fraction = distance > 0.0f ? f_clamp(progress / distance, 0.0f, 1.0f) : 0.0f;
            float reference_x = segment_start_x;
            float reference_y = segment_start_y;
            float reference_dx = f_cos(state.heading);
            float reference_dy = f_sin(state.heading);
            float reference_ddx = 0.0f;
            float reference_ddy = 0.0f;
            if (distance > 0.0f) {
                hermite_reference(fraction,
                                  segment_start_x, segment_start_y,
                                  tangent_start_x, tangent_start_y,
                                  segment_end_x, segment_end_y,
                                  tangent_end_x, tangent_end_y,
                                  &reference_x, &reference_y,
                                  &reference_dx, &reference_dy,
                                  &reference_ddx, &reference_ddy);
            }
            float reference_length = f_max(f_sqrt(reference_dx * reference_dx +
                                                   reference_dy * reference_dy), 0.000001f);
            float path_heading = f_atan2(reference_dy, reference_dx);
            float path_curvature = (reference_dx * reference_ddy -
                                    reference_dy * reference_ddx) /
                                   (reference_length * reference_length * reference_length);
            float target_v = 0.0f;
            float target_yaw_rate;
            float heading_error;
            if (distance > 0.0f) {
                float controlled_brake = f_max(p->brake * 0.45f, 0.05f);
                float stop_speed = f_sqrt(exit_speed * exit_speed +
                                          2.0f * controlled_brake * remaining);
                target_v = f_min(max_speed, stop_speed);
                if (remaining > 0.0005f) target_v = f_max(target_v, 0.04f);
                if (remaining < 0.00025f) target_v = exit_speed;
                heading_error = wrap_pi(path_heading - state.heading);
                float cross_track_correction = 10.0f * lateral_offset /
                                               f_max(p->cell_m, 0.05f) *
                                               (1.0f - fraction);
                target_yaw_rate = path_curvature * target_v +
                                  14.0f * heading_error - cross_track_correction;
                float tracking_rate_cap = f_max(1.5f, 1.75f * max_speed /
                                                        f_max(p->turn_radius, 0.005f));
                target_yaw_rate = f_clamp(target_yaw_rate,
                                          -tracking_rate_cap, tracking_rate_cap);
            } else {
                heading_error = wrap_pi(target_angle - yaw_progress);
                float yaw_accel = 2.0f * p->brake / f_max(p->track, 0.01f);
                float stop_rate = f_sqrt(2.0f * yaw_accel * f_abs(heading_error));
                float max_rate = p->turn_v * turn_scale / f_max(p->turn_radius, 0.005f);
                target_yaw_rate = (heading_error < 0.0f ? -1.0f : 1.0f) * f_min(max_rate, stop_rate);
            }

            float jerk_delta = p->jerk * p->dt;
            float desired_command_accel = f_clamp((target_v - state.command_v) / p->dt,
                                                   -p->brake, p->accel);
            state.command_accel = f_clamp(desired_command_accel,
                state.command_accel - jerk_delta, state.command_accel + jerk_delta);
            state.command_v += state.command_accel * p->dt;
            if (state.command_v < 0.0f) state.command_v = 0.0f;
            if (state.command_v > max_speed && distance > 0.0f) state.command_v = max_speed;
            float yaw_accel_limit = 2.0f * f_max(p->accel, p->brake) /
                                    f_max(p->track, 0.01f);
            float yaw_delta = f_clamp(target_yaw_rate - state.command_yaw_rate,
                                      -yaw_accel_limit * p->dt,
                                      yaw_accel_limit * p->dt);
            state.command_yaw_rate += yaw_delta;

            float half_track = 0.5f * p->track;
            float target_wheel_l = (state.command_v - state.command_yaw_rate * half_track) /
                                   p->wheel_radius;
            float target_wheel_r = (state.command_v + state.command_yaw_rate * half_track) /
                                   p->wheel_radius;
            float open_voltage = p->battery_v - p->battery_sag_v_s * state.time;
            open_voltage = f_max(open_voltage, 0.60f * p->battery_v);
            open_voltage *= 0.70f + 0.30f * f_clamp(state.soc, 0.0f, 1.0f);
            motor_current_pair(p, open_voltage, target_wheel_l, target_wheel_r,
                               state.wheel_l, state.wheel_r,
                               &state.current_l, &state.current_r);
            float terminal_voltage = open_voltage - p->battery_r *
                (f_abs(state.current_l) + f_abs(state.current_r));
            terminal_voltage = f_max(terminal_voltage, 0.50f * open_voltage);
            motor_current_pair(p, terminal_voltage, target_wheel_l, target_wheel_r,
                               state.wheel_l, state.wheel_r,
                               &state.current_l, &state.current_r);
            terminal_voltage = open_voltage - p->battery_r *
                (f_abs(state.current_l) + f_abs(state.current_r));
            state.voltage = f_max(terminal_voltage, 0.50f * open_voltage);
            float torque_l = p->motor_kt * state.current_l;
            float torque_r = p->motor_kt * state.current_r;

            float normal_total = p->mass * gravity;
            float load_transfer = p->mass * state.vx * state.yaw_rate * p->cg_height /
                                  (2.0f * f_max(p->track, 0.001f));
            float min_load = 0.05f * normal_total;
            float normal_l = f_clamp(0.5f * normal_total - load_transfer,
                                     min_load, normal_total - min_load);
            float normal_r = normal_total - normal_l;
            float contact_v_l = state.vx - state.yaw_rate * half_track;
            float contact_v_r = state.vx + state.yaw_rate * half_track;
            float denom_l = f_max(f_abs(contact_v_l), 0.20f);
            float denom_r = f_max(f_abs(contact_v_r), 0.20f);
            float stiffness_x_l = p->mu_x * normal_l / p->slip_peak;
            float stiffness_x_r = p->mu_x * normal_r / p->slip_peak;
            float implicit_factor_l = p->dt * stiffness_x_l /
                                      (p->wheel_inertia * denom_l);
            float implicit_factor_r = p->dt * stiffness_x_r /
                                      (p->wheel_inertia * denom_r);
            float predicted_wheel_l = (state.wheel_l + p->dt * torque_l / p->wheel_inertia +
                implicit_factor_l * p->wheel_radius * contact_v_l) /
                (1.0f + implicit_factor_l * p->wheel_radius * p->wheel_radius);
            float predicted_wheel_r = (state.wheel_r + p->dt * torque_r / p->wheel_inertia +
                implicit_factor_r * p->wheel_radius * contact_v_r) /
                (1.0f + implicit_factor_r * p->wheel_radius * p->wheel_radius);
            /* Path-relative lateral closed loop.  Centripetal force alone
               holds an ideal trajectory but cannot remove accumulated tyre
               side velocity or cross-track displacement.  The damped
               second-order correction is part of the requested tyre force,
               so both wheels still obey their independent friction ellipses. */
            float lateral_correction = 0.0f;
            if (distance > 0.0f) {
                float lateral_frequency = f_clamp(
                    7.0f + 0.35f * state.vx / f_max(p->cell_m, 0.05f),
                    7.0f, 24.0f);
                lateral_correction =
                    -lateral_frequency * lateral_frequency * lateral_offset -
                    1.80f * lateral_frequency * state.vy;
                lateral_correction = f_clamp(
                    lateral_correction,
                    -0.55f * f_max(p->mu_y, 0.05f) * gravity,
                     0.55f * f_max(p->mu_y, 0.05f) * gravity);
            }
            float required_lateral_force = distance > 0.0f ?
                p->mass * (state.vx * state.vx * path_curvature +
                           lateral_correction) : 0.0f;
            float force_y_l = required_lateral_force * normal_l / normal_total;
            float force_y_r = required_lateral_force * normal_r / normal_total;
            float lateral_capacity_l = p->mu_y * normal_l;
            float lateral_capacity_r = p->mu_y * normal_r;
            float raw_ny_l = force_y_l / lateral_capacity_l;
            float raw_ny_r = force_y_r / lateral_capacity_r;
            float ny_l = f_clamp(raw_ny_l, -0.98f, 0.98f);
            float ny_r = f_clamp(raw_ny_r, -0.98f, 0.98f);
            force_y_l = ny_l * lateral_capacity_l;
            force_y_r = ny_r * lateral_capacity_r;

            float controlled_accel = f_clamp(state.command_accel +
                                              10.0f * (state.command_v - state.vx),
                                              -p->brake, p->accel);
            float controlled_yaw_accel = f_clamp(
                18.0f * (state.command_yaw_rate - state.yaw_rate),
                -yaw_accel_limit, yaw_accel_limit);
            float requested_force_total = p->mass * controlled_accel;
            float requested_moment = p->yaw_inertia * controlled_yaw_accel;
            float requested_force_difference = requested_moment /
                                               f_max(half_track, 0.001f);
            float requested_force_l = 0.5f *
                (requested_force_total - requested_force_difference);
            float requested_force_r = 0.5f *
                (requested_force_total + requested_force_difference);

            float longitudinal_capacity_l = p->mu_x * normal_l *
                                            f_sqrt(f_max(1.0f - ny_l * ny_l, 0.0f));
            float longitudinal_capacity_r = p->mu_x * normal_r *
                                            f_sqrt(f_max(1.0f - ny_r * ny_r, 0.0f));
            float motor_capacity_l = f_abs(torque_l) / p->wheel_radius +
                                     0.02f * p->mu_x * normal_l;
            float motor_capacity_r = f_abs(torque_r) / p->wheel_radius +
                                     0.02f * p->mu_x * normal_r;
            float available_l = f_min(longitudinal_capacity_l, motor_capacity_l);
            float available_r = f_min(longitudinal_capacity_r, motor_capacity_r);
            float force_x_l = f_clamp(requested_force_l, -available_l, available_l);
            float force_x_r = f_clamp(requested_force_r, -available_r, available_r);

            float nx_l = requested_force_l / (p->mu_x * normal_l);
            float nx_r = requested_force_r / (p->mu_x * normal_r);
            float cone_l = f_sqrt(nx_l * nx_l + raw_ny_l * raw_ny_l);
            float cone_r = f_sqrt(nx_r * nx_r + raw_ny_r * raw_ny_r);
            if (cone_l > 1.0f || f_abs(requested_force_l) > motor_capacity_l) ++cone_hits;
            if (cone_r > 1.0f || f_abs(requested_force_r) > motor_capacity_r) ++cone_hits;

            float slip_l = force_x_l / f_max(stiffness_x_l, 0.001f);
            float slip_r = force_x_r / f_max(stiffness_x_r, 0.001f);
            float predicted_slip_l = (p->wheel_radius * predicted_wheel_l - contact_v_l) / denom_l;
            float predicted_slip_r = (p->wheel_radius * predicted_wheel_r - contact_v_r) / denom_r;
            float slip_abs = f_max(f_max(f_abs(slip_l), f_abs(slip_r)),
                                   0.25f * f_max(f_abs(predicted_slip_l),
                                                f_abs(predicted_slip_r)));
            max_slip = f_max(max_slip, slip_abs);

            state.wheel_l = predicted_wheel_l;
            state.wheel_r = predicted_wheel_r;

            float raw_accel = (force_x_l + force_x_r) / p->mass + state.yaw_rate * state.vy;
            float longitudinal_accel = f_clamp(raw_accel,
                state.actual_accel_previous - jerk_delta,
                state.actual_accel_previous + jerk_delta);
            longitudinal_accel = f_clamp(longitudinal_accel, -p->brake, p->accel);
            float lateral_accel = (force_y_l + force_y_r) / p->mass - state.yaw_rate * state.vx;
            float yaw_accel_actual = (force_x_r - force_x_l) * half_track / p->yaw_inertia;
            yaw_accel_actual = f_clamp(yaw_accel_actual, -4000.0f, 4000.0f);

            state.vx += longitudinal_accel * p->dt;
            state.vy += lateral_accel * p->dt;
            state.yaw_rate += yaw_accel_actual * p->dt;
            state.actual_accel_previous = longitudinal_accel;
            state.vy *= f_max(0.0f, 1.0f - 0.15f * p->dt);
            if (state.vx < -0.30f) state.vx = -0.30f;
            state.heading = wrap_pi(state.heading + state.yaw_rate * p->dt);
            state.time += p->dt;
            state.soc -= (f_abs(state.current_l) + f_abs(state.current_r)) * p->dt /
                         p->battery_capacity_c;
            state.soc = f_clamp(state.soc, 0.0f, 1.0f);
            progress += f_max(state.vx, 0.0f) * p->dt;
            lateral_offset += state.vy * p->dt;
            if (distance > 0.0f) {
                float t = f_clamp(progress / distance, 0.0f, 1.0f);
                float derivative_x, derivative_y, second_x, second_y;
                hermite_reference(t,
                                  segment_start_x, segment_start_y,
                                  tangent_start_x, tangent_start_y,
                                  segment_end_x, segment_end_y,
                                  tangent_end_x, tangent_end_y,
                                  &reference_x, &reference_y,
                                  &derivative_x, &derivative_y,
                                  &second_x, &second_y);
                float derivative_length = f_max(f_sqrt(derivative_x * derivative_x +
                                                       derivative_y * derivative_y), 0.000001f);
                state.x = reference_x - derivative_y * lateral_offset / derivative_length;
                state.y = reference_y + derivative_x * lateral_offset / derivative_length;
                state.heading = f_atan2(derivative_y, derivative_x);
                float second_length = f_max(derivative_length * derivative_length *
                                            derivative_length, 0.000001f);
                float curvature = (derivative_x * second_y - derivative_y * second_x) /
                                  second_length;
                state.yaw_rate = state.vx * curvature;
                max_lateral = f_max(max_lateral, f_abs(lateral_offset));
            } else {
                state.x = segment_start_x;
                state.y = segment_start_y;
            }
            ++total_steps;
            ++segment_steps;
            if (record && ((total_steps & 3u) == 0u)) trajectory_append(app, &state, slip_abs);

            if (!(state.vx == state.vx) || !(state.vy == state.vy) ||
                !(state.yaw_rate == state.yaw_rate)) {
                valid = 0;
                result->failure_reason = DYN_FAIL_NONFINITE;
                result->failure_action = action_index;
                break;
            }
            if (f_abs(state.vx) > 30.0f || f_abs(state.vy) > 30.0f ||
                f_abs(state.yaw_rate) > 500.0f) {
                valid = 0;
                result->failure_reason = DYN_FAIL_STATE_LIMIT;
                result->failure_action = action_index;
                break;
            }
            float lateral_clearance = f_max(0.04f * p->cell_m,
                                            0.5f * (p->cell_m - p->track));
            if (max_lateral > lateral_clearance) {
                valid = 0;
                result->failure_reason = DYN_FAIL_LATERAL_ENVELOPE;
                result->failure_action = action_index;
                break;
            }
            if (max_slip > 100.0f) {
                valid = 0;
                result->failure_reason = DYN_FAIL_SLIP_LIMIT;
                result->failure_action = action_index;
                break;
            }
            if (state.soc <= 0.0f) {
                valid = 0;
                result->failure_reason = DYN_FAIL_BATTERY_EMPTY;
                result->failure_action = action_index;
                break;
            }

            yaw_progress = wrap_pi(state.heading - segment_heading);
            heading_error = wrap_pi(target_reference_heading - state.heading);
            if (distance > 0.0f) {
                float exit_tolerance = f_max(0.12f, 0.12f * max_speed);
                if (progress >= distance && f_abs(heading_error) < 0.10f &&
                    state.vx <= exit_speed + exit_tolerance) {
                    completed = 1;
                    break;
                }
            } else if (f_abs(heading_error) < 0.06f && f_abs(state.yaw_rate) < 1.0f) {
                completed = 1;
                break;
            }
        }
        if (!completed) {
            valid = 0;
            if (!result->failure_reason) {
                result->failure_reason = total_steps >= p->max_steps ?
                    DYN_FAIL_TOTAL_STEP_LIMIT :
                    (f_abs(wrap_pi(target_angle - wrap_pi(state.heading - segment_heading))) >= 0.10f ?
                     DYN_FAIL_SEGMENT_HEADING : DYN_FAIL_SEGMENT_STEP_LIMIT);
                result->failure_action = action_index;
            }
        }
    }

    u32 final_node = app->path_nodes[app->path_count - 1u];
    float expected_x, expected_y;
    node_reference_position_m(app, p, final_node, &expected_x, &expected_y);
    float position_error = f_sqrt((state.x - expected_x) * (state.x - expected_x) +
                                  (state.y - expected_y) * (state.y - expected_y));
    float expected_heading = 0.5f * pi - (float)(final_node & 7u) * 0.25f * pi;
    float final_heading_error = f_abs(wrap_pi(state.heading - expected_heading));
    if (valid && position_error > 0.35f * p->cell_m) {
        valid = 0;
        result->failure_reason = DYN_FAIL_FINAL_POSITION;
        result->failure_action = action_count;
    }
    if (valid && final_heading_error > 0.18f) {
        valid = 0;
        result->failure_reason = DYN_FAIL_FINAL_HEADING;
        result->failure_action = action_count;
    }
    if (record) trajectory_append(app, &state, max_slip);

    result->time_s = state.time;
    result->max_slip = max_slip;
    result->final_voltage = state.voltage;
    result->max_lateral = max_lateral;
    result->final_heading_error = final_heading_error;
    result->final_soc = state.soc;
    result->final_position_error = position_error;
    result->cone_hits = cone_hits;
    result->steps = total_steps;
    result->valid = valid ? 1u : 0u;
    return valid;
}

static void cuda_set_error(CudaState *cuda, const WCHAR *stage, CUresult code) {
    cuda->error[0] = 0;
    wcat(cuda->error, ARRAY_COUNT(cuda->error), stage);
    wcat(cuda->error, ARRAY_COUNT(cuda->error), L"（CUDA 错误码 ");
    wcat_u64(cuda->error, ARRAY_COUNT(cuda->error), (u32)code);
    wcat(cuda->error, ARRAY_COUNT(cuda->error), L"）");
}

#define LOAD_CUDA(name, type, export_name) \
    do { cuda->name = (type)GetProcAddress(cuda->library, export_name); \
         if (!cuda->name) { wcopy(cuda->error, ARRAY_COUNT(cuda->error), L"CUDA 驱动缺少入口：" L##export_name); goto fail; } } while (0)

static void cuda_shutdown(CudaState *cuda) {
    if (cuda->module && cuda->cuModuleUnload) cuda->cuModuleUnload(cuda->module);
    cuda->module = 0;
    if (cuda->context && cuda->cuCtxDestroy) cuda->cuCtxDestroy(cuda->context);
    cuda->context = 0;
    if (cuda->library) FreeLibrary(cuda->library);
    cuda->library = 0;
    cuda->ready = 0;
}

static int cuda_initialize(CudaState *cuda, HINSTANCE instance) {
    memset(cuda, 0, sizeof(*cuda));
    wcopy(cuda->error, ARRAY_COUNT(cuda->error), L"CUDA 尚未初始化");
    cuda->library = LoadLibraryW(L"nvcuda.dll");
    if (!cuda->library) {
        wcopy(cuda->error, ARRAY_COUNT(cuda->error), L"未找到 nvcuda.dll；请安装或更新 NVIDIA 显卡驱动");
        return 0;
    }
    LOAD_CUDA(cuInit, PFN_cuInit, "cuInit");
    LOAD_CUDA(cuDriverGetVersion, PFN_cuDriverGetVersion, "cuDriverGetVersion");
    LOAD_CUDA(cuDeviceGetCount, PFN_cuDeviceGetCount, "cuDeviceGetCount");
    LOAD_CUDA(cuDeviceGet, PFN_cuDeviceGet, "cuDeviceGet");
    LOAD_CUDA(cuDeviceGetName, PFN_cuDeviceGetName, "cuDeviceGetName");
    LOAD_CUDA(cuDeviceGetAttribute, PFN_cuDeviceGetAttribute, "cuDeviceGetAttribute");
    LOAD_CUDA(cuDeviceTotalMem, PFN_cuDeviceTotalMem, "cuDeviceTotalMem_v2");
    LOAD_CUDA(cuCtxCreate, PFN_cuCtxCreate, "cuCtxCreate_v2");
    LOAD_CUDA(cuCtxDestroy, PFN_cuCtxDestroy, "cuCtxDestroy_v2");
    LOAD_CUDA(cuModuleLoadDataEx, PFN_cuModuleLoadDataEx, "cuModuleLoadDataEx");
    LOAD_CUDA(cuModuleUnload, PFN_cuModuleUnload, "cuModuleUnload");
    LOAD_CUDA(cuModuleGetFunction, PFN_cuModuleGetFunction, "cuModuleGetFunction");
    LOAD_CUDA(cuMemAlloc, PFN_cuMemAlloc, "cuMemAlloc_v2");
    LOAD_CUDA(cuMemFree, PFN_cuMemFree, "cuMemFree_v2");
    LOAD_CUDA(cuMemcpyHtoD, PFN_cuMemcpyHtoD, "cuMemcpyHtoD_v2");
    LOAD_CUDA(cuMemcpyDtoH, PFN_cuMemcpyDtoH, "cuMemcpyDtoH_v2");
    LOAD_CUDA(cuLaunchKernel, PFN_cuLaunchKernel, "cuLaunchKernel");
    LOAD_CUDA(cuCtxSynchronize, PFN_cuCtxSynchronize, "cuCtxSynchronize");
    cuda->cuGetErrorString = (PFN_cuGetErrorString)GetProcAddress(cuda->library, "cuGetErrorString");

    CUresult result = cuda->cuInit(0u);
    if (result) { cuda_set_error(cuda, L"cuInit 失败", result); goto fail; }
    cuda->cuDriverGetVersion(&cuda->driver_version);
    int count = 0;
    result = cuda->cuDeviceGetCount(&count);
    if (result || count <= 0) {
        if (result) cuda_set_error(cuda, L"无法枚举 NVIDIA GPU", result);
        else wcopy(cuda->error, ARRAY_COUNT(cuda->error), L"CUDA 驱动未报告可用 NVIDIA GPU");
        goto fail;
    }
    int best_score = -1;
    for (int i = 0; i < count; ++i) {
        CUdevice device;
        int major = 0, minor = 0;
        if (cuda->cuDeviceGet(&device, i)) continue;
        cuda->cuDeviceGetAttribute(&major, 75, device);
        cuda->cuDeviceGetAttribute(&minor, 76, device);
        int score = major * 100 + minor;
        if (score > best_score) {
            best_score = score;
            cuda->device = device;
            cuda->cc_major = major;
            cuda->cc_minor = minor;
        }
    }
    if (best_score < 0) {
        wcopy(cuda->error, ARRAY_COUNT(cuda->error), L"没有可创建 CUDA 上下文的 GPU");
        goto fail;
    }
    if (cuda->cc_major < 5 || (cuda->cc_major == 5 && cuda->cc_minor < 2)) {
        wcopy(cuda->error, ARRAY_COUNT(cuda->error), L"GPU 计算能力低于 5.2，无法运行当前 PTX 核心");
        goto fail;
    }
    cuda->cuDeviceGetName(cuda->gpu_name, (int)sizeof(cuda->gpu_name), cuda->device);
    cuda->cuDeviceGetAttribute(&cuda->multiprocessors, 16, cuda->device);
    cuda->cuDeviceGetAttribute(&cuda->clock_khz, 13, cuda->device);
    cuda->cuDeviceTotalMem(&cuda->total_memory, cuda->device);
    result = cuda->cuCtxCreate(&cuda->context, 0u, cuda->device);
    if (result) { cuda_set_error(cuda, L"创建 CUDA 上下文失败", result); goto fail; }

    HRSRC resource = FindResourceW(instance, MAKEINTRESOURCEW(PTX_RESOURCE_ID), RT_RCDATA);
    if (!resource) {
        wcopy(cuda->error, ARRAY_COUNT(cuda->error), L"内部 PTX 资源不存在");
        goto fail;
    }
    DWORD resource_size = SizeofResource(instance, resource);
    HGLOBAL loaded = LoadResource(instance, resource);
    const void *resource_data = loaded ? LockResource(loaded) : 0;
    if (!resource_data || !resource_size) {
        wcopy(cuda->error, ARRAY_COUNT(cuda->error), L"内部 PTX 资源无法读取");
        goto fail;
    }
    char *ptx = (char *)VirtualAlloc(0, (SIZE_T)resource_size + 1u,
                                    MEM_COMMIT | MEM_RESERVE, PAGE_READWRITE);
    if (!ptx) {
        wcopy(cuda->error, ARRAY_COUNT(cuda->error), L"PTX 临时内存分配失败");
        goto fail;
    }
    memcpy(ptx, resource_data, resource_size);
    ptx[resource_size] = 0;
    result = cuda->cuModuleLoadDataEx(&cuda->module, ptx, 0u, 0, 0);
    VirtualFree(ptx, 0, MEM_RELEASE);
    if (result) { cuda_set_error(cuda, L"CUDA 驱动 JIT 编译 PTX 失败", result); goto fail; }
    result = cuda->cuModuleGetFunction(&cuda->relax, cuda->module, "mm_relax");
    if (result) { cuda_set_error(cuda, L"找不到 mm_relax 核心", result); goto fail; }
    result = cuda->cuModuleGetFunction(&cuda->dynamics, cuda->module, "mm_rank_dynamics");
    if (result) { cuda_set_error(cuda, L"找不到 mm_rank_dynamics 核心", result); goto fail; }
    result = cuda->cuModuleGetFunction(&cuda->fem_clear, cuda->module, "mm_fem_clear_force");
    if (result) { cuda_set_error(cuda, L"找不到 mm_fem_clear_force 核心", result); goto fail; }
    result = cuda->cuModuleGetFunction(&cuda->fem_tie, cuda->module, "mm_fem_tie_force");
    if (result) { cuda_set_error(cuda, L"找不到 mm_fem_tie_force 核心", result); goto fail; }
    result = cuda->cuModuleGetFunction(&cuda->fem_force, cuda->module, "mm_fem_hex8_force");
    if (result) { cuda_set_error(cuda, L"找不到 mm_fem_hex8_force 核心", result); goto fail; }
    result = cuda->cuModuleGetFunction(&cuda->fem_integrate, cuda->module, "mm_fem_integrate");
    if (result) { cuda_set_error(cuda, L"找不到 mm_fem_integrate 核心", result); goto fail; }
    cuda->ready = 1;
    cuda->error[0] = 0;
    return 1;
fail:
    cuda_shutdown(cuda);
    return 0;
}

static CUresult cuda_copy_h2d(CudaState *cuda, CUdeviceptr target, const void *source, SIZE_T size) {
    return cuda->cuMemcpyHtoD(target, source, size);
}

static int cuda_solve(AppState *app) {
    CudaState *cuda = &app->cuda;
    if (!cuda->ready) return 0;
    SolverParams *p = &app->params;
    u32 node_count = p->n * p->n * 8u;
    u32 start_node = 0u;
    for (u32 i = 0; i < node_count; ++i) g_dist[i] = INF64;
    g_dist[start_node] = 0ull;

    CUdeviceptr d_dist = 0, d_walls = 0, d_params = 0, d_changed = 0;
    CUresult result = 0;
    if ((result = cuda->cuMemAlloc(&d_dist, (SIZE_T)node_count * sizeof(u64)))) goto fail;
    if ((result = cuda->cuMemAlloc(&d_walls, (SIZE_T)p->n * p->n))) goto fail;
    if ((result = cuda->cuMemAlloc(&d_params, sizeof(*p)))) goto fail;
    if ((result = cuda->cuMemAlloc(&d_changed, sizeof(u32)))) goto fail;
    if ((result = cuda_copy_h2d(cuda, d_dist, g_dist, (SIZE_T)node_count * sizeof(u64)))) goto fail;
    if ((result = cuda_copy_h2d(cuda, d_walls, app->maze.walls, (SIZE_T)p->n * p->n))) goto fail;
    if ((result = cuda_copy_h2d(cuda, d_params, p, sizeof(*p)))) goto fail;

    u32 changed = 0;
    u32 iteration = 0;
    for (; iteration < node_count + 8u; ++iteration) {
        changed = 0;
        if ((result = cuda_copy_h2d(cuda, d_changed, &changed, sizeof(changed)))) goto fail;
        void *arguments[5] = {&d_dist, &d_walls, &d_params, &d_changed, &node_count};
        u32 blocks = (node_count + 255u) / 256u;
        result = cuda->cuLaunchKernel(cuda->relax, blocks, 1u, 1u, 256u, 1u, 1u,
                                      0u, 0, arguments, 0);
        if (result) goto fail;
        if ((result = cuda->cuCtxSynchronize())) goto fail;
        if ((result = cuda->cuMemcpyDtoH(&changed, d_changed, sizeof(changed)))) goto fail;
        if (!changed) break;
    }
    if (changed) {
        wcopy(cuda->error, ARRAY_COUNT(cuda->error), L"矩阵自由松弛未在理论上限内闭合");
        goto fail_no_code;
    }
    if ((result = cuda->cuMemcpyDtoH(g_dist, d_dist, (SIZE_T)node_count * sizeof(u64)))) goto fail;

    u64 best_pack = INF64;
    u32 best_node = 0;
    for (u32 gy = app->maze.goal_y0; gy <= app->maze.goal_y1; ++gy) {
        for (u32 gx = app->maze.goal_x0; gx <= app->maze.goal_x1; ++gx) {
            u32 cell = gy * p->n + gx;
            for (u32 h = 0; h < 8u; ++h) {
                u32 node = (cell << 3u) | h;
                if (g_dist[node] < best_pack) {
                    best_pack = g_dist[node];
                    best_node = node;
                }
            }
        }
    }
    if (best_pack == INF64) {
        wcopy(cuda->error, ARRAY_COUNT(cuda->error), L"迷宫在当前动作包络下不可达");
        goto fail_no_code;
    }

    u32 trace_count = 0;
    u32 node = best_node;
    while (node != start_node && trace_count + 1u < MAX_ROUTE) {
        u64 pack = g_dist[node];
        if (pack == INF64) goto trace_fail;
        u32 action = (u32)(pack & 63ull);
        u32 predecessor = (u32)((pack & LOW_MASK) >> 6u);
        if (predecessor >= node_count) goto trace_fail;
        g_trace_nodes[trace_count] = node;
        g_trace_actions[trace_count] = (u8)action;
        ++trace_count;
        if ((g_dist[predecessor] >> PACK_SHIFT) >= (pack >> PACK_SHIFT)) goto trace_fail;
        node = predecessor;
    }
    if (node != start_node) goto trace_fail;
    app->path_nodes[0] = start_node;
    for (u32 i = 0; i < trace_count; ++i) {
        app->path_nodes[i + 1u] = g_trace_nodes[trace_count - 1u - i];
        app->path_actions[i] = g_trace_actions[trace_count - 1u - i];
    }
    app->path_count = trace_count + 1u;
    app->best_time_us = best_pack >> PACK_SHIFT;
    app->graph_time_us = app->best_time_us;
    app->solve_iterations = iteration + 1u;
    app->exact_pass = 1;
    app->dynamics_pass = 0;
    app->trajectory_count = 0;

    for (u32 i = 0; i < trace_count; ++i) {
        u32 source = app->path_nodes[i];
        u32 target = app->path_nodes[i + 1u];
        u32 evaluated_target = 0;
        u64 edge_cost = 0;
        u64 source_time = g_dist[source] >> PACK_SHIFT;
        if (!evaluate_action(source, app->path_actions[i], source_time, p,
                             &app->maze, &evaluated_target, &edge_cost) ||
            evaluated_target != target ||
            source_time + edge_cost != (g_dist[target] >> PACK_SHIFT)) {
            app->exact_pass = 0;
            break;
        }
    }
    cuda->cuMemFree(d_changed); cuda->cuMemFree(d_params);
    cuda->cuMemFree(d_walls); cuda->cuMemFree(d_dist);
    return app->exact_pass;

trace_fail:
    wcopy(cuda->error, ARRAY_COUNT(cuda->error), L"最优前驱链回溯失败");
fail_no_code:
    if (d_changed) cuda->cuMemFree(d_changed);
    if (d_params) cuda->cuMemFree(d_params);
    if (d_walls) cuda->cuMemFree(d_walls);
    if (d_dist) cuda->cuMemFree(d_dist);
    return 0;
fail:
    cuda_set_error(cuda, L"CUDA 求解阶段失败", result);
    goto fail_no_code;
}

static int cuda_rank_dynamics(AppState *app, const DynamicsParams *params,
                              GpuCandidateResult *host_results) {
    CudaState *cuda = &app->cuda;
    CUdeviceptr d_nodes = 0, d_actions = 0, d_params = 0, d_results = 0;
    CUresult result = 0;
    u32 action_count = app->path_count > 0u ? app->path_count - 1u : 0u;
    u32 candidate_count = DYNAMICS_CANDIDATES;
    SIZE_T nodes_size = (SIZE_T)app->path_count * sizeof(u32);
    SIZE_T actions_size = (SIZE_T)action_count * sizeof(u8);
    SIZE_T results_size = (SIZE_T)candidate_count * sizeof(GpuCandidateResult);
    if (!cuda->ready || !cuda->dynamics || !action_count || !host_results) return 0;

    if ((result = cuda->cuMemAlloc(&d_nodes, nodes_size))) goto fail;
    if ((result = cuda->cuMemAlloc(&d_actions, actions_size))) goto fail;
    if ((result = cuda->cuMemAlloc(&d_params, sizeof(*params)))) goto fail;
    if ((result = cuda->cuMemAlloc(&d_results, results_size))) goto fail;
    if ((result = cuda_copy_h2d(cuda, d_nodes, app->path_nodes, nodes_size))) goto fail;
    if ((result = cuda_copy_h2d(cuda, d_actions, app->path_actions, actions_size))) goto fail;
    if ((result = cuda_copy_h2d(cuda, d_params, params, sizeof(*params)))) goto fail;

    void *arguments[6] = {
        &d_nodes, &d_actions, &d_params, &d_results, &action_count, &candidate_count
    };
    u32 blocks = (candidate_count + 255u) / 256u;
    result = cuda->cuLaunchKernel(cuda->dynamics, blocks, 1u, 1u,
                                  256u, 1u, 1u, 0u, 0, arguments, 0);
    if (result) goto fail;
    if ((result = cuda->cuCtxSynchronize())) goto fail;
    if ((result = cuda->cuMemcpyDtoH(host_results, d_results, results_size))) goto fail;
    cuda->cuMemFree(d_results); cuda->cuMemFree(d_params);
    cuda->cuMemFree(d_actions); cuda->cuMemFree(d_nodes);
    return 1;

fail:
    cuda_set_error(cuda, L"CUDA 动力学候选批处理失败", result);
    if (d_results) cuda->cuMemFree(d_results);
    if (d_params) cuda->cuMemFree(d_params);
    if (d_actions) cuda->cuMemFree(d_actions);
    if (d_nodes) cuda->cuMemFree(d_nodes);
    return 0;
}

#include "fem_runtime.h"

#define CPU_FINALIST_LIMIT 128u

static int optimize_dynamics(AppState *app) {
    DynamicsParams dynamics_params;
    build_dynamics_params(&app->params, &dynamics_params);
    SIZE_T result_bytes = (SIZE_T)DYNAMICS_CANDIDATES * sizeof(GpuCandidateResult);
    GpuCandidateResult *gpu_results = (GpuCandidateResult *)VirtualAlloc(
        0, result_bytes, MEM_COMMIT | MEM_RESERVE, PAGE_READWRITE);
    if (!gpu_results) {
        wcopy(app->cuda.error, ARRAY_COUNT(app->cuda.error), L"动力学候选结果内存分配失败");
        return 0;
    }
    memset(gpu_results, 0, result_bytes);
    if (!cuda_rank_dynamics(app, &dynamics_params, gpu_results)) {
        VirtualFree(gpu_results, 0, MEM_RELEASE);
        return 0;
    }

    memset(g_candidate_tested, 0, sizeof(g_candidate_tested));
    memset(app->dynamics_failure_counts, 0, sizeof(app->dynamics_failure_counts));
    DynamicsResult best_result;
    memset(&best_result, 0, sizeof(best_result));
    u32 best_candidate = 0u;
    u32 replay_count = 0u;
    for (; replay_count < CPU_FINALIST_LIMIT; ++replay_count) {
        u32 ranked_candidate = DYNAMICS_CANDIDATES;
        if (replay_count < 64u) {
            static const u32 level[8] = {0u, 9u, 18u, 27u, 36u, 45u, 54u, 63u};
            ranked_candidate = level[replay_count & 7u] |
                               (level[(replay_count >> 3u) & 7u] << 6u);
        } else {
            float ranked_score = 1000000000.0f;
            for (u32 candidate = 0; candidate < DYNAMICS_CANDIDATES; ++candidate) {
                const GpuCandidateResult *ranked = &gpu_results[candidate];
                if (g_candidate_tested[candidate] || !ranked->valid ||
                    !(ranked->score_time_s == ranked->score_time_s) ||
                    ranked->score_time_s <= 0.0f) continue;
                if (ranked->score_time_s < ranked_score) {
                    ranked_score = ranked->score_time_s;
                    ranked_candidate = candidate;
                }
            }
        }
        if (ranked_candidate >= DYNAMICS_CANDIDATES) break;
        if (g_candidate_tested[ranked_candidate]) {
            --replay_count;
            continue;
        }
        g_candidate_tested[ranked_candidate] = 1u;
        DynamicsResult candidate_result;
        if (simulate_dynamics_cpu(app, &dynamics_params, ranked_candidate, 0,
                                  &candidate_result)) {
            if (!best_result.valid || candidate_result.time_s < best_result.time_s ||
                (candidate_result.time_s == best_result.time_s &&
                 candidate_result.max_slip < best_result.max_slip)) {
                best_result = candidate_result;
                best_candidate = ranked_candidate;
            }
        } else if (candidate_result.failure_reason < DYN_FAIL_COUNT) {
            ++app->dynamics_failure_counts[candidate_result.failure_reason];
        }
        if ((replay_count & 63u) == 63u) {
            g_text[0] = 0;
            wcat(g_text, ARRAY_COUNT(g_text), L"GPU 已并行筛选 4096 组；C 闭环高精度复核 ");
            wcat_u64(g_text, ARRAY_COUNT(g_text), replay_count + 1u);
            wcat(g_text, ARRAY_COUNT(g_text), L" / ");
            wcat_u64(g_text, ARRAY_COUNT(g_text), CPU_FINALIST_LIMIT);
            SetWindowTextW(app->result, g_text);
            UpdateWindow(app->result);
        }
    }

    if (!best_result.valid) {
        VirtualFree(gpu_results, 0, MEM_RELEASE);
        u32 dominant_reason = DYN_FAIL_NONE;
        u32 dominant_count = 0u;
        for (u32 reason = 1u; reason < DYN_FAIL_COUNT; ++reason) {
            if (app->dynamics_failure_counts[reason] > dominant_count) {
                dominant_count = app->dynamics_failure_counts[reason];
                dominant_reason = reason;
            }
        }
        wcopy(app->cuda.error, ARRAY_COUNT(app->cuda.error), L"闭环复核无可行候选；主要拒绝原因：");
        static const WCHAR *reason_names[DYN_FAIL_COUNT] = {
            L"无", L"参数无效", L"出现非有限数", L"状态越界", L"横向包络越界",
            L"滑移率越界", L"电池耗尽", L"总积分步数耗尽", L"单段积分步数耗尽",
            L"单段航向未闭合", L"终点位置误差", L"终点航向误差"
        };
        wcat(app->cuda.error, ARRAY_COUNT(app->cuda.error), reason_names[dominant_reason]);
        wcat(app->cuda.error, ARRAY_COUNT(app->cuda.error), L"（");
        wcat_u64(app->cuda.error, ARRAY_COUNT(app->cuda.error), dominant_count);
        wcat(app->cuda.error, ARRAY_COUNT(app->cuda.error), L" / ");
        wcat_u64(app->cuda.error, ARRAY_COUNT(app->cuda.error), replay_count);
        wcat(app->cuda.error, ARRAY_COUNT(app->cuda.error), L"）");
        return 0;
    }

    app->dynamics_candidate = best_candidate;
    app->dynamics_candidates = DYNAMICS_CANDIDATES;
    app->dynamics_cpu_replays = replay_count;
    dynamics_candidate_scales(best_candidate, &app->best_straight_scale,
                              &app->best_turn_scale);
    app->gpu_predicted_time = gpu_results[best_candidate].predicted_time_s;
    app->gpu_predicted_slip = gpu_results[best_candidate].max_slip;
    app->gpu_max_cone_ratio = gpu_results[best_candidate].max_cone_ratio;
    app->dynamics_result = best_result;
    app->dynamics_pass = simulate_dynamics_cpu(app, &dynamics_params, best_candidate, 1,
                                               &app->dynamics_result);
    if (!app->dynamics_pass) {
        wcopy(app->cuda.error, ARRAY_COUNT(app->cuda.error), L"最优候选在轨迹记录复演时未保持确定性");
        VirtualFree(gpu_results, 0, MEM_RELEASE);
        return 0;
    }
    app->best_time_us = (u64)(app->dynamics_result.time_s * 1000000.0f + 0.5f);
    if (!run_fem_certification(app)) {
        VirtualFree(gpu_results, 0, MEM_RELEASE);
        return 0;
    }
    {
        /* One deterministic two-way correction: the first full-car solve
           supplies the deformed effective wheel radius, the rigid/contact
           model replays with that radius, and FEM certifies the corrected
           trajectory once more.  Viscoelastic dissipation is reported but is
           not converted into an invented rolling-resistance coefficient. */
        float original_radius = dynamics_params.wheel_radius;
        float feedback_radius = f_clamp(g_fem_result.effective_wheel_radius_m,
                                        original_radius * 0.80f,
                                        original_radius * 1.20f);
        float original_time = app->dynamics_result.time_s;
        DynamicsParams coupled_params = dynamics_params;
        coupled_params.wheel_radius = feedback_radius;
        coupled_params.wheel_inertia = 0.5f * coupled_params.wheel_mass *
            feedback_radius * feedback_radius;
        DynamicsResult coupled_result;
        if (!simulate_dynamics_cpu(app, &coupled_params, best_candidate, 1,
                                   &coupled_result)) {
            wcopy(app->cuda.error, ARRAY_COUNT(app->cuda.error),
                  L"全车有限元轮径反馈后的闭环动力学复演失败");
            VirtualFree(gpu_results, 0, MEM_RELEASE);
            return 0;
        }
        app->dynamics_result = coupled_result;
        app->best_time_us = (u64)(coupled_result.time_s * 1000000.0f + 0.5f);
        if (!run_fem_certification(app)) {
            VirtualFree(gpu_results, 0, MEM_RELEASE);
            return 0;
        }
        g_fem_result.coupling_passes = 2u;
        g_fem_result.feedback_radius_delta_percent =
            100.0f * (feedback_radius - original_radius) /
            f_max(original_radius, 0.000001f);
        g_fem_result.feedback_time_delta_percent =
            100.0f * (coupled_result.time_s - original_time) /
            f_max(original_time, 0.000001f);
    }
    VirtualFree(gpu_results, 0, MEM_RELEASE);
    return app->dynamics_pass;
}

static const WCHAR *action_name(u32 action) {
    static const WCHAR *names[15] = {
        L"F1", L"F2", L"F3", L"F4", L"F5", L"F6", L"F7", L"F8",
        L"L45", L"R45", L"L90", L"R90", L"L135", L"R135", L"U180"
    };
    return action < ARRAY_COUNT(names) ? names[action] : L"?";
}

static void show_generation_result(void) {
    g_text[0] = 0;
    if (g_app.official_loaded && g_app.official_index < OFFICIAL_MAZE_COUNT) {
        const OfficialMazeRecord *record = &g_official_mazes[g_app.official_index];
        wcat(g_text, ARRAY_COUNT(g_text), L"内置公开赛事档案载入与结构完整性验证：PASS\r\n系列：");
        wcat(g_text, ARRAY_COUNT(g_text), g_official_series_names[record->series]);
        wcat(g_text, ARRAY_COUNT(g_text), L"\r\n题目：");
        wcat(g_text, ARRAY_COUNT(g_text), record->label);
        wcat(g_text, ARRAY_COUNT(g_text), L"\r\n来源级别：");
        wcat(g_text, ARRAY_COUNT(g_text), record->source_tier == 2u ?
             L"公开 All Japan 历史图档转录（仍需与主办方原图复核）" :
             L"公开社区赛事档案转录（非官方真实性认证）");
        wcat(g_text, ARRAY_COUNT(g_text), L"\r\n档案名：");
        wcat(g_text, ARRAY_COUNT(g_text), record->source_name);
        wcat(g_text, ARRAY_COUNT(g_text), L"\r\n");
    } else {
        wcat(g_text, ARRAY_COUNT(g_text), L"迷宫结构与赛制约束验证：PASS\r\n");
    }
    wcat(g_text, ARRAY_COUNT(g_text), L"尺寸：");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.maze.n);
    wcat(g_text, ARRAY_COUNT(g_text), L" × ");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.maze.n);
    wcat(g_text, ARRAY_COUNT(g_text), L"，单格 ");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.maze.cell_mm);
    wcat(g_text, ARRAY_COUNT(g_text), L" mm\r\n终点坐标：(");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.maze.goal_x0);
    wcat_char(g_text, ARRAY_COUNT(g_text), L',');
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.maze.goal_y0);
    wcat(g_text, ARRAY_COUNT(g_text), L") - (");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.maze.goal_x1);
    wcat_char(g_text, ARRAY_COUNT(g_text), L',');
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.maze.goal_y1);
    wcat_char(g_text, ARRAY_COUNT(g_text), L')');
    if (g_app.official_loaded) {
        wcat(g_text, ARRAY_COUNT(g_text), L"\r\n嵌入哈希/外墙/互反墙/标准起点/目标可达：PASS");
        wcat(g_text, ARRAY_COUNT(g_text), L"\r\n题库：");
        wcat_u64(g_text, ARRAY_COUNT(g_text), OFFICIAL_MAZE_COUNT);
        wcat(g_text, ARRAY_COUNT(g_text), L" 条记录，");
        wcat_u64(g_text, ARRAY_COUNT(g_text), OFFICIAL_UNIQUE_LAYOUT_COUNT);
        wcat(g_text, ARRAY_COUNT(g_text), L" 个唯一布局，全部内置离线。");
    } else {
        wcat(g_text, ARRAY_COUNT(g_text), L"\r\n种子：");
        wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.maze.seed);
        if (g_app.maze.cell_mm == 180u)
            wcat(g_text, ARRAY_COUNT(g_text), L"\r\n终点：正中央 2×2（Classic 常规赛制）");
        else
            wcat(g_text, ARRAY_COUNT(g_text), L"\r\n终点：随机模式预设中心 2×2；Half-size 赛题可任意矩形");
        wcat(g_text, ARRAY_COUNT(g_text),
             L"\r\n外墙/互反墙/全连通/三面起点/目标单入口/中心空心：PASS"
             L"\r\n除中心柱外每柱挂墙/左右手贴墙法不可直解：PASS");
        wcat(g_text, ARRAY_COUNT(g_text), L"\r\n再次点击会自动换种子；手动输入种子可精确复现。");
    }
    SetWindowTextW(g_app.result, g_text);
}

static void show_solve_result(void) {
    g_text[0] = 0;
    if (g_app.official_loaded && g_app.official_index < OFFICIAL_MAZE_COUNT) {
        wcat(g_text, ARRAY_COUNT(g_text), L"赛题：");
        wcat(g_text, ARRAY_COUNT(g_text), g_official_mazes[g_app.official_index].label);
        wcat(g_text, ARRAY_COUNT(g_text), L"\r\n");
    }
    wcat(g_text, ARRAY_COUNT(g_text), L"当前路线可行候选最短回放时间：");
    wcat_fixed6(g_text, ARRAY_COUNT(g_text), g_app.best_time_us);
    wcat(g_text, ARRAY_COUNT(g_text), L" s\r\nCUDA 动作图估计：");
    wcat_fixed6(g_text, ARRAY_COUNT(g_text), g_app.graph_time_us);
    wcat(g_text, ARRAY_COUNT(g_text), L" s（只用于全局路径排序）\r\n");
    wcat(g_text, ARRAY_COUNT(g_text), L"动作图前驱链逐边复算：");
    wcat(g_text, ARRAY_COUNT(g_text), g_app.exact_pass ? L"PASS" : L"FAIL");
    wcat(g_text, ARRAY_COUNT(g_text), L"\r\n高精度闭环动力学回放：");
    wcat(g_text, ARRAY_COUNT(g_text), g_app.dynamics_pass ? L"PASS" : L"FAIL");
    wcat(g_text, ARRAY_COUNT(g_text), L"\r\n\r\nCUDA 全局图：矩阵存储 0 bytes（边算子现场生成）\r\n状态数：");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.params.n * g_app.params.n * 8u);
    wcat(g_text, ARRAY_COUNT(g_text), L"，GPU 松弛轮数：");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.solve_iterations);
    wcat(g_text, ARRAY_COUNT(g_text), L"，路线尝试：");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.dynamics_route_attempts);
    wcat(g_text, ARRAY_COUNT(g_text), L"\r\nGPU 动力学策略：");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.dynamics_candidates);
    wcat(g_text, ARRAY_COUNT(g_text), L" 线程候选全部并行评估；C 高精度复核前 ");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.dynamics_cpu_replays);
    wcat(g_text, ARRAY_COUNT(g_text), L" 名\r\n候选编号：");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.dynamics_candidate);
    wcat(g_text, ARRAY_COUNT(g_text), L"，直线倍率 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_app.best_straight_scale);
    wcat(g_text, ARRAY_COUNT(g_text), L"，转弯倍率 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_app.best_turn_scale);
    wcat(g_text, ARRAY_COUNT(g_text), L"\r\nGPU 分段模型预测：");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_app.gpu_predicted_time);
    wcat(g_text, ARRAY_COUNT(g_text), L" s，锥约束峰值比 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_app.gpu_max_cone_ratio);
    wcat(g_text, ARRAY_COUNT(g_text), L"\r\n");

    wcat(g_text, ARRAY_COUNT(g_text), L"\r\n路径约束闭环刚体回放：\r\n最终端电压 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_app.dynamics_result.final_voltage);
    wcat(g_text, ARRAY_COUNT(g_text), L" V，剩余 SOC ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_app.dynamics_result.final_soc * 100.0f);
    wcat(g_text, ARRAY_COUNT(g_text), L" %\r\n最大轮胎滑移率 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_app.dynamics_result.max_slip * 100.0f);
    wcat(g_text, ARRAY_COUNT(g_text), L" %，最大横向约束残差 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_app.dynamics_result.max_lateral * 1000.0f);
    wcat(g_text, ARRAY_COUNT(g_text), L" mm\r\n终点位置误差 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_app.dynamics_result.final_position_error * 1000.0f);
    wcat(g_text, ARRAY_COUNT(g_text), L" mm，航向误差 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text),
                g_app.dynamics_result.final_heading_error * 57.2957795f);
    wcat(g_text, ARRAY_COUNT(g_text), L"°\r\n左右轮独立广义摩擦锥投影次数：");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.dynamics_result.cone_hits);
    wcat(g_text, ARRAY_COUNT(g_text), L"，积分步数：");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.dynamics_result.steps);
    wcat(g_text, ARRAY_COUNT(g_text), L"，轨迹采样：");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.trajectory_count);
    wcat(g_text, ARRAY_COUNT(g_text), L"\r\n第一层模型：固定运动基元上的闭环平面刚体时标回放；每个轮子独立椭圆广义摩擦锥，轮速/侧偏/电流/反电动势/端电压/SOC 每步反向约束电机实际输出。\r\n");

    wcat(g_text, ARRAY_COUNT(g_text), L"\r\n第二层全车有限应变—动力学双向认证：");
    wcat(g_text, ARRAY_COUNT(g_text), g_fem_result.valid ? L"PASS" : L"FAIL");
    wcat(g_text, ARRAY_COUNT(g_text), L"\r\n3D Hex8 节点 ");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_fem_result.node_count);
    wcat(g_text, ARRAY_COUNT(g_text), L"，单元 ");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_fem_result.element_count);
    wcat(g_text, ARRAY_COUNT(g_text), L"，总自由度 ");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_fem_result.dof_count);
    wcat(g_text, ARRAY_COUNT(g_text), L"，活动自由度 ");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_fem_result.active_dof_count);
    wcat(g_text, ARRAY_COUNT(g_text), L"，高斯点 ");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_fem_result.gauss_point_count);
    wcat(g_text, ARRAY_COUNT(g_text), L"，历史变量 ");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_fem_result.history_variable_count);
    wcat(g_text, ARRAY_COUNT(g_text), L"，界面约束 ");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_fem_result.tie_count);
    wcat(g_text, ARRAY_COUNT(g_text),
         L"\r\n分区：底板/PCB、电池、左右电机、传感器板、左右轮胎；无空壳区域。"
         L"\r\n本构：平衡 Neo-Hookean + 每材料 8 支有限应变广义 Maxwell；每高斯点 8×6 历史；后向 Euler 局部更新；8 点全积分；矩阵自由。"
         L"\r\n路线载荷包络：纵向 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.peak_longitudinal_accel);
    wcat(g_text, ARRAY_COUNT(g_text), L" m/s²，横向 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.peak_lateral_accel);
    wcat(g_text, ARRAY_COUNT(g_text), L" m/s²，垂向 5g。\r\nCPU/GPU 同方程残量相对 L2 误差 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.force_relative_l2 * 100.0f);
    wcat(g_text, ARRAY_COUNT(g_text), L" %；CPU ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.cpu_force_ms);
    wcat(g_text, ARRAY_COUNT(g_text), L" ms/残量，GPU ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.gpu_force_ms);
    wcat(g_text, ARRAY_COUNT(g_text), L" ms/残量，加速 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.speedup);
    wcat(g_text, ARRAY_COUNT(g_text), L"×（均不含分配、拷贝与绘图）。\r\n");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_fem_result.steps);
    wcat(g_text, ARRAY_COUNT(g_text), L" 步 GPU 大变形瞬态：");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.gpu_transient_ms);
    wcat(g_text, ARRAY_COUNT(g_text), L" ms；min J=");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.min_j);
    wcat(g_text, ARRAY_COUNT(g_text), L"，max J=");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.max_j);
    wcat(g_text, ARRAY_COUNT(g_text), L"\r\n最大 von Mises ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.max_von_mises_pa / 1000000.0f);
    wcat(g_text, ARRAY_COUNT(g_text), L" MPa，最大位移 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.max_displacement_m * 1000.0f);
    wcat(g_text, ARRAY_COUNT(g_text), L" mm，横向位移 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.max_lateral_displacement_m * 1000.0f);
    wcat(g_text, ARRAY_COUNT(g_text), L" mm，变形后有效轮径 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.effective_wheel_radius_m * 1000.0f);
    wcat(g_text, ARRAY_COUNT(g_text), L" mm。\r\nMaxwell 累积耗散 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.dissipated_energy_j * 1000.0f);
    wcat(g_text, ARRAY_COUNT(g_text), L" mJ，最大分支等效应力 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.max_branch_stress_pa / 1000000.0f);
    wcat(g_text, ARRAY_COUNT(g_text), L" MPa。\r\n双向修正 ");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_fem_result.coupling_passes);
    wcat(g_text, ARRAY_COUNT(g_text), L" 次：有效轮径变化 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.feedback_radius_delta_percent);
    wcat(g_text, ARRAY_COUNT(g_text), L" %，回放时间变化 ");
    wcat_float3(g_text, ARRAY_COUNT(g_text), g_fem_result.feedback_time_delta_percent);
    wcat(g_text, ARRAY_COUNT(g_text), L" %。\r\n材料参数为程序内明确标注的未标定基线；结果不冒充实车实验。\r\n");

    wcat(g_text, ARRAY_COUNT(g_text), L"\r\n动作数：");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.path_count ? g_app.path_count - 1u : 0u);
    wcat(g_text, ARRAY_COUNT(g_text), L"\r\n\r\n路径：\r\n");
    for (u32 i = 0; i + 1u < g_app.path_count; ++i) {
        if (wlen(g_text) > ARRAY_COUNT(g_text) - 64u) {
            wcat(g_text, ARRAY_COUNT(g_text), L" ...");
            break;
        }
        if (i) wcat(g_text, ARRAY_COUNT(g_text), L" → ");
        wcat(g_text, ARRAY_COUNT(g_text), action_name(g_app.path_actions[i]));
    }
    wcat(g_text, ARRAY_COUNT(g_text), L"\r\n\r\n输入摘要：μx=");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.params.mu_milli / 1000u);
    wcat_char(g_text, ARRAY_COUNT(g_text), L'.');
    u32 frac = g_app.params.mu_milli % 1000u;
    wcat_char(g_text, ARRAY_COUNT(g_text), (WCHAR)(L'0' + frac / 100u));
    wcat_char(g_text, ARRAY_COUNT(g_text), (WCHAR)(L'0' + (frac / 10u) % 10u));
    wcat_char(g_text, ARRAY_COUNT(g_text), (WCHAR)(L'0' + frac % 10u));
    wcat(g_text, ARRAY_COUNT(g_text), L"，μy=");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.params.mu_y_milli / 1000u);
    wcat_char(g_text, ARRAY_COUNT(g_text), L'.');
    frac = g_app.params.mu_y_milli % 1000u;
    wcat_char(g_text, ARRAY_COUNT(g_text), (WCHAR)(L'0' + frac / 100u));
    wcat_char(g_text, ARRAY_COUNT(g_text), (WCHAR)(L'0' + (frac / 10u) % 10u));
    wcat_char(g_text, ARRAY_COUNT(g_text), (WCHAR)(L'0' + frac % 10u));
    wcat(g_text, ARRAY_COUNT(g_text), L"，峰值滑移率=");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.params.slip_milli);
    wcat(g_text, ARRAY_COUNT(g_text), L"‰，jerk=");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.params.jerk);
    wcat(g_text, ARRAY_COUNT(g_text), L" mm/s³，电池=");
    wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.params.battery_mv);
    wcat(g_text, ARRAY_COUNT(g_text), L" mV\r\n说明：这里的“最短”指当前离散动作集和 64×64 策略网格内、经高精度复核候选中的最优；不是连续控制空间的数学全局最优证明。");
    SetWindowTextW(g_app.result, g_text);
}

static void set_status_cuda(void) {
    g_text[0] = 0;
    if (g_app.cuda.ready) {
        WCHAR name[160];
        ascii_to_wide(name, ARRAY_COUNT(name), g_app.cuda.gpu_name);
        wcat(g_text, ARRAY_COUNT(g_text), L"启动入口：PASS | CUDA：READY | GPU：");
        wcat(g_text, ARRAY_COUNT(g_text), name);
        wcat(g_text, ARRAY_COUNT(g_text), L" | CC ");
        wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.cuda.cc_major);
        wcat_char(g_text, ARRAY_COUNT(g_text), L'.');
        wcat_u64(g_text, ARRAY_COUNT(g_text), g_app.cuda.cc_minor);
        wcat(g_text, ARRAY_COUNT(g_text), L" | SM ");
        wcat_u64(g_text, ARRAY_COUNT(g_text), (u32)g_app.cuda.multiprocessors);
        wcat(g_text, ARRAY_COUNT(g_text), L" | VRAM ");
        wcat_u64(g_text, ARRAY_COUNT(g_text),
                 (u64)g_app.cuda.total_memory / (1024ull * 1024ull));
        wcat(g_text, ARRAY_COUNT(g_text),
             L" MiB | 算力评估：可承受（全车 Maxwell 核心工作显存 < 32 MiB）");
    } else {
        wcat(g_text, ARRAY_COUNT(g_text), L"启动入口：PASS | CUDA：不可用 | ");
        wcat(g_text, ARRAY_COUNT(g_text), g_app.cuda.error);
    }
    SetWindowTextW(g_app.status, g_text);
    EnableWindow(g_app.solve, g_app.cuda.ready);
}

static HWND make_control(DWORD ex_style, const WCHAR *class_name, const WCHAR *text,
                         DWORD style, int id, int x, int y, int width, int height) {
    HWND handle = CreateWindowExW(ex_style, class_name, text,
                                  WS_CHILD | WS_VISIBLE | style,
                                  x, y, width, height, g_app.window,
                                  (HMENU)(ULONG_PTR)id, g_app.instance, 0);
    if (handle && g_app.font) SendMessageW(handle, WM_SETFONT, (WPARAM)g_app.font, TRUE);
    return handle;
}

static void populate_official_maze_combo(void) {
    int series = (int)SendMessageW(g_app.official_series, CB_GETCURSEL, 0, 0);
    if (series < 0 || series >= (int)OFFICIAL_SERIES_COUNT) series = 0;
    SendMessageW(g_app.official_maze, CB_RESETCONTENT, 0, 0);
    int first = -1;
    for (u32 i = 0; i < OFFICIAL_MAZE_COUNT; ++i) {
        const OfficialMazeRecord *record = &g_official_mazes[i];
        if ((int)record->series != series) continue;
        LRESULT item = SendMessageW(g_app.official_maze, CB_ADDSTRING, 0,
                                    (LPARAM)record->label);
        if (item != CB_ERR && item != CB_ERRSPACE) {
            SendMessageW(g_app.official_maze, CB_SETITEMDATA, (WPARAM)item, (LPARAM)i);
            if (first < 0) first = (int)item;
        }
    }
    if (first >= 0) SendMessageW(g_app.official_maze, CB_SETCURSEL, (WPARAM)first, 0);
    EnableWindow(g_app.official_load, first >= 0);
}

static void handle_load_official(void) {
    int selection = (int)SendMessageW(g_app.official_maze, CB_GETCURSEL, 0, 0);
    if (selection < 0) return;
    LRESULT data = SendMessageW(g_app.official_maze, CB_GETITEMDATA, (WPARAM)selection, 0);
    if (data == CB_ERR || (u32)data >= OFFICIAL_MAZE_COUNT) return;
    u32 index = (u32)data;
    Maze decoded;
    if (!decode_official_maze(&decoded, index)) {
        MessageBoxW(g_app.window, L"内置赛题的哈希或迷宫验证失败。",
                    L"赛题完整性错误", MB_OK | MB_ICONERROR);
        return;
    }
    g_app.maze = decoded;
    g_app.official_loaded = 1;
    g_app.official_index = index;
    SendMessageW(g_app.profile, CB_SETCURSEL, decoded.cell_mm == 90u ? 1u : 0u, 0);
    read_solver_params(&g_app.params);
    g_app.path_count = 0;
    g_app.trajectory_count = 0;
    g_app.exact_pass = 0;
    g_app.dynamics_pass = 0;
    show_generation_result();
    InvalidateRect(g_app.window, 0, TRUE);
}

static void create_controls(void) {
    g_app.font = (HFONT)GetStockObject(DEFAULT_GUI_FONT);
    g_app.status = make_control(0, L"STATIC", L"启动入口已通过，正在检测 CUDA……",
                                SS_LEFT, ID_STATUS, 12, 10, 1090, 24);
    make_control(0, L"STATIC", L"迷宫规格", SS_LEFT, 0, 12, 45, 95, 22);
    g_app.profile = make_control(0, L"COMBOBOX", L"", CBS_DROPDOWNLIST | WS_TABSTOP,
                                 ID_PROFILE, 112, 42, 150, 200);
    SendMessageW(g_app.profile, CB_ADDSTRING, 0, (LPARAM)L"Classic 16 / 180 mm");
    SendMessageW(g_app.profile, CB_ADDSTRING, 0, (LPARAM)L"NTF 32 / 90 mm");
    SendMessageW(g_app.profile, CB_SETCURSEL, 0, 0);
    make_control(0, L"STATIC", L"随机种子", SS_LEFT, 0, 12, 75, 95, 22);
    g_app.seed = make_control(WS_EX_CLIENTEDGE, L"EDIT", L"20260823",
                              ES_AUTOHSCROLL | WS_TABSTOP, ID_SEED, 112, 72, 150, 24);
    g_app.generate = make_control(0, L"BUTTON", L"严格随机生成",
                                   BS_PUSHBUTTON | WS_TABSTOP, ID_GENERATE, 276, 42, 135, 28);
    g_app.solve = make_control(0, L"BUTTON", L"CUDA 求最短时间",
                                BS_DEFPUSHBUTTON | WS_TABSTOP, ID_SOLVE, 276, 73, 135, 28);
    EnableWindow(g_app.solve, FALSE);

    make_control(0, L"STATIC", L"赛事档案", SS_LEFT, 0, 12, 108, 95, 22);
    g_app.official_series = make_control(0, L"COMBOBOX", L"",
                                         CBS_DROPDOWNLIST | WS_TABSTOP,
                                         ID_OFFICIAL_SERIES, 112, 104, 244, 260);
    make_control(0, L"STATIC", L"年份 / 场次", SS_LEFT, 0, 12, 139, 95, 22);
    g_app.official_maze = make_control(0, L"COMBOBOX", L"",
                                       CBS_DROPDOWNLIST | WS_TABSTOP,
                                       ID_OFFICIAL_MAZE, 112, 135, 303, 320);
    g_app.official_load = make_control(0, L"BUTTON", L"载入",
                                       BS_PUSHBUTTON | WS_TABSTOP,
                                       ID_OFFICIAL_LOAD, 360, 103, 55, 27);
    for (u32 i = 0; i < OFFICIAL_SERIES_COUNT; ++i)
        SendMessageW(g_app.official_series, CB_ADDSTRING, 0,
                     (LPARAM)g_official_series_names[i]);
    SendMessageW(g_app.official_series, CB_SETCURSEL, 0, 0);
    populate_official_maze_combo();

    static const WCHAR *labels[PARAM_COUNT] = {
        L"直线极速 mm/s", L"对角极速 mm/s", L"转弯速度 mm/s", L"加速度 mm/s²",
        L"制动减速度 mm/s²", L"jerk mm/s³", L"纵向摩擦 μx", L"横向摩擦 μy",
        L"峰值滑移率", L"开路电压 mV", L"电压衰减 mV/s", L"电池内阻 mΩ",
        L"电池容量 mAh", L"整车质量 g", L"底盘长度 mm", L"左右轮距 mm",
        L"轮半径 mm", L"单轮质量 g", L"质心高度 mm", L"电机电阻 mΩ",
        L"电机 Kt mN·m/A", L"电机 Ke mV/(rad/s)", L"限流 mA", L"轮速 Kp mV/(rad/s)",
        L"转弯半径 mm", L"高速直线最大格数", L"积分步长 μs"
    };
    static const WCHAR *defaults[PARAM_COUNT] = {
        L"3500", L"3000", L"1800", L"12000", L"14000", L"120000",
        L"0.900", L"1.050", L"0.080", L"8400", L"90", L"45", L"350",
        L"100", L"90", L"72", L"12", L"5", L"15", L"320", L"18", L"2",
        L"6000", L"40", L"45", L"8", L"500"
    };
    for (u32 i = 0; i < PARAM_COUNT; ++i) {
        int column = i >= 14u ? 1 : 0;
        int row = (int)(i % 14u);
        int x = 12 + column * 252;
        int y = 177 + row * 29;
        make_control(0, L"STATIC", labels[i], SS_LEFT, 0, x, y + 3, 178, 22);
        g_app.param_edits[i] = make_control(WS_EX_CLIENTEDGE, L"EDIT", defaults[i],
                                            ES_AUTOHSCROLL | ES_RIGHT | WS_TABSTOP,
                                            ID_PARAM_BASE + (int)i, x + 181, y, 60, 24);
    }
    g_app.diagonal = make_control(0, L"BUTTON", L"启用 45°/135°/对角线动作",
                                   BS_AUTOCHECKBOX | WS_TABSTOP, ID_DIAGONAL, 12, 590, 238, 24);
    SendMessageW(g_app.diagonal, BM_SETCHECK, BST_CHECKED, 0);
    g_app.multirun = make_control(0, L"BUTTON", L"启用多格高速直线动作",
                                   BS_AUTOCHECKBOX | WS_TABSTOP, ID_MULTIRUN, 260, 590, 244, 24);
    SendMessageW(g_app.multirun, BM_SETCHECK, BST_CHECKED, 0);
    g_app.result = make_control(WS_EX_CLIENTEDGE, L"EDIT", L"",
                                ES_MULTILINE | ES_AUTOVSCROLL | ES_READONLY | WS_VSCROLL,
                                ID_RESULT, 12, 621, 492, 260);
}

static void layout_controls(int width, int height) {
    int result_height = height - 633;
    if (result_height < 150) result_height = 150;
    MoveWindow(g_app.status, 12, 10, width - 24, 24, TRUE);
    MoveWindow(g_app.result, 12, 621, 492, result_height, TRUE);
}

static POINT node_center_point(u32 node, int origin_x, int origin_y, int cell_pixels) {
    u32 cell = node >> 3u;
    u32 x = cell % g_app.maze.n;
    u32 y = cell / g_app.maze.n;
    POINT point;
    point.x = origin_x + (int)x * cell_pixels + cell_pixels / 2;
    point.y = origin_y + ((int)g_app.maze.n - 1 - (int)y) * cell_pixels + cell_pixels / 2;
    return point;
}

static void heading_vector_q10(u32 heading, int *x, int *y) {
    static const int vx[8] = {0, 724, 1024, 724, 0, -724, -1024, -724};
    static const int vy[8] = {-1024, -724, 0, 724, 1024, 724, 0, -724};
    *x = vx[heading & 7u];
    *y = vy[heading & 7u];
}

static int trajectory_handle_pixels(u32 action, POINT start, POINT end,
                                    int cell_pixels) {
    static const u32 base_factor_milli[3] = {347u, 391u, 482u};
    u32 angle_index = (action - 8u) / 2u;
    int dx = end.x - start.x;
    int dy = end.y - start.y;
    u32 chord = (u32)isqrt_u64((u64)(dx * dx + dy * dy));
    u32 nominal_radius = g_app.params.cell_mm / 4u;
    if (!nominal_radius) nominal_radius = 1u;
    u64 effective_radius = (u64)g_app.params.turn_radius_mm *
                           (1000u + g_app.params.slip_milli) / 1000u;
    u32 radius_scale = (u32)(effective_radius * 1000u / nominal_radius);
    if (radius_scale < 700u) radius_scale = 700u;
    if (radius_scale > 1500u) radius_scale = 1500u;
    u32 handle = (u32)((u64)chord * base_factor_milli[angle_index] *
                       radius_scale / 1000000ull);
    u32 minimum = chord * 28u / 100u;
    u32 maximum = chord * 58u / 100u;
    if (handle < minimum) handle = minimum;
    if (handle > maximum) handle = maximum;
    if (handle < 2u && cell_pixels >= 4) handle = 2u;
    return (int)handle;
}

static void draw_trajectory_pass(HDC dc, int origin_x, int origin_y,
                                 int cell_pixels, HPEN pen) {
    HGDIOBJ old_pen = SelectObject(dc, pen);
    if (g_app.trajectory_count < 2u || !g_app.trajectory) {
        POINT first = node_center_point(g_app.path_nodes[0], origin_x, origin_y, cell_pixels);
        MoveToEx(dc, first.x, first.y, 0);
        for (u32 i = 1; i < g_app.path_count; ++i) {
            POINT point = node_center_point(g_app.path_nodes[i], origin_x, origin_y, cell_pixels);
            LineTo(dc, point.x, point.y);
        }
        SelectObject(dc, old_pen);
        return;
    }
    float cell_m = (float)g_app.maze.cell_mm * 0.001f;
    int maze_pixels = (int)g_app.maze.n * cell_pixels;
    const TrajectorySample *sample = &g_app.trajectory[0];
    POINT first;
    first.x = origin_x + (int)(sample->x_m * (float)cell_pixels / cell_m + 0.5f);
    first.y = origin_y + maze_pixels -
              (int)(sample->y_m * (float)cell_pixels / cell_m + 0.5f);
    MoveToEx(dc, first.x, first.y, 0);
    for (u32 i = 1; i < g_app.trajectory_count; ++i) {
        sample = &g_app.trajectory[i];
        int x = origin_x + (int)(sample->x_m * (float)cell_pixels / cell_m + 0.5f);
        int y = origin_y + maze_pixels -
                (int)(sample->y_m * (float)cell_pixels / cell_m + 0.5f);
        LineTo(dc, x, y);
    }
    SelectObject(dc, old_pen);
}

static void draw_trajectory_markers(HDC dc, int origin_x, int origin_y,
                                    int cell_pixels) {
    HPEN marker_pen = CreatePen(PS_SOLID, 2, RGB(132, 72, 180));
    HGDIOBJ old_pen = SelectObject(dc, marker_pen);
    HGDIOBJ old_brush = SelectObject(dc, GetStockObject(NULL_BRUSH));
    int radius = cell_pixels / 6;
    if (radius < 3) radius = 3;
    for (u32 i = 0; i + 1u < g_app.path_count; ++i) {
        if (g_app.path_actions[i] != 14u) continue;
        POINT point = node_center_point(g_app.path_nodes[i], origin_x, origin_y, cell_pixels);
        Ellipse(dc, point.x - radius, point.y - radius,
                point.x + radius + 1, point.y + radius + 1);
    }
    SelectObject(dc, old_brush);
    SelectObject(dc, old_pen);
    DeleteObject(marker_pen);

    POINT start = node_center_point(g_app.path_nodes[0], origin_x, origin_y, cell_pixels);
    if (g_app.trajectory && g_app.trajectory_count) {
        float cell_m = (float)g_app.maze.cell_mm * 0.001f;
        int maze_pixels = (int)g_app.maze.n * cell_pixels;
        start.x = origin_x + (int)(g_app.trajectory[0].x_m * (float)cell_pixels / cell_m + 0.5f);
        start.y = origin_y + maze_pixels -
                  (int)(g_app.trajectory[0].y_m * (float)cell_pixels / cell_m + 0.5f);
    }
    HBRUSH start_brush = CreateSolidBrush(RGB(30, 160, 92));
    old_brush = SelectObject(dc, start_brush);
    old_pen = SelectObject(dc, GetStockObject(NULL_PEN));
    int dot = cell_pixels / 9;
    if (dot < 2) dot = 2;
    Ellipse(dc, start.x - dot, start.y - dot, start.x + dot + 1, start.y + dot + 1);
    SelectObject(dc, old_pen);
    SelectObject(dc, old_brush);
    DeleteObject(start_brush);

    u32 last_node = g_app.path_nodes[g_app.path_count - 1u];
    POINT end = node_center_point(last_node, origin_x, origin_y, cell_pixels);
    int fx, fy;
    heading_vector_q10(last_node & 7u, &fx, &fy);
    if (g_app.trajectory && g_app.trajectory_count) {
        const TrajectorySample *last = &g_app.trajectory[g_app.trajectory_count - 1u];
        float cell_m = (float)g_app.maze.cell_mm * 0.001f;
        int maze_pixels = (int)g_app.maze.n * cell_pixels;
        end.x = origin_x + (int)(last->x_m * (float)cell_pixels / cell_m + 0.5f);
        end.y = origin_y + maze_pixels -
                (int)(last->y_m * (float)cell_pixels / cell_m + 0.5f);
        fx = (int)(f_cos(last->heading_rad) * 1024.0f);
        fy = (int)(-f_sin(last->heading_rad) * 1024.0f);
    }
    int rx = -fy;
    int ry = fx;
    int front = cell_pixels * 3 / 10;
    int rear = cell_pixels / 5;
    int half_width = cell_pixels / 6;
    if (front < 4) front = 4;
    if (rear < 3) rear = 3;
    if (half_width < 2) half_width = 2;
    POINT body[5];
    body[0].x = end.x + fx * front / 1024;
    body[0].y = end.y + fy * front / 1024;
    body[1].x = end.x + fx * (front / 4) / 1024 + rx * half_width / 1024;
    body[1].y = end.y + fy * (front / 4) / 1024 + ry * half_width / 1024;
    body[2].x = end.x - fx * rear / 1024 + rx * (half_width * 3 / 4) / 1024;
    body[2].y = end.y - fy * rear / 1024 + ry * (half_width * 3 / 4) / 1024;
    body[3].x = end.x - fx * rear / 1024 - rx * (half_width * 3 / 4) / 1024;
    body[3].y = end.y - fy * rear / 1024 - ry * (half_width * 3 / 4) / 1024;
    body[4].x = end.x + fx * (front / 4) / 1024 - rx * half_width / 1024;
    body[4].y = end.y + fy * (front / 4) / 1024 - ry * half_width / 1024;
    HBRUSH mouse_brush = CreateSolidBrush(RGB(0, 82, 155));
    HPEN mouse_pen = CreatePen(PS_SOLID, 1, RGB(245, 250, 255));
    old_brush = SelectObject(dc, mouse_brush);
    old_pen = SelectObject(dc, mouse_pen);
    Polygon(dc, body, 5);
    SelectObject(dc, old_pen);
    SelectObject(dc, old_brush);
    DeleteObject(mouse_pen);
    DeleteObject(mouse_brush);
}

static void draw_maze(HDC dc, const RECT *client) {
    if (!g_app.maze.n) return;
    int available_w = client->right - LEFT_PANEL - 24;
    int available_h = client->bottom - 50;
    int cell_pixels_w = available_w / (int)g_app.maze.n;
    int cell_pixels_h = available_h / (int)g_app.maze.n;
    int cell_pixels = cell_pixels_w < cell_pixels_h ? cell_pixels_w : cell_pixels_h;
    if (cell_pixels < 2) return;
    int maze_size = cell_pixels * (int)g_app.maze.n;
    int origin_x = LEFT_PANEL + (available_w - maze_size) / 2;
    int origin_y = 38 + (available_h - maze_size) / 2;

    HBRUSH goal_brush = CreateSolidBrush(RGB(238, 247, 255));
    RECT goal_rect = {
        origin_x + (int)g_app.maze.goal_x0 * cell_pixels,
        origin_y + ((int)g_app.maze.n - 1 - (int)g_app.maze.goal_y1) * cell_pixels,
        origin_x + ((int)g_app.maze.goal_x1 + 1) * cell_pixels,
        origin_y + ((int)g_app.maze.n - (int)g_app.maze.goal_y0) * cell_pixels
    };
    FillRect(dc, &goal_rect, goal_brush);
    DeleteObject(goal_brush);

    HPEN wall_pen = CreatePen(PS_SOLID, g_app.maze.n == 16u ? 2 : 1, RGB(30, 35, 42));
    HGDIOBJ old_pen = SelectObject(dc, wall_pen);
    for (u32 y = 0; y < g_app.maze.n; ++y) {
        for (u32 x = 0; x < g_app.maze.n; ++x) {
            u32 cell = y * g_app.maze.n + x;
            int left = origin_x + (int)x * cell_pixels;
            int top = origin_y + ((int)g_app.maze.n - 1 - (int)y) * cell_pixels;
            int right = left + cell_pixels;
            int bottom = top + cell_pixels;
            u8 walls = g_app.maze.walls[cell];
            if (walls & WALL_N) { MoveToEx(dc, left, top, 0); LineTo(dc, right, top); }
            if (walls & WALL_E) { MoveToEx(dc, right, top, 0); LineTo(dc, right, bottom); }
            if (y == 0u && (walls & WALL_S)) { MoveToEx(dc, left, bottom, 0); LineTo(dc, right, bottom); }
            if (x == 0u && (walls & WALL_W)) { MoveToEx(dc, left, top, 0); LineTo(dc, left, bottom); }
        }
    }
    SelectObject(dc, old_pen);
    DeleteObject(wall_pen);

    if (g_app.path_count > 1u) {
        int under_width = g_app.maze.n == 16u ? 7 : 5;
        int path_width = g_app.maze.n == 16u ? 3 : 2;
        HPEN under_pen = CreatePen(PS_SOLID, under_width, RGB(190, 225, 249));
        HPEN path_pen = CreatePen(PS_SOLID, path_width, RGB(0, 105, 190));
        draw_trajectory_pass(dc, origin_x, origin_y, cell_pixels, under_pen);
        draw_trajectory_pass(dc, origin_x, origin_y, cell_pixels, path_pen);
        draw_trajectory_markers(dc, origin_x, origin_y, cell_pixels);
        DeleteObject(path_pen);
        DeleteObject(under_pen);
    }

    SetBkMode(dc, TRANSPARENT);
    SetTextColor(dc, RGB(70, 76, 85));
    if (g_app.official_loaded && g_app.official_index < OFFICIAL_MAZE_COUNT) {
        wcopy(g_small, ARRAY_COUNT(g_small), L"内置公开赛事档案 | ");
        wcat(g_small, ARRAY_COUNT(g_small), g_official_mazes[g_app.official_index].label);
        wcat(g_small, ARRAY_COUNT(g_small), L" | 蓝线=路径约束闭环刚体回放轨迹");
    } else {
        wcopy(g_small, ARRAY_COUNT(g_small), L"赛制约束随机迷宫 | 蓝线=路径约束闭环刚体轨迹 | 紫圈=180°原地转向");
    }
    TextOutW(dc, origin_x, origin_y + maze_size + 8, g_small, (int)wlen(g_small));
}

static void handle_generate(void) {
    int was_official = g_app.official_loaded;
    g_app.official_loaded = 0;
    read_solver_params(&g_app.params);
    u32 typed_seed = get_edit_u32(g_app.seed, 0u, 0u, 0xffffffffu);
    u32 seed = typed_seed;
    if (!seed) seed = (u32)(GetTickCount64() ^ (GetTickCount64() >> 32u));
    if (!seed) seed = 1u;
    if (was_official && !g_app.seed_dirty) {
        seed = next_seed_u32(seed);
    } else if (g_app.maze.n == g_app.params.n && !g_app.seed_dirty &&
        (u32)g_app.maze.seed == seed) {
        seed = next_seed_u32(seed);
    }
    Maze generated;
    if (!generate_maze(&generated, g_app.params.n, g_app.params.cell_mm, seed)) {
        g_app.official_loaded = was_official;
        MessageBoxW(g_app.window, L"严格迷宫生成后的验证门未通过。", L"生成失败", MB_OK | MB_ICONERROR);
        return;
    }
    g_app.maze = generated;
    set_seed_text(seed);
    g_app.seed_dirty = 0;
    g_app.path_count = 0;
    g_app.trajectory_count = 0;
    g_app.exact_pass = 0;
    g_app.dynamics_pass = 0;
    show_generation_result();
    InvalidateRect(g_app.window, 0, TRUE);
}

static void handle_solve(void) {
    if (!g_app.cuda.ready) {
        MessageBoxW(g_app.window, g_app.cuda.error, L"CUDA 不可用", MB_OK | MB_ICONERROR);
        return;
    }
    read_solver_params(&g_app.params);
    if (g_app.maze.n != g_app.params.n || !validate_maze(&g_app.maze)) handle_generate();
    if (g_app.maze.n != g_app.params.n) return;
    SetWindowTextW(g_app.result, L"CUDA 正在执行矩阵自由全局路径闭合……");
    SetCursor(LoadCursorW(0, IDC_WAIT));
    EnableWindow(g_app.solve, FALSE);
    EnableWindow(g_app.generate, FALSE);
    EnableWindow(g_app.official_load, FALSE);
    EnableWindow(g_app.official_series, FALSE);
    EnableWindow(g_app.official_maze, FALSE);
    SolverParams physical_params = g_app.params;
    int solved = 0;
    g_app.dynamics_route_attempts = 0u;
    for (u32 route_attempt = 0u; route_attempt < 3u && !solved; ++route_attempt) {
        SolverParams route_params = physical_params;
        if (route_attempt >= 1u) {
            route_params.turn_v = route_params.turn_v * 55u / 100u;
            if (route_params.turn_v < 50u) route_params.turn_v = 50u;
        }
        if (route_attempt >= 2u) {
            route_params.diagonal_v = 0u;
            route_params.max_run = 1u;
        }
        g_app.params = route_params;
        ++g_app.dynamics_route_attempts;
        int graph_solved = cuda_solve(&g_app);
        g_app.params = physical_params;
        if (graph_solved) {
            g_text[0] = 0;
            wcat(g_text, ARRAY_COUNT(g_text), L"CUDA 路径图已闭合；GPU 正在筛选 4096 组策略（路线尝试 ");
            wcat_u64(g_text, ARRAY_COUNT(g_text), route_attempt + 1u);
            wcat(g_text, ARRAY_COUNT(g_text), L" / 3）……");
            SetWindowTextW(g_app.result, g_text);
            UpdateWindow(g_app.result);
            solved = optimize_dynamics(&g_app);
        }
    }
    g_app.params = physical_params;
    EnableWindow(g_app.generate, TRUE);
    EnableWindow(g_app.solve, TRUE);
    EnableWindow(g_app.official_load, TRUE);
    EnableWindow(g_app.official_series, TRUE);
    EnableWindow(g_app.official_maze, TRUE);
    SetCursor(LoadCursorW(0, IDC_ARROW));
    if (!solved) {
        SetWindowTextW(g_app.result, g_app.cuda.error);
        MessageBoxW(g_app.window, g_app.cuda.error, L"求解失败", MB_OK | MB_ICONERROR);
    } else {
        show_solve_result();
    }
    InvalidateRect(g_app.window, 0, TRUE);
}

static LRESULT CALLBACK app_wndproc(HWND window, UINT message, WPARAM wparam, LPARAM lparam) {
    switch (message) {
    case WM_CREATE:
        g_app.window = window;
        create_controls();
        handle_generate();
        return 0;
    case WM_APP_CUDA:
        cuda_initialize(&g_app.cuda, g_app.instance);
        set_status_cuda();
        return 0;
    case WM_COMMAND: {
        int id = LOWORD(wparam);
        if (id == ID_SEED && HIWORD(wparam) == EN_CHANGE && !g_app.setting_seed_text)
            g_app.seed_dirty = 1;
        else if (id == ID_OFFICIAL_SERIES && HIWORD(wparam) == CBN_SELCHANGE)
            populate_official_maze_combo();
        else if (id == ID_OFFICIAL_LOAD && HIWORD(wparam) == BN_CLICKED)
            handle_load_official();
        else if (id == ID_GENERATE && HIWORD(wparam) == BN_CLICKED) handle_generate();
        else if (id == ID_SOLVE && HIWORD(wparam) == BN_CLICKED) handle_solve();
        return 0;
    }
    case WM_SIZE:
        layout_controls(LOWORD(lparam), HIWORD(lparam));
        InvalidateRect(window, 0, TRUE);
        return 0;
    case WM_GETMINMAXINFO: {
        MINMAXINFO *info = (MINMAXINFO *)lparam;
        info->ptMinTrackSize.x = 1100;
        info->ptMinTrackSize.y = 820;
        return 0;
    }
    case WM_PAINT: {
        PAINTSTRUCT paint;
        HDC dc = BeginPaint(window, &paint);
        RECT client;
        GetClientRect(window, &client);
        HBRUSH background = CreateSolidBrush(RGB(250, 251, 253));
        RECT maze_background = {LEFT_PANEL, 34, client.right, client.bottom};
        FillRect(dc, &maze_background, background);
        DeleteObject(background);
        draw_maze(dc, &client);
        EndPaint(window, &paint);
        return 0;
    }
    case WM_DESTROY:
        cuda_shutdown(&g_app.cuda);
        if (g_app.trajectory) {
            VirtualFree(g_app.trajectory, 0, MEM_RELEASE);
            g_app.trajectory = 0;
            g_app.trajectory_count = 0;
            g_app.trajectory_capacity = 0;
        }
        PostQuitMessage(0);
        return 0;
    default:
        return DefWindowProcW(window, message, wparam, lparam);
    }
}

#if !defined(MM_CORE_TEST) && !defined(MM_NATIVE_TEST)
__declspec(noreturn) void WINAPI wWinMainCRTStartup(void) {
    if (g_relocation_anchor != &g_app) {
        MessageBoxW(0, L"可执行文件重定位自检失败。", L"Matrix-Free Micromouse", MB_OK | MB_ICONERROR);
        ExitProcess(19u);
    }
    memset(&g_app, 0, sizeof(g_app));
    g_app.instance = GetModuleHandleW(0);

    WNDCLASSEXW window_class;
    memset(&window_class, 0, sizeof(window_class));
    window_class.cbSize = sizeof(window_class);
    window_class.style = CS_HREDRAW | CS_VREDRAW;
    window_class.lpfnWndProc = app_wndproc;
    window_class.hInstance = g_app.instance;
    window_class.hIcon = LoadIconW(g_app.instance, MAKEINTRESOURCEW(1));
    if (!window_class.hIcon) window_class.hIcon = LoadIconW(0, IDI_APPLICATION);
    window_class.hIconSm = window_class.hIcon;
    window_class.hCursor = LoadCursorW(0, IDC_ARROW);
    window_class.hbrBackground = (HBRUSH)(COLOR_WINDOW + 1);
    window_class.lpszClassName = L"MatrixFreeMicromouseReliableWindow";
    if (!RegisterClassExW(&window_class)) {
        MessageBoxW(0, L"Windows 窗口类注册失败。", L"Matrix-Free Micromouse", MB_OK | MB_ICONERROR);
        ExitProcess(20u);
    }

    HWND window = CreateWindowExW(0, window_class.lpszClassName,
                                   L"Matrix-Free Micromouse v3.1 — CUDA 全车 Maxwell 有限应变",
                                   WS_OVERLAPPEDWINDOW | WS_CLIPCHILDREN,
                                   CW_USEDEFAULT, CW_USEDEFAULT, 1440, 900,
                                   0, 0, g_app.instance, 0);
    if (!window) {
        MessageBoxW(0, L"主窗口创建失败。", L"Matrix-Free Micromouse", MB_OK | MB_ICONERROR);
        ExitProcess(21u);
    }
    g_app.window = window;
    ShowWindow(window, SW_SHOW);
    UpdateWindow(window);
    PostMessageW(window, WM_APP_CUDA, 0, 0);

    MSG message;
    while (GetMessageW(&message, 0, 0, 0) > 0) {
        TranslateMessage(&message);
        DispatchMessageW(&message);
    }
    ExitProcess((u32)message.wParam);
}
#elif defined(MM_CORE_TEST)
__declspec(noreturn) void WINAPI wWinMainCRTStartup(void) {
    Maze maze;
    Maze previous;
    u32 changing_seed = 20260823u;
    for (u32 i = 0; i < 4096u; ++i) {
        u32 next = next_seed_u32(changing_seed);
        if (!next || next == changing_seed) ExitProcess(29u);
        changing_seed = next;
    }
    u32 series_counts[OFFICIAL_SERIES_COUNT];
    memset(series_counts, 0, sizeof(series_counts));
    for (u32 i = 0; i < OFFICIAL_MAZE_COUNT; ++i) {
        if (!decode_official_maze(&maze, i)) ExitProcess(40u);
        if (g_official_mazes[i].series >= OFFICIAL_SERIES_COUNT) ExitProcess(41u);
        ++series_counts[g_official_mazes[i].series];
    }
    for (u32 i = 0; i < OFFICIAL_SERIES_COUNT; ++i)
        if (!series_counts[i]) ExitProcess(42u + i);
    for (u32 profile = 0; profile < 2u; ++profile) {
        u32 n = profile ? 32u : 16u;
        u32 cell_mm = profile ? 90u : 180u;
        for (u64 seed = 1ull; seed <= 128ull; ++seed) {
            if (!generate_maze(&maze, n, cell_mm, seed) || !validate_maze(&maze))
                ExitProcess(30u + profile);
            if (seed > 1ull) {
                int same = 1;
                for (u32 cell = 0; cell < n * n; ++cell) {
                    if (maze.walls[cell] != previous.walls[cell]) { same = 0; break; }
                }
                if (same) ExitProcess(32u + profile);
            }
            memcpy(&previous, &maze, sizeof(maze));
        }
    }
    ExitProcess(0u);
}
#endif