//! MATLAB-compatible `filter2` builtin implementing 2-D correlation and convolution.
//!
//! The implementation builds on the shared `imfilter` infrastructure so that host
//! and GPU behaviour stay perfectly aligned. By default the builtin performs a
//! correlation and returns an output the same size as the input image, matching
//! MathWorks MATLAB semantics.

use runmat_accelerate_api::{
    GpuTensorHandle, HostTensorView, ImfilterMode, ImfilterOptions, ImfilterShape,
};
use runmat_builtins::{
    BuiltinCompletionPolicy, BuiltinDescriptor, BuiltinErrorDescriptor, BuiltinOutputMode,
    BuiltinParamArity, BuiltinParamDescriptor, BuiltinParamType, BuiltinSignatureDescriptor,
    Tensor, Value,
};
use runmat_macros::runtime_builtin;

use super::imfilter::apply_imfilter_tensor;
use crate::builtins::common::spec::{
    BroadcastSemantics, BuiltinFusionSpec, BuiltinGpuSpec, ConstantStrategy, GpuOpKind,
    ProviderHook, ReductionNaN, ResidencyPolicy, ScalarType, ShapeRequirements,
};
use crate::builtins::common::{gpu_helpers, tensor};
use crate::builtins::image::filters::type_resolvers::filter2_type;
use crate::{build_runtime_error, BuiltinResult, RuntimeError};

#[runmat_macros::register_gpu_spec(builtin_path = "crate::builtins::image::filters::filter2")]
pub const GPU_SPEC: BuiltinGpuSpec = BuiltinGpuSpec {
    name: "filter2",
    op_kind: GpuOpKind::Custom("filter2"),
    supported_precisions: &[ScalarType::F32, ScalarType::F64],
    broadcast: BroadcastSemantics::None,
    provider_hooks: &[ProviderHook::Custom("imfilter")],
    constant_strategy: ConstantStrategy::InlineLiteral,
    residency: ResidencyPolicy::NewHandle,
    nan_mode: ReductionNaN::Include,
    two_pass_threshold: None,
    workgroup_size: None,
    accepts_nan_mode: false,
    notes: "Delegates to the provider `imfilter` hook with zero padding and correlation/convolution options.",
};

#[runmat_macros::register_fusion_spec(builtin_path = "crate::builtins::image::filters::filter2")]
pub const FUSION_SPEC: BuiltinFusionSpec = BuiltinFusionSpec {
    name: "filter2",
    shape: ShapeRequirements::Any,
    constant_strategy: ConstantStrategy::InlineLiteral,
    elementwise: None,
    reduction: None,
    emits_nan: false,
    notes: "Not a fusion candidate; executes via the dedicated filtering pipeline.",
};

const FILTER2_BUILTIN: &str = "filter2";

const FILTER2_OUTPUT: [BuiltinParamDescriptor; 1] = [BuiltinParamDescriptor {
    name: "B",
    ty: BuiltinParamType::NumericArray,
    arity: BuiltinParamArity::Required,
    default: None,
    description: "Filtered output image/array.",
}];

const FILTER2_INPUTS_CORE: [BuiltinParamDescriptor; 2] = [
    BuiltinParamDescriptor {
        name: "h",
        ty: BuiltinParamType::Any,
        arity: BuiltinParamArity::Required,
        default: None,
        description: "Filter kernel.",
    },
    BuiltinParamDescriptor {
        name: "X",
        ty: BuiltinParamType::Any,
        arity: BuiltinParamArity::Required,
        default: None,
        description: "Input image/array.",
    },
];

const FILTER2_INPUTS_WITH_OPTIONS: [BuiltinParamDescriptor; 3] = [
    BuiltinParamDescriptor {
        name: "h",
        ty: BuiltinParamType::Any,
        arity: BuiltinParamArity::Required,
        default: None,
        description: "Filter kernel.",
    },
    BuiltinParamDescriptor {
        name: "X",
        ty: BuiltinParamType::Any,
        arity: BuiltinParamArity::Required,
        default: None,
        description: "Input image/array.",
    },
    BuiltinParamDescriptor {
        name: "options",
        ty: BuiltinParamType::Any,
        arity: BuiltinParamArity::Variadic,
        default: None,
        description: "Optional filter2 flags: 'same'|'full'|'valid'|'conv'|'corr'.",
    },
];

