//! An OpenCL kernel.

use std;
use std::ops::Deref;
use std::any::Any;
use std::any::TypeId;
use std::collections::{HashMap, BTreeMap};
use std::marker::PhantomData;
use std::cell::RefCell;
use std::borrow::Borrow;
use core::ffi::c_void;
use core::{self, util, OclPrm, Kernel as KernelCore, CommandQueue as CommandQueueCore, Mem as MemCore,
    ArgVal, KernelInfo, KernelInfoResult, KernelArgInfo, KernelArgInfoResult,
    KernelWorkGroupInfo, KernelWorkGroupInfoResult, AsMem, MemCmdAll, ClVersions};
use core::error::{Result as OclCoreResult, ErrorKind as OclCoreErrorKind};
use error::{Error as OclError, Result as OclResult};
use standard::{SpatialDims, Program, Queue, WorkDims, Sampler, Device, ClNullEventPtrEnum,
    ClWaitListPtrEnum, Buffer, Image};
pub use self::arg_type::{BaseType, Cardinality, ArgType};


/// An error related to a `Kernel`, `KernelBuilder`, or `KernelCmd`.
#[derive(Debug, Fail)]
pub enum KernelError {
    #[fail(display = "No queue specified.")]
    CmdNoQueue,
    #[fail(display = "Global Work Size cannot be left unspecified. Set a default for \
        the kernel or specify one when enqueuing command.")]
    CmdNoGws,
    #[fail(display = "Unable to resolve argument named: '{}'. Ensure that an argument with \
        that name has been declared before building kernel.", _0)]
    NamedArgsInvalidArgName(String),
    #[fail(display = "No named arguments have been declared. Declare named arguments \
        when before building kernel.")]
    NamedArgsNone,
    #[fail(display = "Kernel arg index out of range. (kernel: {}, index: {})", _0, _1)]
    ArgIdxOor(String, u32),
    #[fail(display = "Kernel argument type mismatch. The argument named: '{}' at index: [{}] \
        should be a '{}' ({:?}).", arg_name, idx, ty_name, ty)]
    ArgTypeMismatch { idx: u32, arg_name: String, ty_name: String, ty: ArgType },
    #[fail(display = "No program specified.")]
    BuilderNoProgram,
    #[fail(display = "No kernel name specified.")]
    BuilderNoKernelName,
    #[fail(display = "The wrong number of kernel arguments have been specified \
        (required: {}, specified: {}). Use named arguments with 'None' or zero values to \
        declare arguments you plan to assign a value to at a later time.", required, specified)]
    BuilderWrongArgCount { required: u32, specified: u32 },
}


/// A kernel command builder used to enqueue a kernel with a mix of default
/// and optionally specified arguments.
#[must_use = "commands do nothing unless enqueued"]
pub struct KernelCmd<'k> {
    kernel: &'k KernelCore,
    queue: Option<&'k CommandQueueCore>,
    gwo: SpatialDims,
    gws: SpatialDims,
    lws: SpatialDims,
    wait_events: Option<ClWaitListPtrEnum<'k>>,
    new_event: Option<ClNullEventPtrEnum<'k>>,
}

/// A kernel enqueue command.
impl<'k> KernelCmd<'k> {
    /// Specifies a queue to use for this call only.
    ///
    /// Overrides the kernel's default queue if one is set. If no default
    /// queue is set, this method **must** be called before enqueuing the
    /// kernel.
    pub fn queue<'q, Q>(mut self, queue: &'q Q) -> KernelCmd<'k>
            where 'q: 'k, Q: 'k + AsRef<CommandQueueCore> {
        self.queue = Some(queue.as_ref());
        self
    }

    /// Specifies a global work offset for this call only.
    #[deprecated(since = "0.18.0", note = "Use `::global_work_offset` instead.")]
     pub fn gwo<D: Into<SpatialDims>>(mut self, gwo: D) -> KernelCmd<'k> {
        self.gwo = gwo.into();
        self
    }

    /// Specifies a global work size for this call only.
    #[deprecated(since = "0.18.0", note = "Use `::global_work_size` instead.")]
    pub fn gws<D: Into<SpatialDims>>(mut self, gws: D) -> KernelCmd<'k> {
        self.gws = gws.into();
        self
    }

    /// Specifies a local work size for this call only.
    #[deprecated(since = "0.18.0", note = "Use `::local_work_size` instead.")]
    pub fn lws<D: Into<SpatialDims>>(mut self, lws: D) -> KernelCmd<'k> {
        self.lws = lws.into();
        self
    }

    /// Specifies a global work offset for this call only.
     pub fn global_work_offset<D: Into<SpatialDims>>(mut self, gwo: D) -> KernelCmd<'k> {
        self.gwo = gwo.into();
        self
    }

    /// Specifies a global work size for this call only.
    pub fn global_work_size<D: Into<SpatialDims>>(mut self, gws: D) -> KernelCmd<'k> {
        self.gws = gws.into();
        self
    }

    /// Specifies a local work size for this call only.
    pub fn local_work_size<D: Into<SpatialDims>>(mut self, lws: D) -> KernelCmd<'k> {
        self.lws = lws.into();
        self
    }

    /// Specifies an event or list of events to wait on before the command
    /// will run.
    ///
    /// When events generated using the `::enew` method of **other**,
    /// previously enqueued commands are passed here (either individually or
    /// as part of an [`EventList`]), this command will not execute until
    /// those commands have completed.
    ///
    /// Using events can compliment the use of queues to order commands by
    /// creating temporal dependencies between them (where commands in one
    /// queue must wait for the completion of commands in another). Events can
    /// also supplant queues altogether when, for example, using out-of-order
    /// queues.
    ///
    /// # Example
    ///
    /// ```rust,ignore
    /// // Create an event list:
    /// let mut event_list = EventList::new();
    /// // Enqueue a kernel on `queue_1`, creating an event representing the kernel
    /// // command in our list:
    /// kernel.cmd().queue(&queue_1).enew(&mut event_list).enq()?;
    /// // Read from a buffer using `queue_2`, ensuring the read does not begin until
    /// // after the kernel command has completed:
    /// buffer.read(rwvec.clone()).queue(&queue_2).ewait(&event_list).enq_async()?;
    /// ```
    ///
    /// [`EventList`]: struct.EventList.html
    pub fn ewait<'e, Ewl>(mut self, ewait: Ewl) -> KernelCmd<'k>
            where 'e: 'k, Ewl: Into<ClWaitListPtrEnum<'e>> {
        self.wait_events = Some(ewait.into());
        self
    }

    /// Specifies the destination to store a new, optionally created event
    /// associated with this command.
    ///
    /// The destination can be a mutable reference to an empty event (created
    /// using [`Event::empty`]) or a mutable reference to an event list.
    ///
    /// After this command is enqueued, the event in the destination can be
    /// passed to the `::ewait` method of another command. Doing so will cause
    /// the other command to wait until this command has completed before
    /// executing.
    ///
    /// Using events can compliment the use of queues to order commands by
    /// creating temporal dependencies between them (where commands in one
    /// queue must wait for the completion of commands in another). Events can
    /// also supplant queues altogether when, for example, using out-of-order
    /// queues.
    ///
    /// # Example
    ///
    /// ```rust,ignore
    /// // Create an event list:
    /// let mut event = Event::empty();
    /// // Enqueue a kernel on `queue_1`, creating an event representing the kernel
    /// // command in our list:
    /// kernel.cmd().queue(&queue_1).enew(&mut event).enq()?;
    /// // Read from a buffer using `queue_2`, ensuring the read does not begin until
    /// // after the kernel command has completed:
    /// buffer.read(rwvec.clone()).queue(&queue_2).ewait(&event).enq_async()?;
    /// ```
    ///
    /// [`Event::empty`]: struct.Event.html#method.empty
    pub fn enew<'e, En>(mut self, new_event_dest: En) -> KernelCmd<'k>
            where 'e: 'k, En: Into<ClNullEventPtrEnum<'e>> {
        self.new_event = Some(new_event_dest.into());
        self
    }

    /// Enqueues this kernel command.
    ///
    /// # Safety
    ///
    /// All kernel code must be considered untrusted. Therefore the act of
    /// calling this function contains implied unsafety even though the API
    /// itself is safe.
    pub unsafe fn enq(self) -> OclResult<()> {
        let queue = match self.queue {
            Some(q) => q,
            None => return Err(KernelError::CmdNoQueue.into()),
        };

        let dim_count = self.gws.dim_count();

        let gws = match self.gws.to_work_size() {
            Some(gws) => gws,
            None => return Err(KernelError::CmdNoGws.into()),
        };

        core::enqueue_kernel(queue, &self.kernel, dim_count, self.gwo.to_work_offset(),
            &gws, self.lws.to_work_size(), self.wait_events, self.new_event)
            .map_err(OclError::from)
    }
}


