mirror of
https://gitlab.freedesktop.org/mesa/mesa.git
synced 2026-01-06 19:40:10 +01:00
rusticl/kernel: basic implementation
still not able to run kernels, but most of the boilerplate code is there now Signed-off-by: Karol Herbst <kherbst@redhat.com> Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15439>
This commit is contained in:
Karol Herbst
Marge Bot
parent
129b821638
commit
e8de580998
7 changed files with 350 additions and 24 deletions
|
|
@ -58,14 +58,14 @@ pub static DISPATCH: cl_icd_dispatch = cl_icd_dispatch {
|
|||
clGetProgramBuildInfo: Some(cl_get_program_build_info),
|
||||
clCreateKernel: Some(cl_create_kernel),
|
||||
clCreateKernelsInProgram: Some(cl_create_kernels_in_program),
|
||||
clRetainKernel: None,
|
||||
clReleaseKernel: None,
|
||||
clSetKernelArg: None,
|
||||
clGetKernelInfo: None,
|
||||
clGetKernelWorkGroupInfo: None,
|
||||
clRetainKernel: Some(cl_retain_kernel),
|
||||
clReleaseKernel: Some(cl_release_kernel),
|
||||
clSetKernelArg: Some(cl_set_kernel_arg),
|
||||
clGetKernelInfo: Some(cl_get_kernel_info),
|
||||
clGetKernelWorkGroupInfo: Some(cl_get_kernel_work_group_info),
|
||||
clWaitForEvents: Some(cl_wait_for_events),
|
||||
clGetEventInfo: Some(cl_get_event_info),
|
||||
clRetainEvent: None,
|
||||
clRetainEvent: Some(cl_retain_event),
|
||||
clReleaseEvent: Some(cl_release_event),
|
||||
clGetEventProfilingInfo: Some(cl_get_event_profiling_info),
|
||||
clFlush: Some(cl_flush),
|
||||
|
|
@ -81,7 +81,7 @@ pub static DISPATCH: cl_icd_dispatch = cl_icd_dispatch {
|
|||
clEnqueueMapBuffer: Some(cl_enqueue_map_buffer),
|
||||
clEnqueueMapImage: Some(cl_enqueue_map_image),
|
||||
clEnqueueUnmapMemObject: Some(cl_enqueue_unmap_mem_object),
|
||||
clEnqueueNDRangeKernel: None,
|
||||
clEnqueueNDRangeKernel: Some(cl_enqueue_ndrange_kernel),
|
||||
clEnqueueTask: None,
|
||||
clEnqueueNativeKernel: None,
|
||||
clEnqueueMarker: None,
|
||||
|
|
@ -123,7 +123,7 @@ pub static DISPATCH: cl_icd_dispatch = cl_icd_dispatch {
|
|||
clCompileProgram: Some(cl_compile_program),
|
||||
clLinkProgram: Some(cl_link_program),
|
||||
clUnloadPlatformCompiler: Some(cl_unload_platform_compiler),
|
||||
clGetKernelArgInfo: None,
|
||||
clGetKernelArgInfo: Some(cl_get_kernel_arg_info),
|
||||
clEnqueueFillBuffer: None,
|
||||
clEnqueueFillImage: Some(cl_enqueue_fill_image),
|
||||
clEnqueueMigrateMemObjects: None,
|
||||
|
|
@ -775,6 +775,55 @@ extern "C" fn cl_create_kernels_in_program(
|
|||
CL_OUT_OF_HOST_MEMORY
|
||||
}
|
||||
|
||||
extern "C" fn cl_retain_kernel(kernel: cl_kernel) -> cl_int {
|
||||
match_err!(kernel.retain())
|
||||
}
|
||||
|
||||
extern "C" fn cl_release_kernel(kernel: cl_kernel) -> cl_int {
|
||||
match_err!(kernel.release())
|
||||
}
|
||||
|
||||
extern "C" fn cl_set_kernel_arg(
|
||||
kernel: cl_kernel,
|
||||
arg_index: cl_uint,
|
||||
arg_size: usize,
|
||||
arg_value: *const ::std::os::raw::c_void,
|
||||
) -> cl_int {
|
||||
match_err!(set_kernel_arg(kernel, arg_index, arg_size, arg_value))
|
||||
}
|
||||
|
||||
extern "C" fn cl_get_kernel_info(
|
||||
kernel: cl_kernel,
|
||||
param_name: cl_kernel_info,
|
||||
param_value_size: usize,
|
||||
param_value: *mut ::std::os::raw::c_void,
|
||||
param_value_size_ret: *mut usize,
|
||||
) -> cl_int {
|
||||
match_err!(kernel.get_info(
|
||||
param_name,
|
||||
param_value_size,
|
||||
param_value,
|
||||
param_value_size_ret,
|
||||
))
|
||||
}
|
||||
|
||||
extern "C" fn cl_get_kernel_work_group_info(
|
||||
kernel: cl_kernel,
|
||||
device: cl_device_id,
|
||||
param_name: cl_kernel_work_group_info,
|
||||
param_value_size: usize,
|
||||
param_value: *mut ::std::os::raw::c_void,
|
||||
param_value_size_ret: *mut usize,
|
||||
) -> cl_int {
|
||||
match_err!(kernel.get_info_obj(
|
||||
device,
|
||||
param_name,
|
||||
param_value_size,
|
||||
param_value,
|
||||
param_value_size_ret,
|
||||
))
|
||||
}
|
||||
|
||||
extern "C" fn cl_wait_for_events(_num_events: cl_uint, _event_list: *const cl_event) -> cl_int {
|
||||
println!("cl_wait_for_events not implemented");
|
||||
CL_OUT_OF_HOST_MEMORY
|
||||
|
|
@ -795,6 +844,10 @@ extern "C" fn cl_get_event_info(
|
|||
))
|
||||
}
|
||||
|
||||
extern "C" fn cl_retain_event(event: cl_event) -> cl_int {
|
||||
match_err!(event.retain())
|
||||
}
|
||||
|
||||
extern "C" fn cl_release_event(event: cl_event) -> cl_int {
|
||||
match_err!(event.release())
|
||||
}
|
||||
|
|
@ -1025,6 +1078,21 @@ extern "C" fn cl_enqueue_unmap_mem_object(
|
|||
))
|
||||
}
|
||||
|
||||
extern "C" fn cl_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,
|
||||
) -> cl_int {
|
||||
println!("cl_enqueue_ndrange_kernel not implemented");
|
||||
CL_OUT_OF_HOST_MEMORY
|
||||
}
|
||||
|
||||
extern "C" fn cl_get_extension_function_address(
|
||||
function_name: *const ::std::os::raw::c_char,
|
||||
) -> *mut ::std::ffi::c_void {
|
||||
|
|
@ -1249,6 +1317,23 @@ extern "C" fn cl_unload_platform_compiler(_platform: cl_platform_id) -> cl_int {
|
|||
CL_OUT_OF_HOST_MEMORY
|
||||
}
|
||||
|
||||
extern "C" fn cl_get_kernel_arg_info(
|
||||
kernel: cl_kernel,
|
||||
arg_indx: cl_uint,
|
||||
param_name: cl_kernel_arg_info,
|
||||
param_value_size: usize,
|
||||
param_value: *mut ::std::os::raw::c_void,
|
||||
param_value_size_ret: *mut usize,
|
||||
) -> cl_int {
|
||||
match_err!(kernel.get_info_obj(
|
||||
arg_indx,
|
||||
param_name,
|
||||
param_value_size,
|
||||
param_value,
|
||||
param_value_size_ret,
|
||||
))
|
||||
}
|
||||
|
||||
extern "C" fn cl_enqueue_fill_image(
|
||||
_command_queue: cl_command_queue,
|
||||
_image: cl_mem,
|
||||
|
|
|
|||
|
|
@ -1,24 +1,146 @@
|
|||
extern crate mesa_rust_util;
|
||||
extern crate rusticl_opencl_gen;
|
||||
|
||||
use crate::api::icd::*;
|
||||
use crate::api::util::*;
|
||||
use crate::core::kernel::*;
|
||||
|
||||
use self::mesa_rust_util::string::*;
|
||||
use self::rusticl_opencl_gen::*;
|
||||
|
||||
use std::collections::HashSet;
|
||||
use std::sync::Arc;
|
||||
|
||||
impl CLInfo<cl_kernel_info> for cl_kernel {
|
||||
fn query(&self, q: cl_kernel_info) -> CLResult<Vec<u8>> {
|
||||
let kernel = self.get_ref()?;
|
||||
Ok(match q {
|
||||
CL_KERNEL_CONTEXT => {
|
||||
let ptr = Arc::as_ptr(&kernel.prog.context);
|
||||
cl_prop::<cl_context>(cl_context::from_ptr(ptr))
|
||||
}
|
||||
CL_KERNEL_FUNCTION_NAME => cl_prop::<&str>(&kernel.name),
|
||||
CL_KERNEL_NUM_ARGS => cl_prop::<cl_uint>(kernel.args.len() as cl_uint),
|
||||
CL_KERNEL_PROGRAM => {
|
||||
let ptr = Arc::as_ptr(&kernel.prog);
|
||||
cl_prop::<cl_program>(cl_program::from_ptr(ptr))
|
||||
}
|
||||
CL_KERNEL_REFERENCE_COUNT => cl_prop::<cl_uint>(self.refcnt()?),
|
||||
// CL_INVALID_VALUE if param_name is not one of the supported values
|
||||
_ => return Err(CL_INVALID_VALUE),
|
||||
})
|
||||
}
|
||||
}
|
||||
|
||||
impl CLInfoObj<cl_kernel_arg_info, cl_uint> for cl_kernel {
|
||||
fn query(&self, idx: cl_uint, q: cl_kernel_arg_info) -> CLResult<Vec<u8>> {
|
||||
let kernel = self.get_ref()?;
|
||||
|
||||
// CL_INVALID_ARG_INDEX if arg_index is not a valid argument index.
|
||||
if idx as usize >= kernel.args.len() {
|
||||
return Err(CL_INVALID_ARG_INDEX);
|
||||
}
|
||||
|
||||
Ok(match *q {
|
||||
CL_KERNEL_ARG_ACCESS_QUALIFIER => {
|
||||
cl_prop::<cl_kernel_arg_access_qualifier>(kernel.access_qualifier(idx))
|
||||
}
|
||||
CL_KERNEL_ARG_ADDRESS_QUALIFIER => {
|
||||
cl_prop::<cl_kernel_arg_address_qualifier>(kernel.address_qualifier(idx))
|
||||
}
|
||||
CL_KERNEL_ARG_NAME => cl_prop::<&str>(kernel.arg_name(idx)),
|
||||
CL_KERNEL_ARG_TYPE_NAME => cl_prop::<&str>(kernel.arg_type_name(idx)),
|
||||
CL_KERNEL_ARG_TYPE_QUALIFIER => {
|
||||
cl_prop::<cl_kernel_arg_type_qualifier>(kernel.type_qualifier(idx))
|
||||
}
|
||||
// CL_INVALID_VALUE if param_name is not one of the supported values
|
||||
_ => return Err(CL_INVALID_VALUE),
|
||||
})
|
||||
}
|
||||
}
|
||||
|
||||
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) -> CLResult<Vec<u8>> {
|
||||
let _kernel = self.get_ref()?;
|
||||
let _dev = dev.get_ref()?;
|
||||
Ok(match *q {
|
||||
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE => cl_prop::<usize>(1),
|
||||
// TODO
|
||||
CL_KERNEL_WORK_GROUP_SIZE => cl_prop::<usize>(1),
|
||||
// CL_INVALID_VALUE if param_name is not one of the supported values
|
||||
_ => return Err(CL_INVALID_VALUE),
|
||||
})
|
||||
}
|
||||
}
|
||||
|
||||
pub fn create_kernel(
|
||||
program: cl_program,
|
||||
kernel_name: *const ::std::os::raw::c_char,
|
||||
) -> CLResult<cl_kernel> {
|
||||
let _p = program.get_ref()?;
|
||||
let p = program.get_arc()?;
|
||||
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);
|
||||
}
|
||||
|
||||
println!("create_kernel not implemented");
|
||||
Err(CL_OUT_OF_HOST_MEMORY)
|
||||
// CL_INVALID_PROGRAM_EXECUTABLE if there is no successfully built executable for program.
|
||||
if p.kernels().is_empty() {
|
||||
return Err(CL_INVALID_PROGRAM_EXECUTABLE);
|
||||
}
|
||||
|
||||
//• CL_INVALID_PROGRAM_EXECUTABLE if there is no successfully built executable for program.
|
||||
//• CL_INVALID_KERNEL_NAME if kernel_name is not found in program.
|
||||
//• 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.
|
||||
// CL_INVALID_KERNEL_NAME if kernel_name is not found in program.
|
||||
if !p.kernels().contains(&name) {
|
||||
return Err(CL_INVALID_KERNEL_NAME);
|
||||
}
|
||||
|
||||
// CL_INVALID_PROGRAM_EXECUTABLE if there is no successfully built executable for program.
|
||||
let devs: Vec<_> = p
|
||||
.devs
|
||||
.iter()
|
||||
.filter(|d| p.status(d) == CL_BUILD_SUCCESS as cl_build_status)
|
||||
.collect();
|
||||
if devs.is_empty() {
|
||||
return Err(CL_INVALID_PROGRAM_EXECUTABLE);
|
||||
}
|
||||
|
||||
// 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.
|
||||
let kernel_args: HashSet<_> = devs.iter().map(|d| p.args(d, &name)).collect();
|
||||
if kernel_args.len() != 1 {
|
||||
return Err(CL_INVALID_KERNEL_DEFINITION);
|
||||
}
|
||||
|
||||
Ok(cl_kernel::from_arc(Kernel::new(
|
||||
name,
|
||||
p,
|
||||
kernel_args.into_iter().next().unwrap(),
|
||||
)))
|
||||
}
|
||||
|
||||
pub fn set_kernel_arg(
|
||||
kernel: cl_kernel,
|
||||
arg_index: cl_uint,
|
||||
_arg_size: usize,
|
||||
_arg_value: *const ::std::os::raw::c_void,
|
||||
) -> CLResult<()> {
|
||||
let k = kernel.get_arc()?;
|
||||
|
||||
// CL_INVALID_ARG_INDEX if arg_index is not a valid argument index.
|
||||
if arg_index as usize >= k.args.len() {
|
||||
return Err(CL_INVALID_ARG_INDEX);
|
||||
}
|
||||
|
||||
//• CL_INVALID_ARG_VALUE if arg_value specified is not a valid value.
|
||||
//• CL_INVALID_MEM_OBJECT for an argument declared to be a memory object when the specified arg_value is not a valid memory object.
|
||||
//• CL_INVALID_SAMPLER for an argument declared to be of type sampler_t when the specified arg_value is not a valid sampler object.
|
||||
//• 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_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).
|
||||
//• 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_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.
|
||||
|
||||
println!("set_kernel_arg not implemented");
|
||||
Err(CL_OUT_OF_HOST_MEMORY)
|
||||
}
|
||||
|
|
|
|||
|
|
@ -40,6 +40,7 @@ impl CLInfo<cl_program_info> for cl_program {
|
|||
.collect(),
|
||||
)
|
||||
}
|
||||
CL_PROGRAM_KERNEL_NAMES => cl_prop::<String>(prog.kernels().join(";")),
|
||||
CL_PROGRAM_NUM_DEVICES => cl_prop::<cl_uint>(prog.devs.len() as cl_uint),
|
||||
CL_PROGRAM_NUM_KERNELS => cl_prop::<usize>(prog.kernels().len()),
|
||||
CL_PROGRAM_REFERENCE_COUNT => cl_prop::<cl_uint>(self.refcnt()?),
|
||||
|
|
|
|||
|
|
@ -1,9 +1,13 @@
|
|||
extern crate mesa_rust;
|
||||
extern crate mesa_rust_gen;
|
||||
extern crate rusticl_opencl_gen;
|
||||
|
||||
use crate::api::icd::*;
|
||||
use crate::core::program::*;
|
||||
use crate::impl_cl_type_trait;
|
||||
|
||||
use self::mesa_rust::compiler::clc::*;
|
||||
use self::mesa_rust_gen::*;
|
||||
use self::rusticl_opencl_gen::*;
|
||||
|
||||
use std::sync::Arc;
|
||||
|
|
@ -11,14 +15,82 @@ use std::sync::Arc;
|
|||
#[repr(C)]
|
||||
pub struct Kernel {
|
||||
pub base: CLObjectBase<CL_INVALID_KERNEL>,
|
||||
pub prog: Arc<Program>,
|
||||
pub name: String,
|
||||
pub args: Vec<spirv::SPIRVKernelArg>,
|
||||
}
|
||||
|
||||
impl_cl_type_trait!(cl_kernel, Kernel, CL_INVALID_KERNEL);
|
||||
|
||||
impl Kernel {
|
||||
pub fn new() -> Arc<Kernel> {
|
||||
pub fn new(name: String, prog: Arc<Program>, args: Vec<spirv::SPIRVKernelArg>) -> Arc<Kernel> {
|
||||
Arc::new(Self {
|
||||
base: CLObjectBase::new(),
|
||||
prog: prog,
|
||||
name: name,
|
||||
args: args,
|
||||
})
|
||||
}
|
||||
|
||||
pub fn access_qualifier(&self, idx: cl_uint) -> cl_kernel_arg_access_qualifier {
|
||||
let aq = self.args[idx as usize].access_qualifier;
|
||||
|
||||
if aq
|
||||
== clc_kernel_arg_access_qualifier::CLC_KERNEL_ARG_ACCESS_READ
|
||||
| clc_kernel_arg_access_qualifier::CLC_KERNEL_ARG_ACCESS_WRITE
|
||||
{
|
||||
CL_KERNEL_ARG_ACCESS_READ_WRITE
|
||||
} else if aq == clc_kernel_arg_access_qualifier::CLC_KERNEL_ARG_ACCESS_READ {
|
||||
CL_KERNEL_ARG_ACCESS_READ_ONLY
|
||||
} else if aq == clc_kernel_arg_access_qualifier::CLC_KERNEL_ARG_ACCESS_WRITE {
|
||||
CL_KERNEL_ARG_ACCESS_WRITE_ONLY
|
||||
} else {
|
||||
CL_KERNEL_ARG_ACCESS_NONE
|
||||
}
|
||||
}
|
||||
|
||||
pub fn address_qualifier(&self, idx: cl_uint) -> cl_kernel_arg_address_qualifier {
|
||||
match self.args[idx as usize].address_qualifier {
|
||||
clc_kernel_arg_address_qualifier::CLC_KERNEL_ARG_ADDRESS_PRIVATE => {
|
||||
CL_KERNEL_ARG_ADDRESS_PRIVATE
|
||||
}
|
||||
clc_kernel_arg_address_qualifier::CLC_KERNEL_ARG_ADDRESS_CONSTANT => {
|
||||
CL_KERNEL_ARG_ADDRESS_CONSTANT
|
||||
}
|
||||
clc_kernel_arg_address_qualifier::CLC_KERNEL_ARG_ADDRESS_LOCAL => {
|
||||
CL_KERNEL_ARG_ADDRESS_LOCAL
|
||||
}
|
||||
clc_kernel_arg_address_qualifier::CLC_KERNEL_ARG_ADDRESS_GLOBAL => {
|
||||
CL_KERNEL_ARG_ADDRESS_GLOBAL
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
pub fn type_qualifier(&self, idx: cl_uint) -> cl_kernel_arg_type_qualifier {
|
||||
let tq = self.args[idx as usize].type_qualifier;
|
||||
let zero = clc_kernel_arg_type_qualifier(0);
|
||||
let mut res = CL_KERNEL_ARG_TYPE_NONE;
|
||||
|
||||
if tq & clc_kernel_arg_type_qualifier::CLC_KERNEL_ARG_TYPE_CONST != zero {
|
||||
res |= CL_KERNEL_ARG_TYPE_CONST;
|
||||
}
|
||||
|
||||
if tq & clc_kernel_arg_type_qualifier::CLC_KERNEL_ARG_TYPE_RESTRICT != zero {
|
||||
res |= CL_KERNEL_ARG_TYPE_RESTRICT;
|
||||
}
|
||||
|
||||
if tq & clc_kernel_arg_type_qualifier::CLC_KERNEL_ARG_TYPE_VOLATILE != zero {
|
||||
res |= CL_KERNEL_ARG_TYPE_VOLATILE;
|
||||
}
|
||||
|
||||
res.into()
|
||||
}
|
||||
|
||||
pub fn arg_name(&self, idx: cl_uint) -> &String {
|
||||
&self.args[idx as usize].name
|
||||
}
|
||||
|
||||
pub fn arg_type_name(&self, idx: cl_uint) -> &String {
|
||||
&self.args[idx as usize].type_name
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -107,6 +107,14 @@ impl Program {
|
|||
.clone()
|
||||
}
|
||||
|
||||
pub fn args(&self, dev: &Arc<Device>, kernel: &str) -> Vec<spirv::SPIRVKernelArg> {
|
||||
Self::dev_build_info(&mut self.build_info(), dev)
|
||||
.spirv
|
||||
.as_ref()
|
||||
.unwrap()
|
||||
.args(kernel)
|
||||
}
|
||||
|
||||
pub fn kernels(&self) -> Vec<String> {
|
||||
self.build_info().kernels.clone()
|
||||
}
|
||||
|
|
|
|||
|
|
@ -17,6 +17,15 @@ pub struct SPIRVBin {
|
|||
info: Option<clc_parsed_spirv>,
|
||||
}
|
||||
|
||||
#[derive(PartialEq, Eq, Hash)]
|
||||
pub struct SPIRVKernelArg {
|
||||
pub name: String,
|
||||
pub type_name: String,
|
||||
pub access_qualifier: clc_kernel_arg_access_qualifier,
|
||||
pub address_qualifier: clc_kernel_arg_address_qualifier,
|
||||
pub type_qualifier: clc_kernel_arg_type_qualifier,
|
||||
}
|
||||
|
||||
pub struct CLCHeader<'a> {
|
||||
pub name: CString,
|
||||
pub source: &'a CString,
|
||||
|
|
@ -122,16 +131,40 @@ impl SPIRVBin {
|
|||
(res, msgs.join("\n"))
|
||||
}
|
||||
|
||||
fn kernel_infos(&self) -> &[clc_kernel_info] {
|
||||
match self.info {
|
||||
None => &[],
|
||||
Some(info) => unsafe { slice::from_raw_parts(info.kernels, info.num_kernels as usize) },
|
||||
}
|
||||
}
|
||||
|
||||
fn kernel_info(&self, name: &str) -> Option<&clc_kernel_info> {
|
||||
self.kernel_infos()
|
||||
.iter()
|
||||
.find(|i| c_string_to_string(i.name) == name)
|
||||
}
|
||||
|
||||
pub fn kernels(&self) -> Vec<String> {
|
||||
unsafe {
|
||||
match self.info {
|
||||
None => Vec::new(),
|
||||
Some(info) => slice::from_raw_parts(info.kernels, info.num_kernels as usize)
|
||||
.iter()
|
||||
.map(|i| i.name)
|
||||
.map(c_string_to_string)
|
||||
.collect(),
|
||||
}
|
||||
self.kernel_infos()
|
||||
.iter()
|
||||
.map(|i| i.name)
|
||||
.map(c_string_to_string)
|
||||
.collect()
|
||||
}
|
||||
|
||||
pub fn args(&self, name: &str) -> Vec<SPIRVKernelArg> {
|
||||
match self.kernel_info(name) {
|
||||
None => Vec::new(),
|
||||
Some(info) => unsafe { slice::from_raw_parts(info.args, info.num_args) }
|
||||
.iter()
|
||||
.map(|a| SPIRVKernelArg {
|
||||
name: c_string_to_string(a.name),
|
||||
type_name: c_string_to_string(a.type_name),
|
||||
access_qualifier: clc_kernel_arg_access_qualifier(a.access_qualifier),
|
||||
address_qualifier: a.address_qualifier,
|
||||
type_qualifier: clc_kernel_arg_type_qualifier(a.type_qualifier),
|
||||
})
|
||||
.collect(),
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -127,6 +127,7 @@ rusticl_opencl_bindings_rs = rust.bindgen(
|
|||
'--whitelist-var', 'CL_.*',
|
||||
# some info types need to be strongly typed so we can implement various get_infos
|
||||
'--new-type-alias-deref', 'cl_(mem|image)_info',
|
||||
'--new-type-alias-deref', 'cl_kernel_(arg|work_group)_info',
|
||||
],
|
||||
)
|
||||
|
||||
|
|
@ -177,6 +178,10 @@ rusticl_mesa_bindings_rs = rust.bindgen(
|
|||
'--whitelist-function', 'clc_.*',
|
||||
'--whitelist-function', 'pipe_.*',
|
||||
'--whitelist-type', 'pipe_endian',
|
||||
'--whitelist-type', 'clc_kernel_arg_access_qualifier',
|
||||
'--bitfield-enum', 'clc_kernel_arg_access_qualifier',
|
||||
'--whitelist-type', 'clc_kernel_arg_type_qualifier',
|
||||
'--bitfield-enum', 'clc_kernel_arg_type_qualifier',
|
||||
'--whitelist-var', 'PIPE_.*',
|
||||
'--bitfield-enum', 'pipe_map_flags',
|
||||
],
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue