use std::collections::HashMap;
use std::ffi::c_void;
use half::bf16;
use xserv_kernels::*;
use xserv_tensor::{Device, Tensor};

use crate::config::ModelConfig;
use crate::paged_kv_cache::PagedKVCache;

pub struct GptOss {
    pub config: ModelConfig,
    embed_tokens: Tensor,
    layers: Vec<GptOssBlock>,
    norm: Tensor,
    lm_head_t: Tensor,
    rope_cache: RopeCache,
    tp: Option<std::sync::Arc<xserv_distributed::TpContext>>,
    local_num_heads: usize,
    local_num_kv_heads: usize,
}

struct GptOssBlock {
    input_norm: Tensor,
    // Attention (with bias)
    q_proj_wt: Tensor,
    q_proj_bias: Tensor,
    k_proj_wt: Tensor,
    k_proj_bias: Tensor,
    v_proj_wt: Tensor,
    v_proj_bias: Tensor,
    o_proj_wt: Tensor,
    o_proj_bias: Tensor,
    sinks: Tensor,
    #[allow(dead_code)]
    is_sliding: bool,
    window_size: usize,
    // MoE MLP
    post_norm: Tensor,
    router_wt: Tensor,
    router_bias: Tensor,
    expert_gate_up_wt: Vec<Tensor>,
    expert_gate_up_bias: Vec<Tensor>,
    expert_down_wt: Vec<Tensor>,
    expert_down_bias: Vec<Tensor>,
    // Activation params
    glu_alpha: f32,
    glu_limit: f32,
}

impl GptOss {
    pub fn from_weights(config: ModelConfig, w: HashMap<String, Tensor>) -> Self {
        Self::from_weights_tp(config, w, 0, 1, 0, None)
    }