/// Converts an argument index specifier to `u32`.
#[derive(Clone, Copy, Debug)]
pub enum ArgIdxSpecifier {
    Uint(u32),
    Str(&'static str),
}

impl ArgIdxSpecifier {
    // Resolves this specifier into an integer index using the specified
    // argument map if necessary.
    fn to_idx(&self, named_args: &NamedArgs) -> OclResult<u32> {
        match *self {
            ArgIdxSpecifier::Uint(idx) => Ok(idx),
            ArgIdxSpecifier::Str(s) => named_args.resolve_idx(s),
        }
    }
}

impl From<u32> for ArgIdxSpecifier {
    fn from(idx: u32) -> ArgIdxSpecifier {
        ArgIdxSpecifier::Uint(idx)
    }
}

impl From<i32> for ArgIdxSpecifier {
    fn from(idx: i32) -> ArgIdxSpecifier {
        ArgIdxSpecifier::Uint(idx as u32)
    }
}

impl From<usize> for ArgIdxSpecifier {
    fn from(idx: usize) -> ArgIdxSpecifier {
        ArgIdxSpecifier::Uint(idx as u32)
    }
}

impl From<&'static str> for ArgIdxSpecifier {
    fn from(s: &'static str) -> ArgIdxSpecifier {
        ArgIdxSpecifier::Str(s)
    }
}


/// Contains owned or shared argument values.
#[derive(Debug, Clone)]
enum ArgValKeeper<'b> {
    Shared(ArgVal<'b>),
    Owned(Vec<u8>),
}

impl<'b> ArgValKeeper<'b> {
    pub fn owned<T>(prm: T) -> ArgValKeeper<'b> where T: OclPrm {
        unsafe { ArgValKeeper::Owned(util::into_bytes(prm)) }
    }

    pub fn mem_null() -> ArgValKeeper<'b> {
        ArgValKeeper::Shared(ArgVal::mem_null())
    }

    /// Returns the size (in bytes) and raw pointer to the contained kernel
    /// argument value.
    pub fn to_arg_val(&'b self) -> ArgVal<'b> {
        match *self {
            ArgValKeeper::Shared(ref av) => av.clone(),
            ArgValKeeper::Owned(ref bytes) => {
                unsafe { ArgVal::from_raw(bytes.len(), bytes.as_ptr() as *const c_void) }
            }
        }
    }
}


/// Wraps argument values of different types.
pub struct ArgValConverter<'b, T> where T: OclPrm {
    val: ArgValKeeper<'b>,
    type_id: Option<TypeId>,
    mem: Option<MemCore>,
    _ty: PhantomData<T>,
}

impl<'b, T> From<Option<&'b Buffer<T>>> for ArgValConverter<'b, T> where T: OclPrm {
    /// Converts from an Option<`Buffer`>.
    fn from(buf: Option<&'b Buffer<T>>) -> ArgValConverter<'b, T> {
        let val = match buf {
            Some(b) => ArgValKeeper::Shared(ArgVal::mem(b)),
            None => ArgValKeeper::mem_null(),
        };

        ArgValConverter {
            val,
            type_id: Some(TypeId::of::<T>()),
            mem: buf.map(|b| b.as_mem().clone()),
            _ty: PhantomData,
        }
    }
}

impl<'b, T> From<&'b Buffer<T>> for ArgValConverter<'b, T> where T: OclPrm {
    /// Converts from a `Buffer`.
    fn from(buf: &'b Buffer<T>) -> ArgValConverter<'b, T> {
        ArgValConverter {
            val: ArgValKeeper::Shared(ArgVal::mem(buf)),
            type_id: Some(TypeId::of::<T>()),
            mem: Some(buf.as_mem().clone()),
            _ty: PhantomData,
        }
    }
}

impl<'b, T> From<&'b mut Buffer<T>> for ArgValConverter<'b, T> where T: OclPrm {
    fn from(buf: &'b mut Buffer<T>) -> ArgValConverter<'b, T> {
        ArgValConverter::from(&*buf)
    }
}

// impl<'b, T> From<&'b &'b Buffer<T>> for ArgValConverter<'b, T> where T: OclPrm {
//     fn from(buf: &'b &'b Buffer<T>) -> ArgValConverter<'b, T> {
//         ArgValConverter::from(&**buf)
//     }
// }

impl<'b, T> From<Option<&'b Image<T>>> for ArgValConverter<'b, T> where T: OclPrm {
    /// Converts from an Option<`Image`>.
    fn from(img: Option<&'b Image<T>>) -> ArgValConverter<'b, T> {
        let val = match img {
            Some(i) => ArgValKeeper::Shared(ArgVal::mem(i)),
            None => ArgValKeeper::mem_null(),
        };

        ArgValConverter {
            val,
            type_id: None,
            mem: img.map(|i| i.as_mem().clone()),
            _ty: PhantomData,
        }
    }
}

impl<'b, T> From<&'b Image<T>> for ArgValConverter<'b, T> where T: OclPrm {
    /// Converts from an `Image`.
    fn from(img: &'b Image<T>) -> ArgValConverter<'b, T> {
        ArgValConverter {
            val: ArgValKeeper::Shared(ArgVal::mem(img)),
            type_id: None,
            mem: Some(img.as_mem().clone()),
            _ty: PhantomData,
        }
    }
}

impl<'b, T> From<&'b mut Image<T>> for ArgValConverter<'b, T> where T: OclPrm {
    fn from(img: &'b mut Image<T>) -> ArgValConverter<'b, T> {
        ArgValConverter::from(&*img)
    }
}

// impl<'b, T> From<&'b &'b Image<T>> for ArgValConverter<'b, T> where T: OclPrm {
//     fn from(img: &'b &'b Image<T>) -> ArgValConverter<'b, T> {
//         ArgValConverter::from(&**img)
//     }
// }

impl<'b, T> From<&'b T> for ArgValConverter<'b, T> where T: OclPrm {
    /// Converts from a scalar or vector value.
    fn from(prm: &'b T) -> ArgValConverter<'b, T> {
        ArgValConverter {
            val: ArgValKeeper::Shared(ArgVal::scalar(prm)),
            type_id: Some(TypeId::of::<T>()),
            mem: None,
            _ty: PhantomData,
        }
    }
}

impl<'b, T> From<T> for ArgValConverter<'b, T> where T: OclPrm {
    /// Converts from a scalar or vector value.
    fn from(prm: T) -> ArgValConverter<'b, T> {
        ArgValConverter {
            val: ArgValKeeper::owned(prm),
            type_id: Some(TypeId::of::<T>()),
            mem: None,
            _ty: PhantomData,
        }
    }
}