const FILTER2_SIGNATURES: [BuiltinSignatureDescriptor; 2] = [
    BuiltinSignatureDescriptor {
        label: "B = filter2(h, X)",
        inputs: &FILTER2_INPUTS_CORE,
        outputs: &FILTER2_OUTPUT,
    },
    BuiltinSignatureDescriptor {
        label: "B = filter2(h, X, options...)",
        inputs: &FILTER2_INPUTS_WITH_OPTIONS,
        outputs: &FILTER2_OUTPUT,
    },
];

const FILTER2_ERROR_INVALID_INPUT: BuiltinErrorDescriptor = BuiltinErrorDescriptor {
    code: "RM.FILTER2.INVALID_INPUT",
    identifier: Some("RunMat:filter2:InvalidInput"),
    when: "Input/kernel tensors are invalid or cannot be converted.",
    message: "filter2: invalid input",
};

const FILTER2_ERROR_INVALID_OPTION: BuiltinErrorDescriptor = BuiltinErrorDescriptor {
    code: "RM.FILTER2.INVALID_OPTION",
    identifier: Some("RunMat:filter2:InvalidOption"),
    when: "One or more option flags are invalid.",
    message: "filter2: invalid option",
};

const FILTER2_ERROR_INTERNAL: BuiltinErrorDescriptor = BuiltinErrorDescriptor {
    code: "RM.FILTER2.INTERNAL",
    identifier: Some("RunMat:filter2:Internal"),
    when: "Internal filtering operation fails.",
    message: "filter2: internal operation failed",
};

const FILTER2_ERRORS: [BuiltinErrorDescriptor; 3] = [
    FILTER2_ERROR_INVALID_INPUT,
    FILTER2_ERROR_INVALID_OPTION,
    FILTER2_ERROR_INTERNAL,
];

pub const FILTER2_DESCRIPTOR: BuiltinDescriptor = BuiltinDescriptor {
    signatures: &FILTER2_SIGNATURES,
    output_mode: BuiltinOutputMode::Fixed,
    completion_policy: BuiltinCompletionPolicy::Public,
    errors: &FILTER2_ERRORS,
};

fn filter2_descriptor_error(
    error: &'static BuiltinErrorDescriptor,
    detail: Option<&str>,
) -> RuntimeError {
    let message = match detail {
        Some(detail) => format!("{}: {}", error.message, detail),
        None => error.message.to_string(),
    };
    let mut builder = build_runtime_error(message).with_builtin(FILTER2_BUILTIN);
    if let Some(identifier) = error.identifier {
        builder = builder.with_identifier(identifier);
    }
    builder.build()
}

fn filter2_error_with_detail(
    error: &'static BuiltinErrorDescriptor,
    detail: impl AsRef<str>,
) -> RuntimeError {
    filter2_descriptor_error(error, Some(detail.as_ref()))
}

fn filter2_map_error(err: RuntimeError, fallback: &'static BuiltinErrorDescriptor) -> RuntimeError {
    if err.identifier().is_some() {
        err
    } else {
        filter2_error_with_detail(fallback, err.message())
    }
}

