decuda 0.1.1

//! Integration test: feed the canonical `saxpy.cu` fixture through the full
//! migration pipeline and check that each backend's output mentions the
//! expected CUDA constructs in translated form.

use std::fs;
use std::path::Path;

use decuda::cli::Target;
use decuda::ir::{BuiltinKind, CudaQualifier, IrNode};
use decuda::migrate::{run, MigrateOptions};
use decuda::targets::for_target;

use decuda::parser::translate_source;
use decuda::preprocess::preprocess;
use decuda::walker;
use tempfile::tempdir;

fn fixture_path() -> std::path::PathBuf {
    Path::new(env!("CARGO_MANIFEST_DIR"))
        .join("examples/cu/saxpy.cu")
        .to_path_buf()
}

/// Load a named fixture from `examples/cu/<name>` and parse it.
fn load_fixture(name: &str) -> TranslationUnit {
    let p = Path::new(env!("CARGO_MANIFEST_DIR"))
        .join("examples/cu")
        .join(name);
    let raw = fs::read_to_string(&p)
        .unwrap_or_else(|e| panic!("read {}: {e}", p.display()));
    translate_source(raw, p)
}

use decuda::ir::TranslationUnit;

#[test]
fn parser_extracts_qualifier_for_fixture() {
    let p = fixture_path();
    let u = translate_source(fs::read_to_string(&p).unwrap(), p.clone());
    assert!(u.nodes.iter().any(|n| matches!(
        n,
        decuda::ir::IrNode::QualifierDecl { qualifier: decuda::ir::CudaQualifier::Global, .. }
    )));
}

#[test]
fn hip_backend_renames_runtime_calls() {
    let p = fixture_path();
    let raw = fs::read_to_string(&p).unwrap();
    let u = translate_source(raw.clone(), p);
    let backend = for_target(Target::Hip);
    let out = backend.emit(&u);
    assert!(out.contains("hipMalloc"), "hip output should rename cudaMalloc -> hipMalloc:\n{out}");
    assert!(out.contains("hipFree"), "hip output should rename cudaFree -> hipFree:\n{out}");
    assert!(out.contains("threadIdx"), "HIP keeps threadIdx unchanged:\n{out}");
    // `cudaMalloc` may still appear inside a comment that documents the
    // fixture; check that it no longer appears OUTSIDE comments.
    assert!(
        !contains_outside_comments(&out, "cudaMalloc"),
        "HIP output must not contain cudaMalloc outside comments:\n{out}"
    );
}

/// True if `needle` appears in `haystack` in any non-comment region. Comments
/// (lines starting with `//`) are ignored.
fn contains_outside_comments(haystack: &str, needle: &str) -> bool {
    for line in haystack.lines() {
        let stripped = line.trim_start();
        if stripped.starts_with("//") {
            continue;
        }
        if line.contains(needle) {
            return true;
        }
    }
    false
}

#[test]
fn opencl_backend_emits_cl_calls() {
    let p = fixture_path();
    let u = translate_source(fs::read_to_string(&p).unwrap(), p);
    let backend = for_target(Target::Opencl);
    let out = backend.emit(&u);
    assert!(out.contains("get_local_id"), "OpenCL output should use get_local_id:\n{out}");
    assert!(out.contains("__kernel") || out.contains("clEnqueue"), "OpenCL needs kernel/launch replacement:\n{out}");
}

#[test]
fn sycl_backend_emits_placeholder_block() {
    let p = fixture_path();
    let u = translate_source(fs::read_to_string(&p).unwrap(), p);
    let backend = for_target(Target::Sycl);
    let out = backend.emit(&u);
    assert!(
        out.contains("sycl::") || out.contains("queue.submit"),
        "SYCL output should mention sycl:\n{out}"
    );
}

#[test]
fn rust_backend_emits_cust_launch() {
    let p = fixture_path();
    let u = translate_source(fs::read_to_string(&p).unwrap(), p);
    let backend = for_target(Target::Rust);
    let out = backend.emit(&u);
    assert!(out.contains("cust"), "Rust output should reference `cust`:\n{out}");
    assert!(out.contains("// Generated by decuda"), "Rust banner missing:\n{out}");
}

#[test]
fn end_to_end_migrate_writes_all_targets() {
    let dir = tempdir().unwrap();
    let out_dir = dir.path().join("out");
    let opts = MigrateOptions {
        input: fixture_path(),
        output: out_dir.clone(),
        target: Target::All,
        dry_run: false,
        verbose: false,
        filter: None,
    };
    let report = run(opts).expect("migration should succeed");

    for t in Target::iter_real() {
        let dir = out_dir.join(t.as_str());
        assert!(dir.is_dir(), "missing output dir for {t:?}: {}", dir.display());
        let count = fs::read_dir(&dir).unwrap().count();
        assert!(count > 0, "{t:?} output dir is empty: {}", dir.display());
    }

    // The report should mention each backend at least once.
    let json = serde_json::to_string(&report).unwrap();
    assert!(json.contains("hip"));
    assert!(json.contains("sycl"));
    assert!(json.contains("rust"));
    assert!(json.contains("opencl"));
}

#[test]
fn dry_run_does_not_write_files() {
    let dir = tempdir().unwrap();
    let opts = MigrateOptions {
        input: fixture_path(),
        output: dir.path().join("out"),
        target: Target::Hip,
        dry_run: true,
        verbose: false,
        filter: None,
    };
    let _ = run(opts).unwrap();
    assert!(
        !dir.path().join("out/hip").exists(),
        "dry-run should not create output dir"
    );
}

#[test]
fn header_replacements_applied() {
    let src = "#include <cuda_runtime.h>\n// rest\n";
    let u = translate_source(src.to_string(), "x.cu".into());
    let backend = for_target(Target::Hip);
    let out = backend.emit(&u);
    assert!(out.contains("hip/hip_runtime.h"));
    let backend = for_target(Target::Opencl);
    let out = backend.emit(&u);
    assert!(out.contains("CL/cl.h"));
}

#[test]
fn rich_fixture_translates_atomics_and_shared() {
    let p = Path::new(env!("CARGO_MANIFEST_DIR")).join("examples/cu/rich.cu");
    let raw = fs::read_to_string(&p).unwrap();
    let u = translate_source(raw, p.clone());

    let backend = for_target(Target::Hip);
    let out = backend.emit(&u);
    // Atomic intrinsic is preserved as-is on HIP.
    assert!(out.contains("atomicAdd"), "HIP must keep atomicAdd name:\n{out}");
    // __shared__ preserved in HIP.
    assert!(out.contains("__shared__"));
    // __syncthreads() preserved.
    assert!(out.contains("__syncthreads()"));
    // __constant__ preserved (HIP-compatible).
    assert!(out.contains("__constant__"));
    // __syncwarp call preserved.
    assert!(out.contains("__syncwarp"));

    let backend = for_target(Target::Opencl);
    let out = backend.emit(&u);
    assert!(out.contains("barrier(CLK_LOCAL_MEM_FENCE)"));
    assert!(out.contains("__local"), "OpenCL must rename __shared__ -> __local");
    assert!(out.contains("__constant"), "OpenCL keeps __constant");
}

#[test]
fn rich_fixture_2d_launch_dim3_grid() {
    let p = Path::new(env!("CARGO_MANIFEST_DIR")).join("examples/cu/rich.cu");
    let raw = fs::read_to_string(&p).unwrap();
    let u = translate_source(raw, p);

    let backend = for_target(Target::Hip);
    let out = backend.emit(&u);
    // dim3 grid/block survive unchanged.
    assert!(out.contains("dim3 grid"));
    assert!(out.contains("dim3 block"));
    // hipLaunchKernelGGL is used for every launch.
    assert!(out.contains("hipLaunchKernelGGL"));
}

#[test]
fn rich_fixture_no_cuda_names_leak_in_hip_output() {
    let p = Path::new(env!("CARGO_MANIFEST_DIR")).join("examples/cu/rich.cu");
    let raw = fs::read_to_string(&p).unwrap();
    let u = translate_source(raw, p);
    let backend = for_target(Target::Hip);
    let out = backend.emit(&u);
    assert!(
        !contains_outside_comments(&out, "cudaMalloc"),
        "HIP output must not contain cudaMalloc outside comments"
    );
    assert!(
        !contains_outside_comments(&out, "cudaFree"),
        "HIP output must not contain cudaFree outside comments"
    );
}

#[test]
fn dir_walker_handles_symlinked_root_without_infinite_loop() {
    use std::os::unix::fs::symlink;
    use tempfile::tempdir;
    let dir = tempdir().unwrap();
    fs::write(dir.path().join("a.cu"), "__global__ void k() {}").unwrap();
    let link_dir = dir.path().join("link");
    let _ = symlink(dir.path(), &link_dir);
    // We don't assert the exact set of files (symlink semantics differ across
    // platforms); we only assert the walker does not loop or crash.
    let _ = walker::collect_cuda_files(&link_dir, None).unwrap();
}

