mesa/src/gallium/frontends/rusticl/api/kernel.rs

// Copyright 2020 Red Hat.
// SPDX-License-Identifier: MIT

use crate::api::event::create_and_queue;
use crate::api::icd::*;
use crate::api::util::*;
use crate::core::device::*;
use crate::core::event::*;
use crate::core::kernel::*;
use crate::core::memory::*;
use crate::core::program::*;
use crate::core::queue::*;

use mesa_rust_util::ptr::*;
use mesa_rust_util::string::*;
use rusticl_opencl_gen::*;
use rusticl_proc_macros::cl_entrypoint;
use rusticl_proc_macros::cl_info_entrypoint;

use std::cmp;
use std::ffi::CStr;
use std::os::raw::c_void;
use std::ptr;
use std::slice;
use std::sync::Arc;

#[cl_info_entrypoint(clGetKernelInfo)]
unsafe impl CLInfo<cl_kernel_info> for cl_kernel {
    fn query(&self, q: cl_kernel_info, v: CLInfoValue) -> CLResult<CLInfoRes> {
        let kernel = Kernel::ref_from_raw(*self)?;
        match q {
            CL_KERNEL_ATTRIBUTES => v.write::<&str>(&kernel.kernel_info.attributes_string),
            CL_KERNEL_CONTEXT => {
                let ptr = Arc::as_ptr(&kernel.prog.context);
                v.write::<cl_context>(cl_context::from_ptr(ptr))
            }
            CL_KERNEL_FUNCTION_NAME => v.write::<&str>(&kernel.name),
            CL_KERNEL_NUM_ARGS => v.write::<cl_uint>(kernel.kernel_info.args.len() as cl_uint),
            CL_KERNEL_PROGRAM => {
                let ptr = Arc::as_ptr(&kernel.prog);
                v.write::<cl_program>(cl_program::from_ptr(ptr))
            }
            CL_KERNEL_REFERENCE_COUNT => v.write::<cl_uint>(Kernel::refcnt(*self)?),
            // CL_INVALID_VALUE if param_name is not one of the supported values
            _ => Err(CL_INVALID_VALUE),
        }
    }
}

#[cl_info_entrypoint(clGetKernelArgInfo)]
unsafe impl CLInfoObj<cl_kernel_arg_info, cl_uint> for cl_kernel {
    fn query(&self, idx: cl_uint, q: cl_kernel_arg_info, v: CLInfoValue) -> CLResult<CLInfoRes> {
        let kernel = Kernel::ref_from_raw(*self)?;

        let idx = idx as usize;

        // CL_INVALID_ARG_INDEX if arg_index is not a valid argument index.
        if idx >= kernel.kernel_info.args.len() {
            return Err(CL_INVALID_ARG_INDEX);
        }

        match *q {
            CL_KERNEL_ARG_ACCESS_QUALIFIER => {
                v.write::<cl_kernel_arg_access_qualifier>(kernel.access_qualifier(idx))
            }
            CL_KERNEL_ARG_ADDRESS_QUALIFIER => {
                v.write::<cl_kernel_arg_address_qualifier>(kernel.address_qualifier(idx))
            }
            CL_KERNEL_ARG_NAME => v.write::<&CStr>(
                kernel
                    .arg_name(idx)
                    .ok_or(CL_KERNEL_ARG_INFO_NOT_AVAILABLE)?,
            ),
            CL_KERNEL_ARG_TYPE_NAME => v.write::<&CStr>(
                kernel
                    .arg_type_name(idx)
                    .ok_or(CL_KERNEL_ARG_INFO_NOT_AVAILABLE)?,
            ),
            CL_KERNEL_ARG_TYPE_QUALIFIER => {
                v.write::<cl_kernel_arg_type_qualifier>(kernel.type_qualifier(idx))
            }
            // CL_INVALID_VALUE if param_name is not one of the supported values
            _ => Err(CL_INVALID_VALUE),
        }
    }
}

#[cl_info_entrypoint(clGetKernelWorkGroupInfo)]
unsafe impl CLInfoObj<cl_kernel_work_group_info, cl_device_id> for cl_kernel {
    fn query(
        &self,
        dev: cl_device_id,
        q: cl_kernel_work_group_info,
        v: CLInfoValue,
    ) -> CLResult<CLInfoRes> {
        let kernel = Kernel::ref_from_raw(*self)?;

        // CL_INVALID_DEVICE [..] if device is NULL but there is more than one device associated with kernel.
        let dev = if dev.is_null() {
            if kernel.prog.devs.len() > 1 {
                return Err(CL_INVALID_DEVICE);
            } else {
                kernel.prog.devs[0]
            }
        } else {
            let dev = Device::ref_from_raw(dev)?;

            // CL_INVALID_DEVICE if device is not in the list of devices associated with kernel
            if !kernel.prog.devs.contains(&dev) {
                return Err(CL_INVALID_DEVICE);
            }

            dev
        };

        match *q {
            CL_KERNEL_COMPILE_WORK_GROUP_SIZE => v.write::<[usize; 3]>(kernel.work_group_size()),
            CL_KERNEL_LOCAL_MEM_SIZE => v.write::<cl_ulong>(kernel.local_mem_size(dev)),
            CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE => {
                v.write::<usize>(kernel.preferred_simd_size(dev))
            }
            CL_KERNEL_PRIVATE_MEM_SIZE => v.write::<cl_ulong>(kernel.priv_mem_size(dev)),
            CL_KERNEL_SPILL_MEM_SIZE_INTEL => v.write::<cl_ulong>(kernel.priv_mem_size(dev)),
            CL_KERNEL_WORK_GROUP_SIZE => v.write::<usize>(kernel.max_threads_per_block(dev)),
            // CL_INVALID_VALUE if param_name is not one of the supported values
            _ => Err(CL_INVALID_VALUE),
        }
    }
}