#[runtime_builtin(
    name = "filter2",
    category = "image/filters",
    summary = "Apply 2-D correlation/convolution filters.",
    keywords = "filter2,correlation,convolution,image filtering,gpu",
    accel = "custom-imfilter",
    type_resolver(filter2_type),
    descriptor(crate::builtins::image::filters::filter2::FILTER2_DESCRIPTOR),
    builtin_path = "crate::builtins::image::filters::filter2"
)]
async fn filter2_builtin(
    kernel: Value,
    image: Value,
    rest: Vec<Value>,
) -> crate::BuiltinResult<Value> {
    let options = parse_filter2_options(&rest)?;
    match (kernel, image) {
        (Value::GpuTensor(kernel_handle), Value::GpuTensor(image_handle)) => {
            filter2_gpu(Value::GpuTensor(kernel_handle), image_handle, &options).await
        }
        (Value::GpuTensor(kernel_handle), image_value) => {
            let kernel_tensor = gpu_helpers::gather_tensor_async(&kernel_handle)
                .await
                .map_err(|err| filter2_map_error(err, &FILTER2_ERROR_INVALID_INPUT))?;
            filter2_host(Value::Tensor(kernel_tensor), image_value, &options)
        }
        (kernel_value, Value::GpuTensor(image_handle)) => {
            filter2_gpu(kernel_value, image_handle, &options).await
        }
        (kernel_value, image_value) => filter2_host(kernel_value, image_value, &options),
    }
}

fn filter2_host(
    kernel_value: Value,
    image_value: Value,
    options: &ImfilterOptions,
) -> BuiltinResult<Value> {
    let kernel_tensor = tensor::value_into_tensor_for(FILTER2_BUILTIN, kernel_value)
        .map_err(|err| filter2_error_with_detail(&FILTER2_ERROR_INVALID_INPUT, err))?;
    let image_tensor = tensor::value_into_tensor_for(FILTER2_BUILTIN, image_value)
        .map_err(|err| filter2_error_with_detail(&FILTER2_ERROR_INVALID_INPUT, err))?;
    let result = apply_imfilter_tensor(&image_tensor, &kernel_tensor, options, FILTER2_BUILTIN)
        .map_err(|err| filter2_map_error(err, &FILTER2_ERROR_INVALID_INPUT))?;
    Ok(tensor::tensor_into_value(result))
}

async fn filter2_gpu(
    kernel_value: Value,
    image_handle: GpuTensorHandle,
    options: &ImfilterOptions,
) -> BuiltinResult<Value> {
    let kernel_clone = kernel_value.clone();
    #[cfg(all(test, feature = "wgpu"))]
    {
        let kernel_is_wgpu = matches!(kernel_clone, Value::GpuTensor(ref h) if h.device_id != 0);
        if kernel_is_wgpu || image_handle.device_id != 0 {
            let _ = runmat_accelerate::backend::wgpu::provider::register_wgpu_provider(
                runmat_accelerate::backend::wgpu::provider::WgpuProviderOptions::default(),
            );
        }
    }
    let provider = match runmat_accelerate_api::provider() {
        Some(p) => p,
        None => {
            let image_tensor = gpu_helpers::gather_tensor_async(&image_handle)
                .await
                .map_err(|err| filter2_map_error(err, &FILTER2_ERROR_INVALID_INPUT))?;
            let kernel_tensor = tensor::value_into_tensor_for(FILTER2_BUILTIN, kernel_clone)
                .map_err(|err| filter2_error_with_detail(&FILTER2_ERROR_INVALID_INPUT, err))?;
            let result =
                apply_imfilter_tensor(&image_tensor, &kernel_tensor, options, FILTER2_BUILTIN)
                    .map_err(|err| filter2_map_error(err, &FILTER2_ERROR_INVALID_INPUT))?;
            return Ok(tensor::tensor_into_value(result));
        }
    };

    let mut uploaded_kernel: Option<GpuTensorHandle> = None;
    let mut kernel_tensor_cache: Option<Tensor> = None;

    let kernel_handle = match kernel_value {
        Value::GpuTensor(handle) => handle,
        other => {
            let tensor = tensor::value_into_tensor_for(FILTER2_BUILTIN, other)
                .map_err(|err| filter2_error_with_detail(&FILTER2_ERROR_INVALID_INPUT, err))?;
            let view = HostTensorView {
                data: &tensor.data,
                shape: &tensor.shape,
            };
            match provider.upload(&view) {
                Ok(uploaded) => {
                    kernel_tensor_cache = Some(tensor);
                    uploaded_kernel = Some(uploaded.clone());
                    uploaded
                }
                Err(_) => {
                    let image_tensor = gpu_helpers::gather_tensor_async(&image_handle)
                        .await
                        .map_err(|err| filter2_map_error(err, &FILTER2_ERROR_INVALID_INPUT))?;
                    let result =
                        apply_imfilter_tensor(&image_tensor, &tensor, options, FILTER2_BUILTIN)
                            .map_err(|err| filter2_map_error(err, &FILTER2_ERROR_INVALID_INPUT))?;
                    return Ok(tensor::tensor_into_value(result));
                }
            }
        }
    };

    let kernel_handle_clone = kernel_handle.clone();

    match provider
        .imfilter(&image_handle, &kernel_handle, options)
        .await
    {
        Ok(output) => {
            if let Some(uploaded) = uploaded_kernel {
                let _ = provider.free(&uploaded);
            }
            Ok(Value::GpuTensor(output))
        }
        Err(_) => {
            if let Some(uploaded) = uploaded_kernel {
                let _ = provider.free(&uploaded);
            }
            let image_tensor = gpu_helpers::gather_tensor_async(&image_handle)
                .await
                .map_err(|err| filter2_map_error(err, &FILTER2_ERROR_INVALID_INPUT))?;
            let kernel_tensor = if let Some(ref tensor) = kernel_tensor_cache {
                tensor.clone()
            } else {
                gpu_helpers::gather_tensor_async(&kernel_handle_clone)
                    .await
                    .map_err(|err| filter2_map_error(err, &FILTER2_ERROR_INVALID_INPUT))?
            };
            let result =
                apply_imfilter_tensor(&image_tensor, &kernel_tensor, options, FILTER2_BUILTIN)
                    .map_err(|err| filter2_map_error(err, &FILTER2_ERROR_INVALID_INPUT))?;
            Ok(tensor::tensor_into_value(result))
        }
    }
}