// ---------------------------------------------------------------------------
// Preprocessor invariants: launches inside comments/strings are not captured.
// ---------------------------------------------------------------------------

#[test]
fn preprocessor_skips_launch_in_line_comment() {
    let src = "// foo<<<1, 1>>>(a);\nbar<<<2, 2>>>(b);\n";
    let p = preprocess(src);
    assert_eq!(p.launches.len(), 1, "only the real launch should be captured");
    assert_eq!(p.launches[0].kernel, "bar");
}

#[test]
fn preprocessor_skips_launch_in_block_comment() {
    let src = "/* bar<<<2, 2>>>(b, c); */\nk<<<1, 1>>>(x);\n";
    let p = preprocess(src);
    assert_eq!(p.launches.len(), 1);
    assert_eq!(p.launches[0].kernel, "k");
}

#[test]
fn preprocessor_skips_launch_in_string_literal() {
    let src = "const char* s = \"k<<<1, 1>>>(x);\";\nk<<<1, 1>>>(y);\n";
    let p = preprocess(src);
    assert_eq!(p.launches.len(), 1);
    assert_eq!(p.launches[0].args, vec!["y".to_string()]);
}

#[test]
fn preprocessor_captures_full_launch_with_smem_and_stream() {
    let src = "k<<<g, b, 64, stream>>>(a, b);";
    let p = preprocess(src);
    assert_eq!(p.launches.len(), 1);
    let l = &p.launches[0];
    assert_eq!(l.kernel, "k");
    assert_eq!(l.grid, "g");
    assert_eq!(l.block, "b");
    assert_eq!(l.smem.as_deref(), Some("64"));
    assert_eq!(l.stream.as_deref(), Some("stream"));
    assert_eq!(l.args, vec!["a".to_string(), "b".to_string()]);
}

#[test]
fn launch_in_comment_fixture_only_captures_real_launch() {
    let p = Path::new(env!("CARGO_MANIFEST_DIR"))
        .join("examples/cu/launch_in_comment.cu");
    let raw = fs::read_to_string(&p).unwrap();
    let u = translate_source(raw, p);
    let launches: Vec<_> = u
        .nodes
        .iter()
        .filter_map(|n| match n {
            IrNode::KernelLaunch { kernel, .. } => Some(kernel.clone()),
            _ => None,
        })
        .collect();
    assert_eq!(launches, vec!["real_kernel".to_string()]);
}

// ---------------------------------------------------------------------------
// Parser: built-in field access and kernel-def detection.
// ---------------------------------------------------------------------------

#[test]
fn parser_records_threadidx_field_access() {
    let u = translate_source("int i = threadIdx.x;".to_string(), "x.cu".into());
    let node = u.nodes.iter().find_map(|n| match n {
        n @ IrNode::BuiltinRef { kind: BuiltinKind::ThreadIdx, .. } => Some(n),
        _ => None,
    });
    let node = node.expect("threadIdx builtin ref");
    match node {
        IrNode::BuiltinRef { has_field_access, .. } => {
            assert!(*has_field_access, "threadIdx.x should set has_field_access");
        }
        _ => unreachable!(),
    }
}

#[test]
fn parser_finds_multiple_qualifiers_on_one_kernel() {
    let u = translate_source(
        "__global__ __device__ void k(int* x) { *x = 1; }".to_string(),
        "x.cu".into(),
    );
    let globals: Vec<_> = u
        .nodes
        .iter()
        .filter_map(|n| match n {
            IrNode::QualifierDecl { qualifier: CudaQualifier::Global, .. } => Some(()),
            _ => None,
        })
        .collect();
    assert!(!globals.is_empty(), "expected a __global__ qualifier");
}

// ---------------------------------------------------------------------------
// Emit: byte-shift correctness and overlap handling.
// ---------------------------------------------------------------------------

#[test]
fn emit_applies_multiple_replacements_without_corrupting_positions() {
    // Three distinct rewrites in one source: a header, a runtime call, and a
    // builtin. Each lands at a different byte offset and has a different
    // replacement length. If the cumulative byte shift is wrong, a later
    // edit lands at the wrong position and the output is corrupted.
    let src = "#include <cuda_runtime.h>\n"
        .to_string()
        + "__global__ void k() { int i = threadIdx.x; cudaMalloc((void**)&p, n); }";
    let u = translate_source(src, "x.cu".into());

    let hip = for_target(Target::Hip);
    let out = hip.emit(&u);
    // Header replaced, threadIdx preserved (HIP keeps .x), cudaMalloc renamed.
    assert!(out.contains("hip/hip_runtime.h"), "header rewrite missing:\n{out}");
    assert!(out.contains("threadIdx.x"), "HIP must keep threadIdx.x:\n{out}");
    assert!(out.contains("hipMalloc"), "cudaMalloc -> hipMalloc missing:\n{out}");
    // The original `#include <cuda_runtime.h>` directive must be gone; the
    // header rewriter emits `#include <hip/hip_runtime.h> /* was: cuda_runtime.h */`,
    // so the only remaining `cuda_runtime.h` occurrence is inside that comment.
    assert!(
        !out.contains("#include <cuda_runtime.h>"),
        "original cuda_runtime.h include must be replaced:\n{out}"
    );
    assert!(!contains_outside_comments(&out, "cudaMalloc"));

    let opencl = for_target(Target::Opencl);
    let out = opencl.emit(&u);
    // OpenCL consumes the `.x` field access and uses get_local_id(0).
    assert!(out.contains("get_local_id(0)"), "OpenCL threadIdx rewrite missing:\n{out}");
    assert!(!out.contains("threadIdx.x"), "OpenCL must drop threadIdx.x:\n{out}");
}

#[test]
fn emit_overlapping_edits_keep_earlier_one() {
    // `#include <cuda_runtime.h>` contains the substring `cuda_runtime`,
    // which the runtime-call regex would also match. The HeaderInclude node
    // starts earlier, so it must win and the runtime-call edit must be
    // dropped. Otherwise the output would contain a mangled half-rewritten
    // include line.
    let src = "#include <cuda_runtime.h>\n".to_string();
    let u = translate_source(src, "x.cu".into());
    let hip = for_target(Target::Hip);
    let out = hip.emit(&u);
    assert!(out.contains("hip/hip_runtime.h"));
    // The line should still be a single #include directive, not two.
    let include_lines = out
        .lines()
        .filter(|l| l.trim_start().starts_with("#include"))
        .count();
    assert_eq!(include_lines, 1, "expected exactly one #include line:\n{out}");
}

#[test]
fn emit_empty_fixture_produces_only_banner_and_source() {
    let p = Path::new(env!("CARGO_MANIFEST_DIR")).join("examples/cu/empty.cu");
    let raw = fs::read_to_string(&p).unwrap();
    let u = translate_source(raw.clone(), p);
    assert!(u.nodes.is_empty(), "empty fixture should yield no IR nodes");
    let hip = for_target(Target::Hip);
    let out = hip.emit(&u);
    assert!(out.contains("// Generated by decuda"));
    assert!(out.contains("int main() { return 0; }"));
}

// ---------------------------------------------------------------------------
// Headers: every include-replacement path.
// ---------------------------------------------------------------------------

#[test]
fn header_replacements_for_every_known_header() {
    let p = Path::new(env!("CARGO_MANIFEST_DIR"))
        .join("examples/cu/headers_only.cu");
    let raw = fs::read_to_string(&p).unwrap();
    let u = translate_source(raw, p);

    let hip = for_target(Target::Hip);
    let out = hip.emit(&u);
    assert!(out.contains("hip/hip_runtime.h"));
    assert!(out.contains("hip/device_functions.h"));

    let sycl = for_target(Target::Sycl);
    let out = sycl.emit(&u);
    assert!(out.contains("sycl/sycl.hpp"));

    let opencl = for_target(Target::Opencl);
    let out = opencl.emit(&u);
    assert!(out.contains("CL/cl.h"));

    // Generic cuda_*.h header (cuda_fp16.h) becomes a TODO placeholder on
    // every target.
    let rust = for_target(Target::Rust);
    let out = rust.emit(&u);
    assert!(out.contains("TODO"), "generic cuda_*.h should become a TODO:\n{out}");
}

// ---------------------------------------------------------------------------
// Walker: output-path mapping and error handling.
// ---------------------------------------------------------------------------