/// A map of argument names -> indexes.
#[derive(Clone, Debug)]
struct NamedArgs(Option<HashMap<&'static str, u32>>);

impl NamedArgs {
    /// Inserts a named argument into the map.
    fn insert(&mut self, name: &'static str, arg_idx: u32) {
        if self.0.is_none() {
            self.0 = Some(HashMap::with_capacity(8));
        }
        self.0.as_mut().unwrap().insert(name, arg_idx);
    }

    /// Resolves the index of a named argument with a friendly error message.
    fn resolve_idx(&self, name: &'static str) -> OclResult<u32> {
        match self.0 {
            Some(ref map) => {
                match map.get(name) {
                    Some(&ai) => Ok(ai),
                    None => Err(KernelError::NamedArgsInvalidArgName(name.into()).into()),
                }
            },
            None => Err(KernelError::NamedArgsNone.into()),
        }
    }
}


/// Storage for `Mem` arguments.
#[derive(Clone, Debug)]
struct MemArgs(Option<RefCell<BTreeMap<u32, MemCore>>>);

impl MemArgs {
    /// Inserts a `Mem` argument for storage if possible.
    fn insert(&self, idx: u32, mem: MemCore) {
        if let Some(ref map) = self.0 {
            map.borrow_mut().insert(idx, mem);
        }
    }

    /// Removes a `Mem` argument.
    fn remove(&self, idx: &u32) {
        if let Some(ref map) = self.0 {
            map.borrow_mut().remove(idx);
        }
    }
}


/// A kernel which represents a 'procedure'.
///
/// Corresponds to code which must have already been compiled into a program.
///
/// Set arguments using `::set_arg` or any of the `::set_arg...` methods.
///
/// `Kernel` includes features that a raw OpenCL kernel does not, including:
///
/// 1. Type-checked arguments (not just size-checked)
/// 2. Named arguments (with a `&'static str` name)
/// 3. Prevention of a potential (difficult to debug) segfault if a buffer or
///    image used by a kernel is dropped prematurely.
/// 4. Stored defaults for the:
///     - Queue
///     - Global Work Offset
///     - Global Work Size
///     - Local Work Size
///
/// ### `Clone` and `Send`
///
/// A `Kernel` may not be cloned but may be sent between threads. This ensures
/// that no two threads create a race condition by attempting to set an
/// argument and enqueue a kernel at the same time. Use the `KernelBuilder` to
/// create multiple identical kernels (`KernelBuilder` is clonable and
/// re-usable).
#[derive(Debug)]
pub struct Kernel {
    obj_core: KernelCore,
    named_args: NamedArgs,
    mem_args: MemArgs,
    queue: Option<Queue>,
    gwo: SpatialDims,
    gws: SpatialDims,
    lws: SpatialDims,
    arg_types: Option<Vec<ArgType>>,
}

impl Kernel {
    /// Returns a new `KernelBuilder`.
    pub fn builder<'p>() -> KernelBuilder<'p> {
        KernelBuilder::new()
    }

    /// Verifies that a type matches the kernel arg info:
    ///
    /// This function does nothing and always returns `Ok` if argument type
    /// checking is disabled by the user or for any other reason.
    ///
    /// Argument type checking will automatically be disabled if any of the
    /// devices in use do not support OpenCL version 1.2 or higher or if
    /// argument information is not available on the associated platform.
    fn verify_arg_type<T: OclPrm + Any>(&self, arg_idx: u32) -> OclResult<()> {
        if let Some(ref arg_types) = self.arg_types {
            let arg_type = arg_types.get(arg_idx as usize)
                .ok_or(KernelError::ArgIdxOor(self.name()?, arg_idx))?;

            if arg_type.is_match::<T>() {
                Ok(())
            } else {
                let ty_name = arg_type_name(&self.obj_core, arg_idx)?;
                let arg_name = arg_name(&self.obj_core, arg_idx)?;
                Err(KernelError::ArgTypeMismatch { idx: arg_idx, arg_name,
                    ty_name, ty: arg_type.clone() }.into())
            }
        } else {
            Ok(())
        }
    }

    /// Returns the argument index of a named argument if it exists.
    pub fn named_arg_idx(&self, name: &'static str) -> Option<u32> {
        self.named_args.resolve_idx(name).ok()
    }

    /// Sets an argument by index without checks of any kind.
    ///
    /// Setting buffer or image (`cl_mem`) arguments this way may cause
    /// segfaults or errors if the buffer goes out of scope at any point
    /// before this kernel is dropped.
    ///
    /// This also bypasses the check to determine if the type of the value you
    /// pass here matches the type defined in your kernel.
    pub unsafe fn set_arg_unchecked(&self, arg_idx: u32, arg_val: ArgVal) -> OclResult<()> {
        core::set_kernel_arg(&self.obj_core, arg_idx, arg_val).map_err(OclError::from)
    }

    /// Sets an argument by index.
    fn _set_arg<T: OclPrm>(&self, arg_idx: u32, arg_val: ArgVal) -> OclResult<()> {
        self.verify_arg_type::<T>(arg_idx)?;
        core::set_kernel_arg(&self.obj_core, arg_idx, arg_val).map_err(OclError::from)
    }

    /// Sets a `Buffer`, `Image`, scalar, or vector argument by index or by
    /// name.
    ///
    /// ### Example
    /// ```rust,ignore
    /// // Create a kernel with arguments corresponding to those in the kernel.
    /// // Just for fun, one argument will be 'named':
    /// let kern = ocl_pq.kernel_builder("multiply_by_scalar")
    ///     .arg(&0)
    ///     .arg(None::<&Buffer<f32>>)
    ///     .arg_named("result", None::<&Buffer<f32>>)
    ///     .build()?;
    ///
    /// // Set our named argument. The Option<_> wrapper is, well... optional:
    /// kern.set_arg("result", &result_buffer)?;
    /// // We can also set arguments (named or not) by index. Just for
    /// // demonstration, we'll set one using an option:
    /// kern.set_arg(0, &COEFF)?;
    /// kern.set_arg(1, Some(&source_buffer))?;
    /// kern.set_arg(2, &result_buffer)?;
    /// ```
    pub fn set_arg<'a, T, Ai, Av>(&self, idx: Ai, arg: Av) -> OclResult<()>
            where T: OclPrm, Ai: Into<ArgIdxSpecifier>, Av: Into<ArgValConverter<'a, T>> {
        let arg_idx = idx.into().to_idx(&self.named_args)?;
        self.verify_arg_type::<T>(arg_idx)?;
        let arg: ArgValConverter<T> = arg.into();

        // If the `KernelArg` is a `Mem` variant, clone the `MemCore` it
        // refers to, store it in `self.mem_args`. This prevents a buffer
        // which has gone out of scope from being erroneously referred to when
        // this kernel is enqueued and causing either a misleading error
        // message or a hard to debug segfault depending on the platform.
        match arg.mem {
            Some(mem) => self.mem_args.insert(arg_idx, mem),
            None => self.mem_args.remove(&arg_idx),
        };

        let val = arg.val.to_arg_val();
        self._set_arg::<T>(arg_idx, val)
    }

    /// Modifies the kernel argument named: `name`.
    #[deprecated(since = "0.18.0", note = "Use `::set_arg` instead.")]
    pub fn set_arg_buf_named<'a, T, M>(&'a self, name: &'static str,
            buffer_opt: Option<M>) -> OclResult<()>
            where T: OclPrm, M: AsMem<T> + MemCmdAll {
        let arg_idx = self.named_args.resolve_idx(name)?;
        match buffer_opt {
            Some(buffer) => {
                self.mem_args.insert(arg_idx, buffer.as_mem().clone());
                self._set_arg::<T>(arg_idx, ArgVal::mem(buffer.as_mem()))
            },
            None => {
                self.mem_args.remove(&arg_idx);
                self._set_arg::<T>(arg_idx, ArgVal::mem_null())
            },
        }
    }

    /// Modifies the kernel argument named: `name`.
    #[deprecated(since = "0.18.0", note = "Use `::set_arg` instead.")]
    pub fn set_arg_img_named<'a, T, M>(&'a self, name: &'static str,
            image_opt: Option<M>) -> OclResult<()>
            where T: OclPrm, M: AsMem<T> + MemCmdAll {
        let arg_idx = self.named_args.resolve_idx(name)?;
        match image_opt {
            Some(img) => {
                self.mem_args.insert(arg_idx, img.as_mem().clone());
                self._set_arg::<T>(arg_idx, ArgVal::mem(img.as_mem()))
            },
            None => {
                self.mem_args.remove(&arg_idx);
                self._set_arg::<T>(arg_idx, ArgVal::mem_null())
            },
        }
    }

    /// Sets the value of a named sampler argument.
    #[deprecated(since = "0.18.0", note = "Use `::set_arg_sampler_named` instead.")]
    pub fn set_arg_smp_named<'a>(&'a self, name: &'static str,
            sampler_opt: Option<&Sampler>) -> OclResult<()> {
        let arg_idx = self.named_args.resolve_idx(name)?;
        match sampler_opt {
            Some(sampler) => self._set_arg::<u64>(arg_idx, ArgVal::sampler(sampler)),
            None => self._set_arg::<u64>(arg_idx, ArgVal::sampler_null()),
        }
    }

    /// Modifies the kernel argument named: `name`.
    #[deprecated(since = "0.18.0", note = "Use `::set_arg` instead.")]
    pub fn set_arg_scl_named<'a, T, B>(&'a self, name: &'static str, scalar: B)
            -> OclResult<()>
            where T: OclPrm, B: Borrow<T> {
        let arg_idx = self.named_args.resolve_idx(name)?;
        self._set_arg::<T>(arg_idx, ArgVal::scalar(scalar.borrow()))
    }

    /// Modifies the kernel argument named: `name`.
    #[deprecated(since = "0.18.0", note = "Use `::set_arg` instead.")]
    pub fn set_arg_vec_named<'a, T, B>(&'a self, name: &'static str, vector: B)
            -> OclResult<()>
            where T: OclPrm, B: Borrow<T> {
        let arg_idx = self.named_args.resolve_idx(name)?;
        self._set_arg::<T>(arg_idx, ArgVal::vector(vector.borrow()))
    }

    /// Sets the value of a named sampler argument.
    pub fn set_arg_sampler_named<'a, Ai>(&'a self, idx: Ai,
            sampler_opt: Option<&Sampler>) -> OclResult<()>
            where Ai: Into<ArgIdxSpecifier> {
        // let arg_idx = self.named_args.resolve_idx(name)?;
        let arg_idx = idx.into().to_idx(&self.named_args)?;
        match sampler_opt {
            Some(sampler) => self._set_arg::<u64>(arg_idx, ArgVal::sampler(sampler)),
            None => self._set_arg::<u64>(arg_idx, ArgVal::sampler_null()),
        }
    }

    /// Returns a command builder which is used to chain parameters of an
    /// 'enqueue' command together.
    pub fn cmd(&self) -> KernelCmd {
        KernelCmd {
            kernel: &self.obj_core,
            queue: self.queue.as_ref().map(|q| q.as_ref()),
            gwo: self.gwo,
            gws: self.gws,
            lws: self.lws,
            wait_events: None,
            new_event: None
        }
    }

    /// Enqueues this kernel on the default queue using the default work sizes
    /// and offsets.
    ///
    /// Shorthand for `.cmd().enq()`
    ///
    /// # Safety
    ///
    /// All kernel code must be considered untrusted. Therefore the act of
    /// calling this function contains implied unsafety even though the API
    /// itself is safe.
    pub unsafe fn enq(&self) -> OclResult<()> {
        self.cmd().enq()
    }

    /// Changes the default queue.
    ///
    /// Returns a ref for chaining i.e.:
    ///
    /// `kernel.set_default_queue(queue).enqueue(....);`
    ///
    /// Even when used as above, the queue is changed permanently,
    /// not just for the one call. Changing the queue is cheap so feel free
    /// to change as often as needed.
    ///
    /// If you want to change the queue for only a single call, use:
    /// `::cmd.queue(...)...enq()...`
    ///
    /// The new queue must be associated with a device associated with the
    /// kernel's program.
    ///
    pub fn set_default_queue(&mut self, queue: Queue) -> &mut Kernel {
        self.queue = Some(queue);
        self
    }

    /// Returns the default global work offset.
    #[deprecated(since = "0.18.0", note = "Use `::global_work_offset` instead.")]
    pub fn get_gwo(&self) -> SpatialDims {
        self.gwo
    }

    /// Returns the default global work size.
    #[deprecated(since = "0.18.0", note = "Use `::global_work_size` instead.")]
    pub fn get_gws(&self) -> SpatialDims {
        self.gws
    }

    /// Returns the default local work size.
    #[deprecated(since = "0.18.0", note = "Use `::local_work_size` instead.")]
    pub fn get_lws(&self) -> SpatialDims {
        self.lws
    }

    /// Sets the default global work offset.
    pub fn set_default_global_work_offset(&mut self, gwo: SpatialDims) -> &mut Kernel {
        self.gwo = gwo;
        self
    }

    /// Sets the default global work size.
    pub fn set_default_global_work_size(&mut self, gws: SpatialDims) -> &mut Kernel {
        self.gws = gws;
        self
    }

    /// Sets the default local work size.
    pub fn set_default_local_work_size(&mut self, lws: SpatialDims) -> &mut Kernel {
        self.lws = lws;
        self
    }

    /// Returns the default queue for this kernel if one has been set.
    pub fn default_queue<'a>(&'a self) -> Option<&'a Queue> {
        self.queue.as_ref()
    }

    /// Returns the default global work offset.
    pub fn default_global_work_offset(&self) -> SpatialDims {
        self.gwo
    }

    /// Returns the default global work size.
    pub fn default_global_work_size(&self) -> SpatialDims {
        self.gws
    }

    /// Returns the default local work size.
    pub fn default_local_work_size(&self) -> SpatialDims {
        self.lws
    }

    /// Returns a reference to the core pointer wrapper, usable by functions in
    /// the `core` module.
    #[inline]
    pub fn as_core<'a>(&'a self) -> &'a KernelCore {
        self
    }

    /// Returns information about this kernel.
    pub fn info(&self, info_kind: KernelInfo) -> OclCoreResult<KernelInfoResult> {
        core::get_kernel_info(&self.obj_core, info_kind)
    }

    /// Returns work group information for this kernel.
    pub fn wg_info(&self, device: Device, info_kind: KernelWorkGroupInfo)
            -> OclCoreResult<KernelWorkGroupInfoResult> {
        core::get_kernel_work_group_info(&self.obj_core, device, info_kind)
    }

    /// Returns argument information for this kernel.
    pub fn arg_info(&self, arg_idx: u32, info_kind: KernelArgInfo)
            -> OclCoreResult<KernelArgInfoResult> {
        arg_info(&*self.as_core(), arg_idx, info_kind)
    }

    /// Returns the name of this kernel.
    pub fn name(&self) -> OclCoreResult<String> {
        core::get_kernel_info(&self.obj_core, KernelInfo::FunctionName).map(|r| r.into())
    }

    /// Returns the number of arguments this kernel has.
    pub fn num_args(&self) -> OclCoreResult<u32> {
        match core::get_kernel_info(&self.obj_core, KernelInfo::NumArgs) {
            Ok(KernelInfoResult::NumArgs(num)) => Ok(num),
            Err(err) => Err(err),
            _=> unreachable!(),
        }
    }

    fn fmt_info(&self, f: &mut std::fmt::Formatter) -> std::fmt::Result {
        f.debug_struct("Kernel")
            .field("FunctionName", &self.info(KernelInfo::FunctionName))
            .field("ReferenceCount", &self.info(KernelInfo::ReferenceCount))
            .field("Context", &self.info(KernelInfo::Context))
            .field("Program", &self.info(KernelInfo::Program))
            .field("Attributes", &self.info(KernelInfo::Attributes))
            .finish()
    }

    // fn fmt_arg_info(&self, f: &mut std::fmt::Formatter) -> std::fmt::Result {
    //     f.debug_struct("Kernel")
    //         .field("FunctionName", &self.info(KernelInfo::FunctionName))
    //         .field("ReferenceCount", &self.info(KernelInfo::ReferenceCount))
    //         .field("Context", &self.info(KernelInfo::Context))
    //         .field("Program", &self.info(KernelInfo::Program))
    //         .field("Attributes", &self.info(KernelInfo::Attributes))
    //         .finish()
    // }

    fn fmt_wg_info<D>(&self, f: &mut std::fmt::Formatter, devices: Vec<D>) -> std::fmt::Result
            where D: Into<Device> {
        for device in devices {
            let device = device.into();
            if !device.vendor().unwrap()
                    .contains("NVIDIA") {
                f.debug_struct("WorkGroup")
                    .field("WorkGroupSize", &self.wg_info(device, KernelWorkGroupInfo::WorkGroupSize))
                    .field("CompileWorkGroupSize", &self.wg_info(device, KernelWorkGroupInfo::CompileWorkGroupSize))
                    .field("LocalMemSize", &self.wg_info(device, KernelWorkGroupInfo::LocalMemSize))
                    .field("PreferredWorkGroupSizeMultiple",
                        &self.wg_info(device, KernelWorkGroupInfo::PreferredWorkGroupSizeMultiple))
                    .field("PrivateMemSize", &self.wg_info(device, KernelWorkGroupInfo::PrivateMemSize))
                    // .field("GlobalWorkSize", &self.wg_info(device, KernelWorkGroupInfo::GlobalWorkSize))
                    .finish()?;
            }
        }
        Ok(())
    }
}