unsafe impl CLInfoObj<cl_kernel_sub_group_info, (cl_device_id, usize, *const c_void, usize)>
    for cl_kernel
{
    fn query(
        &self,
        (dev, input_value_size, input_value, output_value_size): (
            cl_device_id,
            usize,
            *const c_void,
            usize,
        ),
        q: cl_program_build_info,
        v: CLInfoValue,
    ) -> CLResult<CLInfoRes> {
        let kernel = Kernel::ref_from_raw(*self)?;

        // CL_INVALID_DEVICE [..] if device is NULL but there is more than one device associated
        // with kernel.
        let dev = if dev.is_null() {
            if kernel.prog.devs.len() > 1 {
                return Err(CL_INVALID_DEVICE);
            } else {
                kernel.prog.devs[0]
            }
        } else {
            Device::ref_from_raw(dev)?
        };

        // CL_INVALID_DEVICE if device is not in the list of devices associated with kernel
        if !kernel.prog.devs.contains(&dev) {
            return Err(CL_INVALID_DEVICE);
        }

        // CL_INVALID_OPERATION if device does not support subgroups.
        if !dev.subgroups_supported() {
            return Err(CL_INVALID_OPERATION);
        }

        let usize_byte = size_of::<usize>();
        // first we have to convert the input to a proper thing
        let input: &[usize] = match q {
            CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE | CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE => {
                // CL_INVALID_VALUE if param_name is CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE,
                // CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE or ... and the size in bytes specified by
                // input_value_size is not valid or if input_value is NULL.
                if ![usize_byte, 2 * usize_byte, 3 * usize_byte].contains(&input_value_size) {
                    return Err(CL_INVALID_VALUE);
                }
                // SAFETY: we verified the size as best as possible, with the rest we trust the client
                unsafe { slice::from_raw_parts(input_value.cast(), input_value_size / usize_byte) }
            }
            CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT => {
                // CL_INVALID_VALUE if param_name is ... CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT
                // and the size in bytes specified by input_value_size is not valid or if
                // input_value is NULL.
                if input_value_size != usize_byte || input_value.is_null() {
                    return Err(CL_INVALID_VALUE);
                }
                // SAFETY: we trust the client here
                unsafe { slice::from_raw_parts(input_value.cast(), 1) }
            }
            _ => &[],
        };

        match q {
            CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE => {
                v.write::<usize>(kernel.subgroups_for_block(dev, input))
            }
            CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE => {
                v.write::<usize>(kernel.subgroup_size_for_block(dev, input))
            }
            CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT => {
                let subgroups = input[0];
                let mut res = [0; 3];

                for subgroup_size in kernel.subgroup_sizes(dev) {
                    let threads = subgroups * subgroup_size;

                    if threads > dev.max_threads_per_block() {
                        continue;
                    }

                    let block = [threads, 1, 1];
                    let real_subgroups = kernel.subgroups_for_block(dev, &block);

                    if real_subgroups == subgroups {
                        res = block;
                        break;
                    }
                }

                v.write_iter::<usize>(res.into_iter().take(output_value_size / usize_byte))
            }
            CL_KERNEL_MAX_NUM_SUB_GROUPS => {
                let threads = kernel.max_threads_per_block(dev);
                let max_groups = dev.max_subgroups();

                let mut result = 0;
                for sgs in kernel.subgroup_sizes(dev) {
                    result = cmp::max(result, threads / sgs);
                    result = cmp::min(result, max_groups as usize);
                }
                v.write::<usize>(result)
            }
            CL_KERNEL_COMPILE_NUM_SUB_GROUPS => v.write::<usize>(kernel.num_subgroups()),
            CL_KERNEL_COMPILE_SUB_GROUP_SIZE_INTEL => v.write::<usize>(kernel.subgroup_size()),
            // CL_INVALID_VALUE if param_name is not one of the supported values
            _ => Err(CL_INVALID_VALUE),
        }
    }
}

const ZERO_ARR: [usize; 3] = [0; 3];

/// # Safety
///
/// This function is only safe when called on an array of `work_dim` length
unsafe fn kernel_work_arr_or_default<'a>(arr: *const usize, work_dim: cl_uint) -> &'a [usize] {
    if !arr.is_null() {
        unsafe { slice::from_raw_parts(arr, work_dim as usize) }
    } else {
        &ZERO_ARR
    }
}