#[test]
fn walker_maps_output_extension_per_target() {
    use std::path::Path;
    let root = Path::new("/proj/src");
    let file = Path::new("/proj/src/kernels/foo.cu");
    assert_eq!(
        walker::map_output_path(root, Path::new("/out"), "hip", file)
            .file_name()
            .unwrap(),
        "foo.hip.cpp"
    );
    assert_eq!(
        walker::map_output_path(root, Path::new("/out"), "sycl", file)
            .file_name()
            .unwrap(),
        "foo.sycl.cpp"
    );
    assert_eq!(
        walker::map_output_path(root, Path::new("/out"), "rust", file)
            .file_name()
            .unwrap(),
        "foo.rs"
    );
    assert_eq!(
        walker::map_output_path(root, Path::new("/out"), "opencl", file)
            .file_name()
            .unwrap(),
        "foo.cl"
    );
}

#[test]
fn walker_preserves_subdirectory_layout() {
    use std::path::Path;
    let root = Path::new("/proj/src");
    let file = Path::new("/proj/src/sub/a/b.cu");
    let out = walker::map_output_path(root, Path::new("/out"), "hip", file);
    assert!(out.ends_with("/out/hip/sub/a/b.hip.cpp"));
}

#[test]
fn walker_rejects_nonexistent_path() {
    let p = Path::new("/this/path/does/not/exist/decuda_test");
    let res = walker::collect_cuda_files(p, None);
    assert!(res.is_err());
}

// ---------------------------------------------------------------------------
// cuda_db: math intrinsics and type aliases.
// ---------------------------------------------------------------------------

#[test]
fn cuda_db_math_intrinsics_map_per_target() {
    let e = decuda::cuda_db::lookup("__sinf").expect("__sinf in db");
    assert_eq!(e.hip, Some("__sinf"));
    assert_eq!(e.rust, Some("__sinf"));
    assert_eq!(e.opencl, Some("sinf"), "opencl strips leading underscores");
}

#[test]
fn cuda_db_dim3_type_alias() {
    let e = decuda::cuda_db::lookup("dim3").expect("dim3 in db");
    assert_eq!(e.hip, Some("dim3"));
    assert_eq!(e.sycl, Some("sycl::range<3>"));
}

#[test]
fn cuda_db_supports_predicate_matches_mapping() {
    let e = decuda::cuda_db::lookup("cudaMalloc").unwrap();
    assert!(e.supports(Target::Hip));
    assert!(!e.supports(Target::Opencl));
}

// ---------------------------------------------------------------------------
// Migrate: filter, report warnings, directory layout.
// ---------------------------------------------------------------------------

#[test]
fn migrate_filter_restricts_files() {
    let dir = tempdir().unwrap();
    fs::write(dir.path().join("keep.cu"), "__global__ void k() {}").unwrap();
    fs::write(dir.path().join("skip.cu"), "__global__ void k() {}").unwrap();
    let out = dir.path().join("out");

    let report = run(MigrateOptions {
        input: dir.path().to_path_buf(),
        output: out.clone(),
        target: Target::Hip,
        dry_run: false,
        verbose: false,
        filter: Some("keep".into()),
    })
    .expect("migration should succeed");

    assert_eq!(report.files.len(), 1, "filter should select exactly one file");
    assert!(report.files[0].source.file_name().unwrap() == "keep.cu");
}

#[test]
fn migrate_directory_input_preserves_layout() {
    let dir = tempdir().unwrap();
    fs::create_dir_all(dir.path().join("sub")).unwrap();
    fs::write(
        dir.path().join("sub/k.cu"),
        "#include <cuda_runtime.h>\n__global__ void k() {}\n",
    )
    .unwrap();
    let out = dir.path().join("out");

    run(MigrateOptions {
        input: dir.path().to_path_buf(),
        output: out.clone(),
        target: Target::Hip,
        dry_run: false,
        verbose: false,
        filter: None,
    })
    .unwrap();

    let hip_out = out.join("hip").join("sub").join("k.hip.cpp");
    assert!(hip_out.is_file(), "expected nested output at {}", hip_out.display());
}

#[test]
fn migrate_records_warning_for_unmapped_api() {
    // cudaMalloc has no OpenCL mapping; the report must record a warning.
    let dir = tempdir().unwrap();
    fs::write(
        dir.path().join("k.cu"),
        "#include <cuda_runtime.h>\nint main(){ cudaMalloc((void**)&p, n); }\n",
    )
    .unwrap();
    let out = dir.path().join("out");

    let report = run(MigrateOptions {
        input: dir.path().join("k.cu"),
        output: out,
        target: Target::Opencl,
        dry_run: false,
        verbose: false,
        filter: None,
    })
    .unwrap();

    let has_unmapped_warning = report
        .by_file
        .values()
        .flat_map(|m| m.values())
        .flatten()
        .any(|(_, msg)| msg.contains("cudaMalloc"));
    assert!(
        has_unmapped_warning,
        "expected a warning about unmapped cudaMalloc for OpenCL"
    );
}

#[test]
fn migrate_writes_json_report_file() {
    let dir = tempdir().unwrap();
    fs::write(dir.path().join("k.cu"), "__global__ void k() {}\n").unwrap();
    let out = dir.path().join("out");

    let report = run(MigrateOptions {
        input: dir.path().join("k.cu"),
        output: out.clone(),
        target: Target::Hip,
        dry_run: false,
        verbose: false,
        filter: None,
    })
    .unwrap();

    assert!(report.report_path.starts_with(&out), "report path under output");
    assert!(
        report.report_path.file_name().unwrap().to_string_lossy().starts_with("migration-report.json."),
        "report file name should be migration-report.json.<timestamp>"
    );
    assert!(report.report_path.is_file(), "json report file should exist on disk");
}

// ---------------------------------------------------------------------------
// Backend banners: every target emits a decuda banner.
// ---------------------------------------------------------------------------

#[test]
fn every_backend_emits_decuda_banner() {
    let src = "__global__ void k() {}\n";
    let u = translate_source(src.to_string(), "x.cu".into());
    for t in Target::iter_real() {
        let backend = for_target(t);
        let out = backend.emit(&u);
        assert!(
            out.contains("decuda"),
            "{t:?} banner should mention decuda:\n{out}"
        );
    }
}

// ===========================================================================
// Complex fixture: histogram.cu
//   atomicAdd/atomicMin/atomicMax, shared bins, grid-stride loop, __syncthreads,
//   __syncwarp, __laneid, warpSize, __constant__ memory, __device__ helper.
// ===========================================================================

#[test]
fn histogram_parser_harvests_all_constructs() {
    let u = load_fixture("histogram.cu");

    // Qualifiers: __global__, __device__, __shared__, __constant__, __forceinline__.
    let mut has_global = false;
    let mut has_device = false;
    let mut has_shared = false;
    let mut has_constant = false;
    for n in &u.nodes {
        if let IrNode::QualifierDecl { qualifier, .. } = n {
            match qualifier {
                CudaQualifier::Global => has_global = true,
                CudaQualifier::Device => has_device = true,
                CudaQualifier::Shared => has_shared = true,
                CudaQualifier::Constant => has_constant = true,
                _ => {}
            }
        }
    }
    assert!(has_global, "expected __global__");
    assert!(has_device, "expected __device__");
    assert!(has_shared, "expected __shared__");
    assert!(has_constant, "expected __constant__");

    // Atomics: atomicAdd, atomicMin, atomicMax.
    let atomic_names: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::AtomicIntrinsic { name, .. } => Some(name.clone()),
        _ => None,
    }).collect();
    assert!(atomic_names.iter().any(|n| n == "atomicAdd"), "expected atomicAdd");
    assert!(atomic_names.iter().any(|n| n == "atomicMin"), "expected atomicMin");
    assert!(atomic_names.iter().any(|n| n == "atomicMax"), "expected atomicMax");

    // Builtins: __syncthreads, __syncwarp, __laneid, warpSize.
    let builtin_kinds: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::BuiltinRef { kind, .. } => Some(*kind),
        _ => None,
    }).collect();
    assert!(builtin_kinds.contains(&BuiltinKind::SyncThreads));
    assert!(builtin_kinds.contains(&BuiltinKind::SyncWarp));
    assert!(builtin_kinds.contains(&BuiltinKind::LaneId));
    assert!(builtin_kinds.contains(&BuiltinKind::WarpSize));

    // Runtime calls: cudaMalloc, cudaFree, cudaMemcpy, cudaMemset.
    let rt_names: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::RuntimeCall { name, .. } => Some(name.clone()),
        _ => None,
    }).collect();
    assert!(rt_names.iter().any(|n| n == "cudaMalloc"));
    assert!(rt_names.iter().any(|n| n == "cudaFree"));
    assert!(rt_names.iter().any(|n| n == "cudaMemcpy"));
    assert!(rt_names.iter().any(|n| n == "cudaMemset"));

    // Two kernel launches.
    let launches: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::KernelLaunch { kernel, .. } => Some(kernel.clone()),
        _ => None,
    }).collect();
    assert!(launches.iter().any(|k| k == "histogram"));
    assert!(launches.iter().any(|k| k == "warp_reduce"));
}