impl std::fmt::Display for Kernel {
    fn fmt(&self, f: &mut std::fmt::Formatter) -> std::fmt::Result {
        self.fmt_info(f)?;
        write!(f, " ")?;
        self.fmt_wg_info(f, self.obj_core.devices().unwrap())
    }
}

impl Deref for Kernel {
    type Target = KernelCore;

    fn deref<'a>(&'a self) -> &'a KernelCore {
        &self.obj_core
    }
}

impl Clone for Kernel {
    /// Creates a new, identical kernel.
    fn clone(&self) -> Kernel {
        unimplemented!();
    }
}


/// A kernel builder.
///
///
/// ### Examples
///
/// ```rust,ignore
/// // Create a kernel:
/// let kernel = Kernel::builder()
///     .program(&program)
///     .name("process")
///     .queue(queue.clone())
///     .global_work_size(&gws_patch_count)
///     .arg(&(patch_size as i32))
///     .arg(&source_buffer)
///     .arg(&destination_buffer)
///     .build()?;
/// ```
///
/// Re-use and clone:
///
/// ```rust,ignore
/// // Create a builder for re-use:
/// let builder = Kernel::builder()
///     .program(&program)
///     .name("process")
///     .queue(queue.clone())
///     .global_work_size([512, 64, 64])
///     .arg(&(patch_size as i32))
///     .arg(&source_buffer)
///     .arg_named("dest", &destination_buffer);
///
/// // Create multiple kernels using the same builder:
/// let kernel_0 = builder.build()?;
/// let kernel_1 = builder.build()?;
///
/// // Clone the builder:
/// let mut builder_clone = builder.clone();
///
/// // Modify a default parameter:
/// builder_clone.global_work_size([1024, 128, 128]);
///
/// // Create another kernel using the cloned builder:
/// let kernel_2 = builder_clone.build()?;
///
/// // Modify one of the arguments using the created kernel directly
/// // (arguments cannot be modified using the builder):
/// kernel_2.set_arg("dest", &alternate_destination_buffer)?;
///
/// // Arguments can also be referred to by index:
/// kernel_2.set_arg(1, &alternate_source_buffer)?;
/// ```
///
#[derive(Debug)]
pub struct KernelBuilder<'b> {
    program: Option<&'b Program>,
    name: Option<String>,
    named_args: NamedArgs,
    mem_args: MemArgs,
    args: Vec<(ArgValKeeper<'b>, Option<TypeId>)>,
    queue: Option<Queue>,
    gwo: SpatialDims,
    gws: SpatialDims,
    lws: SpatialDims,
    disable_arg_check: bool,
}