/// # Safety
///
/// This function is only safe when called on an array of `work_dim` length
unsafe fn kernel_work_arr_mut<'a>(arr: *mut usize, work_dim: cl_uint) -> Option<&'a mut [usize]> {
    if !arr.is_null() {
        unsafe { Some(slice::from_raw_parts_mut(arr, work_dim as usize)) }
    } else {
        None
    }
}

#[cl_entrypoint(clCreateKernel)]
fn create_kernel(
    program: cl_program,
    kernel_name: *const ::std::os::raw::c_char,
) -> CLResult<cl_kernel> {
    let p = Program::arc_from_raw(program)?;
    let name = c_string_to_string(kernel_name);

    // CL_INVALID_VALUE if kernel_name is NULL.
    if kernel_name.is_null() {
        return Err(CL_INVALID_VALUE);
    }

    let build = p.build_info();
    // CL_INVALID_PROGRAM_EXECUTABLE if there is no successfully built executable for program.
    if build.kernels().is_empty() {
        return Err(CL_INVALID_PROGRAM_EXECUTABLE);
    }

    // CL_INVALID_KERNEL_NAME if kernel_name is not found in program.
    if !build.kernels().contains(&name) {
        return Err(CL_INVALID_KERNEL_NAME);
    }

    // CL_INVALID_KERNEL_DEFINITION if the function definition for __kernel function given by
    // kernel_name such as the number of arguments, the argument types are not the same for all
    // devices for which the program executable has been built.
    if !p.has_unique_kernel_signatures(&name) {
        return Err(CL_INVALID_KERNEL_DEFINITION);
    }

    Ok(Kernel::new(name, Arc::clone(&p), &build).into_cl())
}

#[cl_entrypoint(clRetainKernel)]
fn retain_kernel(kernel: cl_kernel) -> CLResult<()> {
    Kernel::retain(kernel)
}

#[cl_entrypoint(clReleaseKernel)]
fn release_kernel(kernel: cl_kernel) -> CLResult<()> {
    Kernel::release(kernel)
}

#[cl_entrypoint(clCreateKernelsInProgram)]
fn create_kernels_in_program(
    program: cl_program,
    num_kernels: cl_uint,
    kernels: *mut cl_kernel,
    num_kernels_ret: *mut cl_uint,
) -> CLResult<()> {
    let p = Program::arc_from_raw(program)?;
    let build = p.build_info();

    // CL_INVALID_PROGRAM_EXECUTABLE if there is no successfully built executable for any device in
    // program.
    if build.kernels().is_empty() {
        return Err(CL_INVALID_PROGRAM_EXECUTABLE);
    }

    // CL_INVALID_VALUE if kernels is not NULL and num_kernels is less than the number of kernels
    // in program.
    if !kernels.is_null() && build.kernels().len() > num_kernels as usize {
        return Err(CL_INVALID_VALUE);
    }

    let mut num_kernels = 0;
    for name in build.kernels() {
        // Kernel objects are not created for any __kernel functions in program that do not have the
        // same function definition across all devices for which a program executable has been
        // successfully built.
        if !p.has_unique_kernel_signatures(name) {
            continue;
        }

        if !kernels.is_null() {
            // we just assume the client isn't stupid
            unsafe {
                kernels
                    .add(num_kernels as usize)
                    .write(Kernel::new(name.clone(), Arc::clone(&p), &build).into_cl());
            }
        }
        num_kernels += 1;
    }

    // SAFETY: Caller is responsible for providing a pointer valid for a write
    // of `size_of::<cl_uint>()`.
    unsafe { num_kernels_ret.write_checked(num_kernels) };

    Ok(())
}