fn parse_filter2_options(args: &[Value]) -> BuiltinResult<ImfilterOptions> {
    let mut options = ImfilterOptions::default();
    for value in args {
        let Some(text) = tensor::value_to_string(value) else {
            return Err(filter2_error_with_detail(
                &FILTER2_ERROR_INVALID_OPTION,
                "expected string option ('same', 'full', 'valid', 'conv', 'corr')",
            ));
        };
        let lowered = text.trim().to_ascii_lowercase();
        match lowered.as_str() {
            "same" => options.shape = ImfilterShape::Same,
            "full" => options.shape = ImfilterShape::Full,
            "valid" => options.shape = ImfilterShape::Valid,
            "conv" => options.mode = ImfilterMode::Convolution,
            "corr" => options.mode = ImfilterMode::Correlation,
            other => {
                return Err(filter2_error_with_detail(
                    &FILTER2_ERROR_INVALID_OPTION,
                    format!(
                        "unknown option '{}' (supported: 'same', 'full', 'valid', 'conv', 'corr')",
                        other
                    ),
                ))
            }
        }
    }
    Ok(options)
}

#[cfg(test)]
pub(crate) mod tests {
    use super::*;
    use crate::builtins::common::test_support;
    use futures::executor::block_on;
    use runmat_accelerate_api::{HostTensorView, ImfilterMode, ImfilterOptions, ImfilterShape};
    use runmat_builtins::LogicalArray;

    fn tensor(data: Vec<f64>, rows: usize, cols: usize) -> Tensor {
        Tensor::new(data, vec![rows, cols]).expect("tensor construction")
    }