impl<'b> KernelBuilder<'b> {
    /// Returns a new kernel builder.
    pub fn new() -> KernelBuilder<'b> {
        KernelBuilder {
            program: None,
            name: None,
            named_args: NamedArgs(None),
            args: Vec::with_capacity(16),
            mem_args: MemArgs(Some(RefCell::new(BTreeMap::new()))),
            queue: None,
            gwo: SpatialDims::Unspecified,
            gws: SpatialDims::Unspecified,
            lws: SpatialDims::Unspecified,
            disable_arg_check: false,
        }
    }

    /// Specifies a program object with a successfully built executable.
    pub fn program<'s>(&'s mut self, program: &'b Program) -> &'s mut KernelBuilder<'b> {
        self.program = Some(program);
        self
    }

    /// Specifies a function name in the program declared with the `__kernel`
    /// qualifier (e.g. `__kernel void add_values(...`).
    pub fn name<'s, S>(&'s mut self, name: S) -> &'s mut KernelBuilder<'b>
            where S: Into<String> {
        self.name = Some(name.into());
        self
    }

    /// Sets the default queue to be used by all subsequent enqueue commands
    /// unless otherwise changed (with `::set_default_queue`) or overridden
    /// (by `::cmd().queue(...)...`).
    ///
    /// The queue must be associated with a device associated with the
    /// kernel's program.
    pub fn queue<'s>(&'s mut self, queue: Queue) -> &'s mut KernelBuilder<'b> {
        self.queue = Some(queue);
        self
    }

    /// Sets the default global work offset.
    ///
    /// Used when enqueuing kernel commands. Superseded if specified while
    /// building a queue command with `::cmd`.
    #[deprecated(since = "0.18.0", note = "Use `::global_work_offset` instead.")]
    pub fn gwo<'s, D: Into<SpatialDims>>(&'s mut self, gwo: D) -> &'s mut KernelBuilder<'b> {
        self.gwo = gwo.into();
        self
    }

    /// Sets the default global work size.
    ///
    /// Used when enqueuing kernel commands. Superseded if specified while
    /// building a queue command with `::cmd`.
    #[deprecated(since = "0.18.0", note = "Use `::global_work_size` instead.")]
    pub fn gws<'s, D: Into<SpatialDims>>(&'s mut self, gws: D) -> &'s mut KernelBuilder<'b> {
        self.gws = gws.into();
        self
    }

    /// Sets the default local work size.
    ///
    /// Used when enqueuing kernel commands. Superseded if specified while
    /// building a queue command with `::cmd`.
    #[deprecated(since = "0.18.0", note = "Use `::local_work_size` instead.")]
    pub fn lws<'s, D: Into<SpatialDims>>(&'s mut self, lws: D) -> &'s mut KernelBuilder<'b> {
        self.lws = lws.into();
        self
    }

    /// Sets the default global work offset.
    ///
    /// Used when enqueuing kernel commands. Superseded if specified while
    /// building a queue command with `::cmd`.
    pub fn global_work_offset<'s, D: Into<SpatialDims>>(&'s mut self, gwo: D) -> &'s mut KernelBuilder<'b> {
        self.gwo = gwo.into();
        self
    }

    /// Sets the default global work size.
    ///
    /// Used when enqueuing kernel commands. Superseded if specified while
    /// building a queue command with `::cmd`.
    pub fn global_work_size<'s, D: Into<SpatialDims>>(&'s mut self, gws: D) -> &'s mut KernelBuilder<'b> {
        self.gws = gws.into();
        self
    }

    /// Sets the default local work size.
    ///
    /// Used when enqueuing kernel commands. Superseded if specified while
    /// building a queue command with `::cmd`.
    pub fn local_work_size<'s, D: Into<SpatialDims>>(&'s mut self, lws: D) -> &'s mut KernelBuilder<'b> {
        self.lws = lws.into();
        self
    }

    /// Adds a new argument to the kernel and returns the index.
    fn new_arg(&mut self, arg_val: ArgValKeeper<'b>, type_id: Option<TypeId>, mem: Option<MemCore>) -> u32 {
        let arg_idx = self.args.len() as u32;

        // If the `KernelArg` is a `Mem` variant, clone the `MemCore` it
        // refers to, store it in `self.mem_args`. This prevents a buffer
        // which has gone out of scope from being erroneously referred to when
        // this kernel is enqueued and causing either a misleading error
        // message or a hard to debug segfault depending on the platform.
        match mem {
            Some(mem) => self.mem_args.insert(arg_idx, mem),
            None => self.mem_args.remove(&arg_idx),
        };

        self.args.push((arg_val, type_id));
        arg_idx
    }

    /// Non-builder-style version of `::arg_buf()`.
    fn new_arg_buf<T, M>(&mut self, buffer_opt: Option<&'b M>) -> u32
            where T: OclPrm, M: 'b + AsMem<T> + MemCmdAll {
        match buffer_opt {
            Some(buffer) => {
                self.new_arg(ArgValKeeper::Shared(ArgVal::mem(buffer.as_mem())),
                    Some(TypeId::of::<T>()),
                    Some(buffer.as_mem().clone()))
            },
            None => {
                self.new_arg(ArgValKeeper::mem_null(),
                    Some(TypeId::of::<T>()), None)
            },
        }
    }

    /// Non-builder-style version of `::arg_img()`.
    fn new_arg_img<T, M>(&mut self, image_opt: Option<&'b M>) -> u32
            where T: OclPrm, M: 'b + AsMem<T> + MemCmdAll {
        match image_opt {
            Some(image) => {
                // Type is ignored:
                self.new_arg(ArgValKeeper::Shared(ArgVal::mem(image.as_mem())), None,
                    Some(image.as_mem().clone()))
            },
            None => {
                self.new_arg(ArgValKeeper::mem_null(), None, None)
            },
        }
    }

    /// Non-builder-style version of `::arg_img()`.
    fn new_arg_smp(&mut self, sampler_opt: Option<&'b Sampler>) -> u32 {
        match sampler_opt {
            Some(sampler) => {
                // Type is ignored:
                self.new_arg(ArgValKeeper::Shared(ArgVal::sampler(sampler)), None, None)
            },
            None => {
                self.new_arg(ArgValKeeper::Shared(ArgVal::sampler_null()), None, None)
            },
        }
    }

    /// Non-builder-style version of `::arg_scl()`.
    fn new_arg_scl<T>(&mut self, scalar: T) -> u32
            where T: OclPrm {
        self.new_arg(ArgValKeeper::owned(scalar), Some(TypeId::of::<T>()), None)
    }

    /// Non-builder-style version of `::arg_vec()`.
    fn new_arg_vec<T>(&mut self, vector: T) -> u32
            where T: OclPrm {
        self.new_arg(ArgValKeeper::owned(vector), Some(TypeId::of::<T>()), None)
    }

    /// Non-builder-style version of `::arg_loc()`.
    fn new_arg_loc<T>(&mut self, length: usize) -> u32
            where T: OclPrm {
        self.new_arg(ArgValKeeper::Shared(ArgVal::local::<T>(&length)), None, None)
    }

    /// Adds a new `Buffer`, `Image`, scalar, or vector argument to the
    /// kernel.
    ///
    /// The argument is added to the bottom of the argument order.
    ///
    /// ### Example
    ///
    /// ```rust,ignore
    /// let kern = ocl_pq.kernel_builder("multiply_by_scalar")
    ///     .arg(&100.0f32)
    ///     .arg(&source_buffer)
    ///     .arg(&result_buffer)
    ///     .build()?;
    /// ```
    pub fn arg<'s, T, A>(&'s mut self, arg: A) -> &'s mut KernelBuilder<'b>
            where T: OclPrm, A: Into<ArgValConverter<'b, T>> {
        let arg = arg.into();
        self.new_arg(arg.val, arg.type_id, arg.mem);
        self
    }

    /// Adds a new argument to the kernel specifying the buffer object represented
    /// by 'buffer'.
    ///
    /// The argument is added to the bottom of the argument order.
    #[deprecated(since = "0.18", note = "Use ::arg instead.")]
    pub fn arg_buf<'s, T, M>(&'s mut self, buffer: &'b M) -> &'s mut KernelBuilder<'b>
            where T: OclPrm, M: 'b + AsMem<T> + MemCmdAll {
        self.new_arg_buf::<T, _>(Some(buffer));
        self
    }

    /// Adds a new argument to the kernel specifying the image object represented
    /// by 'image'.
    ///
    /// The argument is added to the bottom of the argument order.
    #[deprecated(since = "0.18", note = "Use ::arg instead.")]
    pub fn arg_img<'s, T, M>(&'s mut self, image: &'b M) -> &'s mut KernelBuilder<'b>
            where T: OclPrm, M: 'b + AsMem<T> + MemCmdAll {
        self.new_arg_img::<T, _>(Some(image));
        self
    }

    /// Adds a new argument to the kernel specifying the sampler object represented
    /// by 'sampler'. Argument is added to the bottom of the argument
    /// order.
    #[deprecated(since = "0.18", note = "Use ::arg_sampler instead.")]
    pub fn arg_smp<'s>(&'s mut self, sampler: &'b Sampler) -> &'s mut KernelBuilder<'b> {
        self.new_arg_smp(Some(sampler));
        self
    }

    /// Adds a new argument specifying the value: `scalar`.
    ///
    /// The argument is added to the bottom of the argument order.
    #[deprecated(since = "0.18", note = "Use ::arg instead.")]
    pub fn arg_scl<'s, T>(&'s mut self, scalar: T) -> &'s mut KernelBuilder<'b>
            where T: OclPrm {
        self.new_arg_scl(scalar);
        self
    }

    /// Adds a new argument specifying the value: `vector`.
    ///
    /// The argument is added to the bottom of the argument order.
    #[deprecated(since = "0.18", note = "Use ::arg instead.")]
    pub fn arg_vec<'s, T>(&'s mut self, vector: T) -> &'s mut KernelBuilder<'b>
            where T: OclPrm {
        self.new_arg_vec(vector);
        self
    }

    /// Adds a new argument specifying the allocation of a local variable of size
    /// `length * sizeof(T)` bytes (builder_style).
    ///
    /// The argument is added to the bottom of the argument order.
    ///
    /// Local variables are used to share data between work items in the same
    /// workgroup.
    #[deprecated(since = "0.18", note = "Use ::arg_local instead.")]
    pub fn arg_loc<'s, T>(&'s mut self, length: usize) -> &'s mut KernelBuilder<'b>
            where T: OclPrm {
        self.new_arg_loc::<T>(length);
        self
    }

    /// Adds a new argument to the kernel specifying the sampler object represented
    /// by 'sampler'. Argument is added to the bottom of the argument
    /// order.
    pub fn arg_sampler<'s>(&'s mut self, sampler: &'b Sampler) -> &'s mut KernelBuilder<'b> {
        self.new_arg_smp(Some(sampler));
        self
    }

    /// Adds a new argument specifying the allocation of a local variable of size
    /// `length * sizeof(T)` bytes (builder_style).
    ///
    /// The argument is added to the bottom of the argument order.
    ///
    /// Local variables are used to share data between work items in the same
    /// workgroup.
    pub fn arg_local<'s, T>(&'s mut self, length: usize) -> &'s mut KernelBuilder<'b>
            where T: OclPrm {
        self.new_arg_loc::<T>(length);
        self
    }

    /// Adds a new *named* `Buffer`, `Image`, scalar, or vector argument to the
    /// kernel.
    ///
    /// The argument is added to the bottom of the argument order.
    ///
    /// To set a `Buffer` or `Image` argument to `None` (null), you must use a
    /// type annotation (e.g. `None::<&Buffer<f32>>`). Scalar and vector
    /// arguments may not be null; use zero (e.g. `&0`) instead.
    ///
    /// Named arguments can be modified later using `::set_arg()`.
    ///
    /// ### Example
    ///
    /// ```rust,ignore
    /// // Create a kernel with arguments corresponding to those in the kernel.
    /// // One argument will be 'named':
    /// let kern = ocl_pq.kernel_builder("multiply_by_scalar")
    ///     .arg(&COEFF)
    ///     .arg(&source_buffer)
    ///     .arg_named("result", None::<&Buffer<f32>>)
    ///     .build()?;
    ///
    /// // Set our named argument. The Option<_> wrapper is, well... optional:
    /// kern.set_arg("result", &result_buffer)?;
    /// // We can also set arguments (named or not) by index:
    /// kern.set_arg(2, &result_buffer)?;
    /// ```
    pub fn arg_named<'s, T, A>(&'s mut self, name: &'static str, arg: A) -> &'s mut KernelBuilder<'b>
            where T: OclPrm, A: Into<ArgValConverter<'b, T>> {
        let arg = arg.into();
        let arg_idx = self.new_arg(arg.val, arg.type_id, arg.mem);
        self.named_args.insert(name, arg_idx);
        self
    }

    /// Adds a new *named* argument specifying the buffer object represented by
    /// 'buffer'.
    ///
    /// The argument is added to the bottom of the argument order.
    ///
    /// Named arguments can be easily modified later using `::set_arg_buf_named()`.
    #[deprecated(since = "0.18", note = "Use ::arg_named instead.")]
    pub fn arg_buf_named<'s, T, M>(&'s mut self, name: &'static str, buffer_opt: Option<&'b M>) -> &'s mut KernelBuilder<'b>
            where T: OclPrm, M: 'b + AsMem<T> + MemCmdAll {
        let arg_idx = self.new_arg_buf::<T, _>(buffer_opt);
        self.named_args.insert(name, arg_idx);
        self
    }

    /// Adds a new *named* argument specifying the image object represented by
    /// 'image'.
    ///
    /// The argument is added to the bottom of the argument order.
    ///
    /// Named arguments can be easily modified later using `::set_arg_img_named()`.
    #[deprecated(since = "0.18", note = "Use ::arg_named instead.")]
    pub fn arg_img_named<'s, T, M>(&'s mut self, name: &'static str, image_opt: Option<&'b M>) -> &'s mut KernelBuilder<'b>
            where T: OclPrm, M: 'b + AsMem<T> + MemCmdAll {
        let arg_idx = self.new_arg_img::<T, _>(image_opt);
        self.named_args.insert(name, arg_idx);
        self
    }

    /// Adds a new *named* argument specifying the sampler object represented by
    /// 'sampler'.
    ///
    /// The argument is added to the bottom of the argument order.
    ///
    /// Named arguments can be easily modified later using `::set_arg_smp_named()`.
    #[deprecated(since = "0.18", note = "Use ::arg_sampler_named instead.")]
    pub fn arg_smp_named<'s>(&'s mut self, name: &'static str, sampler_opt: Option<&'b Sampler>) -> &'s mut KernelBuilder<'b> {
        let arg_idx = self.new_arg_smp(sampler_opt);
        self.named_args.insert(name, arg_idx);
        self
    }

    /// Adds a new *named* argument  specifying the value: `scalar`.
    ///
    /// The argument is added to the bottom of the argument order.
    ///
    /// Scalar arguments may not be null, use zero (e.g. `&0`) instead.
    ///
    /// Named arguments can be easily modified later using `::set_arg_scl_named()`.
    #[deprecated(since = "0.18", note = "Use ::arg_named instead.")]
    pub fn arg_scl_named<'s, T>(&'s mut self, name: &'static str, scalar: T) -> &'s mut KernelBuilder<'b>
            where T: OclPrm {
        let arg_idx = self.new_arg_scl(scalar);
        self.named_args.insert(name, arg_idx);
        self
    }

    /// Adds a new *named* argument specifying the value: `vector`.
    ///
    /// The argument is added to the bottom of the argument order.
    ///
    /// Vector arguments may not be null, use zero (e.g. `&0`) instead.
    ///
    /// Named arguments can be easily modified later using `::set_arg_vec_named()`.
    #[deprecated(since = "0.18", note = "Use ::arg_named instead.")]
    pub fn arg_vec_named<'s, T>(&'s mut self, name: &'static str, vector: T) -> &'s mut KernelBuilder<'b>
            where T: OclPrm {
        let arg_idx = self.new_arg_vec(vector);
        self.named_args.insert(name, arg_idx);
        self
    }

    /// Adds a new *named* argument specifying the sampler object represented by
    /// 'sampler'.
    ///
    /// The argument is added to the bottom of the argument order.
    ///
    /// Named arguments can be easily modified later using `::set_arg_smp_named()`.
    pub fn arg_sampler_named<'s>(&'s mut self, name: &'static str, sampler_opt: Option<&'b Sampler>) -> &'s mut KernelBuilder<'b> {
        let arg_idx = self.new_arg_smp(sampler_opt);
        self.named_args.insert(name, arg_idx);
        self
    }

    /// Specifies whether or not to store a copy of memory objects (`Buffer`
    /// and `Image`).
    ///
    /// ### Safety
    ///
    /// Disabling memory object argument retention can lead to a misleading
    /// error message or a difficult to debug segfault (depending on the
    /// platform) *if* a memory object is dropped before a kernel referring to
    /// it is enqueued. Only disable this if you are certain all of the memory
    /// objects set as kernel arguments will outlive the `Kernel` itself.
    pub unsafe fn disable_mem_arg_retention<'s>(&'s mut self) -> &'s mut KernelBuilder<'b> {
        self.mem_args = MemArgs(None);
        self
    }

    /// Disables argument type checking when setting arguments.
    ///
    /// Because the performance cost of argument type checking is negligible,
    /// disabling is not recommended.
    ///
    /// Argument type checking will automatically be disabled if any of the
    /// devices in use do not support OpenCL version 1.2 or higher or if
    /// argument information is not available on the associated platform.
    pub unsafe fn disable_arg_type_check<'s>(&'s mut self) -> &'s mut KernelBuilder<'b> {
        self.disable_arg_check = true;
        self
    }

    /// Builds and returns a new `Kernel`
    pub fn build(&self) -> OclResult<Kernel> {
        let program = self.program.ok_or(KernelError::BuilderNoProgram)?;
        let name = self.name.as_ref().ok_or(KernelError::BuilderNoKernelName)?;

        let obj_core = core::create_kernel(program, name)?;

        let num_args = match core::get_kernel_info(&obj_core, KernelInfo::NumArgs) {
            Ok(KernelInfoResult::NumArgs(num)) => num,
            Err(err) => return Err(OclError::from(err)),
            _=> unreachable!(),
        };

        if self.args.len() as u32 != num_args {
            return Err(KernelError::BuilderWrongArgCount {
                required: num_args,
                specified: self.args.len() as u32,
             }.into())
        }

        let mut arg_types = Vec::with_capacity(num_args as usize);
        let mut all_arg_types_unknown = true;
        let mut disable_arg_check = false;

        // Cache argument types for later use, bypassing if the OpenCL version
        // is too low (v1.1).
        for arg_idx in 0..num_args {
            let arg_type = match ArgType::from_kern_and_idx(&obj_core, arg_idx) {
                Ok(at) => {
                    if !at.is_unknown() { all_arg_types_unknown = false; }
                    at
                },
                Err(e) => {
                    if let OclCoreErrorKind::VersionLow { .. } = *e.kind() {
                        disable_arg_check = true;
                        break;
                    }
                    return Err(OclError::from(e));
                },
            };
            arg_types.push(arg_type);
        }

        // Check argument types then set arguments.
        for (arg_idx, &(ref arg, ref type_id_opt)) in self.args.iter().enumerate() {
            if !disable_arg_check {
                if let Some(type_id) = *type_id_opt {
                    if !arg_types[arg_idx].matches(type_id) {
                        let ty_name = arg_type_name(&obj_core, arg_idx as u32)?;
                        let arg_name = arg_name(&obj_core, arg_idx as u32)?;
                        return Err(KernelError::ArgTypeMismatch { idx: arg_idx as u32,
                            arg_name, ty_name, ty: arg_types[arg_idx].clone() }.into());
                    }
                }
            }

            let val = arg.to_arg_val();
            core::set_kernel_arg(&obj_core, arg_idx as u32, val)?;
        }

        let arg_types = if all_arg_types_unknown || disable_arg_check {
            None
        } else {
            Some(arg_types)
        };

        Ok(Kernel {
            obj_core: obj_core,
            named_args: self.named_args.clone(),
            mem_args: self.mem_args.clone(),
            queue: self.queue.clone(),
            gwo: self.gwo,
            gws: self.gws,
            lws: self.lws,
            arg_types,
        })
    }
}