    pub fn from_weights_tp(
        config: ModelConfig,
        mut w: HashMap<String, Tensor>,
        rank: usize,
        world: usize,
        device: u32,
        tp: Option<std::sync::Arc<xserv_distributed::TpContext>>,
    ) -> Self {
        crate::init_kernels();
        let dev = Device::Cuda(device);
        let take = |w: &mut HashMap<String, Tensor>, name: &str| -> Tensor {
            w.remove(name).unwrap_or_else(|| panic!("missing weight: {name}"))
        };
        let repl = |t: Tensor| -> Tensor { t.to_device(dev) };
        // column-parallel: shard rows of [out, in], transpose → [in, out/world]
        let col = |t: Tensor| -> Tensor {
            shard_rows(&t, rank, world).to_device(dev).transpose(0, 1).contiguous()
        };
        // row-parallel: shard cols of [out, in], transpose → [in/world, out]
        let row = |t: Tensor| -> Tensor {
            shard_cols(&t, rank, world).to_device(dev).transpose(0, 1).contiguous()
        };
        // Bias sharding helpers
        let col_bias = |t: Tensor| -> Tensor { shard_1d(&t, rank, world).to_device(dev) };
        let repl_bias = |t: Tensor| -> Tensor { t.to_device(dev) };

        let embed_tokens = repl(take(&mut w, "model.embed_tokens.weight"));
        let norm = repl(take(&mut w, "model.norm.weight"));
        let lm_head_t = repl(take(&mut w, "lm_head.weight")).transpose(0, 1).contiguous();

        let head_dim = config.head_dim();
        let rope_theta = config.rope_theta.unwrap_or(150000.0);
        let max_seq_len = config.max_seq_len().min(8192); // cap for memory

        let rope_cache = if let Some(ref rs) = config.rope_scaling {
            if rs.rope_type.as_deref() == Some("yarn") {
                RopeCache::new_yarn(
                    max_seq_len,
                    head_dim,
                    rope_theta,
                    rs.factor.unwrap_or(1.0),
                    rs.original_max_position_embeddings.unwrap_or(4096),
                    rs.beta_fast.unwrap_or(32.0),
                    rs.beta_slow.unwrap_or(1.0),
                )
            } else {
                RopeCache::new(max_seq_len, head_dim, rope_theta as f32)
            }
        } else {
            RopeCache::new(max_seq_len, head_dim, rope_theta as f32)
        };

        let num_layers = config.num_layers();
        let num_experts = config.num_experts();
        let glu_alpha = 1.702f32;
        let glu_limit = config.swiglu_limit.unwrap_or(7.0) as f32;

        let mut layers = Vec::with_capacity(num_layers);
        if rank == 0 {
            eprintln!(
                "Loading gpt-oss weights: {} layers, {} experts, world={world}...",
                num_layers, num_experts
            );
        }

        for i in 0..num_layers {
            let p = format!("model.layers.{i}");

            // Attention weights — column-parallel for Q/K/V, row-parallel for O
            let q_proj_wt = col(take(&mut w, &format!("{p}.self_attn.q_proj.weight")));
            let q_proj_bias = col_bias(take(&mut w, &format!("{p}.self_attn.q_proj.bias")));
            let k_proj_wt = col(take(&mut w, &format!("{p}.self_attn.k_proj.weight")));
            let k_proj_bias = col_bias(take(&mut w, &format!("{p}.self_attn.k_proj.bias")));
            let v_proj_wt = col(take(&mut w, &format!("{p}.self_attn.v_proj.weight")));
            let v_proj_bias = col_bias(take(&mut w, &format!("{p}.self_attn.v_proj.bias")));
            let o_proj_wt = row(take(&mut w, &format!("{p}.self_attn.o_proj.weight")));
            let o_proj_bias = repl_bias(take(&mut w, &format!("{p}.self_attn.o_proj.bias")));

            // Sinks: shard per-head across TP ranks
            let sinks_full = take(&mut w, &format!("{p}.self_attn.sinks"));
            let sinks = shard_1d(&sinks_full, rank, world).to_device(dev);

            let is_sliding = config.is_sliding_layer(i);
            let window_size = if is_sliding { config.window_size() } else { 0 };

            // MoE weights — router replicated, experts split across TP ranks
            let router_wt_raw = take(&mut w, &format!("{p}.mlp.router.weight"));
            let router_wt = router_wt_raw.to_device(dev).transpose(0, 1).contiguous();
            let router_bias = repl_bias(take(&mut w, &format!("{p}.mlp.router.bias")));

            // Expert weights: [num_experts, hidden, 2*inter] — stored as 3D tensors
            // Expert parallelism: rank owns experts [rank*E/world .. (rank+1)*E/world)
            let gate_up_3d = take(&mut w, &format!("{p}.mlp.experts.gate_up_proj"));
            let gate_up_bias_2d = take(&mut w, &format!("{p}.mlp.experts.gate_up_proj_bias"));
            let down_3d = take(&mut w, &format!("{p}.mlp.experts.down_proj"));
            let down_bias_2d = take(&mut w, &format!("{p}.mlp.experts.down_proj_bias"));

            let local_experts = num_experts / world;
            let expert_start = rank * local_experts;

            let mut expert_gate_up_wt = Vec::with_capacity(local_experts);
            let mut expert_gate_up_bias = Vec::with_capacity(local_experts);
            let mut expert_down_wt = Vec::with_capacity(local_experts);
            let mut expert_down_bias = Vec::with_capacity(local_experts);

            let inter2 = gate_up_3d.shape()[2]; // 2 * intermediate_size
            let hidden = gate_up_3d.shape()[1];
            let inter = down_3d.shape()[1]; // intermediate_size

            for local_e in 0..local_experts {
                let e = expert_start + local_e;
                let gu_slice = slice_expert_3d(&gate_up_3d, e, hidden, inter2);
                expert_gate_up_wt.push(gu_slice.to_device(dev));

                let gu_bias = slice_expert_2d(&gate_up_bias_2d, e, inter2);
                expert_gate_up_bias.push(gu_bias.to_device(dev));

                let d_slice = slice_expert_3d(&down_3d, e, inter, hidden);
                expert_down_wt.push(d_slice.to_device(dev));

                let d_bias = slice_expert_2d(&down_bias_2d, e, hidden);
                expert_down_bias.push(d_bias.to_device(dev));
            }

            xserv_cuda::allocator::cached_trim();

            layers.push(GptOssBlock {
                input_norm: repl(take(&mut w, &format!("{p}.input_layernorm.weight"))),
                q_proj_wt,
                q_proj_bias,
                k_proj_wt,
                k_proj_bias,
                v_proj_wt,
                v_proj_bias,
                o_proj_wt,
                o_proj_bias,
                sinks,
                is_sliding,
                window_size,
                post_norm: repl(take(&mut w, &format!("{p}.post_attention_layernorm.weight"))),
                router_wt,
                router_bias,
                expert_gate_up_wt,
                expert_gate_up_bias,
                expert_down_wt,
                expert_down_bias,
                glu_alpha,
                glu_limit,
            });
        }

        let local_num_heads = config.num_heads() / world;
        let local_num_kv_heads = config.num_kv_heads() / world;

        Self {
            config,
            embed_tokens,
            layers,
            norm,
            lm_head_t,
            rope_cache,
            tp,
            local_num_heads,
            local_num_kv_heads,
        }
    }