#[cl_entrypoint(clSetKernelArg)]
fn set_kernel_arg(
    kernel: cl_kernel,
    arg_index: cl_uint,
    arg_size: usize,
    arg_value: *const ::std::os::raw::c_void,
) -> CLResult<()> {
    // SAFETY: Per spec, "OpenCL calls that modify the state of `cl_kernel`
    // objects are not thread-safe: clSetKernelArg, ...", so no other threads
    // must hold a reference to this `cl_kernel`. Additionally, we may require
    // callers to pass either null or a valid CL object, satisfying validity
    // constraints.
    let k = unsafe { Kernel::mut_ref_from_raw(kernel) }?;
    let arg_index = arg_index as usize;

    // CL_INVALID_ARG_INDEX if arg_index is not a valid argument index.
    if let Some(arg) = k.kernel_info.args.get(arg_index) {
        // CL_INVALID_ARG_SIZE if arg_size does not match the size of the data type for an argument
        // that is not a memory object or if the argument is a memory object and
        // arg_size != sizeof(cl_mem) or if arg_size is zero and the argument is declared with the
        // local qualifier or if the argument is a sampler and arg_size != sizeof(cl_sampler).
        match arg.kind {
            KernelArgType::MemLocal => {
                if arg_size == 0 {
                    return Err(CL_INVALID_ARG_SIZE);
                }
            }
            KernelArgType::MemGlobal
            | KernelArgType::MemConstant
            | KernelArgType::Image
            | KernelArgType::RWImage
            | KernelArgType::Texture => {
                if arg_size != size_of::<cl_mem>() {
                    return Err(CL_INVALID_ARG_SIZE);
                }
            }

            KernelArgType::Sampler => {
                if arg_size != size_of::<cl_sampler>() {
                    return Err(CL_INVALID_ARG_SIZE);
                }
            }

            KernelArgType::Constant(size) => {
                if size as usize != arg_size {
                    return Err(CL_INVALID_ARG_SIZE);
                }
            }
        }

        // CL_INVALID_ARG_VALUE if arg_value specified is not a valid value.
        match arg.kind {
            // If the argument is declared with the local qualifier, the arg_value entry must be
            // NULL.
            KernelArgType::MemLocal => {
                if !arg_value.is_null() {
                    return Err(CL_INVALID_ARG_VALUE);
                }
            }
            KernelArgType::Constant(_) => {
                if arg_value.is_null() {
                    return Err(CL_INVALID_ARG_VALUE);
                }
            }
            KernelArgType::Sampler => {
                // CL_INVALID_SAMPLER for an argument declared to be of type sampler_t when the
                // specified arg_value is not a valid sampler object.
                if arg_value.is_null() {
                    return Err(CL_INVALID_SAMPLER);
                }
            }
            _ => {}
        };

        // let's create the arg now
        let arg = match arg.kind {
            KernelArgType::Constant(_) if !arg.dead => KernelArgValue::Constant(
                // SAFETY: for constant args arg_value points to raw data.
                unsafe { slice::from_raw_parts(arg_value.cast(), arg_size) }.to_vec(),
            ),
            KernelArgType::MemConstant | KernelArgType::MemGlobal => {
                let ptr: *const cl_mem = arg_value.cast();
                // SAFETY: for buffer args arg_value points to a cl_mem object or NULL.
                if ptr.is_null() || unsafe { *ptr }.is_null() {
                    KernelArgValue::None
                } else {
                    // SAFETY: as above
                    let buffer = Buffer::arc_from_raw(unsafe { *ptr })?;
                    // We are required to prevent mutable access to immutable memory objects,
                    // however no explicit error code has been specified yet.
                    if arg.kind == KernelArgType::MemGlobal
                        && bit_check(buffer.flags, CL_MEM_IMMUTABLE_EXT)
                    {
                        return Err(CL_INVALID_ARG_VALUE);
                    }

                    KernelArgValue::Buffer(Arc::downgrade(&buffer))
                }
            }
            KernelArgType::MemLocal if !arg.dead => KernelArgValue::LocalMem(arg_size),
            KernelArgType::Image | KernelArgType::RWImage | KernelArgType::Texture if !arg.dead => {
                let img: *const cl_mem = arg_value.cast();
                // SAFETY: for image args arg_value pointers to a cl_mem object.
                let img = Image::arc_from_raw(unsafe { *img })?;

                // CL_INVALID_ARG_VALUE if the argument is an image declared with the read_only
                // qualifier and arg_value refers to an image object created with cl_mem_flags
                // of CL_MEM_WRITE_ONLY or if the image argument is declared with the write_only
                // qualifier and arg_value refers to an image object created with cl_mem_flags
                // of CL_MEM_READ_ONLY.
                // We are required to prevent mutable access to immutable memory objects, however no
                // explicit error code has been specified yet.
                if arg.kind == KernelArgType::Texture && bit_check(img.flags, CL_MEM_WRITE_ONLY)
                    || arg.kind == KernelArgType::Image
                        && bit_check(img.flags, CL_MEM_READ_ONLY | CL_MEM_IMMUTABLE_EXT)
                    || arg.kind == KernelArgType::RWImage
                        && bit_check(img.flags, CL_MEM_IMMUTABLE_EXT)
                {
                    return Err(CL_INVALID_ARG_VALUE);
                }

                KernelArgValue::Image(Arc::downgrade(&img))
            }
            KernelArgType::Sampler if !arg.dead => {
                let ptr: *const cl_sampler = arg_value.cast();
                // SAFETY: for a sampler arg arg_value points to a cl_sampler object.
                KernelArgValue::Sampler(Sampler::arc_from_raw(unsafe { *ptr })?)
            }
            _ => {
                debug_assert!(
                    arg.dead
                        || matches!(
                            arg.kind,
                            KernelArgType::MemConstant | KernelArgType::MemGlobal
                        )
                );
                KernelArgValue::None
            }
        };
        k.set_kernel_arg(arg_index, arg)
    } else {
        Err(CL_INVALID_ARG_INDEX)
    }

    //• CL_INVALID_DEVICE_QUEUE for an argument declared to be of type queue_t when the specified arg_value is not a valid device queue object. This error code is missing before version 2.0.
    //• CL_MAX_SIZE_RESTRICTION_EXCEEDED if the size in bytes of the memory object (if the argument is a memory object) or arg_size (if the argument is declared with local qualifier) exceeds a language- specified maximum size restriction for this argument, such as the MaxByteOffset SPIR-V decoration. This error code is missing before version 2.2.
}