/// Returns argument information for a kernel.
pub fn arg_info(core: &KernelCore, arg_idx: u32, info_kind: KernelArgInfo)
        -> OclCoreResult<KernelArgInfoResult> {
    let device_versions = match core.device_versions() {
        Ok(vers) => vers,
        Err(e) => return Err(e.into()),
    };

    core::get_kernel_arg_info(core, arg_idx, info_kind,
        Some(&device_versions))
}

/// Returns the type name for a kernel argument at the specified index.
pub fn arg_type_name(core: &KernelCore, arg_idx: u32) -> OclCoreResult<String> {
    match arg_info(core, arg_idx, KernelArgInfo::TypeName) {
        Ok(KernelArgInfoResult::TypeName(type_name)) => Ok(type_name),
        Err(err) => Err(err.into()),
        _ => unreachable!(),
    }
}

/// Returns the type name for a kernel argument at the specified index.
pub fn arg_name(core: &KernelCore, arg_idx: u32) -> OclCoreResult<String> {
    match arg_info(core, arg_idx, KernelArgInfo::Name) {
        Ok(KernelArgInfoResult::Name(name)) => Ok(name),
        Err(err) => Err(err.into()),
        _ => unreachable!(),
    }
}


pub mod arg_type {
    #![allow(unused_imports)]
    use std::any::{Any, TypeId};
    use ffi::{cl_char, cl_uchar, cl_short, cl_ushort, cl_int, cl_uint, cl_long, cl_ulong,
        cl_half, cl_float, cl_double, cl_bool, cl_bitfield};
    use core::{Error as OclCoreError, Result as OclCoreResult, Status, OclPrm, Kernel as KernelCore};
    use error::{Error as OclError, Result as OclResult};
    use standard::Sampler;
    use super::{arg_info, arg_type_name};