#[test]
fn histogram_hip_keeps_atomics_and_builtins() {
    let u = load_fixture("histogram.cu");
    let out = for_target(Target::Hip).emit(&u);
    // HIP keeps atomic names unchanged.
    assert!(out.contains("atomicAdd"), "HIP keeps atomicAdd:\n{out}");
    assert!(out.contains("atomicMin"), "HIP keeps atomicMin:\n{out}");
    assert!(out.contains("atomicMax"), "HIP keeps atomicMax:\n{out}");
    // HIP keeps threadIdx/blockIdx/blockDim with .x and __syncthreads/__syncwarp.
    assert!(out.contains("threadIdx.x"));
    assert!(out.contains("blockIdx.x"));
    assert!(out.contains("blockDim.x"));
    assert!(out.contains("gridDim.x"));
    assert!(out.contains("__syncthreads()"));
    assert!(out.contains("__syncwarp"));
    // HIP keeps __shared__ and __constant__.
    assert!(out.contains("__shared__"));
    assert!(out.contains("__constant__"));
    // Runtime calls renamed.
    assert!(out.contains("hipMalloc"));
    assert!(out.contains("hipFree"));
    assert!(out.contains("hipMemcpy"));
    assert!(out.contains("hipMemset"));
    // Header replaced.
    assert!(out.contains("hip/hip_runtime.h"));
    // Launches rewritten to hipLaunchKernelGGL.
    assert!(out.contains("hipLaunchKernelGGL"));
    // No cudaMalloc outside comments.
    assert!(!contains_outside_comments(&out, "cudaMalloc"));
}

#[test]
fn histogram_opencl_rewrites_builtins_and_qualifiers() {
    let u = load_fixture("histogram.cu");
    let out = for_target(Target::Opencl).emit(&u);
    // OpenCL builtins.
    assert!(out.contains("get_local_id(0)"));
    assert!(out.contains("get_group_id(0)"));
    assert!(out.contains("get_local_size(0)"));
    assert!(out.contains("get_num_groups(0)"));
    assert!(out.contains("barrier(CLK_LOCAL_MEM_FENCE)"));
    // OpenCL qualifiers.
    assert!(out.contains("__kernel"));
    assert!(out.contains("__local"));
    assert!(out.contains("__constant"));
    // OpenCL keeps atomic names.
    assert!(out.contains("atomicAdd"));
    assert!(out.contains("atomicMin"));
    assert!(out.contains("atomicMax"));
    // Header replaced.
    assert!(out.contains("CL/cl.h"));
    // Launches rewritten to clEnqueueNDRangeKernel.
    assert!(out.contains("clEnqueueNDRangeKernel"));
    // threadIdx.x must be consumed (OpenCL uses scalar get_local_id(0)).
    assert!(!out.contains("threadIdx.x"));
}

#[test]
fn histogram_sycl_and_rust_emit_placeholders() {
    let u = load_fixture("histogram.cu");
    let sycl = for_target(Target::Sycl).emit(&u);
    assert!(sycl.contains("sycl::") || sycl.contains("queue.submit"));
    assert!(sycl.contains("item.get_local_id()"));
    assert!(sycl.contains("item.barrier"));

    let rust = for_target(Target::Rust).emit(&u);
    assert!(rust.contains("thread_idx"));
    assert!(rust.contains("block_idx"));
    assert!(rust.contains("group.sync()"));
    assert!(rust.contains("cust"));
}

#[test]
fn histogram_end_to_end_migrate_all_targets() {
    let dir = tempdir().unwrap();
    let out_dir = dir.path().join("out");
    let report = run(MigrateOptions {
        input: Path::new(env!("CARGO_MANIFEST_DIR")).join("examples/cu/histogram.cu"),
        output: out_dir.clone(),
        target: Target::All,
        dry_run: false,
        verbose: false,
        filter: None,
    }).expect("migration should succeed");

    for t in Target::iter_real() {
        let dir = out_dir.join(t.as_str());
        assert!(dir.is_dir(), "missing {t:?} output dir");
        assert!(fs::read_dir(&dir).unwrap().count() > 0, "{t:?} output empty");
    }
    // __launch_bounds__ is NOT in this fixture, but the report should still
    // record warnings for unmapped runtime APIs (cudaMemcpyToSymbol etc.).
    assert!(!report.files.is_empty());
}

// ===========================================================================
// Complex fixture: transpose.cu
//   2D dim3 grid/block, shared-memory tile, __syncthreads, blockIdx.{x,y}.
// ===========================================================================

#[test]
fn transpose_parser_captures_2d_indices_and_launch() {
    let u = load_fixture("transpose.cu");
    // Builtins with .x and .y field access.
    let field_accesses: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::BuiltinRef { kind, has_field_access: true, .. } => Some(*kind),
        _ => None,
    }).collect();
    assert!(field_accesses.contains(&BuiltinKind::BlockIdx));
    assert!(field_accesses.contains(&BuiltinKind::ThreadIdx));
    // __syncthreads.
    assert!(u.nodes.iter().any(|n| matches!(
        n,
        IrNode::BuiltinRef { kind: BuiltinKind::SyncThreads, .. }
    )));
    // __shared__ tile.
    assert!(u.nodes.iter().any(|n| matches!(
        n,
        IrNode::QualifierDecl { qualifier: CudaQualifier::Shared, .. }
    )));
    // One launch with a 2D dim3 grid/block.
    let launches: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::KernelLaunch { kernel, grid, block, .. } => Some((kernel.clone(), grid.clone(), block.clone())),
        _ => None,
    }).collect();
    assert_eq!(launches.len(), 1);
    assert_eq!(launches[0].0, "transpose");
    // The launch uses the `grid` and `block` variable names; decuda keeps
    // them verbatim (dim3 expressions are not inlined into the launch).
    assert_eq!(launches[0].1, "grid");
    assert_eq!(launches[0].2, "block");
}

#[test]
fn transpose_hip_preserves_2d_field_access() {
    let u = load_fixture("transpose.cu");
    let out = for_target(Target::Hip).emit(&u);
    // HIP keeps blockIdx.x/blockIdx.y and threadIdx.x/threadIdx.y.
    assert!(out.contains("blockIdx.x"));
    assert!(out.contains("blockIdx.y"));
    assert!(out.contains("threadIdx.x"));
    assert!(out.contains("threadIdx.y"));
    assert!(out.contains("blockDim.x"));
    assert!(out.contains("blockDim.y"));
    // __shared__ and __syncthreads preserved.
    assert!(out.contains("__shared__"));
    assert!(out.contains("__syncthreads()"));
    // dim3 grid/block survive.
    assert!(out.contains("dim3 grid"));
    assert!(out.contains("dim3 block"));
    // Launch rewritten.
    assert!(out.contains("hipLaunchKernelGGL"));
}

#[test]
fn transpose_opencl_drops_2d_field_access() {
    let u = load_fixture("transpose.cu");
    let out = for_target(Target::Opencl).emit(&u);
    // OpenCL replaces blockIdx.x -> get_group_id(0) and consumes the .x.
    assert!(out.contains("get_group_id(0)"));
    assert!(out.contains("get_local_id(0)"));
    assert!(out.contains("get_local_size(0)"));
    // No threadIdx.x / blockIdx.x should remain.
    assert!(!out.contains("threadIdx.x"));
    assert!(!out.contains("blockIdx.x"));
    // __shared__ -> __local, __syncthreads -> barrier.
    assert!(out.contains("__local"));
    assert!(out.contains("barrier(CLK_LOCAL_MEM_FENCE)"));
}

// ===========================================================================
// Complex fixture: stream_pipeline.cu
//   streams, events, async memcpy, launch with smem+stream args.
// ===========================================================================

#[test]
fn stream_pipeline_parser_captures_streams_events_and_launch_args() {
    let u = load_fixture("stream_pipeline.cu");
    // Runtime calls for streams and events.
    let rt_names: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::RuntimeCall { name, .. } => Some(name.clone()),
        _ => None,
    }).collect();
    for expected in [
        "cudaStreamCreate", "cudaStreamDestroy", "cudaStreamSynchronize",
        "cudaEventCreate", "cudaEventRecord", "cudaEventSynchronize", "cudaEventDestroy",
        "cudaMemcpyAsync", "cudaMalloc", "cudaFree",
    ] {
        assert!(
            rt_names.iter().any(|n| n == expected),
            "expected runtime call `{expected}` in IR, got: {rt_names:?}"
        );
    }

    // Two launches; the second one carries smem and stream args.
    let launches: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::KernelLaunch { kernel, smem, stream, .. } => Some((kernel.clone(), smem.clone(), stream.clone())),
        _ => None,
    }).collect();
    assert_eq!(launches.len(), 2);
    // The `add` launch has smem=128 and stream=s2.
    let add_launch = launches.iter().find(|(k, _, _)| k == "add").expect("add launch");
    assert_eq!(add_launch.1.as_deref(), Some("128"), "smem arg for add launch");
    assert_eq!(add_launch.2.as_deref(), Some("s2"), "stream arg for add launch");
}