#[cl_entrypoint(clSetKernelArgSVMPointer)]
fn set_kernel_arg_svm_pointer(
    kernel: cl_kernel,
    arg_index: cl_uint,
    arg_value: *const ::std::os::raw::c_void,
) -> CLResult<()> {
    // SAFETY: Per spec, "OpenCL calls that modify the state of `cl_kernel`
    // objects are not thread-safe: ..., clSetKernelArgSVMPointer, ...", so no
    // other threads must hold a reference to this `cl_kernel`. Additionally, we
    // may require callers to pass either null or a valid CL object, satisfying
    // validity constraints.
    let kernel = unsafe { Kernel::mut_ref_from_raw(kernel) }?;
    let arg_index = arg_index as usize;
    let arg_value = arg_value as usize;

    if !kernel.has_svm_devs() {
        return Err(CL_INVALID_OPERATION);
    }

    if let Some(arg) = kernel.kernel_info.args.get(arg_index) {
        if !matches!(
            arg.kind,
            KernelArgType::MemConstant | KernelArgType::MemGlobal
        ) {
            return Err(CL_INVALID_ARG_INDEX);
        }

        let arg_value = KernelArgValue::SVM(arg_value);
        kernel.set_kernel_arg(arg_index, arg_value)
    } else {
        Err(CL_INVALID_ARG_INDEX)
    }

    // CL_INVALID_ARG_VALUE if arg_value specified is not a valid value.
}

#[cl_entrypoint(clSetKernelArgDevicePointerEXT)]
fn set_kernel_arg_device_pointer(
    kernel: cl_kernel,
    arg_index: cl_uint,
    arg_value: cl_mem_device_address_ext,
) -> CLResult<()> {
    // SAFETY: Per spec, "OpenCL calls that modify the state of `cl_kernel`
    // objects are not thread-safe...", so no other threads must hold a
    // reference to this `cl_kernel`. Additionally, we may require callers to
    // pass either null or a valid CL object, satisfying validity constraints.
    let kernel = unsafe { Kernel::mut_ref_from_raw(kernel) }?;
    let arg_index = arg_index as usize;
    let devs = &kernel.prog.context.devs;

    // CL_INVALID_OPERATION if no devices in the context associated with kernel support the device
    // pointer.
    if devs.iter().any(|dev| !dev.bda_supported()) {
        return Err(CL_INVALID_OPERATION);
    }

    // CL_INVALID_ARG_INDEX if arg_index is not a valid argument index.
    let Some(arg) = kernel.kernel_info.args.get(arg_index) else {
        return Err(CL_INVALID_ARG_INDEX);
    };

    // The device pointer can only be used for arguments that are declared to be a pointer to global
    // memory allocated with clCreateBufferWithProperties with the CL_MEM_DEVICE_PRIVATE_ADDRESS_EXT
    // property.
    if arg.kind != KernelArgType::MemGlobal {
        return Err(CL_INVALID_ARG_INDEX);
    }

    // we set the arg also when it's a dead argument, as we need to ensure the buffer gets migrated.
    kernel.set_kernel_arg(arg_index, KernelArgValue::BDA(arg_value))
}