    fn error_message(err: &crate::RuntimeError) -> String {
        err.message().to_string()
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn default_same_matches_imfilter_reference() {
        let kernel = tensor(vec![1.0, 2.0, 3.0, 4.0], 2, 2);
        let image = tensor(vec![4.0, 1.0, 2.0, 0.0], 2, 2);
        let value = block_on(filter2_builtin(
            Value::Tensor(kernel.clone()),
            Value::Tensor(image.clone()),
            Vec::new(),
        ))
        .expect("filter2");
        let gathered = test_support::gather(value).expect("gather");
        let options = ImfilterOptions {
            shape: ImfilterShape::Same,
            mode: ImfilterMode::Correlation,
            ..Default::default()
        };
        let expected =
            apply_imfilter_tensor(&image, &kernel, &options, FILTER2_BUILTIN).expect("reference");
        assert_eq!(gathered.shape, expected.shape);
        assert_eq!(gathered.data, expected.data);
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn full_shape_option_expands_output() {
        let kernel = tensor(vec![1.0, 2.0, 3.0, 4.0], 2, 2);
        let image = tensor(vec![1.0, 2.0, 3.0, 4.0], 2, 2);
        let value = block_on(filter2_builtin(
            Value::Tensor(kernel.clone()),
            Value::Tensor(image.clone()),
            vec![Value::from("full")],
        ))
        .expect("filter2");
        let gathered = test_support::gather(value).expect("gather");
        let options = ImfilterOptions {
            shape: ImfilterShape::Full,
            ..Default::default()
        };
        let expected =
            apply_imfilter_tensor(&image, &kernel, &options, FILTER2_BUILTIN).expect("reference");
        assert_eq!(gathered.shape, expected.shape);
        assert_eq!(gathered.data, expected.data);
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn valid_shape_option_matches_reference() {
        let kernel = tensor(vec![1.0, 0.0, 0.0, -1.0], 2, 2);
        let image = tensor(vec![1.0, 2.0, 3.0, 4.0], 2, 2);
        let value = block_on(filter2_builtin(
            Value::Tensor(kernel.clone()),
            Value::Tensor(image.clone()),
            vec![Value::from("valid")],
        ))
        .expect("filter2");
        let gathered = test_support::gather(value).expect("gather");
        let options = ImfilterOptions {
            shape: ImfilterShape::Valid,
            ..Default::default()
        };
        let expected =
            apply_imfilter_tensor(&image, &kernel, &options, FILTER2_BUILTIN).expect("reference");
        assert_eq!(gathered.shape, expected.shape);
        assert_eq!(gathered.data, expected.data);
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn convolution_mode_rotates_kernel() {
        let kernel = tensor(vec![1.0, 2.0, 3.0, 4.0], 2, 2);
        let image = tensor(vec![1.0, 2.0, 3.0, 4.0], 2, 2);
        let value = block_on(filter2_builtin(
            Value::Tensor(kernel.clone()),
            Value::Tensor(image.clone()),
            vec![Value::from("conv")],
        ))
        .expect("filter2");
        let gathered = test_support::gather(value).expect("gather");
        let options = ImfilterOptions {
            mode: ImfilterMode::Convolution,
            ..Default::default()
        };
        let expected =
            apply_imfilter_tensor(&image, &kernel, &options, FILTER2_BUILTIN).expect("reference");
        assert_eq!(gathered.shape, expected.shape);
        assert_eq!(gathered.data, expected.data);
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn explicit_corr_option_matches_default() {
        let kernel = tensor(vec![0.0, 1.0, 1.0, 0.0], 2, 2);
        let image = tensor(vec![1.0, 0.0, 0.0, 1.0], 2, 2);

        let default_value = block_on(filter2_builtin(
            Value::Tensor(kernel.clone()),
            Value::Tensor(image.clone()),
            Vec::new(),
        ))
        .expect("filter2 default");
        let corr_value = block_on(filter2_builtin(
            Value::Tensor(kernel.clone()),
            Value::Tensor(image.clone()),
            vec![Value::from("corr")],
        ))
        .expect("filter2 corr");

        let default_tensor = test_support::gather(default_value).expect("gather default");
        let corr_tensor = test_support::gather(corr_value).expect("gather corr");
        assert_eq!(default_tensor.shape, corr_tensor.shape);
        assert_eq!(default_tensor.data, corr_tensor.data);
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn conv_and_shape_order_independent() {
        let kernel = tensor(vec![1.0, -1.0, 2.0, -2.0], 2, 2);
        let image = tensor(vec![1.0, 2.0, 3.0, 4.0], 2, 2);

        let conv_first = block_on(filter2_builtin(
            Value::Tensor(kernel.clone()),
            Value::Tensor(image.clone()),
            vec![Value::from("conv"), Value::from("full")],
        ))
        .expect("filter2 conv,full");
        let shape_first = block_on(filter2_builtin(
            Value::Tensor(kernel.clone()),
            Value::Tensor(image.clone()),
            vec![Value::from("full"), Value::from("conv")],
        ))
        .expect("filter2 full,conv");

        let conv_first_tensor = test_support::gather(conv_first).expect("gather conv first");
        let shape_first_tensor = test_support::gather(shape_first).expect("gather shape first");
        assert_eq!(conv_first_tensor.shape, shape_first_tensor.shape);
        assert_eq!(conv_first_tensor.data, shape_first_tensor.data);
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn non_string_option_is_rejected() {
        let kernel = tensor(vec![1.0], 1, 1);
        let image = tensor(vec![1.0], 1, 1);
        let err = block_on(filter2_builtin(
            Value::Tensor(kernel),
            Value::Tensor(image),
            vec![Value::Tensor(tensor(vec![2.0], 1, 1))],
        ))
        .expect_err("expected option parsing error");
        let message = error_message(&err);
        assert!(
            message.contains("expected string option"),
            "unexpected error message: {message}"
        );
        assert_eq!(err.identifier(), FILTER2_ERROR_INVALID_OPTION.identifier);
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn logical_input_promoted_to_double() {
        let kernel = tensor(vec![1.0, 1.0, 1.0, 1.0], 2, 2);
        let logical = LogicalArray::new(vec![1, 0, 1, 0], vec![2, 2]).expect("logical array");

        let value = block_on(filter2_builtin(
            Value::Tensor(kernel.clone()),
            Value::LogicalArray(logical.clone()),
            Vec::new(),
        ))
        .expect("filter2 logical");

        let gathered = test_support::gather(value).expect("gather logical result");
        let image_tensor = Tensor::new(
            logical
                .data
                .iter()
                .map(|&b| if b != 0 { 1.0 } else { 0.0 })
                .collect(),
            logical.shape.clone(),
        )
        .expect("logical->tensor");
        let expected = apply_imfilter_tensor(
            &image_tensor,
            &kernel,
            &ImfilterOptions::default(),
            FILTER2_BUILTIN,
        )
        .expect("expected logical reference");

        assert_eq!(gathered.shape, expected.shape);
        assert_eq!(gathered.data, expected.data);
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn rejects_unknown_option() {
        let kernel = tensor(vec![1.0], 1, 1);
        let image = tensor(vec![1.0], 1, 1);
        let err = block_on(filter2_builtin(
            Value::Tensor(kernel),
            Value::Tensor(image),
            vec![Value::from("replicate")],
        ))
        .expect_err("filter2 should error");
        let message = error_message(&err);
        assert!(
            message.contains("filter2"),
            "expected error mentioning builtin, got {message}"
        );
        assert_eq!(err.identifier(), FILTER2_ERROR_INVALID_OPTION.identifier);
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn filter2_gpu_matches_cpu() {
        test_support::with_test_provider(|provider| {
            let kernel = tensor(vec![1.0, 0.0, -1.0, 2.0], 2, 2);
            let image = tensor(vec![1.0, 2.0, 3.0, 4.0], 2, 2);
            let image_view = HostTensorView {
                data: &image.data,
                shape: &image.shape,
            };
            let image_handle = provider.upload(&image_view).expect("upload image");
            let result = block_on(filter2_builtin(
                Value::Tensor(kernel.clone()),
                Value::GpuTensor(image_handle.clone()),
                Vec::new(),
            ))
            .expect("filter2 gpu");
            let gathered = test_support::gather(result).expect("gather");
            let expected = apply_imfilter_tensor(
                &image,
                &kernel,
                &ImfilterOptions::default(),
                FILTER2_BUILTIN,
            )
            .expect("ref");
            assert_eq!(gathered.shape, expected.shape);
            assert_eq!(gathered.data, expected.data);
        });
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn filter2_gpu_kernel_and_image() {
        test_support::with_test_provider(|provider| {
            let kernel = tensor(vec![0.0, 1.0, -1.0, 0.0], 2, 2);
            let image = tensor(vec![1.0, 2.0, 3.0, 4.0], 2, 2);
            let kernel_view = HostTensorView {
                data: &kernel.data,
                shape: &kernel.shape,
            };
            let image_view = HostTensorView {
                data: &image.data,
                shape: &image.shape,
            };
            let kernel_handle = provider.upload(&kernel_view).expect("upload kernel");
            let image_handle = provider.upload(&image_view).expect("upload image");
            let result = block_on(filter2_builtin(
                Value::GpuTensor(kernel_handle.clone()),
                Value::GpuTensor(image_handle.clone()),
                Vec::new(),
            ))
            .expect("filter2 gpu");
            let gathered = test_support::gather(result).expect("gather");
            let expected = apply_imfilter_tensor(
                &image,
                &kernel,
                &ImfilterOptions::default(),
                FILTER2_BUILTIN,
            )
            .expect("reference");
            assert_eq!(gathered.shape, expected.shape);
            assert_eq!(gathered.data, expected.data);
        });
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    #[cfg(feature = "wgpu")]
    fn filter2_wgpu_same_matches_cpu() {
        std::env::set_var("RUNMAT_WGPU_SKIP_WARMUP", "1");
        std::env::set_var("RUNMAT_WGPU_DISABLE_IMFILTER", "1");
        let _ = runmat_accelerate::backend::wgpu::provider::register_wgpu_provider(
            runmat_accelerate::backend::wgpu::provider::WgpuProviderOptions::default(),
        );
        let provider = runmat_accelerate_api::provider().expect("wgpu provider registered");

        let kernel = tensor(vec![0.0, 1.0, 0.0, 1.0, -4.0, 1.0, 0.0, 1.0, 0.0], 3, 3);
        let image = tensor(
            vec![
                1.0, 2.0, 3.0, 4.0, //
                5.0, 6.0, 7.0, 8.0, //
                9.0, 10.0, 11.0, 12.0, //
                13.0, 14.0, 15.0, 16.0,
            ],
            4,
            4,
        );

        let kernel_view = HostTensorView {
            data: &kernel.data,
            shape: &kernel.shape,
        };
        let image_view = HostTensorView {
            data: &image.data,
            shape: &image.shape,
        };

        let kernel_handle = provider.upload(&kernel_view).expect("upload kernel");
        let image_handle = provider.upload(&image_view).expect("upload image");

        let gpu_value = block_on(filter2_builtin(
            Value::GpuTensor(kernel_handle.clone()),
            Value::GpuTensor(image_handle.clone()),
            Vec::new(),
        ))
        .expect("filter2 wgpu same");
        let gpu_tensor = test_support::gather(gpu_value).expect("gather gpu result");

        let expected = apply_imfilter_tensor(
            &image,
            &kernel,
            &ImfilterOptions::default(),
            FILTER2_BUILTIN,
        )
        .expect("expected");

        assert_eq!(gpu_tensor.shape, expected.shape);
        assert_eq!(gpu_tensor.data.len(), expected.data.len());
        let tol = match runmat_accelerate_api::provider().unwrap().precision() {
            runmat_accelerate_api::ProviderPrecision::F64 => 1e-12,
            runmat_accelerate_api::ProviderPrecision::F32 => 1e-5,
        };
        for (gpu, cpu) in gpu_tensor.data.iter().zip(expected.data.iter()) {
            assert!((gpu - cpu).abs() < tol, "{gpu} vs {cpu}");
        }

        let _ = provider.free(&kernel_handle);
        let _ = provider.free(&image_handle);
        std::env::remove_var("RUNMAT_WGPU_SKIP_WARMUP");
        std::env::remove_var("RUNMAT_WGPU_DISABLE_IMFILTER");
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    #[cfg(feature = "wgpu")]
    fn filter2_wgpu_full_conv_matches_cpu() {
        std::env::set_var("RUNMAT_WGPU_SKIP_WARMUP", "1");
        std::env::set_var("RUNMAT_WGPU_DISABLE_IMFILTER", "1");
        let _ = runmat_accelerate::backend::wgpu::provider::register_wgpu_provider(
            runmat_accelerate::backend::wgpu::provider::WgpuProviderOptions::default(),
        );
        let provider = runmat_accelerate_api::provider().expect("wgpu provider registered");

        let kernel = tensor(vec![1.0, 2.0, 3.0, 0.0], 2, 2);
        let image = tensor(
            vec![
                3.0, 1.0, 4.0, //
                1.0, 5.0, 9.0, //
                2.0, 6.0, 5.0,
            ],
            3,
            3,
        );

        let kernel_view = HostTensorView {
            data: &kernel.data,
            shape: &kernel.shape,
        };
        let image_view = HostTensorView {
            data: &image.data,
            shape: &image.shape,
        };
        let kernel_handle = provider.upload(&kernel_view).expect("upload kernel");
        let image_handle = provider.upload(&image_view).expect("upload image");

        let gpu_value = block_on(filter2_builtin(
            Value::GpuTensor(kernel_handle.clone()),
            Value::GpuTensor(image_handle.clone()),
            vec![Value::from("full"), Value::from("conv")],
        ))
        .expect("filter2 wgpu full conv");
        let gpu_tensor = test_support::gather(gpu_value).expect("gather gpu result");

        let options = ImfilterOptions {
            shape: ImfilterShape::Full,
            mode: ImfilterMode::Convolution,
            ..Default::default()
        };
        let expected = apply_imfilter_tensor(&image, &kernel, &options, FILTER2_BUILTIN)
            .expect("expected host result");

        assert_eq!(gpu_tensor.shape, expected.shape);
        assert_eq!(gpu_tensor.data.len(), expected.data.len());
        let tol = match runmat_accelerate_api::provider().unwrap().precision() {
            runmat_accelerate_api::ProviderPrecision::F64 => 1e-12,
            runmat_accelerate_api::ProviderPrecision::F32 => 1e-5,
        };
        for (gpu, cpu) in gpu_tensor.data.iter().zip(expected.data.iter()) {
            assert!((gpu - cpu).abs() < tol, "{gpu} vs {cpu}");
        }

        let _ = provider.free(&kernel_handle);
        let _ = provider.free(&image_handle);
        std::env::remove_var("RUNMAT_WGPU_SKIP_WARMUP");
        std::env::remove_var("RUNMAT_WGPU_DISABLE_IMFILTER");
    }

    #[test]
    fn filter2_descriptor_signatures_cover_surface() {
        let labels: Vec<&str> = FILTER2_DESCRIPTOR
            .signatures
            .iter()
            .map(|signature| signature.label)
            .collect();
        assert!(labels.contains(&"B = filter2(h, X)"));
        assert!(labels.contains(&"B = filter2(h, X, options...)"));
    }

    #[test]
    fn filter2_descriptor_errors_have_stable_codes() {
        let codes: Vec<&str> = FILTER2_DESCRIPTOR
            .errors
            .iter()
            .map(|error| error.code)
            .collect();
        assert!(codes.contains(&"RM.FILTER2.INVALID_INPUT"));
        assert!(codes.contains(&"RM.FILTER2.INVALID_OPTION"));
        assert!(codes.contains(&"RM.FILTER2.INTERNAL"));
    }
}