#[test]
fn stream_pipeline_hip_renames_streams_events_and_async() {
    let u = load_fixture("stream_pipeline.cu");
    let out = for_target(Target::Hip).emit(&u);
    // Stream/event APIs renamed to hip* equivalents.
    assert!(out.contains("hipStreamCreate"), "missing hipStreamCreate:\n{out}");
    assert!(out.contains("hipStreamDestroy"));
    assert!(out.contains("hipStreamSynchronize"));
    assert!(out.contains("hipEventCreate"));
    assert!(out.contains("hipEventRecord"));
    assert!(out.contains("hipEventSynchronize"));
    assert!(out.contains("hipEventDestroy"));
    assert!(out.contains("hipMemcpyAsync"));
    // Launch with smem+stream -> hipLaunchKernelGGL with smem and stream slots.
    assert!(out.contains("hipLaunchKernelGGL"));
    // The add launch should carry smem=128 and stream=s2 in the rewritten form.
    assert!(out.contains("128") && out.contains("s2"));
    // No cuda* stream/event names leak outside comments.
    assert!(!contains_outside_comments(&out, "cudaStreamCreate"));
    assert!(!contains_outside_comments(&out, "cudaEventRecord"));
}

#[test]
fn stream_pipeline_opencl_keeps_unmapped_calls_as_warnings() {
    let u = load_fixture("stream_pipeline.cu");
    let out = for_target(Target::Opencl).emit(&u);
    // OpenCL has no mapping for streams/events/async memcpy; the calls are
    // left verbatim and the migration report flags them. We only assert the
    // output still parses through decuda (no panic) and the header is replaced.
    assert!(out.contains("CL/cl.h"));
    // Launches rewritten to clEnqueueNDRangeKernel.
    assert!(out.contains("clEnqueueNDRangeKernel"));
}

// ===========================================================================
// Complex fixture: device_helpers.cu
//   __device__ helpers, __forceinline__/__noinline__, __constant__, __launch_bounds__,
//   __laneid, __syncwarp, __syncthreads, device_functions.h header.
// ===========================================================================

#[test]
fn device_helpers_parser_harvests_inline_hints_and_launch_bounds() {
    let u = load_fixture("device_helpers.cu");
    // __forceinline__ and __noinline__ qualifiers.
    let quals: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::QualifierDecl { qualifier, .. } => Some(*qualifier),
        _ => None,
    }).collect();
    assert!(quals.contains(&CudaQualifier::ForceInline), "expected __forceinline__");
    assert!(quals.contains(&CudaQualifier::NoInline), "expected __noinline__");
    assert!(quals.contains(&CudaQualifier::Device), "expected __device__");
    assert!(quals.contains(&CudaQualifier::Constant), "expected __constant__");

    // Two headers: cuda_runtime.h and device_functions.h.
    let headers: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::HeaderInclude { header, .. } => Some(header.clone()),
        _ => None,
    }).collect();
    assert!(headers.iter().any(|h| h == "cuda_runtime.h"));
    assert!(headers.iter().any(|h| h == "device_functions.h"));

    // __laneid and __syncwarp builtins.
    assert!(u.nodes.iter().any(|n| matches!(
        n,
        IrNode::BuiltinRef { kind: BuiltinKind::LaneId, .. }
    )));
    assert!(u.nodes.iter().any(|n| matches!(
        n,
        IrNode::BuiltinRef { kind: BuiltinKind::SyncWarp, .. }
    )));
}

#[test]
fn device_helpers_hip_emits_inline_hints() {
    let u = load_fixture("device_helpers.cu");
    let out = for_target(Target::Hip).emit(&u);
    // HIP keeps __forceinline__ / __noinline__ / __device__.
    assert!(out.contains("__forceinline__"), "HIP keeps __forceinline__:\n{out}");
    assert!(out.contains("__noinline__"));
    assert!(out.contains("__device__"));
    // __constant__ preserved.
    assert!(out.contains("__constant__"));
    // device_functions.h -> hip/device_functions.h.
    assert!(out.contains("hip/device_functions.h"));
    // __laneid and __syncwarp preserved.
    assert!(out.contains("__laneid"));
    assert!(out.contains("__syncwarp"));
}

#[test]
fn device_helpers_rust_emits_inline_attributes() {
    let u = load_fixture("device_helpers.cu");
    let out = for_target(Target::Rust).emit(&u);
    // Rust maps __forceinline__ -> #[inline(always)] and __noinline__ -> #[inline(never)].
    assert!(out.contains("#[inline(always)]"), "Rust inline(always) missing:\n{out}");
    assert!(out.contains("#[inline(never)]"));
    // __device__ -> /* device function */.
    assert!(out.contains("/* device function */"));
    // __constant__ -> /* constant */.
    assert!(out.contains("/* constant */"));
}

#[test]
fn device_helpers_opencl_drops_inline_hints() {
    let u = load_fixture("device_helpers.cu");
    let out = for_target(Target::Opencl).emit(&u);
    // OpenCL maps __forceinline__/__noinline__ to empty string. The tokens
    // still appear in the fixture's comment header, so check outside comments.
    assert!(
        !contains_outside_comments(&out, "__forceinline__"),
        "OpenCL must drop __forceinline__ outside comments:\n{out}"
    );
    assert!(
        !contains_outside_comments(&out, "__noinline__"),
        "OpenCL must drop __noinline__ outside comments:\n{out}"
    );
    // __device__ -> __device.
    assert!(out.contains("__device"));
    // __constant__ -> __constant.
    assert!(out.contains("__constant"));
    // device_functions.h has no OpenCL mapping -> not replaced (left verbatim).
    // cuda_runtime.h -> CL/cl.h.
    assert!(out.contains("CL/cl.h"));
}

#[test]
fn device_helpers_end_to_end_migrate_records_launch_bounds_warning() {
    let dir = tempdir().unwrap();
    let out_dir = dir.path().join("out");
    let report = run(MigrateOptions {
        input: Path::new(env!("CARGO_MANIFEST_DIR")).join("examples/cu/device_helpers.cu"),
        output: out_dir,
        target: Target::Hip,
        dry_run: false,
        verbose: false,
        filter: None,
    }).expect("migration should succeed");

    // __launch_bounds__ is parsed but not in the from_token list (it's not
    // auto-translated). The fixture uses it; the migrate run must still
    // succeed and produce output. We assert the run completed and produced
    // at least one output file.
    assert!(!report.files.is_empty());
    assert!(!report.files[0].outputs.is_empty());
}

// ---------------------------------------------------------------------------
// Regenerate the persisted example outputs under examples/out/ so they stay
// in sync with the current implementation. Unlike the tempdir-based tests
// above, this one writes to a real on-disk location that survives the test
// run and can be inspected / committed.
// ---------------------------------------------------------------------------

#[test]
fn regenerate_examples_out_dir() {
    let manifest = Path::new(env!("CARGO_MANIFEST_DIR"));
    let input = manifest.join("examples/cu");
    let output = manifest.join("examples/out");

    // Start from a clean slate so stale files from a previous run (e.g. a
    // renamed fixture) don't linger.
    if output.exists() {
        fs::remove_dir_all(&output).expect("clean examples/out before regen");
    }

    let report = run(MigrateOptions {
        input,
        output: output.clone(),
        target: Target::All,
        dry_run: false,
        verbose: false,
        filter: None,
    })
    .expect("regenerate examples/out should succeed");

    // Every backend must have written its subdirectory with at least one file.
    for t in Target::iter_real() {
        let dir = output.join(t.as_str());
        assert!(dir.is_dir(), "missing {t:?} output dir: {}", dir.display());
        let count = fs::read_dir(&dir).unwrap().count();
        assert!(count > 0, "{t:?} output dir is empty: {}", dir.display());
    }

    // A single consolidated migration report must be written next to the
    // target subdirectories.
    let report_files: Vec<_> = fs::read_dir(&output)
        .unwrap()
        .filter_map(|e| e.ok())
        .filter(|e| {
            e.file_name()
                .to_string_lossy()
                .starts_with("migration-report.json.")
        })
        .collect();
    assert!(
        !report_files.is_empty(),
        "expected a migration-report.json.<ts> file under examples/out"
    );

    // Sanity: the report must cover the example CUDA inputs.
    assert!(report.files.len() >= 5, "expected >=5 files in report");

    // Visible confirmation so `cargo test` output shows what was written.
    let targets = Target::iter_real();
    let total: usize = targets
        .iter()
        .map(|t| fs::read_dir(output.join(t.as_str())).unwrap().count())
        .sum();
    eprintln!(
        "\n[regenerate_examples_out_dir] wrote {total} files to examples/out/ \
         ({} CUDA inputs x {} targets) + 1 migration report",
        report.files.len(),
        targets.len()
    );
}