#[cl_entrypoint(clSetKernelExecInfo)]
fn set_kernel_exec_info(
    kernel: cl_kernel,
    param_name: cl_kernel_exec_info,
    param_value_size: usize,
    param_value: *const ::std::os::raw::c_void,
) -> CLResult<()> {
    // SAFETY: Per spec, "OpenCL calls that modify the state of `cl_kernel`
    // objects are not thread-safe: ..., clSetKernelExecInfo, ...", so no other
    // threads must hold a reference to this `cl_kernel`. Additionally, we may
    // require callers to pass either null or a valid CL object, satisfying
    // validity constraints.
    let k = unsafe { Kernel::mut_ref_from_raw(kernel) }?;
    let devs = &k.prog.devs;

    // CL_INVALID_OPERATION for CL_KERNEL_EXEC_INFO_DEVICE_PTRS_EXT if no device in the context
    // associated with kernel support the cl_ext_buffer_device_address extension.
    let check_bda_support = || {
        if devs.iter().all(|dev| !dev.bda_supported()) {
            Err(CL_INVALID_OPERATION)
        } else {
            Ok(())
        }
    };

    // CL_INVALID_OPERATION for CL_KERNEL_EXEC_INFO_SVM_PTRS and
    // CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM if no devices in the context associated with kernel
    // support SVM.
    let check_svm_support = || {
        if devs.iter().all(|dev| !dev.api_svm_supported()) {
            Err(CL_INVALID_OPERATION)
        } else {
            Ok(())
        }
    };

    // CL_INVALID_VALUE ... if the size specified by param_value_size is not valid.
    match param_name {
        CL_KERNEL_EXEC_INFO_DEVICE_PTRS_EXT => {
            check_bda_support()?;
            let handles = unsafe {
                cl_slice::from_raw_parts_bytes_len::<cl_mem_device_address_ext>(
                    param_value,
                    param_value_size,
                )?
            };

            handles.clone_into(&mut k.bdas);
        }
        CL_KERNEL_EXEC_INFO_SVM_PTRS | CL_KERNEL_EXEC_INFO_SVM_PTRS_ARM => {
            check_svm_support()?;

            // To specify that no SVM allocations will be accessed by a kernel other than those set
            // as kernel arguments, specify an empty set by passing param_value_size equal to zero
            // and param_value equal to NULL.
            if param_value_size == 0 && param_value.is_null() {
                k.svms.clear();
            } else {
                let pointers = unsafe {
                    cl_slice::from_raw_parts_bytes_len::<*const c_void>(
                        param_value,
                        param_value_size,
                    )?
                };

                // We need to clear _after_ the error checking above. We could just assign a new
                // container, however we also want to reuse the allocations.
                k.svms.clear();
                pointers
                    .iter()
                    // Each of the pointers can be the pointer returned by clSVMAlloc or can be a
                    // pointer to the middle of an SVM allocation. It is sufficient to specify one
                    // pointer for each SVM allocation.
                    //
                    // So we'll simply fetch the base and store that one.
                    .filter_map(|&handle| k.prog.context.find_svm_alloc(handle as usize))
                    .for_each(|(base, _)| {
                        k.svms.insert(base as usize);
                    });
            }
        }
        CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM
        | CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM_ARM => {
            check_svm_support()?;
            let val = unsafe {
                cl_slice::from_raw_parts_bytes_len::<cl_bool>(param_value, param_value_size)?
            };

            // we must explicitly check that we only got one element
            if val.len() != 1 {
                return Err(CL_INVALID_VALUE);
            }

            // CL_INVALID_OPERATION if param_name is CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM and
            // param_value is CL_TRUE but no devices in context associated with kernel support
            // fine-grain system SVM allocations.
            if val[0] == CL_TRUE && devs.iter().all(|dev| !dev.system_svm_supported()) {
                return Err(CL_INVALID_OPERATION);
            }
        }
        // CL_INVALID_VALUE if param_name is not valid
        _ => return Err(CL_INVALID_VALUE),
    }

    Ok(())
}