    #[inline]
    fn all_reduce(&self, t: &Tensor) {
        if let Some(tp) = &self.tp {
            if tp.world > 1 {
                let ptr = t.storage().gpu_buffer().as_ptr() as *mut c_void;
                tp.all_reduce_sum_bf16_ptr(ptr, t.numel());
            }
        }
    }

    /// Paged decode: process one token per sequence using paged KV cache.
    pub fn forward_decode_paged(
        &self,
        tokens: &[u32],
        positions: &[usize],
        seq_slots: &[usize],
        paged_cache: &mut PagedKVCache,
    ) -> Tensor {
        let batch = tokens.len();
        assert_eq!(positions.len(), batch);
        assert_eq!(seq_slots.len(), batch);
        assert!(batch > 0);

        let num_heads = self.local_num_heads;
        let num_kv_heads = self.local_num_kv_heads;
        let head_dim = self.config.head_dim();
        let eps = self.config.rms_norm_eps.unwrap_or(1e-5) as f32;

        let kv_lens: Vec<i32> = positions.iter().map(|&p| (p + 1) as i32).collect();
        for (b, &slot) in seq_slots.iter().enumerate() {
            paged_cache.ensure_capacity(slot, positions[b] + 1);
        }
        paged_cache.sync_active_batch_with_lens(seq_slots, &kv_lens);

        let bt_ptr = paged_cache.block_table_gpu().as_ptr() as *const i32;
        let cl_ptr = paged_cache.context_lens_gpu().as_ptr() as *const i32;
        let max_blocks = paged_cache.max_blocks_per_seq();

        let positions_u32: Vec<u32> = positions.iter().map(|&p| p as u32).collect();

        let mut x = embedding(&self.embed_tokens, tokens);

        for (layer_idx, layer) in self.layers.iter().enumerate() {
            let residual = x.clone();
            let normed = rmsnorm(&x, &layer.input_norm, eps);

            // Q/K/V projections with bias
            let q_all = add_bias(&matmul_2d(&normed, &layer.q_proj_wt), &layer.q_proj_bias);
            let k_all = add_bias(&matmul_2d(&normed, &layer.k_proj_wt), &layer.k_proj_bias);
            let v_all = add_bias(&matmul_2d(&normed, &layer.v_proj_wt), &layer.v_proj_bias);


            // Reshape for RoPE: [B, H*D] → [B, H, D]
            let q_3d = q_all.reshape(&[batch, num_heads, head_dim]);
            let k_3d = k_all.reshape(&[batch, num_kv_heads, head_dim]);

            // RoPE (no QK-norm for gpt-oss)
            rope_inplace(&q_3d, &self.rope_cache, &positions_u32);
            rope_inplace(&k_3d, &self.rope_cache, &positions_u32);

            let v_3d = v_all.reshape(&[batch, num_kv_heads, head_dim]);

            // KV cache scatter
            paged_cache.append_tokens_batched(layer_idx, &k_3d, &v_3d, batch);

            // Paged attention with sinks + sliding window
            let q_4d = q_3d.reshape(&[batch, num_heads, 1, head_dim]);
            let k_pool_ptr = paged_cache.k_pool(layer_idx).as_ptr() as *const c_void;
            let v_pool_ptr = paged_cache.v_pool(layer_idx).as_ptr() as *const c_void;
            let sinks_ptr = layer.sinks.data_ptr() as *const c_void;

            let attn_out = paged_decode_attention_sinks(
                &q_4d, k_pool_ptr, v_pool_ptr, bt_ptr, cl_ptr,
                sinks_ptr,
                batch, num_heads, num_kv_heads, head_dim, max_blocks,
                layer.window_size,
            );

            let attn_merged = attn_out.reshape(&[batch, num_heads * head_dim]);
            let attn_proj = matmul_2d(&attn_merged, &layer.o_proj_wt);
            self.all_reduce(&attn_proj);
            let attn_proj = add_bias(&attn_proj, &layer.o_proj_bias);


            // Residual + post-norm
            let (normed, x_new) = add_rmsnorm(&attn_proj, &residual, &layer.post_norm, eps);


            let residual = x_new;
            let normed = normed.contiguous();


            // MoE MLP
            let moe_out = self.moe_forward(&normed, layer, batch);
            x = xserv_kernels::add(&residual, &moe_out);
        }

        // Advance KV cache
        for &slot in seq_slots {
            paged_cache.advance_seq_len(slot, 1);
        }

        unsafe { xserv_cuda::ffi::cudaDeviceSynchronize(); }
        let x = rmsnorm(&x, &self.norm, eps);
        unsafe { xserv_cuda::ffi::cudaDeviceSynchronize(); }
        let logits = matmul_2d(&x, &self.lm_head_t);
        unsafe { xserv_cuda::ffi::cudaDeviceSynchronize(); }
        logits
    }