// ===========================================================================
// Advanced fixture: reduction.cu
//   Warp-shuffle reduction (__shfl_sync flagged), atomicAdd, tree reduction,
//   __syncthreads, warpSize, __laneid, grid-stride loop, multiple kernels.
// ===========================================================================

#[test]
fn reduction_parser_harvests_warpsize_laneid_and_atomics() {
    let u = load_fixture("reduction.cu");
    // Known builtins: __syncthreads, __syncwarp (not used), warpSize, __laneid.
    let kinds: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::BuiltinRef { kind, .. } => Some(*kind),
        _ => None,
    }).collect();
    assert!(kinds.contains(&BuiltinKind::SyncThreads), "expected __syncthreads");
    assert!(kinds.contains(&BuiltinKind::LaneId), "expected __laneid");

    // Atomics: atomicAdd.
    let atomics: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::AtomicIntrinsic { name, .. } => Some(name.clone()),
        _ => None,
    }).collect();
    assert!(atomics.contains(&"atomicAdd".to_string()));

    // Two kernel launches.
    let launches: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::KernelLaunch { kernel, .. } => Some(kernel.clone()),
        _ => None,
    }).collect();
    assert!(launches.contains(&"reduce_block".to_string()));
    assert!(launches.contains(&"reduce_final".to_string()));

    // __device__ and __forceinline__ qualifiers.
    let quals: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::QualifierDecl { qualifier, .. } => Some(*qualifier),
        _ => None,
    }).collect();
    assert!(quals.contains(&CudaQualifier::Device));
    assert!(quals.contains(&CudaQualifier::ForceInline));
}

#[test]
fn reduction_hip_rewrites_known_builtins_keeps_shfl() {
    let u = load_fixture("reduction.cu");
    let out = for_target(Target::Hip).emit(&u);
    // HIP keeps threadIdx.x, blockIdx.x, etc.
    assert!(out.contains("threadIdx.x"));
    assert!(out.contains("blockIdx.x"));
    // HIP keeps __syncthreads, __laneid, warpSize.
    assert!(out.contains("__syncthreads()"));
    assert!(out.contains("__laneid"));
    assert!(out.contains("warpSize"));
    // HIP keeps atomicAdd.
    assert!(out.contains("atomicAdd"));
    // __shfl_sync is now a known builtin — HIP keeps it as-is (no-op name
    // swap), arguments preserved.
    assert!(out.contains("__shfl_sync("), "HIP keeps __shfl_sync:\n{out}");
    // Runtime calls renamed.
    assert!(out.contains("hipMalloc"));
    assert!(out.contains("hipFree"));
    assert!(out.contains("hipMemcpy"));
    assert!(out.contains("hipMemset"));
    // Launches rewritten.
    assert!(out.contains("hipLaunchKernelGGL"));
}

#[test]
fn reduction_opencl_rewrites_builtins_and_shfl() {
    let u = load_fixture("reduction.cu");
    let out = for_target(Target::Opencl).emit(&u);
    // OpenCL rewrites known builtins.
    assert!(out.contains("get_local_id(0)"));
    assert!(out.contains("get_group_id(0)"));
    assert!(out.contains("barrier(CLK_LOCAL_MEM_FENCE)"));
    // __shfl_sync is now rewritten to sub_group_shuffle (args preserved).
    assert!(out.contains("sub_group_shuffle("), "OpenCL __shfl_sync -> sub_group_shuffle:\n{out}");
    // Original __shfl_sync should NOT appear outside comments.
    assert!(
        !contains_outside_comments(&out, "__shfl_sync"),
        "OpenCL should rewrite __shfl_sync:\n{out}"
    );
    // OpenCL keeps atomicAdd.
    assert!(out.contains("atomicAdd"));
    // threadIdx.x consumed outside comments (it appears in the fixture's
    // comment header which is preserved verbatim).
    assert!(
        !contains_outside_comments(&out, "threadIdx.x"),
        "OpenCL must drop threadIdx.x outside comments:\n{out}"
    );
}

#[test]
fn reduction_migrate_flags_unknown_warpsize_in_report() {
    let dir = tempdir().unwrap();
    let out_dir = dir.path().join("out");
    let report = run(MigrateOptions {
        input: Path::new(env!("CARGO_MANIFEST_DIR")).join("examples/cu/reduction.cu"),
        output: out_dir,
        target: Target::Sycl,
        dry_run: false,
        verbose: false,
        filter: None,
    }).expect("migration should succeed");

    // reduction.cu uses cudaMemset which has a HIP mapping but no SYCL
    // mapping — the report should flag it for SYCL.
    let has_warning = report
        .by_file
        .values()
        .flat_map(|m| m.values())
        .flatten()
        .any(|(_, msg)| msg.contains("no automatic mapping"));
    assert!(has_warning, "expected at least one unmapped-API warning for SYCL");
}

// ===========================================================================
// Advanced fixture: stencil_3d.cu
//   3D dim3 grid/block, shared-memory halo, threadIdx.{x,y,z},
//   blockIdx.{x,y,z}, blockDim.{x,y,z}, gridDim.{x,y,z}.
// ===========================================================================

#[test]
fn stencil_3d_parser_captures_xyz_field_access() {
    let u = load_fixture("stencil_3d.cu");
    // All four built-in kinds with field access.
    let field_kinds: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::BuiltinRef { kind, has_field_access: true, .. } => Some(*kind),
        _ => None,
    }).collect();
    assert!(field_kinds.contains(&BuiltinKind::ThreadIdx), "expected threadIdx.x/y/z");
    assert!(field_kinds.contains(&BuiltinKind::BlockIdx), "expected blockIdx.x/y/z");
    assert!(field_kinds.contains(&BuiltinKind::BlockDim), "expected blockDim.x/y/z");
    // gridDim is not used in the stencil kernel body (no grid-stride loop).

    // __shared__ tile.
    assert!(u.nodes.iter().any(|n| matches!(
        n,
        IrNode::QualifierDecl { qualifier: CudaQualifier::Shared, .. }
    )));
    // __syncthreads.
    assert!(u.nodes.iter().any(|n| matches!(
        n,
        IrNode::BuiltinRef { kind: BuiltinKind::SyncThreads, .. }
    )));
    // One launch with 3D dim3.
    let launches: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::KernelLaunch { kernel, grid, block, .. } => Some((kernel.clone(), grid.clone(), block.clone())),
        _ => None,
    }).collect();
    assert_eq!(launches.len(), 1);
    assert_eq!(launches[0].0, "stencil_3d");
}

#[test]
fn stencil_3d_hip_preserves_xyz_access() {
    let u = load_fixture("stencil_3d.cu");
    let out = for_target(Target::Hip).emit(&u);
    // HIP keeps all .x, .y, .z field accesses (gridDim not used in kernel).
    for v in ["threadIdx.x", "threadIdx.y", "threadIdx.z",
              "blockIdx.x", "blockIdx.y", "blockIdx.z",
              "blockDim.x", "blockDim.y", "blockDim.z"] {
        assert!(out.contains(v), "HIP should keep {v}:\n{out}");
    }
    assert!(out.contains("__shared__"));
    assert!(out.contains("__syncthreads()"));
}

#[test]
fn stencil_3d_opencl_rewrites_all_xyz_to_scalar() {
    let u = load_fixture("stencil_3d.cu");
    let out = for_target(Target::Opencl).emit(&u);
    // OpenCL replaces all .x/.y/.z field accesses with scalar get_* calls.
    assert!(out.contains("get_local_id(0)"));
    assert!(out.contains("get_group_id(0)"));
    assert!(out.contains("get_local_size(0)"));
    // gridDim is not used in the kernel, so get_num_groups is absent.
    // No threadIdx.x / blockIdx.x should remain outside comments.
    assert!(!contains_outside_comments(&out, "threadIdx.x"));
    assert!(!contains_outside_comments(&out, "blockIdx.x"));
    // __shared__ -> __local.
    assert!(out.contains("__local"));
    // __syncthreads -> barrier.
    assert!(out.contains("barrier(CLK_LOCAL_MEM_FENCE)"));
}