#[cl_entrypoint(clEnqueueNDRangeKernel)]
fn enqueue_ndrange_kernel(
    command_queue: cl_command_queue,
    kernel: cl_kernel,
    work_dim: cl_uint,
    global_work_offset: *const usize,
    global_work_size: *const usize,
    local_work_size: *const usize,
    num_events_in_wait_list: cl_uint,
    event_wait_list: *const cl_event,
    event: *mut cl_event,
) -> CLResult<()> {
    let q = Queue::arc_from_raw(command_queue)?;
    let k = Kernel::arc_from_raw(kernel)?;
    let evs = event_list_from_cl(&q, num_events_in_wait_list, event_wait_list)?;

    // CL_INVALID_CONTEXT if context associated with command_queue and kernel are not the same
    if q.context != k.prog.context {
        return Err(CL_INVALID_CONTEXT);
    }

    // CL_INVALID_PROGRAM_EXECUTABLE if there is no successfully built program executable available
    // for device associated with command_queue.
    if k.prog.status(q.device) != CL_BUILD_SUCCESS as cl_build_status {
        return Err(CL_INVALID_PROGRAM_EXECUTABLE);
    }

    // CL_INVALID_KERNEL_ARGS if the kernel argument values have not been specified.
    if k.arg_values().iter().any(|v| v.is_none()) {
        return Err(CL_INVALID_KERNEL_ARGS);
    }

    // CL_INVALID_WORK_DIMENSION if work_dim is not a valid value (i.e. a value between 1 and
    // CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS).
    if work_dim == 0 || work_dim > q.device.max_grid_dimensions() {
        return Err(CL_INVALID_WORK_DIMENSION);
    }

    // we assume the application gets it right and doesn't pass shorter arrays then actually needed.
    let global_work_size = unsafe { kernel_work_arr_or_default(global_work_size, work_dim) };
    let local_work_size = unsafe { kernel_work_arr_or_default(local_work_size, work_dim) };
    let global_work_offset = unsafe { kernel_work_arr_or_default(global_work_offset, work_dim) };

    let device_bits = q.device.address_bits();
    let device_max = u64::MAX >> (u64::BITS - device_bits);

    let mut threads = 1;
    for i in 0..work_dim as usize {
        let lws = local_work_size[i];
        let gws = global_work_size[i];
        let gwo = global_work_offset[i];

        threads *= lws;

        // CL_INVALID_WORK_ITEM_SIZE if the number of work-items specified in any of
        // local_work_size[0], … local_work_size[work_dim - 1] is greater than the corresponding
        // values specified by
        // CL_DEVICE_MAX_WORK_ITEM_SIZES[0], …, CL_DEVICE_MAX_WORK_ITEM_SIZES[work_dim - 1].
        if lws > q.device.max_block_sizes()[i] {
            return Err(CL_INVALID_WORK_ITEM_SIZE);
        }

        // CL_INVALID_WORK_GROUP_SIZE if the work-group size must be uniform and the
        // local_work_size is not NULL, [...] if the global_work_size is not evenly divisible by
        // the local_work_size.
        if lws != 0 && gws % lws != 0 {
            return Err(CL_INVALID_WORK_GROUP_SIZE);
        }

        // CL_INVALID_WORK_GROUP_SIZE if local_work_size is specified and does not match the
        // required work-group size for kernel in the program source.
        if lws != 0 && k.work_group_size()[i] != 0 && lws != k.work_group_size()[i] {
            return Err(CL_INVALID_WORK_GROUP_SIZE);
        }

        // CL_INVALID_GLOBAL_WORK_SIZE if any of the values specified in global_work_size[0], …
        // global_work_size[work_dim - 1] exceed the maximum value representable by size_t on
        // the device on which the kernel-instance will be enqueued.
        if gws as u64 > device_max {
            return Err(CL_INVALID_GLOBAL_WORK_SIZE);
        }

        // CL_INVALID_GLOBAL_OFFSET if the value specified in global_work_size + the
        // corresponding values in global_work_offset for any dimensions is greater than the
        // maximum value representable by size t on the device on which the kernel-instance
        // will be enqueued
        if u64::checked_add(gws as u64, gwo as u64)
            .filter(|&x| x <= device_max)
            .is_none()
        {
            return Err(CL_INVALID_GLOBAL_OFFSET);
        }
    }

    // CL_INVALID_WORK_GROUP_SIZE if local_work_size is specified and the total number of work-items
    // in the work-group computed as local_work_size[0] × … local_work_size[work_dim - 1] is greater
    // than the value specified by CL_KERNEL_WORK_GROUP_SIZE in the Kernel Object Device Queries
    // table.
    if threads != 0 && threads > k.max_threads_per_block(q.device) {
        return Err(CL_INVALID_WORK_GROUP_SIZE);
    }

    // If global_work_size is NULL, or the value in any passed dimension is 0 then the kernel
    // command will trivially succeed after its event dependencies are satisfied and subsequently
    // update its completion event.
    let cb: EventSig = if global_work_size.contains(&0) {
        Box::new(|_, _| Ok(()))
    } else {
        k.launch(
            &q,
            work_dim,
            local_work_size,
            global_work_size,
            global_work_offset,
        )?
    };

    create_and_queue(q, CL_COMMAND_NDRANGE_KERNEL, evs, event, false, cb)

    //• CL_INVALID_WORK_GROUP_SIZE if local_work_size is specified and is not consistent with the required number of sub-groups for kernel in the program source.
    //• CL_MISALIGNED_SUB_BUFFER_OFFSET if a sub-buffer object is specified as the value for an argument that is a buffer object and the offset specified when the sub-buffer object is created is not aligned to CL_DEVICE_MEM_BASE_ADDR_ALIGN value for device associated with queue. This error code
    //• CL_INVALID_IMAGE_SIZE if an image object is specified as an argument value and the image dimensions (image width, height, specified or compute row and/or slice pitch) are not supported by device associated with queue.
    //• CL_IMAGE_FORMAT_NOT_SUPPORTED if an image object is specified as an argument value and the image format (image channel order and data type) is not supported by device associated with queue.
    //• CL_OUT_OF_RESOURCES if there is a failure to queue the execution instance of kernel on the command-queue because of insufficient resources needed to execute the kernel. For example, the explicitly specified local_work_size causes a failure to execute the kernel because of insufficient resources such as registers or local memory. Another example would be the number of read-only image args used in kernel exceed the CL_DEVICE_MAX_READ_IMAGE_ARGS value for device or the number of write-only and read-write image args used in kernel exceed the CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS value for device or the number of samplers used in kernel exceed CL_DEVICE_MAX_SAMPLERS for device.
    //• CL_MEM_OBJECT_ALLOCATION_FAILURE if there is a failure to allocate memory for data store associated with image or buffer objects specified as arguments to kernel.
    //• CL_INVALID_OPERATION if SVM pointers are passed as arguments to a kernel and the device does not support SVM or if system pointers are passed as arguments to a kernel and/or stored inside SVM allocations passed as kernel arguments and the device does not support fine grain system SVM allocations.
}

#[cl_entrypoint(clEnqueueTask)]
fn enqueue_task(
    command_queue: cl_command_queue,
    kernel: cl_kernel,
    num_events_in_wait_list: cl_uint,
    event_wait_list: *const cl_event,
    event: *mut cl_event,
) -> CLResult<()> {
    // clEnqueueTask is equivalent to calling clEnqueueNDRangeKernel with work_dim set to 1,
    // global_work_offset set to NULL, global_work_size[0] set to 1, and local_work_size[0] set to
    // 1.
    enqueue_ndrange_kernel(
        command_queue,
        kernel,
        1,
        ptr::null(),
        [1, 1, 1].as_ptr(),
        [1, 0, 0].as_ptr(),
        num_events_in_wait_list,
        event_wait_list,
        event,
    )
}

