opencl3/
program.rs

1// Copyright (c) 2020-2025 Via Technology Ltd. All Rights Reserved.
2//
3// Licensed under the Apache License, Version 2.0 (the "License");
4// you may not use this file except in compliance with the License.
5// You may obtain a copy of the License at
6//
7//    http://www.apache.org/licenses/LICENSE-2.0
8//
9// Unless required by applicable law or agreed to in writing, software
10// distributed under the License is distributed on an "AS IS" BASIS,
11// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12// See the License for the specific language governing permissions and
13// limitations under the License.
14
15#![allow(clippy::missing_safety_doc)]
16
17pub use cl3::program::*;
18
19use super::context::Context;
20
21use super::Result;
22#[allow(unused_imports)]
23use cl3::error_codes::CL_BUILD_PROGRAM_FAILURE;
24#[allow(unused_imports)]
25use cl3::ext;
26#[allow(unused_imports)]
27use libc::{c_void, intptr_t, size_t};
28#[allow(unused_imports)]
29use std::ffi::{CStr, CString};
30use std::ptr;
31use std::result;
32
33// Compile, link and build options.
34// These options can be passed to Program::compile, Program::link or Program::build, see:
35// [Compiler Options](https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_API.html#compiler-options)
36// [Linker Options](https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_API.html#linker-options)
37// [Build Options](https://man.opencl.org/clBuildProgram.html)
38
39// Note: the options have a trailing space so that they can be concatenated.
40
41// Math Intrinsics Options
42pub const CL_SINGLE_RECISION_CONSTANT: &str = "-cl-single-precision-constant ";
43pub const CL_DENORMS_ARE_ZERO: &str = "-cl-denorms-are-zero ";
44pub const CL_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT: &str = "-cl-fp32-correctly-rounded-divide-sqrt ";
45
46// Optimization Options
47pub const CL_OPT_DISABLE: &str = "-cl-opt-disable ";
48pub const CL_STRICT_ALIASING: &str = "-cl-strict-aliasing ";
49pub const CL_UNIFORM_WORK_GROUP_SIZE: &str = "-cl-uniform-work-group-size ";
50pub const CL_NO_SUBGROUP_INFO: &str = "-cl-no-subgroup-ifp ";
51pub const CL_MAD_ENABLE: &str = "-cl-mad-enable ";
52pub const CL_NO_SIGNED_ZEROS: &str = "-cl-no-signed-zeros ";
53pub const CL_UNSAFE_MATH_OPTIMIZATIONS: &str = "-cl-unsafe-math-optimizations ";
54pub const CL_FINITE_MATH_ONLY: &str = "-cl-finite-math-only ";
55pub const CL_FAST_RELAXED_MATH: &str = "-cl-fast-relaxed-math ";
56
57// OpenCL C version Options
58
59/// Applications are required to specify the -cl-std=CL2.0 build option to
60/// compile or build programs with OpenCL C 2.0.
61pub const CL_STD_2_0: &str = "-cl-std=CL2.0 ";
62
63/// Applications are required to specify the -cl-std=CL3.0 build option to
64/// compile or build programs with OpenCL C 3.0.
65pub const CL_STD_3_0: &str = "-cl-std=CL3.0 ";
66
67/// This option allows the compiler to store information about the
68/// arguments of kernels in the program executable.
69pub const CL_KERNEL_ARG_INFO: &str = "-cl-kernel-arg-info ";
70
71pub const DEBUG_OPTION: &str = "-g ";
72
73// Options enabled by the cl_khr_spir extension
74pub const BUILD_OPTION_X_SPIR: &str = "-x spir ";
75pub const BUILD_OPTION_SPIR_STD_1_2: &str = "-spir-std=1.2 ";
76
77// Link and build options.
78pub const CREATE_LIBRARY: &str = "-create-library ";
79pub const ENABLE_LINK_OPTIONS: &str = "-enable-link-options ";
80
81/// An OpenCL program object.  
82/// Stores the names of the OpenCL kernels in the program.
83/// Implements the Drop trait to call release_program when the object is dropped.
84#[derive(Debug)]
85pub struct Program {
86    program: cl_program,
87    kernel_names: String,
88}
89
90impl From<Program> for cl_program {
91    fn from(value: Program) -> Self {
92        value.program as Self
93    }
94}
95
96impl Drop for Program {
97    fn drop(&mut self) {
98        unsafe { release_program(self.program).expect("Error: clReleaseProgram") };
99    }
100}
101
102unsafe impl Send for Program {}
103unsafe impl Sync for Program {}
104
105impl Program {
106    fn new(program: cl_program, kernel_names: &str) -> Self {
107        Self {
108            program,
109            kernel_names: kernel_names.to_owned(),
110        }
111    }
112
113    /// Get the underlying OpenCL cl_program.
114    pub const fn get(&self) -> cl_program {
115        self.program
116    }
117
118    /// Get the names of the OpenCL kernels in the Program, in a string
119    /// separated by semicolons.
120    #[allow(clippy::missing_const_for_fn)]
121    pub fn kernel_names(&self) -> &str {
122        &self.kernel_names
123    }
124
125    /// Create a Program for a context and load source code into that object.  
126    ///
127    /// * `context` - a valid OpenCL context.
128    /// * `sources` - an array of strs containing the source code strings.
129    ///
130    /// returns a Result containing the new Program
131    /// or the error code from the OpenCL C API function.
132    pub fn create_from_sources(context: &Context, sources: &[&str]) -> Result<Self> {
133        Ok(Self::new(
134            create_program_with_source(context.get(), sources)?,
135            "",
136        ))
137    }
138
139    /// Create a Program for a context and load a source code string into that object.  
140    ///
141    /// * `context` - a valid OpenCL context.
142    /// * `src` - a str containing a source code string.
143    ///
144    /// returns a Result containing the new Program
145    /// or the error code from the OpenCL C API function.
146    pub fn create_from_source(context: &Context, src: &str) -> Result<Self> {
147        let sources = [src];
148        Ok(Self::new(
149            create_program_with_source(context.get(), &sources)?,
150            "",
151        ))
152    }
153
154    /// Create a Program for a context and load binary bits into that object.  
155    ///
156    /// * `context` - a valid OpenCL context.
157    /// * `devices` - a slice of devices that are in context.
158    /// * `binaries` - a slice of program binaries slices.
159    ///
160    /// returns a Result containing the new Program
161    /// or the error code from the OpenCL C API function.
162    ///
163    /// # Safety
164    ///
165    /// This is unsafe when a device is not a member of context.
166    pub unsafe fn create_from_binary(
167        context: &Context,
168        devices: &[cl_device_id],
169        binaries: &[&[u8]],
170    ) -> Result<Self> {
171        unsafe {
172            Ok(Self::new(
173                create_program_with_binary(context.get(), devices, binaries)?,
174                "",
175            ))
176        }
177    }
178
179    /// Create a Program for a context and  loads the information related to
180    /// the built-in kernels into that object.  
181    ///
182    /// * `context` - a valid OpenCL context.
183    /// * `devices` - a slice of devices that are in context.
184    /// * `kernel_names` - a semi-colon separated list of built-in kernel names.
185    ///
186    /// returns a Result containing the new Program
187    /// or the error code from the OpenCL C API function.
188    ///
189    /// # Safety
190    ///
191    /// This is unsafe when a device is not a member of context.
192    #[cfg(any(feature = "CL_VERSION_1_2", feature = "dynamic"))]
193    pub unsafe fn create_from_builtin_kernels(
194        context: &Context,
195        devices: &[cl_device_id],
196        kernel_names: &str,
197    ) -> Result<Self> {
198        unsafe {
199            // Ensure options string is null terminated
200            let c_names = CString::new(kernel_names)
201                .expect("Program::create_from_builtin_kernels, invalid kernel_names");
202            Ok(Self::new(
203                create_program_with_builtin_kernels(context.get(), devices, &c_names)?,
204                kernel_names,
205            ))
206        }
207    }
208
209    /// Create a Program for a context and load code in an intermediate language
210    /// into that object.  
211    /// CL_VERSION_2_1
212    ///
213    /// * `context` - a valid OpenCL context.
214    /// * `il` - a slice of program intermediate language code.
215    ///
216    /// returns a Result containing the new Program
217    /// or the error code from the OpenCL C API function.
218    #[cfg(any(feature = "CL_VERSION_2_1", feature = "dynamic"))]
219    pub fn create_from_il(context: &Context, il: &[u8]) -> Result<Self> {
220        Ok(Self::new(create_program_with_il(context.get(), il)?, ""))
221    }
222
223    #[cfg(any(feature = "cl_khr_il_program", feature = "dynamic"))]
224    pub fn create_from_il_khr(context: &Context, il: &[u8]) -> Result<Self> {
225        Ok(Self::new(
226            ext::create_program_with_il_khr(context.get(), il)?,
227            "",
228        ))
229    }
230
231    /// Build (compile & link) a Program.
232    ///
233    /// * `devices` - a slice of devices that are in context.
234    /// * `options` - the build options in a null-terminated string.
235    /// * `pfn_notify` - an optional function pointer to a notification routine.
236    /// * `user_data` - passed as an argument when pfn_notify is called, or ptr::null_mut().
237    ///
238    /// returns a null Result
239    /// or the error code from the OpenCL C API function.
240    pub fn build(&mut self, devices: &[cl_device_id], options: &str) -> Result<()> {
241        // Ensure options string is null terminated
242        let c_options = CString::new(options).expect("Program::build, invalid options");
243        build_program(self.program, devices, &c_options, None, ptr::null_mut())?;
244        self.kernel_names = self.get_kernel_names()?;
245        Ok(())
246    }
247
248    /// Create and build an OpenCL Program from an array of source code strings
249    /// with the given options.  
250    ///
251    /// * `context` - a valid OpenCL context.
252    /// * `sources` - an array of strs containing the source code strings.
253    /// * `options` - the build options in a null-terminated string.
254    ///
255    /// returns a Result containing the new Program, the name of the error code
256    /// from the OpenCL C API function or the build log, if the build failed.
257    pub fn create_and_build_from_sources(
258        context: &Context,
259        sources: &[&str],
260        options: &str,
261    ) -> result::Result<Self, String> {
262        let mut program = Self::create_from_sources(context, sources).map_err(String::from)?;
263        match program.build(context.devices(), options) {
264            Ok(_) => Ok(program),
265            Err(e) => {
266                if CL_BUILD_PROGRAM_FAILURE == e.0 {
267                    let log = program
268                        .get_build_log(context.devices()[0])
269                        .map_err(String::from)?;
270                    Err(String::from(e) + ", build log: " + &log)
271                } else {
272                    Err(String::from(e))
273                }
274            }
275        }
276    }
277
278    /// Create and build an OpenCL Program from source code with the given options.  
279    ///
280    /// * `context` - a valid OpenCL context.
281    /// * `src` - a str containing a source code string.
282    /// * `options` - the build options in a null-terminated string.
283    ///
284    /// returns a Result containing the new Program, the name of the error code
285    /// from the OpenCL C API function or the build log, if the build failed.
286    pub fn create_and_build_from_source(
287        context: &Context,
288        src: &str,
289        options: &str,
290    ) -> result::Result<Self, String> {
291        let sources = [src];
292        Self::create_and_build_from_sources(context, &sources, options)
293    }
294
295    /// Create and build an OpenCL Program from binaries with the given options.  
296    ///
297    /// * `context` - a valid OpenCL context.
298    /// * `binaries` - a slice of program binaries slices.
299    /// * `options` - the build options in a null-terminated string.
300    ///
301    /// returns a Result containing the new Program
302    /// or the error code from the OpenCL C API function.
303    pub fn create_and_build_from_binary(
304        context: &Context,
305        binaries: &[&[u8]],
306        options: &str,
307    ) -> Result<Self> {
308        let mut program =
309            unsafe { Self::create_from_binary(context, context.devices(), binaries)? };
310        program.build(context.devices(), options)?;
311        Ok(program)
312    }
313
314    /// Create and build an OpenCL Program from intermediate language with the
315    /// given options.  
316    /// CL_VERSION_2_1
317    ///
318    /// * `context` - a valid OpenCL context.
319    /// * `il` - a slice of program intermediate language code.
320    /// * `options` - the build options in a null-terminated string.
321    ///
322    /// returns a Result containing the new `Program`
323    /// or the error code from the OpenCL C API function.
324    #[cfg(any(feature = "CL_VERSION_2_1", feature = "dynamic"))]
325    pub fn create_and_build_from_il(context: &Context, il: &[u8], options: &str) -> Result<Self> {
326        let mut program = Self::create_from_il(context, il)?;
327        program.build(context.devices(), options)?;
328        Ok(program)
329    }
330
331    /// Compile a program’s source for the devices the OpenCL context associated
332    /// with the program.
333    ///
334    /// * `devices` - a slice of devices that are in context.
335    /// * `options` - the compilation options in a null-terminated string.
336    /// * `input_headers` - a slice of programs that describe headers in the input_headers.
337    /// * `header_include_names` - an array that has a one to one correspondence with
338    ///   input_headers.
339    ///
340    /// returns a null Result
341    /// or the error code from the OpenCL C API function.
342    #[cfg(any(feature = "CL_VERSION_1_2", feature = "dynamic"))]
343    pub fn compile(
344        &mut self,
345        devices: &[cl_device_id],
346        options: &str,
347        input_headers: &[cl_program],
348        header_include_names: &[&CStr],
349    ) -> Result<()> {
350        // Ensure options string is null terminated
351        let c_options = CString::new(options).expect("Program::compile, invalid options");
352        Ok(compile_program(
353            self.program,
354            devices,
355            &c_options,
356            input_headers,
357            header_include_names,
358            None,
359            ptr::null_mut(),
360        )?)
361    }
362
363    /// Link a set of compiled program objects and libraries for the devices in the
364    /// OpenCL context associated with the program.
365    ///
366    /// * `devices` - a slice of devices that are in context.
367    /// * `options` - the link options in a null-terminated string.
368    /// * `input_programs` - a slice of programs that describe headers in the input_headers.
369    ///
370    /// returns a null Result
371    /// or the error code from the OpenCL C API function.
372    ///
373    /// # Safety
374    ///
375    /// This is unsafe when a device is not a member of context.
376    #[cfg(any(feature = "CL_VERSION_1_2", feature = "dynamic"))]
377    pub unsafe fn link(
378        &mut self,
379        devices: &[cl_device_id],
380        options: &str,
381        input_programs: &[cl_program],
382    ) -> Result<()> {
383        unsafe {
384            // Ensure options string is null terminated
385            let c_options = CString::new(options).expect("Program::link, invalid options");
386            self.program = link_program(
387                self.program,
388                devices,
389                &c_options,
390                input_programs,
391                None,
392                ptr::null_mut(),
393            )?;
394            self.kernel_names = self.get_kernel_names()?;
395            Ok(())
396        }
397    }
398
399    /// Set the value of a specialization constant.
400    /// CL_VERSION_2_2
401    ///
402    /// * `spec_id` - the specialization constant whose value will be set.
403    /// * `spec_size` - size in bytes of the data pointed to by spec_value.
404    /// * `spec_value` - pointer to the memory location that contains the value
405    ///   of the specialization constant.
406    ///
407    /// returns an empty Result or the error code from the OpenCL C API function.
408    #[cfg(any(feature = "CL_VERSION_2_2", feature = "dynamic"))]
409    pub unsafe fn set_specialization_constant(
410        &self,
411        spec_id: cl_uint,
412        spec_size: size_t,
413        spec_value: *const c_void,
414    ) -> Result<()> {
415        unsafe {
416            Ok(set_program_specialization_constant(
417                self.program,
418                spec_id,
419                spec_size,
420                spec_value,
421            )?)
422        }
423    }
424
425    pub fn get_reference_count(&self) -> Result<cl_uint> {
426        Ok(get_program_info(self.program, CL_PROGRAM_REFERENCE_COUNT)?.into())
427    }
428
429    pub fn get_context(&self) -> Result<cl_context> {
430        Ok(intptr_t::from(get_program_info(self.program, CL_PROGRAM_CONTEXT)?) as cl_context)
431    }
432
433    pub fn get_num_devices(&self) -> Result<cl_uint> {
434        Ok(get_program_info(self.program, CL_PROGRAM_NUM_DEVICES)?.into())
435    }
436
437    pub fn get_devices(&self) -> Result<Vec<intptr_t>> {
438        Ok(get_program_info(self.program, CL_PROGRAM_DEVICES)?.into())
439    }
440
441    pub fn get_source(&self) -> Result<String> {
442        Ok(get_program_info(self.program, CL_PROGRAM_SOURCE)?.into())
443    }
444
445    pub fn get_binary_sizes(&self) -> Result<Vec<size_t>> {
446        Ok(get_program_info(self.program, CL_PROGRAM_BINARY_SIZES)?.into())
447    }
448
449    pub fn get_binaries(&self) -> Result<Vec<Vec<cl_uchar>>> {
450        Ok(get_program_info(self.program, CL_PROGRAM_BINARIES)?.into())
451    }
452
453    pub fn get_num_kernels(&self) -> Result<size_t> {
454        Ok(get_program_info(self.program, CL_PROGRAM_NUM_KERNELS)?.into())
455    }
456
457    pub fn get_kernel_names(&self) -> Result<String> {
458        Ok(get_program_info(self.program, CL_PROGRAM_KERNEL_NAMES)?.into())
459    }
460
461    /// CL_VERSION_2_1
462    pub fn get_program_il(&self) -> Result<String> {
463        Ok(get_program_info(self.program, CL_PROGRAM_IL)?.into())
464    }
465
466    /// CL_VERSION_2_2
467    pub fn get_program_scope_global_ctors_present(&self) -> Result<bool> {
468        Ok(cl_uint::from(get_program_info(
469            self.program,
470            CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT,
471        )?) != CL_FALSE)
472    }
473
474    /// CL_VERSION_2_2
475    pub fn get_program_scope_global_dtors_present(&self) -> Result<bool> {
476        Ok(cl_uint::from(get_program_info(
477            self.program,
478            CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT,
479        )?) != CL_FALSE)
480    }
481
482    pub fn get_build_status(&self, device: cl_device_id) -> Result<cl_int> {
483        Ok(get_program_build_info(self.program, device, CL_PROGRAM_BUILD_STATUS)?.into())
484    }
485
486    pub fn get_build_options(&self, device: cl_device_id) -> Result<String> {
487        Ok(get_program_build_info(self.program, device, CL_PROGRAM_BUILD_OPTIONS)?.into())
488    }
489
490    pub fn get_build_log(&self, device: cl_device_id) -> Result<String> {
491        Ok(get_program_build_info(self.program, device, CL_PROGRAM_BUILD_LOG)?.into())
492    }
493
494    pub fn get_build_binary_type(&self, device: cl_device_id) -> Result<cl_uint> {
495        Ok(get_program_build_info(self.program, device, CL_PROGRAM_BINARY_TYPE)?.into())
496    }
497
498    /// CL_VERSION_2_0
499    pub fn get_build_global_variable_total_size(&self, device: cl_device_id) -> Result<size_t> {
500        Ok(get_program_build_info(
501            self.program,
502            device,
503            CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE,
504        )?
505        .into())
506    }
507}
508
509#[cfg(test)]
510mod tests {
511    use super::*;
512    use crate::context::Context;
513    use crate::device::Device;
514    use crate::platform::get_platforms;
515    use cl3::device::CL_DEVICE_TYPE_GPU;
516    use std::collections::HashSet;
517
518    const PROGRAM_SOURCE: &str = r#"
519        kernel void add(global float* buffer, float scalar) {
520            buffer[get_global_id(0)] += scalar;
521        }
522
523        kernel void subtract(global float* buffer, float scalar) {
524            buffer[get_global_id(0)] -= scalar;
525        }
526    "#;
527
528    #[test]
529    fn test_create_and_build_from_source() {
530        let platforms = get_platforms().unwrap();
531        assert!(0 < platforms.len());
532
533        // Get the first platform
534        let platform = &platforms[0];
535
536        let devices = platform.get_devices(CL_DEVICE_TYPE_GPU).unwrap();
537        assert!(0 < devices.len());
538
539        // Get the first device
540        let device = Device::new(devices[0]);
541        let context = Context::from_device(&device).unwrap();
542
543        let program =
544            Program::create_and_build_from_source(&context, PROGRAM_SOURCE, CL_DENORMS_ARE_ZERO)
545                .expect("Program::create_and_build_from_source failed");
546
547        let names: HashSet<&str> = program.kernel_names().split(';').collect();
548        println!("OpenCL Program kernel_names len: {}", names.len());
549        println!("OpenCL Program kernel_names: {:?}", names);
550
551        let value = program.get_reference_count().unwrap();
552        println!("program.get_reference_count(): {}", value);
553        assert_eq!(1, value);
554
555        let value = program.get_context().unwrap();
556        assert!(context.get() == value);
557
558        let value = program.get_num_devices().unwrap();
559        println!("program.get_num_devices(): {}", value);
560        assert_eq!(1, value);
561
562        let value = program.get_devices().unwrap();
563        assert!(device.id() == value[0] as cl_device_id);
564
565        let value = program.get_source().unwrap();
566        println!("program.get_source(): {}", value);
567        assert!(!value.is_empty());
568
569        let value = program.get_binary_sizes().unwrap();
570        println!("program.get_binary_sizes(): {:?}", value);
571        assert!(0 < value[0]);
572
573        let value = program.get_binaries().unwrap();
574        // println!("program.get_binaries(): {:?}", value);
575        assert!(!value[0].is_empty());
576
577        let value = program.get_num_kernels().unwrap();
578        println!("program.get_num_kernels(): {}", value);
579        assert_eq!(2, value);
580
581        // let value = program.get_program_il().unwrap();
582        // println!("program.get_program_il(): {:?}", value);
583        // assert!(!value.is_empty());
584
585        let value = program.get_build_status(device.id()).unwrap();
586        println!("program.get_build_status(): {}", value);
587        assert!(CL_BUILD_SUCCESS == value);
588
589        let value = program.get_build_options(device.id()).unwrap();
590        println!("program.get_build_options(): {}", value);
591        assert!(!value.is_empty());
592
593        let value = program.get_build_log(device.id()).unwrap();
594        println!("program.get_build_log(): {}", value);
595        // assert!(!value.is_empty());
596
597        let value = program.get_build_binary_type(device.id()).unwrap();
598        println!("program.get_build_binary_type(): {}", value);
599        assert_eq!(CL_PROGRAM_BINARY_TYPE_EXECUTABLE as u32, value);
600
601        // CL_VERSION_2_0 value
602        match program.get_build_global_variable_total_size(device.id()) {
603            Ok(value) => println!("program.get_build_global_variable_total_size(): {}", value),
604            Err(e) => println!(
605                "OpenCL error, program.get_build_global_variable_total_size(): {}",
606                e
607            ),
608        };
609    }
610}
opencl3/program.rs

opencl3/
program.rs