// ===========================================================================
// Advanced fixture: device_management.cu
//   Multi-GPU, error handling, cudaGetDeviceCount/cudaSetDevice/cudaGetDevice,
//   cudaGetLastError/cudaGetErrorString, cudaHostAlloc, cudaDeviceSynchronize
//   (flagged), error-check macro.
// ===========================================================================

#[test]
fn device_management_parser_harvests_device_and_error_apis() {
    let u = load_fixture("device_management.cu");
    let rt_names: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::RuntimeCall { name, .. } => Some(name.clone()),
        _ => None,
    }).collect();
    // APIs in the DB.
    assert!(rt_names.contains(&"cudaGetDeviceCount".to_string()));
    assert!(rt_names.contains(&"cudaSetDevice".to_string()));
    assert!(rt_names.contains(&"cudaGetDevice".to_string()));
    assert!(rt_names.contains(&"cudaGetErrorString".to_string()));
    assert!(rt_names.contains(&"cudaGetLastError".to_string()));
    assert!(rt_names.contains(&"cudaHostAlloc".to_string()));
    assert!(rt_names.contains(&"cudaMalloc".to_string()));
    assert!(rt_names.contains(&"cudaFree".to_string()));
    assert!(rt_names.contains(&"cudaMemcpy".to_string()));
    // cudaDeviceSynchronize is also harvested (regex catches cuda*), but has
    // no mapping — it will be flagged in the report.
    assert!(rt_names.contains(&"cudaDeviceSynchronize".to_string()));
    // cudaFreeHost is harvested but not in DB.
    assert!(rt_names.contains(&"cudaFreeHost".to_string()));
}

#[test]
fn device_management_hip_renames_known_apis() {
    let u = load_fixture("device_management.cu");
    let out = for_target(Target::Hip).emit(&u);
    // Known APIs renamed.
    assert!(out.contains("hipGetDeviceCount"));
    assert!(out.contains("hipSetDevice"));
    assert!(out.contains("hipGetDevice"));
    assert!(out.contains("hipGetErrorString"));
    assert!(out.contains("hipGetLastError"));
    assert!(out.contains("hipHostMalloc"), "cudaHostAlloc -> hipHostMalloc:\n{out}");
    assert!(out.contains("hipMalloc"));
    assert!(out.contains("hipFree"));
    assert!(out.contains("hipMemcpy"));
    // cudaDeviceSynchronize -> hipDeviceSynchronize (now in DB).
    assert!(out.contains("hipDeviceSynchronize"), "cudaDeviceSynchronize -> hipDeviceSynchronize:\n{out}");
    // cudaFreeHost -> hipFreeHost (now in DB).
    assert!(out.contains("hipFreeHost"), "cudaFreeHost -> hipFreeHost:\n{out}");
    // No CUDA API names should leak into HIP output outside comments.
    assert!(
        !contains_outside_comments(&out, "cudaDeviceSynchronize"),
        "cudaDeviceSynchronize should be rewritten, not preserved:\n{out}"
    );
    assert!(
        !contains_outside_comments(&out, "cudaFreeHost"),
        "cudaFreeHost should be rewritten, not preserved:\n{out}"
    );
}

#[test]
fn device_management_migrate_flags_unmapped_apis() {
    let dir = tempdir().unwrap();
    let out_dir = dir.path().join("out");
    let report = run(MigrateOptions {
        input: Path::new(env!("CARGO_MANIFEST_DIR")).join("examples/cu/device_management.cu"),
        output: out_dir,
        target: Target::Sycl,
        dry_run: false,
        verbose: false,
        filter: None,
    }).expect("migration should succeed");

    // For SYCL, cudaDeviceSynchronize and cudaFreeHost have no mapping.
    let warnings: Vec<_> = report
        .by_file
        .values()
        .flat_map(|m| m.values())
        .flatten()
        .map(|(_, msg)| msg.clone())
        .collect();
    assert!(
        warnings.iter().any(|w| w.contains("cudaDeviceSynchronize")),
        "expected warning for cudaDeviceSynchronize on SYCL: {warnings:?}"
    );
    assert!(
        warnings.iter().any(|w| w.contains("cudaFreeHost")),
        "expected warning for cudaFreeHost on SYCL: {warnings:?}"
    );
}

// ===========================================================================
// Advanced fixture: managed_memory.cu
//   cudaMallocManaged (flagged), __managed__ qualifier, cudaHostAlloc,
//   cudaMallocHost, cudaMemcpy, cudaDeviceSynchronize (flagged).
// ===========================================================================

#[test]
fn managed_memory_parser_harvests_managed_qualifier() {
    let u = load_fixture("managed_memory.cu");
    // __managed__ qualifier.
    assert!(u.nodes.iter().any(|n| matches!(
        n,
        IrNode::QualifierDecl { qualifier: CudaQualifier::Managed, .. }
    )));
    // __device__ and __forceinline__.
    assert!(u.nodes.iter().any(|n| matches!(
        n,
        IrNode::QualifierDecl { qualifier: CudaQualifier::Device, .. }
    )));
    assert!(u.nodes.iter().any(|n| matches!(
        n,
        IrNode::QualifierDecl { qualifier: CudaQualifier::ForceInline, .. }
    )));
    // Runtime calls.
    let rt_names: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::RuntimeCall { name, .. } => Some(name.clone()),
        _ => None,
    }).collect();
    assert!(rt_names.contains(&"cudaMallocManaged".to_string()));
    assert!(rt_names.contains(&"cudaMallocHost".to_string()));
    assert!(rt_names.contains(&"cudaMemcpy".to_string()));
    assert!(rt_names.contains(&"cudaFree".to_string()));
    assert!(rt_names.contains(&"cudaDeviceSynchronize".to_string()));
}

#[test]
fn managed_memory_hip_rewrites_managed_qualifier() {
    let u = load_fixture("managed_memory.cu");
    let out = for_target(Target::Hip).emit(&u);
    // __managed__ -> __managed__ (HIP keeps it).
    assert!(out.contains("__managed__"));
    // cudaMallocHost -> hipHostMalloc.
    assert!(out.contains("hipHostMalloc"));
    // cudaMemcpy -> hipMemcpy.
    assert!(out.contains("hipMemcpy"));
    // cudaFree -> hipFree.
    assert!(out.contains("hipFree"));
    // cudaMallocManaged -> hipMallocManaged (now in DB).
    assert!(out.contains("hipMallocManaged"), "cudaMallocManaged -> hipMallocManaged:\n{out}");
    // cudaDeviceSynchronize -> hipDeviceSynchronize (now in DB).
    assert!(out.contains("hipDeviceSynchronize"), "cudaDeviceSynchronize -> hipDeviceSynchronize:\n{out}");
    // No CUDA API names should leak into HIP output outside comments.
    assert!(
        !contains_outside_comments(&out, "cudaMallocManaged"),
        "cudaMallocManaged should be rewritten, not preserved:\n{out}"
    );
    assert!(
        !contains_outside_comments(&out, "cudaDeviceSynchronize"),
        "cudaDeviceSynchronize should be rewritten, not preserved:\n{out}"
    );
}

#[test]
fn managed_memory_opencl_rewrites_managed_to_svm_comment() {
    let u = load_fixture("managed_memory.cu");
    let out = for_target(Target::Opencl).emit(&u);
    // __managed__ -> /* use SVM: clSVMAlloc */.
    assert!(out.contains("clSVMAlloc"), "OpenCL __managed__ -> SVM comment:\n{out}");
    // __forceinline__ dropped (empty string).
    assert!(!contains_outside_comments(&out, "__forceinline__"));
}

// ===========================================================================
// Advanced fixture: warp_primitives.cu
//   __shfl_sync, __ballot_sync, __any_sync, __all_sync, __activemask (all
//   flagged as unknown), __laneid, __syncwarp, warpSize, __syncthreads (all
//   rewritten), atomicAdd/atomicMin/atomicMax.
// ===========================================================================