#[cl_entrypoint(clCloneKernel)]
fn clone_kernel(source_kernel: cl_kernel) -> CLResult<cl_kernel> {
    let k = Kernel::ref_from_raw(source_kernel)?;
    Ok(Arc::new(k.clone()).into_cl())
}

#[cl_entrypoint(clGetKernelSuggestedLocalWorkSizeKHR)]
fn get_kernel_suggested_local_work_size_khr(
    command_queue: cl_command_queue,
    kernel: cl_kernel,
    work_dim: cl_uint,
    global_work_offset: *const usize,
    global_work_size: *const usize,
    suggested_local_work_size: *mut usize,
) -> CLResult<()> {
    // CL_INVALID_GLOBAL_WORK_SIZE if global_work_size is NULL or if any of the values specified in
    // global_work_size are 0.
    if global_work_size.is_null() {
        return Err(CL_INVALID_GLOBAL_WORK_SIZE);
    }

    // CL_INVALID_VALUE if suggested_local_work_size is NULL.
    if suggested_local_work_size.is_null() {
        return Err(CL_INVALID_VALUE);
    }

    // CL_INVALID_COMMAND_QUEUE if command_queue is not a valid host command-queue.
    let queue = Queue::ref_from_raw(command_queue)?;

    // CL_INVALID_KERNEL if kernel is not a valid kernel object.
    let kernel = Kernel::ref_from_raw(kernel)?;

    // CL_INVALID_CONTEXT if the context associated with kernel is not the same as the context
    // associated with command_queue.
    if queue.context != kernel.prog.context {
        return Err(CL_INVALID_CONTEXT);
    }

    // CL_INVALID_PROGRAM_EXECUTABLE if there is no successfully built program executable available
    // for kernel for the device associated with command_queue.
    if kernel.prog.status(queue.device) != CL_BUILD_SUCCESS as cl_build_status {
        return Err(CL_INVALID_PROGRAM_EXECUTABLE);
    }

    // CL_INVALID_KERNEL_ARGS if all argument values for kernel have not been set.
    if kernel.arg_values().iter().any(|v| v.is_none()) {
        return Err(CL_INVALID_KERNEL_ARGS);
    }

    // CL_INVALID_WORK_DIMENSION if work_dim is not a valid value (i.e. a value between 1 and
    // CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS).
    if work_dim == 0 || work_dim > queue.device.max_grid_dimensions() {
        return Err(CL_INVALID_WORK_DIMENSION);
    }

    let mut global_work_size =
        unsafe { kernel_work_arr_or_default(global_work_size, work_dim).to_vec() };

    let suggested_local_work_size = unsafe {
        kernel_work_arr_mut(suggested_local_work_size, work_dim).ok_or(CL_INVALID_VALUE)?
    };

    let global_work_offset = unsafe { kernel_work_arr_or_default(global_work_offset, work_dim) };

    let device_bits = queue.device.address_bits();
    let device_max = u64::MAX >> (u64::BITS - device_bits);
    for i in 0..work_dim as usize {
        let gws = global_work_size[i];
        let gwo = global_work_offset[i];

        // CL_INVALID_GLOBAL_WORK_SIZE if global_work_size is NULL or if any of the values specified
        // in global_work_size are 0.
        if gws == 0 {
            return Err(CL_INVALID_GLOBAL_WORK_SIZE);
        }
        // CL_INVALID_GLOBAL_WORK_SIZE if any of the values specified in global_work_size exceed the
        // maximum value representable by size_t on the device associated with command_queue.
        if gws as u64 > device_max {
            return Err(CL_INVALID_GLOBAL_WORK_SIZE);
        }
        // CL_INVALID_GLOBAL_OFFSET if the value specified in global_work_size plus the
        // corresponding value in global_work_offset for dimension exceeds the maximum value
        // representable by size_t on the device associated with command_queue.
        if u64::checked_add(gws as u64, gwo as u64)
            .filter(|&x| x <= device_max)
            .is_none()
        {
            return Err(CL_INVALID_GLOBAL_OFFSET);
        }
    }

    kernel.suggest_local_size(
        queue.device,
        work_dim as usize,
        &mut global_work_size,
        suggested_local_work_size,
    );

    Ok(())

    // CL_MISALIGNED_SUB_BUFFER_OFFSET if a sub-buffer object is set as an argument to kernel and the offset specified when the sub-buffer object was created is not aligned to CL_DEVICE_MEM_BASE_ADDR_ALIGN for the device associated with command_queue.
    // CL_INVALID_IMAGE_SIZE if an image object is set as an argument to kernel and the image dimensions are not supported by device associated with command_queue.
    // CL_IMAGE_FORMAT_NOT_SUPPORTED if an image object is set as an argument to kernel and the image format is not supported by the device associated with command_queue.
    // CL_INVALID_OPERATION if an SVM pointer is set as an argument to kernel and the device associated with command_queue does not support SVM or the required SVM capabilities for the SVM pointer.
    // CL_OUT_OF_RESOURCES if there is a failure to allocate resources required by the OpenCL implementation on the device.
    // CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required by the OpenCL implementation on the host.
}