    /// Paged prefill: process full prompt tokens.
    pub fn forward_prefill_paged(
        &self,
        token_ids: &[u32],
        slot: usize,
        paged_cache: &mut PagedKVCache,
    ) -> Tensor {
        let new_tokens = token_ids.len();
        let pos_offset = paged_cache.seq_len(slot);
        let num_heads = self.local_num_heads;
        let num_kv_heads = self.local_num_kv_heads;
        let head_dim = self.config.head_dim();
        let eps = self.config.rms_norm_eps.unwrap_or(1e-5) as f32;

        paged_cache.ensure_capacity(slot, pos_offset + new_tokens);
        paged_cache.advance_seq_len(slot, new_tokens);

        let mut x = embedding(&self.embed_tokens, token_ids);
        let positions: Vec<u32> = (pos_offset..pos_offset + new_tokens).map(|p| p as u32).collect();

        for (layer_idx, layer) in self.layers.iter().enumerate() {
            let residual = x.clone();
            let normed = rmsnorm(&x, &layer.input_norm, eps);

            let q = add_bias(&matmul_2d(&normed, &layer.q_proj_wt), &layer.q_proj_bias);
            let k = add_bias(&matmul_2d(&normed, &layer.k_proj_wt), &layer.k_proj_bias);
            let v = add_bias(&matmul_2d(&normed, &layer.v_proj_wt), &layer.v_proj_bias);

            let q = reshape_heads_gpu(&q, new_tokens, num_heads, head_dim);
            let k = reshape_heads_gpu(&k, new_tokens, num_kv_heads, head_dim);
            let v = reshape_heads_gpu(&v, new_tokens, num_kv_heads, head_dim);

            // RoPE
            let q = transpose_for_rope_gpu(&q, new_tokens, num_heads, head_dim);
            let k = transpose_for_rope_gpu(&k, new_tokens, num_kv_heads, head_dim);
            rope_inplace(&q, &self.rope_cache, &positions);
            rope_inplace(&k, &self.rope_cache, &positions);
            let q = transpose_from_rope_gpu(&q, new_tokens, num_heads, head_dim);
            let k = transpose_from_rope_gpu(&k, new_tokens, num_kv_heads, head_dim);

            // KV cache
            paged_cache.append_tokens(slot, layer_idx, &k, &v, new_tokens, pos_offset);
            let (k_full, v_full) = paged_cache.gather_kv_contiguous(slot, layer_idx);

            // Flash attention with gpt-oss sinks + (per-layer) sliding window.
            let attn_out = flash_attention_sinks(&q, &k_full, &v_full, &layer.sinks, layer.window_size);

            let attn_merged = merge_heads_gpu(&attn_out, new_tokens, num_heads, head_dim);
            let attn_proj = matmul_2d(&attn_merged, &layer.o_proj_wt);
            self.all_reduce(&attn_proj);
            let attn_proj = add_bias(&attn_proj, &layer.o_proj_bias);

            let (normed, x_new) = add_rmsnorm(&attn_proj, &residual, &layer.post_norm, eps);
            let residual = x_new;

            // MoE MLP
            let moe_out = self.moe_forward(&normed, layer, new_tokens);
            x = xserv_kernels::add(&residual, &moe_out);
        }

        let x = rmsnorm(&x, &self.norm, eps);
        let logits = matmul_2d(&x, &self.lm_head_t);
        unsafe { xserv_cuda::ffi::cudaDeviceSynchronize(); }
        logits
    }