#[test]
fn warp_primitives_parser_harvests_all_builtins() {
    let u = load_fixture("warp_primitives.cu");
    // All builtins are now harvested, including warp primitives.
    let kinds: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::BuiltinRef { kind, .. } => Some(*kind),
        _ => None,
    }).collect();
    // Pre-existing known builtins.
    assert!(kinds.contains(&BuiltinKind::LaneId));
    assert!(kinds.contains(&BuiltinKind::SyncWarp));
    assert!(kinds.contains(&BuiltinKind::SyncThreads));
    // Warp primitives now harvested as BuiltinRef.
    assert!(kinds.contains(&BuiltinKind::ShflSync), "expected __shfl_sync");
    assert!(kinds.contains(&BuiltinKind::BallotSync), "expected __ballot_sync");
    assert!(kinds.contains(&BuiltinKind::AnySync), "expected __any_sync");
    assert!(kinds.contains(&BuiltinKind::AllSync), "expected __all_sync");
    assert!(kinds.contains(&BuiltinKind::ActiveMask), "expected __activemask");

    // Atomics: atomicAdd, atomicMin, atomicMax.
    let atomics: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::AtomicIntrinsic { name, .. } => Some(name.clone()),
        _ => None,
    }).collect();
    assert!(atomics.contains(&"atomicAdd".to_string()));
    assert!(atomics.contains(&"atomicMin".to_string()));
    assert!(atomics.contains(&"atomicMax".to_string()));
}

#[test]
fn warp_primitives_hip_rewrites_all_builtins() {
    let u = load_fixture("warp_primitives.cu");
    let out = for_target(Target::Hip).emit(&u);
    // HIP: warp primitives are identical to CUDA (no-op name swap).
    // The function name is rewritten to the same name, args preserved.
    assert!(out.contains("__shfl_sync("), "HIP keeps __shfl_sync:\n{out}");
    assert!(out.contains("__ballot_sync("), "HIP keeps __ballot_sync:\n{out}");
    assert!(out.contains("__any_sync("), "HIP keeps __any_sync:\n{out}");
    assert!(out.contains("__all_sync("), "HIP keeps __all_sync:\n{out}");
    assert!(out.contains("__activemask"), "HIP keeps __activemask:\n{out}");
    // Known builtins rewritten (HIP keeps them as-is).
    assert!(out.contains("__laneid"));
    assert!(out.contains("__syncwarp"));
    assert!(out.contains("__syncthreads()"));
    // Atomics preserved.
    assert!(out.contains("atomicAdd"));
    assert!(out.contains("atomicMin"));
    assert!(out.contains("atomicMax"));
}

#[test]
fn warp_primitives_opencl_rewrites_all_builtins() {
    let u = load_fixture("warp_primitives.cu");
    let out = for_target(Target::Opencl).emit(&u);
    // Known builtins rewritten to OpenCL equivalents.
    assert!(out.contains("get_sub_group_id()"));
    assert!(out.contains("barrier(CLK_LOCAL_MEM_FENCE)"));
    // Warp primitives rewritten to OpenCL sub_group equivalents.
    // Arguments are preserved as-is (may need manual fixup for arg order).
    assert!(out.contains("sub_group_shuffle("), "OpenCL __shfl_sync -> sub_group_shuffle:\n{out}");
    assert!(out.contains("sub_group_ballot("), "OpenCL __ballot_sync -> sub_group_ballot:\n{out}");
    assert!(out.contains("sub_group_any("), "OpenCL __any_sync -> sub_group_any:\n{out}");
    assert!(out.contains("sub_group_all("), "OpenCL __all_sync -> sub_group_all:\n{out}");
    assert!(out.contains("get_sub_group_size"), "OpenCL __activemask -> get_sub_group_size:\n{out}");
    // Original CUDA warp primitive names should NOT appear outside comments.
    assert!(
        !contains_outside_comments(&out, "__shfl_sync"),
        "OpenCL should rewrite __shfl_sync:\n{out}"
    );
    assert!(
        !contains_outside_comments(&out, "__ballot_sync"),
        "OpenCL should rewrite __ballot_sync:\n{out}"
    );
    // Atomics preserved.
    assert!(out.contains("atomicAdd"));
    assert!(out.contains("atomicMin"));
    assert!(out.contains("atomicMax"));
}

#[test]
fn warp_primitives_end_to_end_all_targets() {
    let dir = tempdir().unwrap();
    let out_dir = dir.path().join("out");
    let report = run(MigrateOptions {
        input: Path::new(env!("CARGO_MANIFEST_DIR")).join("examples/cu/warp_primitives.cu"),
        output: out_dir.clone(),
        target: Target::All,
        dry_run: false,
        verbose: false,
        filter: None,
    }).expect("migration should succeed");

    for t in Target::iter_real() {
        let dir = out_dir.join(t.as_str());
        assert!(dir.is_dir(), "missing {t:?} output dir");
        assert!(fs::read_dir(&dir).unwrap().count() > 0, "{t:?} output empty");
    }
    // cudaDeviceSynchronize is mapped for HIP but not for SYCL/Rust/OpenCL.
    // Since Target::All runs all four, the report should still contain
    // warnings from the non-HIP targets.
    let has_unmapped = report
        .by_file
        .values()
        .flat_map(|m| m.values())
        .flatten()
        .any(|(_, msg)| msg.contains("cudaDeviceSynchronize"));
    assert!(has_unmapped, "expected cudaDeviceSynchronize warning");
}

// ===========================================================================
// Advanced fixture: ptx_inline.cu
//   Inline PTX assembly (asm/asm volatile) — flagged as warnings, preserved
//   verbatim in all target outputs.
// ===========================================================================

#[test]
fn ptx_inline_parser_harvests_asm_warnings() {
    let u = load_fixture("ptx_inline.cu");
    // Four asm constructs in the fixture.
    let ptx_warnings: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::Warning { message, .. } if message.contains("PTX") => Some(message.clone()),
        _ => None,
    }).collect();
    assert!(ptx_warnings.len() >= 4, "expected >=4 PTX warnings, got {}: {ptx_warnings:?}", ptx_warnings.len());

    // The fixture also has standard CUDA constructs.
    let kinds: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::BuiltinRef { kind, .. } => Some(*kind),
        _ => None,
    }).collect();
    assert!(kinds.contains(&BuiltinKind::ThreadIdx));
    assert!(kinds.contains(&BuiltinKind::BlockIdx));
    assert!(kinds.contains(&BuiltinKind::BlockDim));

    // Four kernel launches.
    let launches: Vec<_> = u.nodes.iter().filter_map(|n| match n {
        IrNode::KernelLaunch { kernel, .. } => Some(kernel.clone()),
        _ => None,
    }).collect();
    assert!(launches.contains(&"ptx_bswap".to_string()));
    assert!(launches.contains(&"ptx_membar".to_string()));
    assert!(launches.contains(&"ptx_clock".to_string()));
    assert!(launches.contains(&"ptx_lanemask".to_string()));
}

#[test]
fn ptx_inline_hip_preserves_asm_verbatim() {
    let u = load_fixture("ptx_inline.cu");
    let out = for_target(Target::Hip).emit(&u);
    // PTX asm is preserved verbatim (HIP can actually use some PTX, but
    // decuda flags it for manual review regardless).
    assert!(out.contains("asm("), "HIP preserves asm() verbatim:\n{out}");
    assert!(out.contains("asm volatile("), "HIP preserves asm volatile() verbatim:\n{out}");
    assert!(out.contains("prmt.b32"), "HIP preserves PTX instructions:\n{out}");
    assert!(out.contains("membar.gl"), "HIP preserves PTX membar:\n{out}");
    // Standard CUDA constructs are rewritten for HIP.
    assert!(out.contains("hipMalloc"));
    assert!(out.contains("hipFree"));
    assert!(out.contains("hipDeviceSynchronize"));
    assert!(out.contains("hipLaunchKernelGGL"));
}

#[test]
fn ptx_inline_opencl_preserves_asm_verbatim() {
    let u = load_fixture("ptx_inline.cu");
    let out = for_target(Target::Opencl).emit(&u);
    // PTX asm is preserved verbatim — OpenCL has no inline asm.
    assert!(out.contains("asm("), "OpenCL preserves asm() verbatim:\n{out}");
    assert!(out.contains("asm volatile("), "OpenCL preserves asm volatile() verbatim:\n{out}");
    // Standard CUDA constructs are rewritten for OpenCL.
    assert!(out.contains("get_local_id(0)"));
    assert!(out.contains("get_group_id(0)"));
}

#[test]
fn ptx_inline_migrate_flags_ptx_in_report() {
    let dir = tempdir().unwrap();
    let out_dir = dir.path().join("out");
    let report = run(MigrateOptions {
        input: Path::new(env!("CARGO_MANIFEST_DIR")).join("examples/cu/ptx_inline.cu"),
        output: out_dir,
        target: Target::Hip,
        dry_run: false,
        verbose: false,
        filter: None,
    }).expect("migration should succeed");

    // The report must contain PTX warnings.
    let ptx_warnings: Vec<_> = report
        .by_file
        .values()
        .flat_map(|m| m.values())
        .flatten()
        .filter(|(_, msg)| msg.contains("PTX"))
        .map(|(_, msg)| msg.clone())
        .collect();
    assert!(ptx_warnings.len() >= 4, "expected >=4 PTX warnings in report, got {}", ptx_warnings.len());
}