    pub use core::{
        Char, Char2, Char3, Char4, Char8, Char16,
        Uchar, Uchar2, Uchar3, Uchar4, Uchar8, Uchar16,
        Short, Short2, Short3, Short4, Short8, Short16,
        Ushort, Ushort2, Ushort3, Ushort4, Ushort8, Ushort16,
        Int, Int2, Int3, Int4, Int8, Int16,
        Uint, Uint2, Uint3, Uint4, Uint8, Uint16,
        Long, Long2, Long3, Long4, Long8, Long16,
        Ulong, Ulong2, Ulong3, Ulong4, Ulong8, Ulong16,
        Float, Float2, Float3, Float4, Float8, Float16,
        Double, Double2, Double3, Double4, Double8, Double16};

    /// The base type of an OpenCL primitive.
    #[derive(Clone, Debug, Copy, PartialEq, Eq)]
    pub enum BaseType {
        Char,
        Uchar,
        Short,
        Ushort,
        Int,
        Uint,
        Long,
        Ulong,
        Float,
        Double,
        Sampler,
        Image,
        Unknown,
    }

    /// The cardinality of an OpenCL primitive.
    ///
    /// (ex. `float4`: 4)
    #[derive(Clone, Debug, Copy, PartialEq, Eq)]
    pub enum Cardinality {
        One,
        Two,
        Three,
        Four,
        Eight,
        Sixteen,
    }

    /// The type of a kernel argument derived from its string representation.
    #[allow(dead_code)]
    #[derive(Clone, Debug)]
    pub struct ArgType {
        base_type: BaseType,
        cardinality: Cardinality,
        is_ptr: bool,
    }

    impl ArgType {
        /// Returns an `ArgType` that will always return `true` when calling
        /// `::is_match`.
        pub fn unknown() -> OclCoreResult<ArgType> {
            Ok(ArgType {
                base_type: BaseType::Unknown,
                cardinality: Cardinality::One,
                is_ptr: false,
            })
        }

        /// Ascertains a `KernelArgType` from the contents of a
        /// `KernelArgInfoResult::TypeName`.
        ///
        /// * TODO: Optimize or outsource this if possible. Is `::contains`
        /// the fastest way to parse these in this situation? Should
        /// `::starts_with` be used for base type names instead?
        pub fn from_str(type_name: &str) -> OclCoreResult<ArgType> {
            let is_ptr = type_name.contains("*");

            let card = if type_name.contains("16") {
                Cardinality::Sixteen
            } else if type_name.contains("8") {
                Cardinality::Eight
            } else if type_name.contains("4") {
                Cardinality::Four
            } else if type_name.contains("3") {
                Cardinality::Three
            } else if type_name.contains("2") {
                Cardinality::Two
            } else {
                Cardinality::One
            };

            let base = if type_name.contains("uchar") {
                BaseType::Uchar
            } else if type_name.contains("char") {
                BaseType::Char
            } else if type_name.contains("ushort") {
                BaseType::Ushort
            } else if type_name.contains("short") {
                BaseType::Short
            } else if type_name.contains("uint") {
                BaseType::Uint
            } else if type_name.contains("int") {
                BaseType::Int
            } else if type_name.contains("ulong") {
                BaseType::Ulong
            } else if type_name.contains("long") {
                BaseType::Long
            } else if type_name.contains("float") {
                BaseType::Float
            } else if type_name.contains("double") {
                BaseType::Double
            } else if type_name.contains("sampler") {
                BaseType::Sampler
            } else if type_name.contains("image") {
                BaseType::Image
            } else {
                BaseType::Unknown
            };

            Ok(ArgType {
                base_type: base,
                cardinality: card,
                is_ptr: is_ptr,
            })
        }