    /// MoE forward pass for one layer with expert parallelism.
    /// Each rank owns `num_experts / world` experts. Tokens routed to non-local
    /// experts get zero contribution from this rank; AllReduce sums all ranks.
    /// Input: [tokens, hidden], Output: [tokens, hidden]
    fn moe_forward(&self, x: &Tensor, layer: &GptOssBlock, num_tokens: usize) -> Tensor {
        let hidden = self.config.hidden();
        let num_experts = self.config.num_experts();
        let top_k = self.config.experts_per_token();
        let world = self.tp.as_ref().map(|tp| tp.world).unwrap_or(1);
        let rank = self.tp.as_ref().map(|tp| tp.rank).unwrap_or(0);
        let local_experts = num_experts / world;
        let expert_start = rank * local_experts;

        // Router: [tokens, hidden] @ [hidden, num_experts] + bias → [tokens, num_experts]
        unsafe { xserv_cuda::ffi::cudaDeviceSynchronize(); }
        let router_logits = add_bias(
            &matmul_2d(x, &layer.router_wt),
            &layer.router_bias,
        );
        unsafe { xserv_cuda::ffi::cudaDeviceSynchronize(); }
        let router_cpu = router_logits.to_device(Device::Cpu);
        let router_data = router_cpu.as_slice::<bf16>();
        let x_cpu = x.to_device(Device::Cpu);
        let x_data = x_cpu.as_slice::<bf16>();

        let mut output_acc = vec![0.0f32; num_tokens * hidden];

        for t in 0..num_tokens {
            let row = &router_data[t * num_experts..(t + 1) * num_experts];

            // Find top-k expert indices (global)
            let mut indices: Vec<(usize, f32)> = row.iter()
                .enumerate()
                .map(|(i, &v)| (i, v.to_f32()))
                .collect();
            indices.sort_by(|a, b| b.1.partial_cmp(&a.1).unwrap_or(std::cmp::Ordering::Equal));
            let top_indices: Vec<(usize, f32)> = indices[..top_k].to_vec();

            // Softmax over top-k logits
            let max_val = top_indices.iter().map(|x| x.1).fold(f32::NEG_INFINITY, f32::max);
            let exp_sum: f32 = top_indices.iter().map(|x| (x.1 - max_val).exp()).sum();
            let weights: Vec<f32> = top_indices.iter()
                .map(|x| (x.1 - max_val).exp() / exp_sum)
                .collect();

            // Fresh GPU upload of token data — immune to cached allocator buffer reuse
            let token_slice = &x_data[t * hidden..(t + 1) * hidden];
            let token_tensor = Tensor::from_slice(token_slice, &[1, hidden]).to_device(x.device());


            for (k_idx, &(expert_id, _)) in top_indices.iter().enumerate() {
                // Only process experts owned by this rank
                if expert_id < expert_start || expert_id >= expert_start + local_experts {
                    continue;
                }
                let local_id = expert_id - expert_start;
                let weight = weights[k_idx];

                let gate_up_raw = matmul_2d(&token_tensor, &layer.expert_gate_up_wt[local_id]);
                let gate_up = add_bias(&gate_up_raw, &layer.expert_gate_up_bias[local_id]);

                let activated = gpt_oss_glu(&gate_up, layer.glu_alpha, layer.glu_limit);

                let down_raw = matmul_2d(&activated, &layer.expert_down_wt[local_id]);
                let down = add_bias(&down_raw, &layer.expert_down_bias[local_id]);


                let down_cpu = down.to_device(Device::Cpu);
                let down_data = down_cpu.as_slice::<bf16>();
                for d in 0..hidden {
                    output_acc[t * hidden + d] += weight * down_data[d].to_f32();
                }
            }
        }

        // Convert accumulated output to BF16 tensor on GPU
        let output_bf16: Vec<bf16> = output_acc.iter().map(|&v| bf16::from_f32(v)).collect();
        let moe_out = Tensor::from_slice(&output_bf16, &[num_tokens, hidden]).to_device(x.device());

        self.all_reduce(&moe_out);
        moe_out
    }
}

// --- Helpers ---

fn matmul_2d(a: &Tensor, b: &Tensor) -> Tensor {
    assert_eq!(a.ndim(), 2);
    assert_eq!(b.ndim(), 2);
    matmul(a, b, GemmBackend::CuBlas)
}

/// Add bias to a 2D tensor: [rows, cols] + [cols] → [rows, cols]
fn add_bias(x: &Tensor, bias: &Tensor) -> Tensor {
    assert_eq!(x.ndim(), 2);
    assert_eq!(bias.ndim(), 1);
    let rows = x.shape()[0];
    let cols = x.shape()[1];
    assert_eq!(bias.shape()[0], cols, "bias size {} != cols {}", bias.shape()[0], cols);

    // Broadcast bias to each row using GPU kernels.
    // Tile bias [cols] into [rows, cols] by repeating rows, then add element-wise.
    let bias_cpu = bias.to_device(Device::Cpu);
    let bias_data = bias_cpu.as_slice::<bf16>();
    let mut tiled = Vec::with_capacity(rows * cols);
    for _ in 0..rows {
        tiled.extend_from_slice(bias_data);
    }
    let bias_tiled = Tensor::from_slice(&tiled, &[rows, cols]).to_device(x.device());
    let x_c = x.contiguous();
    xserv_kernels::add(&x_c, &bias_tiled)
}

fn shard_rows(t: &Tensor, rank: usize, world: usize) -> Tensor {
    if world == 1 { return t.clone(); }
    let shape = t.shape();
    assert_eq!(shape.len(), 2);
    let (rows, cols) = (shape[0], shape[1]);
    assert!(rows % world == 0, "rows {rows} not divisible by world {world}");
    let local = rows / world;
    let host = t.to_device(Device::Cpu);
    let data = host.as_slice::<bf16>();
    let start = rank * local * cols;
    let shard = data[start..start + local * cols].to_vec();
    Tensor::from_slice(&shard, &[local, cols])
}

fn shard_cols(t: &Tensor, rank: usize, world: usize) -> Tensor {
    if world == 1 { return t.clone(); }
    let shape = t.shape();
    assert_eq!(shape.len(), 2);
    let (rows, cols) = (shape[0], shape[1]);
    assert!(cols % world == 0, "cols {cols} not divisible by world {world}");
    let local = cols / world;
    let c0 = rank * local;
    let host = t.to_device(Device::Cpu);
    let data = host.as_slice::<bf16>();
    let mut shard = Vec::with_capacity(rows * local);
    for r in 0..rows {
        let base = r * cols + c0;
        shard.extend_from_slice(&data[base..base + local]);
    }
    Tensor::from_slice(&shard, &[rows, local])
}

fn shard_1d(t: &Tensor, rank: usize, world: usize) -> Tensor {
    if world == 1 { return t.clone(); }
    let shape = t.shape();
    assert_eq!(shape.len(), 1);
    let total = shape[0];
    assert!(total % world == 0, "dim {total} not divisible by world {world}");
    let local = total / world;
    let host = t.to_device(Device::Cpu);
    let data = host.as_slice::<bf16>();
    let start = rank * local;
    let shard = data[start..start + local].to_vec();
    Tensor::from_slice(&shard, &[local])
}

/// Extract expert `e` from a [num_experts, rows, cols] 3D tensor → [rows, cols] 2D
fn slice_expert_3d(t: &Tensor, e: usize, rows: usize, cols: usize) -> Tensor {
    assert_eq!(t.ndim(), 3);
    let host = t.to_device(Device::Cpu);
    let data = host.as_slice::<bf16>();
    let stride = rows * cols;
    let start = e * stride;
    let slice = data[start..start + stride].to_vec();
    Tensor::from_slice(&slice, &[rows, cols])
}

/// Extract expert `e` from a [num_experts, dim] 2D tensor → [dim] 1D
fn slice_expert_2d(t: &Tensor, e: usize, dim: usize) -> Tensor {
    assert_eq!(t.ndim(), 2);
    let host = t.to_device(Device::Cpu);
    let data = host.as_slice::<bf16>();
    let start = e * dim;
    let slice = data[start..start + dim].to_vec();
    Tensor::from_slice(&slice, &[dim])
}