        /// Returns a new argument type specifier.
        ///
        /// This function calls `core::get_kernel_arg_info`. Some platforms
        /// (Apple, NVIDIA) either do not implement
        /// `core::get_kernel_arg_info` or error in irregular ways. Known
        /// irregular errors are checked for here. The result of a call to
        /// `ArgType::unknown()` (which matches any argument type) is returned
        /// if any are found.
        pub fn from_kern_and_idx(core: &KernelCore, arg_idx: u32) -> OclCoreResult<ArgType> {
            use core::EmptyInfoResultError;
            use core::ErrorKind as OclCoreErrorKind;

            match arg_type_name(core, arg_idx) {
                Ok(type_name) => ArgType::from_str(type_name.as_str()),
                Err(err) => {
                    // Escape hatches for known, platform-specific errors.
                    match *err.kind() {
                        OclCoreErrorKind::Api(ref api_err) => {
                            if api_err.status() == Status::CL_KERNEL_ARG_INFO_NOT_AVAILABLE {
                                return ArgType::unknown();
                            }
                        }
                        OclCoreErrorKind::EmptyInfoResult(EmptyInfoResultError::KernelArg) => {
                            return ArgType::unknown();
                        },
                        _ => (),
                    }

                    Err(err)
                },
            }
        }

        /// Returns true if `type_id` matches the base type of this `ArgType`.
        pub fn matches(&self, type_id: TypeId) -> bool {
            match self.base_type {
                BaseType::Char => {
                    let card_match = match self.cardinality {
                        Cardinality::One => TypeId::of::<cl_char>() == type_id,
                        Cardinality::Two => TypeId::of::<Char2>() == type_id,
                        Cardinality::Three => TypeId::of::<Char3>() == type_id,
                        Cardinality::Four => TypeId::of::<Char4>() == type_id,
                        Cardinality::Eight => TypeId::of::<Char8>() == type_id,
                        Cardinality::Sixteen => TypeId::of::<Char16>() == type_id,
                    };

                    (self.is_ptr && (TypeId::of::<cl_char>() == type_id || card_match)) ||
                        (!self.is_ptr && card_match)
                },
                BaseType::Uchar => {
                    let card_match = match self.cardinality {
                        Cardinality::One => TypeId::of::<cl_uchar>() == type_id,
                        Cardinality::Two => TypeId::of::<Uchar2>() == type_id,
                        Cardinality::Three => TypeId::of::<Uchar3>() == type_id,
                        Cardinality::Four => TypeId::of::<Uchar4>() == type_id,
                        Cardinality::Eight => TypeId::of::<Uchar8>() == type_id,
                        Cardinality::Sixteen => TypeId::of::<Uchar16>() == type_id,
                    };

                    (self.is_ptr && (TypeId::of::<cl_uchar>() == type_id || card_match)) ||
                        (!self.is_ptr && card_match)
                },
                BaseType::Short => {
                    let card_match = match self.cardinality {
                        Cardinality::One => TypeId::of::<cl_short>() == type_id,
                        Cardinality::Two => TypeId::of::<Short2>() == type_id,
                        Cardinality::Three => TypeId::of::<Short3>() == type_id,
                        Cardinality::Four => TypeId::of::<Short4>() == type_id,
                        Cardinality::Eight => TypeId::of::<Short8>() == type_id,
                        Cardinality::Sixteen => TypeId::of::<Short16>() == type_id,
                    };

                    (self.is_ptr && (TypeId::of::<cl_short>() == type_id || card_match)) ||
                        (!self.is_ptr && card_match)
                },
                BaseType::Ushort => {
                    let card_match = match self.cardinality {
                        Cardinality::One => TypeId::of::<cl_ushort>() == type_id,
                        Cardinality::Two => TypeId::of::<Ushort2>() == type_id,
                        Cardinality::Three => TypeId::of::<Ushort3>() == type_id,
                        Cardinality::Four => TypeId::of::<Ushort4>() == type_id,
                        Cardinality::Eight => TypeId::of::<Ushort8>() == type_id,
                        Cardinality::Sixteen => TypeId::of::<Ushort16>() == type_id,
                    };

                    (self.is_ptr && (TypeId::of::<cl_ushort>() == type_id || card_match)) ||
                        (!self.is_ptr && card_match)
                },
                BaseType::Int => {
                    let card_match = match self.cardinality {
                        Cardinality::One => TypeId::of::<cl_int>() == type_id,
                        Cardinality::Two => TypeId::of::<Int2>() == type_id,
                        Cardinality::Three => TypeId::of::<Int3>() == type_id,
                        Cardinality::Four => TypeId::of::<Int4>() == type_id,
                        Cardinality::Eight => TypeId::of::<Int8>() == type_id,
                        Cardinality::Sixteen => TypeId::of::<Int16>() == type_id,
                    };

                    (self.is_ptr && (TypeId::of::<cl_int>() == type_id || card_match)) ||
                        (!self.is_ptr && card_match)
                },
                BaseType::Uint => {
                    let card_match = match self.cardinality {
                        Cardinality::One => TypeId::of::<cl_uint>() == type_id,
                        Cardinality::Two => TypeId::of::<Uint2>() == type_id,
                        Cardinality::Three => TypeId::of::<Uint3>() == type_id,
                        Cardinality::Four => TypeId::of::<Uint4>() == type_id,
                        Cardinality::Eight => TypeId::of::<Uint8>() == type_id,
                        Cardinality::Sixteen => TypeId::of::<Uint16>() == type_id,
                    };

                    (self.is_ptr && (TypeId::of::<cl_uint>() == type_id || card_match)) ||
                        (!self.is_ptr && card_match)
                },
                BaseType::Long => {
                    let card_match = match self.cardinality {
                        Cardinality::One => TypeId::of::<cl_long>() == type_id,
                        Cardinality::Two => TypeId::of::<Long2>() == type_id,
                        Cardinality::Three => TypeId::of::<Long3>() == type_id,
                        Cardinality::Four => TypeId::of::<Long4>() == type_id,
                        Cardinality::Eight => TypeId::of::<Long8>() == type_id,
                        Cardinality::Sixteen => TypeId::of::<Long16>() == type_id,
                    };

                    (self.is_ptr && (TypeId::of::<cl_long>() == type_id || card_match)) ||
                        (!self.is_ptr && card_match)
                },
                BaseType::Ulong => {
                    let card_match = match self.cardinality {
                        Cardinality::One => TypeId::of::<cl_ulong>() == type_id,
                        Cardinality::Two => TypeId::of::<Ulong2>() == type_id,
                        Cardinality::Three => TypeId::of::<Ulong3>() == type_id,
                        Cardinality::Four => TypeId::of::<Ulong4>() == type_id,
                        Cardinality::Eight => TypeId::of::<Ulong8>() == type_id,
                        Cardinality::Sixteen => TypeId::of::<Ulong16>() == type_id,
                    };

                    (self.is_ptr && (TypeId::of::<cl_ulong>() == type_id || card_match)) ||
                        (!self.is_ptr && card_match)
                },
                BaseType::Float => {
                    let card_match = match self.cardinality {
                        Cardinality::One => TypeId::of::<cl_float>() == type_id,
                        Cardinality::Two => TypeId::of::<Float2>() == type_id,
                        Cardinality::Three => TypeId::of::<Float3>() == type_id,
                        Cardinality::Four => TypeId::of::<Float4>() == type_id,
                        Cardinality::Eight => TypeId::of::<Float8>() == type_id,
                        Cardinality::Sixteen => TypeId::of::<Float16>() == type_id,
                    };

                    (self.is_ptr && (TypeId::of::<cl_float>() == type_id || card_match)) ||
                        (!self.is_ptr && card_match)
                },
                BaseType::Double => {
                    let card_match = match self.cardinality {
                        Cardinality::One => TypeId::of::<cl_double>() == type_id,
                        Cardinality::Two => TypeId::of::<Double2>() == type_id,
                        Cardinality::Three => TypeId::of::<Double3>() == type_id,
                        Cardinality::Four => TypeId::of::<Double4>() == type_id,
                        Cardinality::Eight => TypeId::of::<Double8>() == type_id,
                        Cardinality::Sixteen => TypeId::of::<Double16>() == type_id,
                    };

                    (self.is_ptr && (TypeId::of::<cl_double>() == type_id || card_match)) ||
                        (!self.is_ptr && card_match)
                },
                BaseType::Sampler => TypeId::of::<u64>() == type_id,
                BaseType::Image => TypeId::of::<u64>() == type_id,
                // Everything matches if type was undetermined (escape hatch):
                BaseType::Unknown => true,
            }
        }

        /// Returns true if the type of `T` matches the base type of this `ArgType`.
        pub fn is_match<T: OclPrm + Any + 'static>(&self) -> bool {
            self.matches(TypeId::of::<T>())
        }

        #[allow(dead_code)]
        pub fn is_ptr(&self) -> bool {
            self.is_ptr
        }

        pub fn is_unknown(&self) -> bool {
            match self.base_type {
                BaseType::Unknown => true,
                _ => false,
            }
        }
    }

    impl<'a> From<&'a str> for ArgType {
        /// Ascertains a `KernelArgType` from the contents of a
        /// `KernelArgInfoResult::TypeName`.
        ///
        /// Panics if the string is not valid.
        fn from(type_name: &'a str) -> ArgType {
            ArgType::from_str(type_name).unwrap()
        }
    }
}