From 8cb22f070dada329990b778e5d79bae3b01ec18b Mon Sep 17 00:00:00 2001 From: "github-actions[bot]" <41898282+github-actions[bot]@users.noreply.github.com> Date: Tue, 22 Sep 2026 18:17:56 +0000 Subject: [PATCH 1/2] Update generated OpenCL bindings --- API_DESIGN.md | 7 +- GENERATOR_COMMIT | 2 +- OWNERSHIP.md | 30 ++--- convenience_test.v | 138 +++++++++++++++----- event.v | 37 ++++-- examples/image_svm/main.v | 32 ++--- examples/vector_add/main.v | 16 +-- examples/vulkan_particles/compute.v | 43 +++--- examples/vulkan_particles/live_sync.v | 35 ++--- examples/vulkan_particles/main.v | 5 +- examples/vulkan_particles/particle_buffer.v | 33 ++--- examples/vulkan_particles/window_smoke.v | 8 +- examples/vulkan_particles/zero_copy_smoke.v | 48 +++---- external_interop.v | 89 ++++++++----- image.v | 54 ++++++-- ownership.v | 69 ++++++++-- program.v | 47 +++++-- svm.v | 17 +-- test/pointer_abi_test.v | 8 +- 19 files changed, 473 insertions(+), 245 deletions(-) diff --git a/API_DESIGN.md b/API_DESIGN.md index da9386f..88cc43f 100644 --- a/API_DESIGN.md +++ b/API_DESIGN.md @@ -7,10 +7,13 @@ can continue calling the generated functions directly. ## Shared conventions - Preserve native result codes in typed errors and include the failed operation. -- Return `!T` from convenience operations instead of discarding status codes. +- Return `!&T` for owning convenience operations instead of discarding status + codes or copying resource wrappers across module boundaries. - Hide count-then-fill enumeration without hiding the returned native handles. - Use `close()` for reference-counted OpenCL ownership wrappers. -- Owning wrappers are copyable V values; copying does not transfer ownership. Exactly one copy may close the native handle until a future reference-backed ownership redesign. +- Owning wrappers are `@[nocopy]` and constructors return owned pointers. + `clone_ref()` is the only supported way to create another native reference; + SVM remains uniquely owned because OpenCL provides no retain operation. - Keep constructors explicit about device choice and requested capabilities. - Never enable an extension or feature merely because headers declare it; query runtime support. - Keep unsafe pointers at the low-level boundary and expose slices or strings where ownership is clear. diff --git a/GENERATOR_COMMIT b/GENERATOR_COMMIT index 588f89b..dfaa0ca 100644 --- a/GENERATOR_COMMIT +++ b/GENERATOR_COMMIT @@ -1 +1 @@ -ebbdf9433cd3ee2d7ec416bbe079634ed32b0174 +185ac1cdff29e24c39a1fb0dffdbe21c184a31ae diff --git a/OWNERSHIP.md b/OWNERSHIP.md index d14dd0c..ffa852e 100644 --- a/OWNERSHIP.md +++ b/OWNERSHIP.md @@ -3,25 +3,25 @@ The convenience API uses explicit `close()` methods. Close events, kernels, programs, memory objects, and queues before their parent context. -V structs are values and can be copied. Each copy of an `OwnedContext`, -`OwnedCommandQueue`, `Buffer`, `Image2D`, `OwnedSampler`, `SvmAllocation`, -`OwnedProgram`, `OwnedKernel`, `OwnedEvent`, or external-interoperability owner -refers to the same OpenCL resource. Closing one value clears that value's -handle, but it does not clear copies made earlier. +`OwnedContext`, `OwnedCommandQueue`, `Buffer`, `Image2D`, `OwnedSampler`, +`SvmAllocation`, `OwnedProgram`, `OwnedKernel`, `OwnedEvent`, and +`OwnedExternalSemaphore` are `@[nocopy]`. Constructors return owned pointers so +resources cross module boundaries without copying. Pass those pointers directly +to convenience functions; do not add another `&`. -Until a breaking ownership redesign, follow these rules: +`close()` is mutable and idempotent: it releases one native reference and +clears the wrapper's handle. There are no implicit finalizers. For OpenCL +objects which support native reference counting, `clone_ref()` performs the +matching `clRetain*` operation and returns a separate owned pointer which must +also be closed. SVM has no retain operation and remains uniquely owned. -1. Treat owning wrappers as move-only by convention. -2. Pass references or raw handles to helpers instead of copying owners. -3. Register cleanup immediately and close children before parents. -4. Never close more than one copy of the same owned reference. +1. Keep each returned owner pointer and register cleanup immediately. +2. Pass owner pointers directly; copy only raw handles and discovery values. +3. Close children before parents. +4. Use `clone_ref()` only when a second independently retained reference is + required, and close both owners. For coarse-grained SVM, call `map()` before accessing `SvmAllocation.handle` from the host and `unmap()` before submitting device work. Wait for the unmap event before using the allocation from a kernel. Blocking and asynchronous `write`/`read` helpers perform OpenCL SVM copies and do not expose host access. - -The intended post-`0.x` design is a reference-backed control block. All copies -would observe one closed state and the native reference would be released at -most once. Explicit retain/clone operations would remain separate when the -caller actually wants another OpenCL reference. diff --git a/convenience_test.v b/convenience_test.v index 179a342..7c5aa30 100644 --- a/convenience_test.v +++ b/convenience_test.v @@ -23,6 +23,78 @@ fn test_generated_error_name_covers_full_core_range() { assert cl.error_code_name(cl.invalid_event_wait_list) == 'invalid_event_wait_list' } +fn context_reference_count(handle cl.Context) !u32 { + mut count := u32(0) + cl.check(cl.get_context_info(handle, cl.context_reference_count, sizeof(count), &count, + unsafe { nil }), 'query OpenCL context reference count')! + return count +} + +fn queue_reference_count(handle cl.CommandQueue) !u32 { + mut count := u32(0) + cl.check(cl.get_command_queue_info(handle, cl.queue_reference_count, sizeof(count), &count, + unsafe { nil }), 'query OpenCL queue reference count')! + return count +} + +fn memory_reference_count(handle cl.Mem) !u32 { + mut count := u32(0) + cl.check(cl.get_mem_object_info(handle, cl.mem_reference_count, sizeof(count), &count, + unsafe { nil }), 'query OpenCL memory reference count')! + return count +} + +fn event_reference_count(handle cl.Event) !u32 { + mut count := u32(0) + cl.check(cl.get_event_info(handle, cl.event_reference_count, sizeof(count), &count, + unsafe { nil }), 'query OpenCL event reference count')! + return count +} + +fn test_clone_ref_retains_independently_owned_native_references() ! { + available_platforms := cl.platforms()! + if available_platforms.len == 0 { + return + } + available_devices := cl.devices(available_platforms[0], cl.device_type_all)! + if available_devices.len == 0 { + return + } + device := available_devices[0] + mut context := cl.new_context(device)! + context_refs := context_reference_count(context.handle)! + mut retained_context := context.clone_ref()! + assert context_reference_count(context.handle)! == context_refs + 1 + retained_context.close()! + assert context_reference_count(context.handle)! == context_refs + + mut queue := context.command_queue(device, cl.CommandQueueProperties(0))! + queue_refs := queue_reference_count(queue.handle)! + mut retained_queue := queue.clone_ref()! + assert queue_reference_count(queue.handle)! == queue_refs + 1 + retained_queue.close()! + assert queue_reference_count(queue.handle)! == queue_refs + + mut buffer := cl.new_buffer[u32](context, cl.mem_read_write, 4)! + buffer_refs := memory_reference_count(buffer.handle)! + mut retained_buffer := buffer.clone_ref()! + assert memory_reference_count(buffer.handle)! == buffer_refs + 1 + retained_buffer.close()! + assert memory_reference_count(buffer.handle)! == buffer_refs + + mut event := queue.marker([]cl.Event{})! + event_refs := event_reference_count(event.handle)! + mut retained_event := event.clone_ref()! + assert event_reference_count(event.handle)! == event_refs + 1 + retained_event.close()! + assert event_reference_count(event.handle)! == event_refs + + event.close()! + buffer.close()! + queue.close()! + context.close()! +} + fn test_enqueue_1d_after_validates_handles_and_global_size_before_opencl_call() { mut kernel_storage := u8(0) mut queue_storage := u8(0) @@ -68,7 +140,7 @@ fn test_typed_buffer_rejects_byte_size_overflow_before_opencl_call() { context := cl.OwnedContext{ handle: cl.Context(&context_storage) } - cl.new_buffer[u64](&context, cl.mem_read_write, max_int) or { + cl.new_buffer[u64](context, cl.mem_read_write, max_int) or { assert err is cl.OpenCLError if err is cl.OpenCLError { assert err.status == cl.invalid_buffer_size @@ -87,7 +159,7 @@ fn test_typed_image_rejects_mismatched_pixel_layout_before_opencl_call() { image_channel_order: cl.rgba image_channel_data_type: cl.unorm_int8 } - cl.new_image_2d[u8](&context, cl.mem_read_write, format, 2, 2) or { + cl.new_image_2d[u8](context, cl.mem_read_write, format, 2, 2) or { assert err is cl.OpenCLError if err is cl.OpenCLError { assert err.status == cl.invalid_image_format_descriptor @@ -124,7 +196,7 @@ fn test_typed_image_rejects_out_of_bounds_region_before_opencl_call() { queue := cl.OwnedCommandQueue{ handle: cl.CommandQueue(&queue_storage) } - image.write_region(&queue, 1, 0, 2, 2, [u32(1), 2, 3, 4]) or { + image.write_region(queue, 1, 0, 2, 2, [u32(1), 2, 3, 4]) or { assert err is cl.OpenCLError if err is cl.OpenCLError { assert err.status == cl.invalid_value @@ -143,7 +215,7 @@ fn test_svm_rejects_byte_size_overflow_before_opencl_call() { context := cl.OwnedContext{ handle: cl.Context(&context_storage) } - cl.new_svm[u64](&context, cl.mem_read_write, max_int, 0) or { + cl.new_svm[u64](context, cl.mem_read_write, max_int, 0) or { assert err is cl.OpenCLError if err is cl.OpenCLError { assert err.status == cl.invalid_buffer_size @@ -163,7 +235,7 @@ fn test_svm_rejects_out_of_bounds_transfer_before_opencl_call() { queue := cl.OwnedCommandQueue{ handle: cl.CommandQueue(&queue_storage) } - allocation.write(&queue, 3, [u32(1), 2]) or { + allocation.write(queue, 3, [u32(1), 2]) or { assert err is cl.OpenCLError if err is cl.OpenCLError { assert err.status == cl.invalid_value @@ -185,7 +257,7 @@ fn test_external_buffer_rejects_byte_size_overflow_before_opencl_call() { handle: cl.Context(&context_storage) } interop := cl.ExternalMemoryInterop{} - interop.import_opaque_fd_buffer[u64](&context, 0, max_int, cl.mem_read_write) or { + interop.import_opaque_fd_buffer[u64](context, 0, max_int, cl.mem_read_write) or { assert err is cl.OpenCLError if err is cl.OpenCLError { assert err.status == cl.invalid_buffer_size @@ -283,10 +355,10 @@ fn test_typed_buffer_round_trip() ! { } mut context := cl.new_context(available_devices[0])! mut queue := context.command_queue(available_devices[0], cl.CommandQueueProperties(0))! - mut buffer := cl.new_buffer[u32](&context, cl.mem_read_write, 4)! - buffer.write(&queue, 0, [u32(3), 5, 8, 13])! + mut buffer := cl.new_buffer[u32](context, cl.mem_read_write, 4)! + buffer.write(queue, 0, [u32(3), 5, 8, 13])! mut result := []u32{len: 4} - buffer.read(&queue, 0, mut result)! + buffer.read(queue, 0, mut result)! assert result == [u32(3), 5, 8, 13] buffer.close()! queue.close()! @@ -303,7 +375,7 @@ fn test_typed_image_round_trip_and_sampler_lifecycle() ! { return } mut context := cl.new_context(available_devices[0])! - formats := cl.supported_image2d_formats(&context, cl.mem_read_write)! + formats := cl.supported_image2d_formats(context, cl.mem_read_write)! format := cl.ImageFormat{ image_channel_order: cl.rgba image_channel_data_type: cl.unorm_int8 @@ -321,22 +393,22 @@ fn test_typed_image_round_trip_and_sampler_lifecycle() ! { return } mut queue := context.command_queue(available_devices[0], cl.CommandQueueProperties(0))! - mut source_image := cl.new_image_2d[u32](&context, cl.mem_read_only, format, 2, 2)! - mut destination_image := cl.new_image_2d[u32](&context, cl.mem_write_only, format, 2, 2)! - mut sampler := cl.new_sampler(&context, false, cl.address_clamp_to_edge, cl.filter_nearest)! - mut program := cl.build_source_program(&context, available_devices[0], + mut source_image := cl.new_image_2d[u32](context, cl.mem_read_only, format, 2, 2)! + mut destination_image := cl.new_image_2d[u32](context, cl.mem_write_only, format, 2, 2)! + mut sampler := cl.new_sampler(context, false, cl.address_clamp_to_edge, cl.filter_nearest)! + mut program := cl.build_source_program(context, available_devices[0], '__kernel void copy_image(read_only image2d_t source, write_only image2d_t destination, sampler_t image_sampler) { int2 p = (int2)(get_global_id(0), get_global_id(1)); write_imagef(destination, p, read_imagef(source, image_sampler, p)); }', '')! mut kernel := program.kernel('copy_image')! - source_image.set_kernel_arg(&kernel, 0)! - destination_image.set_kernel_arg(&kernel, 1)! - kernel.set_sampler_arg(2, &sampler)! + source_image.set_kernel_arg(kernel, 0)! + destination_image.set_kernel_arg(kernel, 1)! + kernel.set_sampler_arg(2, sampler)! values := [u32(0xff0000ff), 0xff00ff00, 0xffff0000, 0xffffffff] - mut uploaded := source_image.write_async(&queue, values, []cl.Event{})! - mut copied := kernel.enqueue_nd_after(&queue, [usize(2), 2], []usize{}, [ + mut uploaded := source_image.write_async(queue, values, []cl.Event{})! + mut copied := kernel.enqueue_nd_after(queue, [usize(2), 2], []usize{}, [ uploaded.handle, ])! mut result := []u32{len: 4} - mut downloaded := destination_image.read_async(&queue, mut result, [ + mut downloaded := destination_image.read_async(queue, mut result, [ copied.handle, ])! downloaded.wait()! @@ -372,21 +444,21 @@ fn test_svm_kernel_round_trip() ! { } mut context := cl.new_context(device)! mut queue := context.command_queue(device, cl.CommandQueueProperties(0))! - mut allocation := cl.new_svm[u32](&context, cl.mem_read_write, 4, 0)! - allocation.map(&queue, cl.map_write)! - mut unmapped := allocation.unmap(&queue, []cl.Event{})! + mut allocation := cl.new_svm[u32](context, cl.mem_read_write, 4, 0)! + allocation.map(queue, cl.map_write)! + mut unmapped := allocation.unmap(queue, []cl.Event{})! unmapped.wait()! - mut program := cl.build_source_program(&context, device, + mut program := cl.build_source_program(context, device, '__kernel void add(__global uint *values, uint amount) { values[get_global_id(0)] += amount; }', '')! mut kernel := program.kernel('add')! - allocation.set_kernel_arg(&kernel, 0)! + allocation.set_kernel_arg(kernel, 0)! amount := u32(7) kernel.set_arg(1, &amount)! values := [u32(1), 2, 3, 4] - mut uploaded := allocation.write_async(&queue, 0, values, [unmapped.handle])! - mut computed := kernel.enqueue_1d_after(&queue, 4, 0, [uploaded.handle])! + mut uploaded := allocation.write_async(queue, 0, values, [unmapped.handle])! + mut computed := kernel.enqueue_1d_after(queue, 4, 0, [uploaded.handle])! mut result := []u32{len: 4} - mut downloaded := allocation.read_async(&queue, 0, mut result, [computed.handle])! + mut downloaded := allocation.read_async(queue, 0, mut result, [computed.handle])! downloaded.wait()! assert result == [u32(8), 9, 10, 11] downloaded.close()! @@ -413,20 +485,20 @@ fn test_program_kernel_and_typed_argument() ! { device := available_devices[0] mut context := cl.new_context(device)! mut queue := context.command_queue(device, cl.queue_profiling_enable)! - mut buffer := cl.new_buffer[u32](&context, cl.mem_read_write, 4)! - mut program := cl.build_source_program(&context, device, + mut buffer := cl.new_buffer[u32](context, cl.mem_read_write, 4)! + mut program := cl.build_source_program(context, device, '__kernel void add(__global uint *values, uint amount) { size_t index = get_global_id(0) + get_global_size(0) * get_global_id(1); values[index] += amount; }', '')! mut kernel := program.kernel('add')! kernel.set_buffer_arg(0, buffer.handle)! amount := u32(7) kernel.set_arg(1, &amount)! values := [u32(1), 2, 3, 4] - mut write_event := buffer.write_async(&queue, 0, values, []cl.Event{})! - mut kernel_event := kernel.enqueue_nd_after(&queue, [usize(2), 2], []usize{}, [ + mut write_event := buffer.write_async(queue, 0, values, []cl.Event{})! + mut kernel_event := kernel.enqueue_nd_after(queue, [usize(2), 2], []usize{}, [ write_event.handle, ])! mut result := []u32{len: 4} - mut read_event := buffer.read_async(&queue, 0, mut result, [kernel_event.handle])! + mut read_event := buffer.read_async(queue, 0, mut result, [kernel_event.handle])! read_event.wait()! assert result == [u32(8), 9, 10, 11] profile := read_event.profile()! diff --git a/event.v b/event.v index 221fc2f..70ad52e 100644 --- a/event.v +++ b/event.v @@ -2,11 +2,26 @@ module opencl // OwnedEvent owns one OpenCL event reference. The command queue and context // which created it must remain valid until the event has completed. +@[nocopy] pub struct OwnedEvent { pub mut: handle Event } +// clone_ref retains the native event and returns an independently owned reference. +pub fn (event &OwnedEvent) clone_ref() !&OwnedEvent { + if isnil(event.handle) { + return OpenCLError{ + operation: 'retain closed OpenCL event' + status: invalid_event + } + } + check(retain_event(event.handle), 'retain OpenCL event')! + return &OwnedEvent{ + handle: event.handle + } +} + // EventProfile contains device timestamps in nanoseconds. Values are valid for // commands from queues created with queue_profiling_enable. pub struct EventProfile { @@ -27,7 +42,7 @@ pub fn (event &OwnedEvent) wait() ! { if isnil(event.handle) { return OpenCLError{ operation: 'wait for closed OpenCL event' - status: invalid_event + status: invalid_event } } check(wait_for_events(1, &event.handle), 'wait for OpenCL event')! @@ -39,7 +54,7 @@ pub fn (event &OwnedEvent) execution_status() !i32 { if isnil(event.handle) { return OpenCLError{ operation: 'query closed OpenCL event' - status: invalid_event + status: invalid_event } } mut status := i32(0) @@ -52,7 +67,7 @@ pub fn (event &OwnedEvent) profiling_timestamp(parameter ProfilingInfo) !u64 { if isnil(event.handle) { return OpenCLError{ operation: 'profile closed OpenCL event' - status: invalid_event + status: invalid_event } } mut timestamp := u64(0) @@ -68,8 +83,8 @@ pub fn (event &OwnedEvent) profile() !EventProfile { return EventProfile{ queued: event.profiling_timestamp(profiling_command_queued)! submit: event.profiling_timestamp(profiling_command_submit)! - start: event.profiling_timestamp(profiling_command_start)! - end: event.profiling_timestamp(profiling_command_end)! + start: event.profiling_timestamp(profiling_command_start)! + end: event.profiling_timestamp(profiling_command_end)! } } @@ -84,11 +99,11 @@ pub fn (mut event OwnedEvent) close() ! { // marker enqueues a marker after every event in wait_events and returns its // completion event. An empty wait list depends on earlier commands in queue. -pub fn (queue &OwnedCommandQueue) marker(wait_events []Event) !OwnedEvent { +pub fn (queue &OwnedCommandQueue) marker(wait_events []Event) !&OwnedEvent { if isnil(queue.handle) { return OpenCLError{ operation: 'enqueue marker on closed OpenCL queue' - status: invalid_command_queue + status: invalid_command_queue } } mut wait_pointer := &Event(unsafe { nil }) @@ -97,18 +112,18 @@ pub fn (queue &OwnedCommandQueue) marker(wait_events []Event) !OwnedEvent { } mut handle := Event(unsafe { nil }) check(enqueue_marker_with_wait_list(queue.handle, u32(wait_events.len), wait_pointer, &handle), 'enqueue OpenCL marker')! - return OwnedEvent{ + return &OwnedEvent{ handle: handle } } // barrier enqueues a barrier after every event in wait_events and returns its // completion event. -pub fn (queue &OwnedCommandQueue) barrier(wait_events []Event) !OwnedEvent { +pub fn (queue &OwnedCommandQueue) barrier(wait_events []Event) !&OwnedEvent { if isnil(queue.handle) { return OpenCLError{ operation: 'enqueue barrier on closed OpenCL queue' - status: invalid_command_queue + status: invalid_command_queue } } mut wait_pointer := &Event(unsafe { nil }) @@ -117,7 +132,7 @@ pub fn (queue &OwnedCommandQueue) barrier(wait_events []Event) !OwnedEvent { } mut handle := Event(unsafe { nil }) check(enqueue_barrier_with_wait_list(queue.handle, u32(wait_events.len), wait_pointer, &handle), 'enqueue OpenCL barrier')! - return OwnedEvent{ + return &OwnedEvent{ handle: handle } } diff --git a/examples/image_svm/main.v b/examples/image_svm/main.v index f5ce773..81ee468 100644 --- a/examples/image_svm/main.v +++ b/examples/image_svm/main.v @@ -23,21 +23,21 @@ fn run() ! { image_channel_order: cl.rgba image_channel_data_type: cl.unorm_int8 } - mut source_image := cl.new_image_2d[u32](&context, cl.mem_read_only, format, 2, 2)! - mut destination_image := cl.new_image_2d[u32](&context, cl.mem_write_only, format, 2, 2)! - mut sampler := cl.new_sampler(&context, false, cl.address_clamp_to_edge, cl.filter_nearest)! - mut image_program := cl.build_source_program(&context, device, image_source, '')! + mut source_image := cl.new_image_2d[u32](context, cl.mem_read_only, format, 2, 2)! + mut destination_image := cl.new_image_2d[u32](context, cl.mem_write_only, format, 2, 2)! + mut sampler := cl.new_sampler(context, false, cl.address_clamp_to_edge, cl.filter_nearest)! + mut image_program := cl.build_source_program(context, device, image_source, '')! mut image_kernel := image_program.kernel('copy_image')! - source_image.set_kernel_arg(&image_kernel, 0)! - destination_image.set_kernel_arg(&image_kernel, 1)! - image_kernel.set_sampler_arg(2, &sampler)! + source_image.set_kernel_arg(image_kernel, 0)! + destination_image.set_kernel_arg(image_kernel, 1)! + image_kernel.set_sampler_arg(2, sampler)! pixels := [u32(0xff0000ff), 0xff00ff00, 0xffff0000, 0xffffffff] - mut uploaded := source_image.write_async(&queue, pixels, []cl.Event{})! - mut copied := image_kernel.enqueue_nd_after(&queue, [usize(2), 2], []usize{}, [ + mut uploaded := source_image.write_async(queue, pixels, []cl.Event{})! + mut copied := image_kernel.enqueue_nd_after(queue, [usize(2), 2], []usize{}, [ uploaded.handle, ])! mut image_result := []u32{len: pixels.len} - mut downloaded := destination_image.read_async(&queue, mut image_result, [ + mut downloaded := destination_image.read_async(queue, mut image_result, [ copied.handle, ])! downloaded.wait()! @@ -54,16 +54,16 @@ fn run() ! { capabilities := cl.device_svm_support(device)! if capabilities & (cl.device_svm_coarse_grain_buffer | cl.device_svm_fine_grain_buffer) != 0 { - mut allocation := cl.new_svm[u32](&context, cl.mem_read_write, 4, 0)! - mut program := cl.build_source_program(&context, device, svm_source, '')! + mut allocation := cl.new_svm[u32](context, cl.mem_read_write, 4, 0)! + mut program := cl.build_source_program(context, device, svm_source, '')! mut kernel := program.kernel('brighten')! - allocation.set_kernel_arg(&kernel, 0)! + allocation.set_kernel_arg(kernel, 0)! amount := u32(5) kernel.set_arg(1, &amount)! - allocation.write(&queue, 0, [u32(1), 2, 3, 4])! - kernel.enqueue_1d(&queue, 4, 0)! + allocation.write(queue, 0, [u32(1), 2, 3, 4])! + kernel.enqueue_1d(queue, 4, 0)! mut svm_result := []u32{len: 4} - allocation.read(&queue, 0, mut svm_result)! + allocation.read(queue, 0, mut svm_result)! if svm_result != [u32(6), 7, 8, 9] { return error('typed SVM kernel round trip failed: ${svm_result}') } diff --git a/examples/vector_add/main.v b/examples/vector_add/main.v index c494c13..9eaea19 100644 --- a/examples/vector_add/main.v +++ b/examples/vector_add/main.v @@ -16,10 +16,10 @@ fn run() ! { device := devices[0] mut context := cl.new_context(device)! mut queue := context.command_queue(device, cl.queue_profiling_enable)! - mut a := cl.new_buffer[f32](&context, cl.mem_read_only, 4)! - mut b := cl.new_buffer[f32](&context, cl.mem_read_only, 4)! - mut result_buffer := cl.new_buffer[f32](&context, cl.mem_write_only, 4)! - mut program := cl.build_source_program(&context, device, source, '')! + mut a := cl.new_buffer[f32](context, cl.mem_read_only, 4)! + mut b := cl.new_buffer[f32](context, cl.mem_read_only, 4)! + mut result_buffer := cl.new_buffer[f32](context, cl.mem_write_only, 4)! + mut program := cl.build_source_program(context, device, source, '')! mut kernel := program.kernel('add')! kernel.set_buffer_arg(0, a.handle)! kernel.set_buffer_arg(1, b.handle)! @@ -27,11 +27,11 @@ fn run() ! { left := [f32(1), 2, 3, 4] right := [f32(10), 20, 30, 40] - mut left_ready := a.write_async(&queue, 0, left, []cl.Event{})! - mut right_ready := b.write_async(&queue, 0, right, []cl.Event{})! - mut computed := kernel.enqueue_1d_after(&queue, 4, 0, [left_ready.handle, right_ready.handle])! + mut left_ready := a.write_async(queue, 0, left, []cl.Event{})! + mut right_ready := b.write_async(queue, 0, right, []cl.Event{})! + mut computed := kernel.enqueue_1d_after(queue, 4, 0, [left_ready.handle, right_ready.handle])! mut result := []f32{len: 4} - mut downloaded := result_buffer.read_async(&queue, 0, mut result, [ + mut downloaded := result_buffer.read_async(queue, 0, mut result, [ computed.handle, ])! profile := downloaded.profile()! diff --git a/examples/vulkan_particles/compute.v b/examples/vulkan_particles/compute.v index 7487901..88d24ed 100644 --- a/examples/vulkan_particles/compute.v +++ b/examples/vulkan_particles/compute.v @@ -4,20 +4,21 @@ import antono2.opencl as cl const particle_stride = usize(8 * sizeof(f32)) +@[nocopy] struct Compute { platform cl.PlatformId device cl.DeviceId count usize mut: - context cl.OwnedContext - queue cl.OwnedCommandQueue - program cl.OwnedProgram - reset cl.OwnedKernel - step cl.OwnedKernel - buffer cl.Buffer[f32] + context &cl.OwnedContext = unsafe { nil } + queue &cl.OwnedCommandQueue = unsafe { nil } + program &cl.OwnedProgram = unsafe { nil } + reset &cl.OwnedKernel = unsafe { nil } + step &cl.OwnedKernel = unsafe { nil } + buffer &cl.Buffer[f32] = unsafe { nil } } -fn new_compute(count usize) !Compute { +fn new_compute(count usize) !&Compute { platforms := cl.platforms()! if platforms.len == 0 { return error('no OpenCL platforms found') @@ -48,7 +49,7 @@ fn new_compute(count usize) !Compute { return err } source := $embed_file('particles.cl').to_string() - mut program := cl.build_source_program(&context, selected_device, source, '') or { + mut program := cl.build_source_program(context, selected_device, source, '') or { queue.close() or {} context.close() or {} return err @@ -66,7 +67,7 @@ fn new_compute(count usize) !Compute { context.close() or {} return err } - mut buffer := cl.new_buffer[f32](&context, cl.mem_read_write, int(count * 8)) or { + mut buffer := cl.new_buffer[f32](context, cl.mem_read_write, int(count * 8)) or { step.close() or {} reset.close() or {} program.close() or {} @@ -74,16 +75,16 @@ fn new_compute(count usize) !Compute { context.close() or {} return err } - mut compute := Compute{ + mut compute := &Compute{ platform: selected_platform - device: selected_device - context: context - queue: queue - program: program - reset: reset - step: step - buffer: buffer - count: count + device: selected_device + context: context + queue: queue + program: program + reset: reset + step: step + buffer: buffer + count: count } compute.reset_particles(1)! return compute @@ -96,7 +97,7 @@ fn (compute &Compute) reset_particles(seed u32) ! { fn (compute &Compute) reset_buffer(buffer cl.Mem, seed u32) ! { compute.reset.set_buffer_arg(0, buffer)! compute.reset.set_arg(1, &seed)! - compute.reset.enqueue_1d(&compute.queue, compute.count, 0)! + compute.reset.enqueue_1d(compute.queue, compute.count, 0)! cl.check(cl.finish(compute.queue.handle), 'finish reset kernel')! } @@ -112,7 +113,7 @@ fn (compute &Compute) update_buffer(buffer cl.Mem, dt f32, elapsed f32, pointer_ compute.step.set_arg(2, &elapsed)! compute.step.set_slice_arg(3, pointer)! compute.step.set_arg(4, &attraction)! - compute.step.enqueue_1d(&compute.queue, compute.count, 0)! + compute.step.enqueue_1d(compute.queue, compute.count, 0)! } fn (compute &Compute) read_particles(mut destination []f32) ! { @@ -120,7 +121,7 @@ fn (compute &Compute) read_particles(mut destination []f32) ! { if destination.len < required { return error('particle destination is too small') } - compute.buffer.read(&compute.queue, 0, mut destination[..required])! + compute.buffer.read(compute.queue, 0, mut destination[..required])! } fn (mut compute Compute) close() { diff --git a/examples/vulkan_particles/live_sync.v b/examples/vulkan_particles/live_sync.v index 3c06f2b..bb1ec13 100644 --- a/examples/vulkan_particles/live_sync.v +++ b/examples/vulkan_particles/live_sync.v @@ -3,17 +3,18 @@ module main import antono2.opencl as cl import antono2.vulkan as vk +@[nocopy] struct LiveInteropSync { vk_to_cl vk.Semaphore cl_to_vk vk.Semaphore memory cl.ExternalMemoryInterop mut: - cl_wait cl.OwnedExternalSemaphore - cl_signal cl.OwnedExternalSemaphore + cl_wait &cl.OwnedExternalSemaphore = unsafe { nil } + cl_signal &cl.OwnedExternalSemaphore = unsafe { nil } } fn create_live_interop_sync(compute &Compute, memory cl.ExternalMemoryInterop, - semaphores cl.ExternalSemaphoreInterop, device vk.Device, queue vk.Queue) !LiveInteropSync { + semaphores cl.ExternalSemaphoreInterop, device vk.Device, queue vk.Queue) !&LiveInteropSync { mut export_info := vk.ExportSemaphoreCreateInfo{ handleTypes: u32(vk.ExternalSemaphoreHandleTypeFlagBits.opaque_fd) } @@ -29,18 +30,18 @@ fn create_live_interop_sync(compute &Compute, memory cl.ExternalMemoryInterop, mut vk_to_cl_fd := -1 mut cl_to_vk_fd := -1 vk_wait_fd_info := vk.SemaphoreGetFdInfoKHR{ - semaphore: vk_to_cl + semaphore: vk_to_cl handleType: .opaque_fd } cl_signal_fd_info := vk.SemaphoreGetFdInfoKHR{ - semaphore: cl_to_vk + semaphore: cl_to_vk handleType: .opaque_fd } vk_check(vk.get_semaphore_fd_khr(device, &vk_wait_fd_info, &vk_to_cl_fd), 'export live Vulkan-to-OpenCL semaphore')! vk_check(vk.get_semaphore_fd_khr(device, &cl_signal_fd_info, &cl_to_vk_fd), 'export live OpenCL-to-Vulkan semaphore')! - mut cl_wait := semaphores.import_opaque_fd(&compute.context, vk_to_cl_fd)! - cl_signal := semaphores.import_opaque_fd(&compute.context, cl_to_vk_fd) or { + mut cl_wait := semaphores.import_opaque_fd(compute.context, vk_to_cl_fd)! + mut cl_signal := semaphores.import_opaque_fd(compute.context, cl_to_vk_fd) or { cl_wait.close() or {} return err } @@ -48,21 +49,21 @@ fn create_live_interop_sync(compute &Compute, memory cl.ExternalMemoryInterop, // There is no preceding render for frame zero, so seed the binary handshake once. prime := vk.SubmitInfo{ signalSemaphoreCount: 1 - pSignalSemaphores: &vk_to_cl + pSignalSemaphores: &vk_to_cl } vk_check(vk.queue_submit(queue, 1, &prime, unsafe { nil }), 'prime live interop semaphore')! - return LiveInteropSync{ - vk_to_cl: vk_to_cl - cl_to_vk: cl_to_vk - memory: memory - cl_wait: cl_wait + return &LiveInteropSync{ + vk_to_cl: vk_to_cl + cl_to_vk: cl_to_vk + memory: memory + cl_wait: cl_wait cl_signal: cl_signal } } fn (sync &LiveInteropSync) begin_compute(compute &Compute, buffer cl.Mem) ! { - mut wait_event := sync.cl_wait.wait(&compute.queue, [])! - mut acquire_event := sync.memory.acquire(&compute.queue, [buffer], [ + mut wait_event := sync.cl_wait.wait(compute.queue, [])! + mut acquire_event := sync.memory.acquire(compute.queue, [buffer], [ wait_event.handle, ])! acquire_event.close()! @@ -70,8 +71,8 @@ fn (sync &LiveInteropSync) begin_compute(compute &Compute, buffer cl.Mem) ! { } fn (sync &LiveInteropSync) end_compute(compute &Compute, buffer cl.Mem) ! { - mut release_event := sync.memory.release(&compute.queue, [buffer], [])! - mut signal_event := sync.cl_signal.signal(&compute.queue, [release_event.handle])! + mut release_event := sync.memory.release(compute.queue, [buffer], [])! + mut signal_event := sync.cl_signal.signal(compute.queue, [release_event.handle])! signal_event.close()! release_event.close()! } diff --git a/examples/vulkan_particles/main.v b/examples/vulkan_particles/main.v index 5ee281d..cdeb6f0 100644 --- a/examples/vulkan_particles/main.v +++ b/examples/vulkan_particles/main.v @@ -28,13 +28,14 @@ fn main() { panic('zero-copy required: ${interop.describe()}') } if interop.zero_copy_available { - zero_copy_memory_smoke(&compute) or { panic(err) } + zero_copy_memory_smoke(compute) or { panic(err) } } if options.window { if options.force_staged { println('Renderer override: staged transfer path') } - window_device_loop(&compute, count, interop.zero_copy_available && !options.force_staged, options.frame_limit) or { panic(err) } + window_device_loop(compute, count, interop.zero_copy_available && !options.force_staged, + options.frame_limit) or { panic(err) } } // After the renderer closes, or immediately in headless mode, exercise the exact simulation diff --git a/examples/vulkan_particles/particle_buffer.v b/examples/vulkan_particles/particle_buffer.v index e2771a1..426140f 100644 --- a/examples/vulkan_particles/particle_buffer.v +++ b/examples/vulkan_particles/particle_buffer.v @@ -3,16 +3,17 @@ module main import antono2.opencl as cl import antono2.vulkan as vk +@[nocopy] struct ParticleBuffer { handle vk.Buffer memory vk.DeviceMemory size usize zero_copy bool mut: - cl_buffer cl.Buffer[f32] + cl_buffer &cl.Buffer[f32] = unsafe { nil } } -fn create_staged_particle_buffer(physical vk.PhysicalDevice, device vk.Device, particles []f32) !ParticleBuffer { +fn create_staged_particle_buffer(physical vk.PhysicalDevice, device vk.Device, particles []f32) !&ParticleBuffer { size := usize(particles.len) * sizeof(f32) info := vk.BufferCreateInfo{ size: size, usage: u32(vk.BufferUsageFlagBits.vertex_buffer), sharingMode: .exclusive } mut buffer := vk.Buffer(unsafe { nil }) @@ -29,23 +30,23 @@ fn create_staged_particle_buffer(physical vk.PhysicalDevice, device vk.Device, p vk_check(vk.map_memory(device, memory, 0, size, 0, &mapped), 'map particle vertex memory')! unsafe { vmemcpy(mapped, particles.data, size) } vk.unmap_memory(device, memory) - return ParticleBuffer{ + return &ParticleBuffer{ handle: buffer memory: memory - size: size + size: size } } fn create_zero_copy_particle_buffer(compute &Compute, memory_interop cl.ExternalMemoryInterop, - physical vk.PhysicalDevice, device vk.Device) !ParticleBuffer { + physical vk.PhysicalDevice, device vk.Device) !&ParticleBuffer { size := compute.count * particle_stride mut external_info := vk.ExternalMemoryBufferCreateInfo{ handleTypes: u32(vk.ExternalMemoryHandleTypeFlagBits.opaque_fd) } info := vk.BufferCreateInfo{ - pNext: &external_info - size: size - usage: u32(vk.BufferUsageFlagBits.vertex_buffer) | u32(vk.BufferUsageFlagBits.storage_buffer) + pNext: &external_info + size: size + usage: u32(vk.BufferUsageFlagBits.vertex_buffer) | u32(vk.BufferUsageFlagBits.storage_buffer) sharingMode: .exclusive } mut buffer := vk.Buffer(unsafe { nil }) @@ -60,8 +61,8 @@ fn create_zero_copy_particle_buffer(compute &Compute, memory_interop cl.External handleTypes: u32(vk.ExternalMemoryHandleTypeFlagBits.opaque_fd) } allocate := vk.MemoryAllocateInfo{ - pNext: &export_info - allocationSize: requirements.size + pNext: &export_info + allocationSize: requirements.size memoryTypeIndex: memory_type } mut memory := vk.DeviceMemory(unsafe { nil }) @@ -75,7 +76,7 @@ fn create_zero_copy_particle_buffer(compute &Compute, memory_interop cl.External return err } fd_info := vk.MemoryGetFdInfoKHR{ - memory: memory + memory: memory handleType: .opaque_fd } mut fd := -1 @@ -84,15 +85,15 @@ fn create_zero_copy_particle_buffer(compute &Compute, memory_interop cl.External vk.destroy_buffer(device, buffer, unsafe { nil }) return err } - cl_buffer := memory_interop.import_opaque_fd_buffer[f32](&compute.context, fd, int(compute.count * 8), cl.mem_read_write) or { + cl_buffer := memory_interop.import_opaque_fd_buffer[f32](compute.context, fd, int(compute.count * 8), cl.mem_read_write) or { vk.free_memory(device, memory, unsafe { nil }) vk.destroy_buffer(device, buffer, unsafe { nil }) return err } - return ParticleBuffer{ - handle: buffer - memory: memory - size: size + return &ParticleBuffer{ + handle: buffer + memory: memory + size: size cl_buffer: cl_buffer zero_copy: true } diff --git a/examples/vulkan_particles/window_smoke.v b/examples/vulkan_particles/window_smoke.v index e1efaa2..a283032 100644 --- a/examples/vulkan_particles/window_smoke.v +++ b/examples/vulkan_particles/window_smoke.v @@ -138,11 +138,11 @@ fn window_device_loop(compute &Compute, particle_count usize, use_zero_copy bool } defer { particle_buffer.destroy(device) } if particle_buffer.zero_copy { - mut acquire_event := memory_interop.acquire(&compute.queue, [ + mut acquire_event := memory_interop.acquire(compute.queue, [ particle_buffer.cl_buffer.handle, ], [])! compute.reset_buffer(particle_buffer.cl_buffer.handle, 1)! - mut release_event := memory_interop.release(&compute.queue, [ + mut release_event := memory_interop.release(compute.queue, [ particle_buffer.cl_buffer.handle, ], [])! release_event.close()! @@ -152,7 +152,7 @@ fn window_device_loop(compute &Compute, particle_count usize, use_zero_copy bool mut interop_sync := if particle_buffer.zero_copy { create_live_interop_sync(compute, memory_interop, semaphore_interop, device, queue)! } else { - LiveInteropSync{} + &LiveInteropSync{} } if particle_buffer.zero_copy { println('Live synchronization: external Vulkan/OpenCL semaphores') @@ -220,7 +220,7 @@ fn window_device_loop(compute &Compute, particle_count usize, use_zero_copy bool compute.read_particles(mut particles)! particle_buffer.upload(device, particles)! } - frame_ok := frames.draw_particles(device, queue, &swapchain, &pipeline, &particle_buffer, + frame_ok := frames.draw_particles(device, queue, &swapchain, &pipeline, particle_buffer, particle_count, elapsed, interop_sync.cl_to_vk, interop_sync.vk_to_cl, particle_buffer.zero_copy, trails)! if !frame_ok { diff --git a/examples/vulkan_particles/zero_copy_smoke.v b/examples/vulkan_particles/zero_copy_smoke.v index e96c4be..8fd679a 100644 --- a/examples/vulkan_particles/zero_copy_smoke.v +++ b/examples/vulkan_particles/zero_copy_smoke.v @@ -19,7 +19,7 @@ fn zero_copy_memory_smoke(compute &Compute) ! { mut instance := vk.Instance(unsafe { nil }) mut app_info := vk.ApplicationInfo{ pApplicationName: c'V OpenCL zero-copy smoke' - apiVersion: vk.api_version_1_1 + apiVersion: vk.api_version_1_1 } instance_info := vk.InstanceCreateInfo{ pApplicationInfo: &app_info } vk_check(vk.create_instance(&instance_info, unsafe { nil }, &instance), 'create instance')! @@ -31,15 +31,15 @@ fn zero_copy_memory_smoke(compute &Compute) ! { mut priority := f32(1) mut queue_info := vk.DeviceQueueCreateInfo{ queueFamilyIndex: queue_family - queueCount: 1 + queueCount: 1 pQueuePriorities: &priority } extensions := [vk.khr_external_memory_extension_name, vk.khr_external_memory_fd_extension_name, vk.khr_external_semaphore_extension_name, vk.khr_external_semaphore_fd_extension_name] device_info := vk.DeviceCreateInfo{ - queueCreateInfoCount: 1 - pQueueCreateInfos: &queue_info - enabledExtensionCount: u32(extensions.len) + queueCreateInfoCount: 1 + pQueueCreateInfos: &queue_info + enabledExtensionCount: u32(extensions.len) ppEnabledExtensionNames: extensions.data } mut device := vk.Device(unsafe { nil }) @@ -53,9 +53,9 @@ fn zero_copy_memory_smoke(compute &Compute) ! { handleTypes: u32(vk.ExternalMemoryHandleTypeFlagBits.opaque_fd) } buffer_info := vk.BufferCreateInfo{ - pNext: &external_info - size: compute.count * particle_stride - usage: u32(vk.BufferUsageFlagBits.vertex_buffer) | u32(vk.BufferUsageFlagBits.storage_buffer) + pNext: &external_info + size: compute.count * particle_stride + usage: u32(vk.BufferUsageFlagBits.vertex_buffer) | u32(vk.BufferUsageFlagBits.storage_buffer) sharingMode: .exclusive } mut buffer := vk.Buffer(unsafe { nil }) @@ -68,8 +68,8 @@ fn zero_copy_memory_smoke(compute &Compute) ! { handleTypes: u32(vk.ExternalMemoryHandleTypeFlagBits.opaque_fd) } allocation_info := vk.MemoryAllocateInfo{ - pNext: &export_info - allocationSize: requirements.size + pNext: &export_info + allocationSize: requirements.size memoryTypeIndex: memory_type } mut memory := vk.DeviceMemory(unsafe { nil }) @@ -77,24 +77,24 @@ fn zero_copy_memory_smoke(compute &Compute) ! { defer { vk.free_memory(device, memory, unsafe { nil }) } vk_check(vk.bind_buffer_memory(device, buffer, memory, 0), 'bind exportable buffer')! fd_info := vk.MemoryGetFdInfoKHR{ - memory: memory + memory: memory handleType: .opaque_fd } mut fd := -1 vk_check(vk.get_memory_fd_khr(device, &fd_info, &fd), 'export memory FD')! - mut imported_buffer := memory_interop.import_opaque_fd_buffer[f32](&compute.context, fd, int(compute.count * 8), cl.mem_read_write)! + mut imported_buffer := memory_interop.import_opaque_fd_buffer[f32](compute.context, fd, int(compute.count * 8), cl.mem_read_write)! defer { imported_buffer.close() or {} } - mut acquire_event := memory_interop.acquire(&compute.queue, [ + mut acquire_event := memory_interop.acquire(compute.queue, [ imported_buffer.handle, ], [])! seed := u32(7) compute.reset.set_buffer_arg(0, imported_buffer.handle)! compute.reset.set_arg(1, &seed)! - compute.reset.enqueue_1d(&compute.queue, compute.count, 0)! + compute.reset.enqueue_1d(compute.queue, compute.count, 0)! mut sample := []f32{len: 8} cl_check(cl.enqueue_read_buffer(compute.queue.handle, imported_buffer.handle, cl._true, 0, particle_stride, sample.data, 0, unsafe { nil }, unsafe { nil }), 'verify shared particle buffer')! - mut release_event := memory_interop.release(&compute.queue, [ + mut release_event := memory_interop.release(compute.queue, [ imported_buffer.handle, ], [])! release_event.close()! @@ -120,36 +120,36 @@ fn zero_copy_semaphore_smoke(compute &Compute, interop cl.ExternalSemaphoreInter mut vk_to_cl_fd := -1 mut cl_to_vk_fd := -1 vk_to_cl_fd_info := vk.SemaphoreGetFdInfoKHR{ - semaphore: vk_to_cl + semaphore: vk_to_cl handleType: .opaque_fd } cl_to_vk_fd_info := vk.SemaphoreGetFdInfoKHR{ - semaphore: cl_to_vk + semaphore: cl_to_vk handleType: .opaque_fd } vk_check(vk.get_semaphore_fd_khr(device, &vk_to_cl_fd_info, &vk_to_cl_fd), 'export Vulkan-to-OpenCL semaphore FD')! vk_check(vk.get_semaphore_fd_khr(device, &cl_to_vk_fd_info, &cl_to_vk_fd), 'export OpenCL-to-Vulkan semaphore FD')! - mut cl_wait := interop.import_opaque_fd(&compute.context, vk_to_cl_fd)! + mut cl_wait := interop.import_opaque_fd(compute.context, vk_to_cl_fd)! defer { cl_wait.close() or {} } - mut cl_signal := interop.import_opaque_fd(&compute.context, cl_to_vk_fd)! + mut cl_signal := interop.import_opaque_fd(compute.context, cl_to_vk_fd)! defer { cl_signal.close() or {} } vk_signal_submit := vk.SubmitInfo{ signalSemaphoreCount: 1 - pSignalSemaphores: &vk_to_cl + pSignalSemaphores: &vk_to_cl } vk_check(vk.queue_submit(queue, 1, &vk_signal_submit, unsafe { nil }), 'signal Vulkan-to-OpenCL semaphore')! - mut wait_event := cl_wait.wait(&compute.queue, [])! - mut signal_event := cl_signal.signal(&compute.queue, [wait_event.handle])! + mut wait_event := cl_wait.wait(compute.queue, [])! + mut signal_event := cl_signal.signal(compute.queue, [wait_event.handle])! signal_event.close()! wait_event.close()! mut stage := vk.PipelineStageFlags(vk.PipelineStageFlagBits.all_commands) vk_wait_submit := vk.SubmitInfo{ waitSemaphoreCount: 1 - pWaitSemaphores: &cl_to_vk - pWaitDstStageMask: &stage + pWaitSemaphores: &cl_to_vk + pWaitDstStageMask: &stage } mut fence := vk.Fence(unsafe { nil }) fence_info := vk.FenceCreateInfo{} diff --git a/external_interop.v b/external_interop.v index 3038f7a..bad1963 100644 --- a/external_interop.v +++ b/external_interop.v @@ -23,7 +23,7 @@ pub fn load_external_memory_interop(platform PlatformId, capabilities DeviceCapa if !capabilities.has_all(['cl_khr_external_memory', 'cl_khr_external_memory_opaque_fd']) { return OpenCLError{ operation: 'load opaque-FD OpenCL external-memory interoperability' - status: invalid_operation + status: invalid_operation } } acquire_address := get_extension_function_address_for_platform(platform, c'clEnqueueAcquireExternalMemObjectsKHR') @@ -31,7 +31,7 @@ pub fn load_external_memory_interop(platform PlatformId, capabilities DeviceCapa if isnil(acquire_address) || isnil(release_address) { return OpenCLError{ operation: 'resolve OpenCL external-memory entry points' - status: invalid_operation + status: invalid_operation } } return ExternalMemoryInterop{ @@ -43,23 +43,23 @@ pub fn load_external_memory_interop(platform PlatformId, capabilities DeviceCapa // import_opaque_fd_buffer imports an externally allocated buffer. The caller // remains responsible for the exporting API's handle-ownership requirements. pub fn (interop ExternalMemoryInterop) import_opaque_fd_buffer[T](context &OwnedContext, fd int, - count int, flags MemFlags) !Buffer[T] { + count int, flags MemFlags) !&Buffer[T] { if isnil(context.handle) { return OpenCLError{ operation: 'import external buffer into closed OpenCL context' - status: invalid_context + status: invalid_context } } if fd < 0 { return OpenCLError{ operation: 'import OpenCL external buffer with invalid file descriptor' - status: invalid_property + status: invalid_property } } if count <= 0 { return OpenCLError{ operation: 'import OpenCL external buffer with non-positive element count' - status: invalid_buffer_size + status: invalid_buffer_size } } byte_size := checked_element_bytes[T](count, 'import OpenCL external buffer with overflowing element count')! @@ -72,40 +72,40 @@ pub fn (interop ExternalMemoryInterop) import_opaque_fd_buffer[T](context &Owned } $else { return OpenCLError{ operation: 'import opaque-FD OpenCL buffer on unsupported operating system' - status: invalid_operation + status: invalid_operation } } check(status, 'import opaque-FD OpenCL buffer')! - return Buffer[T]{ + return &Buffer[T]{ handle: handle - count: count + count: count } } // acquire enqueues ownership acquisition for external memory objects. pub fn (interop ExternalMemoryInterop) acquire(queue &OwnedCommandQueue, objects []Mem, - wait_events []Event) !OwnedEvent { + wait_events []Event) !&OwnedEvent { return interop.enqueue_memory_command(interop.acquire_command, queue, objects, wait_events, 'acquire OpenCL external memory') } // release enqueues ownership release for external memory objects. pub fn (interop ExternalMemoryInterop) release(queue &OwnedCommandQueue, objects []Mem, - wait_events []Event) !OwnedEvent { + wait_events []Event) !&OwnedEvent { return interop.enqueue_memory_command(interop.release_command, queue, objects, wait_events, 'release OpenCL external memory') } fn (interop ExternalMemoryInterop) enqueue_memory_command(command ExternalMemoryCommand, - queue &OwnedCommandQueue, objects []Mem, wait_events []Event, operation string) !OwnedEvent { + queue &OwnedCommandQueue, objects []Mem, wait_events []Event, operation string) !&OwnedEvent { if isnil(queue.handle) { return OpenCLError{ operation: operation - status: invalid_command_queue + status: invalid_command_queue } } if objects.len == 0 { return OpenCLError{ operation: operation - status: invalid_value + status: invalid_value } } mut wait_pointer := &Event(unsafe { nil }) @@ -114,7 +114,7 @@ fn (interop ExternalMemoryInterop) enqueue_memory_command(command ExternalMemory } mut event := Event(unsafe { nil }) check(command(queue.handle, u32(objects.len), objects.data, u32(wait_events.len), wait_pointer, &event), operation)! - return OwnedEvent{ + return &OwnedEvent{ handle: event } } @@ -122,9 +122,10 @@ fn (interop ExternalMemoryInterop) enqueue_memory_command(command ExternalMemory // ExternalSemaphoreInterop holds platform-specific opaque-FD semaphore entry points. pub struct ExternalSemaphoreInterop { create_command PFN_clCreateSemaphoreWithPropertiesKHR = unsafe { nil } - wait_command ExternalSemaphoreCommand = unsafe { nil } - signal_command ExternalSemaphoreCommand = unsafe { nil } - release_command PFN_clReleaseSemaphoreKHR = unsafe { nil } + wait_command ExternalSemaphoreCommand = unsafe { nil } + signal_command ExternalSemaphoreCommand = unsafe { nil } + retain_command PFN_clRetainSemaphoreKHR = unsafe { nil } + release_command PFN_clReleaseSemaphoreKHR = unsafe { nil } } // load_external_semaphore_interop validates support and resolves entry points. @@ -134,48 +135,66 @@ pub fn load_external_semaphore_interop(platform PlatformId, 'cl_khr_external_semaphore_opaque_fd']) { return OpenCLError{ operation: 'load opaque-FD OpenCL external-semaphore interoperability' - status: invalid_operation + status: invalid_operation } } create_address := get_extension_function_address_for_platform(platform, c'clCreateSemaphoreWithPropertiesKHR') wait_address := get_extension_function_address_for_platform(platform, c'clEnqueueWaitSemaphoresKHR') signal_address := get_extension_function_address_for_platform(platform, c'clEnqueueSignalSemaphoresKHR') + retain_address := get_extension_function_address_for_platform(platform, c'clRetainSemaphoreKHR') release_address := get_extension_function_address_for_platform(platform, c'clReleaseSemaphoreKHR') if isnil(create_address) || isnil(wait_address) || isnil(signal_address) - || isnil(release_address) { + || isnil(retain_address) || isnil(release_address) { return OpenCLError{ operation: 'resolve OpenCL external-semaphore entry points' - status: invalid_operation + status: invalid_operation } } return ExternalSemaphoreInterop{ - create_command: unsafe { PFN_clCreateSemaphoreWithPropertiesKHR(create_address) } - wait_command: unsafe { ExternalSemaphoreCommand(wait_address) } - signal_command: unsafe { ExternalSemaphoreCommand(signal_address) } + create_command: unsafe { PFN_clCreateSemaphoreWithPropertiesKHR(create_address) } + wait_command: unsafe { ExternalSemaphoreCommand(wait_address) } + signal_command: unsafe { ExternalSemaphoreCommand(signal_address) } + retain_command: unsafe { PFN_clRetainSemaphoreKHR(retain_address) } release_command: unsafe { PFN_clReleaseSemaphoreKHR(release_address) } } } // OwnedExternalSemaphore owns one imported cl_semaphore_khr. +@[nocopy] pub struct OwnedExternalSemaphore { interop ExternalSemaphoreInterop pub mut: handle SemaphoreKhr } +// clone_ref retains the native semaphore and returns an independently owned reference. +pub fn (semaphore &OwnedExternalSemaphore) clone_ref() !&OwnedExternalSemaphore { + if isnil(semaphore.handle) { + return OpenCLError{ + operation: 'retain closed OpenCL external semaphore' + status: invalid_semaphore_khr + } + } + check(semaphore.interop.retain_command(semaphore.handle), 'retain OpenCL external semaphore')! + return &OwnedExternalSemaphore{ + interop: semaphore.interop + handle: semaphore.handle + } +} + // import_opaque_fd imports a binary opaque-FD semaphore. pub fn (interop ExternalSemaphoreInterop) import_opaque_fd(context &OwnedContext, - fd int) !OwnedExternalSemaphore { + fd int) !&OwnedExternalSemaphore { if isnil(context.handle) { return OpenCLError{ operation: 'import semaphore into closed OpenCL context' - status: invalid_context + status: invalid_context } } if fd < 0 { return OpenCLError{ operation: 'import OpenCL semaphore with invalid file descriptor' - status: invalid_property + status: invalid_property } } properties := [SemaphorePropertiesKhr(semaphore_type_khr), @@ -185,36 +204,36 @@ pub fn (interop ExternalSemaphoreInterop) import_opaque_fd(context &OwnedContext mut status := success handle := interop.create_command(context.handle, properties.data, &status) check(status, 'import opaque-FD OpenCL semaphore')! - return OwnedExternalSemaphore{ + return &OwnedExternalSemaphore{ interop: interop - handle: handle + handle: handle } } // wait enqueues a binary semaphore wait and returns its completion event. pub fn (semaphore &OwnedExternalSemaphore) wait(queue &OwnedCommandQueue, - wait_events []Event) !OwnedEvent { + wait_events []Event) !&OwnedEvent { return semaphore.enqueue_semaphore_command(semaphore.interop.wait_command, queue, wait_events, 'wait for OpenCL external semaphore') } // signal enqueues a binary semaphore signal and returns its completion event. pub fn (semaphore &OwnedExternalSemaphore) signal(queue &OwnedCommandQueue, - wait_events []Event) !OwnedEvent { + wait_events []Event) !&OwnedEvent { return semaphore.enqueue_semaphore_command(semaphore.interop.signal_command, queue, wait_events, 'signal OpenCL external semaphore') } fn (semaphore &OwnedExternalSemaphore) enqueue_semaphore_command(command ExternalSemaphoreCommand, - queue &OwnedCommandQueue, wait_events []Event, operation string) !OwnedEvent { + queue &OwnedCommandQueue, wait_events []Event, operation string) !&OwnedEvent { if isnil(semaphore.handle) { return OpenCLError{ operation: operation - status: invalid_semaphore_khr + status: invalid_semaphore_khr } } if isnil(queue.handle) { return OpenCLError{ operation: operation - status: invalid_command_queue + status: invalid_command_queue } } mut wait_pointer := &Event(unsafe { nil }) @@ -223,7 +242,7 @@ fn (semaphore &OwnedExternalSemaphore) enqueue_semaphore_command(command Externa } mut event := Event(unsafe { nil }) check(command(queue.handle, 1, &semaphore.handle, unsafe { nil }, u32(wait_events.len), wait_pointer, &event), operation)! - return OwnedEvent{ + return &OwnedEvent{ handle: event } } diff --git a/image.v b/image.v index 6bc11da..573ea3a 100644 --- a/image.v +++ b/image.v @@ -2,6 +2,7 @@ module opencl // Image2D owns a typed two-dimensional OpenCL image. T represents one complete // image element (pixel), not one channel. Its size must match format exactly. +@[nocopy] pub struct Image2D[T] { pub mut: handle Mem @@ -12,12 +13,45 @@ pub: pixel_bytes usize } +// clone_ref retains the native image and returns an independently owned wrapper. +pub fn (image &Image2D[T]) clone_ref() !&Image2D[T] { + if isnil(image.handle) { + return OpenCLError{ + operation: 'retain closed OpenCL image' + status: invalid_mem_object + } + } + check(retain_mem_object(image.handle), 'retain OpenCL image')! + return &Image2D[T]{ + handle: image.handle + width: image.width + height: image.height + format: image.format + pixel_bytes: image.pixel_bytes + } +} + // OwnedSampler owns one OpenCL sampler reference. +@[nocopy] pub struct OwnedSampler { pub mut: handle Sampler } +// clone_ref retains the native sampler and returns an independently owned reference. +pub fn (sampler &OwnedSampler) clone_ref() !&OwnedSampler { + if isnil(sampler.handle) { + return OpenCLError{ + operation: 'retain closed OpenCL sampler' + status: invalid_sampler + } + } + check(retain_sampler(sampler.handle), 'retain OpenCL sampler')! + return &OwnedSampler{ + handle: sampler.handle + } +} + // supported_image2d_formats returns every format supported for a 2D image with flags. pub fn supported_image2d_formats(context &OwnedContext, flags MemFlags) ![]ImageFormat { if isnil(context.handle) { @@ -141,7 +175,7 @@ fn checked_image_dimensions(width int, height int, operation string) !int { // new_image_2d allocates a tightly packed 2D image without a host pointer. pub fn new_image_2d[T](context &OwnedContext, flags MemFlags, format ImageFormat, width int, - height int) !Image2D[T] { + height int) !&Image2D[T] { if isnil(context.handle) { return OpenCLError{ operation: 'create OpenCL image from closed context' @@ -172,7 +206,7 @@ pub fn new_image_2d[T](context &OwnedContext, flags MemFlags, format ImageFormat status: mem_object_allocation_failure } } - return Image2D[T]{ + return &Image2D[T]{ handle: handle width: width height: height @@ -221,7 +255,7 @@ pub fn (image &Image2D[T]) write(queue &OwnedCommandQueue, values []T) ! { // write_async enqueues replacement of every pixel. values must remain allocated // and unchanged until the returned event completes. pub fn (image &Image2D[T]) write_async(queue &OwnedCommandQueue, values []T, - wait_events []Event) !OwnedEvent { + wait_events []Event) !&OwnedEvent { return image.write_region_async(queue, 0, 0, image.width, image.height, values, wait_events) } @@ -239,7 +273,7 @@ pub fn (image &Image2D[T]) write_region(queue &OwnedCommandQueue, x int, y int, // write_region_async enqueues a tightly packed image write. values must remain // allocated and unchanged until the returned event completes. pub fn (image &Image2D[T]) write_region_async(queue &OwnedCommandQueue, x int, y int, - width int, height int, values []T, wait_events []Event) !OwnedEvent { + width int, height int, values []T, wait_events []Event) !&OwnedEvent { image.validate_region(queue, x, y, width, height, values.len, 'write')! origin := [usize(x), usize(y), usize(0)] region := [usize(width), usize(height), usize(1)] @@ -251,7 +285,7 @@ pub fn (image &Image2D[T]) write_region_async(queue &OwnedCommandQueue, x int, y check(enqueue_write_image(queue.handle, image.handle, non_blocking, origin.data, region.data, usize(width) * image.pixel_bytes, 0, values.data, u32(wait_events.len), wait_pointer, &event), 'write OpenCL image asynchronously')! - return OwnedEvent{ + return &OwnedEvent{ handle: event } } @@ -264,7 +298,7 @@ pub fn (image &Image2D[T]) read(queue &OwnedCommandQueue, mut destination []T) ! // read_async enqueues a copy of every pixel. destination must remain allocated // and unread until the returned event completes. pub fn (image &Image2D[T]) read_async(queue &OwnedCommandQueue, mut destination []T, - wait_events []Event) !OwnedEvent { + wait_events []Event) !&OwnedEvent { return image.read_region_async(queue, 0, 0, image.width, image.height, mut destination, wait_events) } @@ -283,7 +317,7 @@ pub fn (image &Image2D[T]) read_region(queue &OwnedCommandQueue, x int, y int, w // read_region_async enqueues a tightly packed image read. destination must // remain allocated and unread until the returned event completes. pub fn (image &Image2D[T]) read_region_async(queue &OwnedCommandQueue, x int, y int, - width int, height int, mut destination []T, wait_events []Event) !OwnedEvent { + width int, height int, mut destination []T, wait_events []Event) !&OwnedEvent { image.validate_region(queue, x, y, width, height, destination.len, 'read')! origin := [usize(x), usize(y), usize(0)] region := [usize(width), usize(height), usize(1)] @@ -295,7 +329,7 @@ pub fn (image &Image2D[T]) read_region_async(queue &OwnedCommandQueue, x int, y check(enqueue_read_image(queue.handle, image.handle, non_blocking, origin.data, region.data, usize(width) * image.pixel_bytes, 0, destination.data, u32(wait_events.len), wait_pointer, &event), 'read OpenCL image asynchronously')! - return OwnedEvent{ + return &OwnedEvent{ handle: event } } @@ -322,7 +356,7 @@ pub fn (image &Image2D[T]) set_kernel_arg(kernel &OwnedKernel, index u32) ! { // new_sampler creates an owned sampler for image kernel arguments. pub fn new_sampler(context &OwnedContext, normalized_coordinates bool, - addressing_mode AddressingMode, filter_mode FilterMode) !OwnedSampler { + addressing_mode AddressingMode, filter_mode FilterMode) !&OwnedSampler { if isnil(context.handle) { return OpenCLError{ operation: 'create OpenCL sampler from closed context' @@ -339,7 +373,7 @@ pub fn new_sampler(context &OwnedContext, normalized_coordinates bool, status: out_of_host_memory } } - return OwnedSampler{ + return &OwnedSampler{ handle: handle } } diff --git a/ownership.v b/ownership.v index 090aa47..f84a624 100644 --- a/ownership.v +++ b/ownership.v @@ -1,6 +1,7 @@ module opencl // OwnedContext owns one reference to an OpenCL context. Call close when done. +@[nocopy] pub struct OwnedContext { pub mut: handle Context @@ -8,8 +9,24 @@ pub: device DeviceId } +// clone_ref retains the native context and returns an independently owned +// reference. Close both wrappers when they are no longer needed. +pub fn (context &OwnedContext) clone_ref() !&OwnedContext { + if isnil(context.handle) { + return OpenCLError{ + operation: 'retain closed OpenCL context' + status: invalid_context + } + } + check(retain_context(context.handle), 'retain OpenCL context')! + return &OwnedContext{ + handle: context.handle + device: context.device + } +} + // new_context creates a context containing exactly one explicitly selected device. -pub fn new_context(device DeviceId) !OwnedContext { +pub fn new_context(device DeviceId) !&OwnedContext { mut status := success handle := create_context(unsafe { nil }, 1, &device, unsafe { nil }, unsafe { nil }, &status) check(status, 'create OpenCL context')! @@ -19,7 +36,7 @@ pub fn new_context(device DeviceId) !OwnedContext { status: out_of_host_memory } } - return OwnedContext{ + return &OwnedContext{ handle: handle device: device } @@ -28,7 +45,7 @@ pub fn new_context(device DeviceId) !OwnedContext { // command_queue creates a legacy-compatible command queue for a device in this context. // The properties argument accepts flags such as queue_profiling_enable. pub fn (context &OwnedContext) command_queue(device DeviceId, - properties CommandQueueProperties) !OwnedCommandQueue { + properties CommandQueueProperties) !&OwnedCommandQueue { if isnil(context.handle) { return OpenCLError{ operation: 'create OpenCL command queue from closed context' @@ -44,7 +61,7 @@ pub fn (context &OwnedContext) command_queue(device DeviceId, status: out_of_host_memory } } - return OwnedCommandQueue{ + return &OwnedCommandQueue{ handle: handle } } @@ -60,11 +77,26 @@ pub fn (mut context OwnedContext) close() ! { } // OwnedCommandQueue owns one reference to an OpenCL command queue. Call close when done. +@[nocopy] pub struct OwnedCommandQueue { pub mut: handle CommandQueue } +// clone_ref retains the native queue and returns an independently owned reference. +pub fn (queue &OwnedCommandQueue) clone_ref() !&OwnedCommandQueue { + if isnil(queue.handle) { + return OpenCLError{ + operation: 'retain closed OpenCL command queue' + status: invalid_command_queue + } + } + check(retain_command_queue(queue.handle), 'retain OpenCL command queue')! + return &OwnedCommandQueue{ + handle: queue.handle + } +} + // close releases the owned queue reference. It is safe to call more than once. pub fn (mut queue OwnedCommandQueue) close() ! { if isnil(queue.handle) { @@ -76,6 +108,7 @@ pub fn (mut queue OwnedCommandQueue) close() ! { // Buffer owns a typed OpenCL buffer containing count elements of T. T must be // a plain C-layout value without V-managed references. +@[nocopy] pub struct Buffer[T] { pub mut: handle Mem @@ -83,6 +116,22 @@ pub: count int } +// clone_ref retains the native memory object and returns an independently +// owned buffer wrapper with the same element count. +pub fn (buffer &Buffer[T]) clone_ref() !&Buffer[T] { + if isnil(buffer.handle) { + return OpenCLError{ + operation: 'retain closed OpenCL buffer' + status: invalid_mem_object + } + } + check(retain_mem_object(buffer.handle), 'retain OpenCL buffer')! + return &Buffer[T]{ + handle: buffer.handle + count: buffer.count + } +} + fn checked_element_bytes[T](count int, operation string) !usize { if count < 0 { return OpenCLError{ @@ -101,7 +150,7 @@ fn checked_element_bytes[T](count int, operation string) !usize { } // new_buffer allocates storage for count elements of T without a host pointer. -pub fn new_buffer[T](context &OwnedContext, flags MemFlags, count int) !Buffer[T] { +pub fn new_buffer[T](context &OwnedContext, flags MemFlags, count int) !&Buffer[T] { if isnil(context.handle) { return OpenCLError{ operation: 'create OpenCL buffer from closed context' @@ -125,7 +174,7 @@ pub fn new_buffer[T](context &OwnedContext, flags MemFlags, count int) !Buffer[T status: mem_object_allocation_failure } } - return Buffer[T]{ + return &Buffer[T]{ handle: handle count: count } @@ -147,7 +196,7 @@ pub fn (buffer &Buffer[T]) write(queue &OwnedCommandQueue, offset int, values [] // write_async enqueues a non-blocking copy and returns its completion event. // values must remain allocated and unchanged until the returned event completes. pub fn (buffer &Buffer[T]) write_async(queue &OwnedCommandQueue, offset int, values []T, - wait_events []Event) !OwnedEvent { + wait_events []Event) !&OwnedEvent { buffer.validate_transfer(queue, offset, values.len, 'write')! if values.len == 0 { return queue.marker(wait_events) @@ -163,7 +212,7 @@ pub fn (buffer &Buffer[T]) write_async(queue &OwnedCommandQueue, offset int, val check(enqueue_write_buffer(queue.handle, buffer.handle, non_blocking, byte_offset, byte_size, values.data, u32(wait_events.len), wait_pointer, &event), 'write OpenCL buffer asynchronously')! - return OwnedEvent{ + return &OwnedEvent{ handle: event } } @@ -184,7 +233,7 @@ pub fn (buffer &Buffer[T]) read(queue &OwnedCommandQueue, offset int, mut destin // read_async enqueues a non-blocking copy and returns its completion event. // destination must remain allocated and must not be read until the event completes. pub fn (buffer &Buffer[T]) read_async(queue &OwnedCommandQueue, offset int, mut destination []T, - wait_events []Event) !OwnedEvent { + wait_events []Event) !&OwnedEvent { buffer.validate_transfer(queue, offset, destination.len, 'read')! if destination.len == 0 { return queue.marker(wait_events) @@ -200,7 +249,7 @@ pub fn (buffer &Buffer[T]) read_async(queue &OwnedCommandQueue, offset int, mut check(enqueue_read_buffer(queue.handle, buffer.handle, non_blocking, byte_offset, byte_size, destination.data, u32(wait_events.len), wait_pointer, &event), 'read OpenCL buffer asynchronously')! - return OwnedEvent{ + return &OwnedEvent{ handle: event } } diff --git a/program.v b/program.v index 6544881..486efb1 100644 --- a/program.v +++ b/program.v @@ -23,6 +23,7 @@ pub fn (err ProgramBuildError) code() int { } // OwnedProgram owns one compiled OpenCL program reference. +@[nocopy] pub struct OwnedProgram { pub mut: handle Program @@ -30,9 +31,24 @@ pub: device DeviceId } +// clone_ref retains the native program and returns an independently owned reference. +pub fn (program &OwnedProgram) clone_ref() !&OwnedProgram { + if isnil(program.handle) { + return OpenCLError{ + operation: 'retain closed OpenCL program' + status: invalid_program + } + } + check(retain_program(program.handle), 'retain OpenCL program')! + return &OwnedProgram{ + handle: program.handle + device: program.device + } +} + // build_source_program creates and synchronously builds source for one device. pub fn build_source_program(context &OwnedContext, device DeviceId, source string, - options string) !OwnedProgram { + options string) !&OwnedProgram { if isnil(context.handle) { return OpenCLError{ operation: 'build program from closed OpenCL context' @@ -63,14 +79,14 @@ pub fn build_source_program(context &OwnedContext, device DeviceId, source strin log: log } } - return OwnedProgram{ + return &OwnedProgram{ handle: handle device: device } } // kernel creates an owned kernel by name. -pub fn (program &OwnedProgram) kernel(name string) !OwnedKernel { +pub fn (program &OwnedProgram) kernel(name string) !&OwnedKernel { if isnil(program.handle) { return OpenCLError{ operation: 'create kernel from closed OpenCL program' @@ -80,7 +96,7 @@ pub fn (program &OwnedProgram) kernel(name string) !OwnedKernel { mut status := success handle := create_kernel(program.handle, name.str, &status) check(status, 'create OpenCL kernel `${name}`')! - return OwnedKernel{ + return &OwnedKernel{ handle: handle } } @@ -95,11 +111,26 @@ pub fn (mut program OwnedProgram) close() ! { } // OwnedKernel owns one OpenCL kernel reference. +@[nocopy] pub struct OwnedKernel { pub mut: handle Kernel } +// clone_ref retains the native kernel and returns an independently owned reference. +pub fn (kernel &OwnedKernel) clone_ref() !&OwnedKernel { + if isnil(kernel.handle) { + return OpenCLError{ + operation: 'retain closed OpenCL kernel' + status: invalid_kernel + } + } + check(retain_kernel(kernel.handle), 'retain OpenCL kernel')! + return &OwnedKernel{ + handle: kernel.handle + } +} + // set_arg copies one scalar or plain-value argument into the kernel. pub fn (kernel &OwnedKernel) set_arg[T](index u32, value &T) ! { if isnil(kernel.handle) { @@ -180,7 +211,7 @@ pub fn (kernel &OwnedKernel) enqueue_1d(queue &OwnedCommandQueue, global_size us // enqueue_1d_after submits a one-dimensional kernel after wait_events and // returns an owned completion event. A local size of zero lets the runtime choose. pub fn (kernel &OwnedKernel) enqueue_1d_after(queue &OwnedCommandQueue, global_size usize, - local_size usize, wait_events []Event) !OwnedEvent { + local_size usize, wait_events []Event) !&OwnedEvent { if isnil(kernel.handle) { return OpenCLError{ operation: 'enqueue closed OpenCL kernel' @@ -213,7 +244,7 @@ pub fn (kernel &OwnedKernel) enqueue_1d_after(queue &OwnedCommandQueue, global_s check(enqueue_nd_range_kernel(queue.handle, kernel.handle, 1, no_global_offset, &global_size, local_pointer, u32(wait_events.len), wait_pointer, &event), 'enqueue OpenCL kernel with event')! - return OwnedEvent{ + return &OwnedEvent{ handle: event } } @@ -222,7 +253,7 @@ pub fn (kernel &OwnedKernel) enqueue_1d_after(queue &OwnedCommandQueue, global_s // wait_events and returns an owned completion event. An empty local_sizes slice // lets the runtime select work-group dimensions. pub fn (kernel &OwnedKernel) enqueue_nd_after(queue &OwnedCommandQueue, global_sizes []usize, - local_sizes []usize, wait_events []Event) !OwnedEvent { + local_sizes []usize, wait_events []Event) !&OwnedEvent { if isnil(kernel.handle) { return OpenCLError{ operation: 'enqueue closed OpenCL kernel' @@ -276,7 +307,7 @@ pub fn (kernel &OwnedKernel) enqueue_nd_after(queue &OwnedCommandQueue, global_s check(enqueue_nd_range_kernel(queue.handle, kernel.handle, u32(global_sizes.len), no_global_offset, global_sizes.data, local_pointer, u32(wait_events.len), wait_pointer, &event), 'enqueue OpenCL kernel with event')! - return OwnedEvent{ + return &OwnedEvent{ handle: event } } diff --git a/svm.v b/svm.v index 36a6e46..42b4fca 100644 --- a/svm.v +++ b/svm.v @@ -53,6 +53,7 @@ fn platform_set_kernel_arg_svm_pointer(kernel Kernel, index u32, pointer voidptr // SvmAllocation owns a typed OpenCL shared virtual memory allocation. T must // be a plain C-layout value without V-managed references. +@[nocopy] pub struct SvmAllocation[T] { pub mut: handle voidptr @@ -82,7 +83,7 @@ pub fn device_svm_support(device DeviceId) !DeviceSvmCapabilities { // new_svm allocates count elements of OpenCL shared virtual memory. flags accepts // CL_MEM_* values, including mem_svm_fine_grain_buffer and mem_svm_atomics. pub fn new_svm[T](context &OwnedContext, flags MemFlags, count int, - alignment u32) !SvmAllocation[T] { + alignment u32) !&SvmAllocation[T] { if isnil(context.handle) { return OpenCLError{ operation: 'allocate SVM from closed OpenCL context' @@ -128,7 +129,7 @@ pub fn new_svm[T](context &OwnedContext, flags MemFlags, count int, status: mem_object_allocation_failure } } - return SvmAllocation[T]{ + return &SvmAllocation[T]{ handle: handle count: count context: context.handle @@ -179,7 +180,7 @@ pub fn (allocation &SvmAllocation[T]) write(queue &OwnedCommandQueue, offset int // write_async enqueues a copy into SVM. values must remain allocated and // unchanged until the returned event completes. pub fn (allocation &SvmAllocation[T]) write_async(queue &OwnedCommandQueue, offset int, - values []T, wait_events []Event) !OwnedEvent { + values []T, wait_events []Event) !&OwnedEvent { byte_size := allocation.validate_transfer(queue, offset, values.len, 'write')! if values.len == 0 { return queue.marker(wait_events) @@ -192,7 +193,7 @@ pub fn (allocation &SvmAllocation[T]) write_async(queue &OwnedCommandQueue, offs destination := allocation.pointer_at(offset, 'write')! check(platform_enqueue_svm_memcpy(queue.handle, non_blocking, destination, values.data, byte_size, u32(wait_events.len), wait_pointer, &event), 'write OpenCL SVM asynchronously')! - return OwnedEvent{ + return &OwnedEvent{ handle: event } } @@ -212,7 +213,7 @@ pub fn (allocation &SvmAllocation[T]) read(queue &OwnedCommandQueue, offset int, // read_async enqueues a copy from SVM. destination must remain allocated and // unread until the returned event completes. pub fn (allocation &SvmAllocation[T]) read_async(queue &OwnedCommandQueue, offset int, - mut destination []T, wait_events []Event) !OwnedEvent { + mut destination []T, wait_events []Event) !&OwnedEvent { byte_size := allocation.validate_transfer(queue, offset, destination.len, 'read')! if destination.len == 0 { return queue.marker(wait_events) @@ -225,7 +226,7 @@ pub fn (allocation &SvmAllocation[T]) read_async(queue &OwnedCommandQueue, offse source := allocation.pointer_at(offset, 'read')! check(platform_enqueue_svm_memcpy(queue.handle, non_blocking, destination.data, source, byte_size, u32(wait_events.len), wait_pointer, &event), 'read OpenCL SVM asynchronously')! - return OwnedEvent{ + return &OwnedEvent{ handle: event } } @@ -239,7 +240,7 @@ pub fn (allocation &SvmAllocation[T]) map(queue &OwnedCommandQueue, flags MapFla // unmap relinquishes host access and returns an event for the device-visible transition. pub fn (allocation &SvmAllocation[T]) unmap(queue &OwnedCommandQueue, - wait_events []Event) !OwnedEvent { + wait_events []Event) !&OwnedEvent { allocation.validate_transfer(queue, 0, allocation.count, 'unmap')! mut wait_pointer := &Event(unsafe { nil }) if wait_events.len > 0 { @@ -248,7 +249,7 @@ pub fn (allocation &SvmAllocation[T]) unmap(queue &OwnedCommandQueue, mut event := Event(unsafe { nil }) check(platform_enqueue_svm_unmap(queue.handle, allocation.handle, u32(wait_events.len), wait_pointer, &event), 'unmap OpenCL SVM')! - return OwnedEvent{ + return &OwnedEvent{ handle: event } } diff --git a/test/pointer_abi_test.v b/test/pointer_abi_test.v index 7b64627..6d807c6 100644 --- a/test/pointer_abi_test.v +++ b/test/pointer_abi_test.v @@ -39,12 +39,12 @@ fn test_blocking_transfers_pass_null_event_pointers() ! { } buffer := cl.Buffer[u32]{ handle: cl.Mem(&buffer_storage) - count: 2 + count: 2 } values := [u32(3), 5] - buffer.write(&queue, 0, values)! + buffer.write(queue, 0, values)! mut destination := []u32{len: 2} - buffer.read(&queue, 0, mut destination)! + buffer.read(queue, 0, mut destination)! assert C.opencl_pointer_shim_blocking_write_calls() == 1 assert C.opencl_pointer_shim_blocking_read_calls() == 1 } @@ -59,6 +59,6 @@ fn test_kernel_enqueue_passes_null_offset_and_event_pointers() ! { kernel := cl.OwnedKernel{ handle: cl.Kernel(&kernel_storage) } - kernel.enqueue_1d(&queue, 64, 0)! + kernel.enqueue_1d(queue, 64, 0)! assert C.opencl_pointer_shim_kernel_enqueue_calls() == 1 } From 34b0a7555a6a5611830c366e8ff729c1649946d5 Mon Sep 17 00:00:00 2001 From: Anton Oreskin Date: Tue, 22 Sep 2026 20:19:58 +0200 Subject: [PATCH 2/2] docs: explain explicit OpenCL ownership --- CHANGELOG.md | 7 +++++++ README.md | 47 +++++++++++++++++++++++++++-------------------- 2 files changed, 34 insertions(+), 20 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index b456034..2c214ba 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -2,6 +2,13 @@ ## Unreleased +- Prevent implicit copying of every owned OpenCL wrapper with `@[nocopy]` and + provide explicit `clone_ref()` operations backed by the corresponding native + retain calls. SVM allocations remain uniquely owned because OpenCL provides + no retain operation for them. +- Return owned pointers from constructors, asynchronous operations, and + `clone_ref()` so owners cross module boundaries without hidden copies on + released V and strict V3. - Pin Vulkan and GLFW example dependencies to immutable revisions in CI rather than relying on mutable VPM installs. - Ignore local compiler products and caches so building the bundled examples diff --git a/README.md b/README.md index e3cc145..78d2f8b 100644 --- a/README.md +++ b/README.md @@ -86,9 +86,9 @@ defer { context.close() or {} } mut queue := context.command_queue(device, cl.CommandQueueProperties(0))! defer { queue.close() or {} } -mut buffer := cl.new_buffer[f32](&context, cl.mem_read_write, 1024)! +mut buffer := cl.new_buffer[f32](context, cl.mem_read_write, 1024)! defer { buffer.close() or {} } -buffer.write(&queue, 0, []f32{len: 1024, init: f32(index)})! +buffer.write(queue, 0, []f32{len: 1024, init: f32(index)})! ``` The element type used by `Buffer[T]`, typed transfers, and kernel arguments @@ -100,23 +100,23 @@ Source compilation preserves compiler diagnostics through `ProgramBuildError`. O kernels support typed scalar and buffer arguments plus one-dimensional dispatch: ```v -mut program := cl.build_source_program(&context, device, source, '')! +mut program := cl.build_source_program(context, device, source, '')! defer { program.close() or {} } mut kernel := program.kernel('transform')! defer { kernel.close() or {} } kernel.set_buffer_arg(0, buffer.handle)! kernel.set_slice_arg(1, [f32(0.5), 1.0])! // e.g. an OpenCL float2 -kernel.enqueue_1d(&queue, usize(buffer.count), 0)! +kernel.enqueue_1d(queue, usize(buffer.count), 0)! ``` Non-blocking transfers and dispatch return owned events and accept native event dependency lists. Host slices must remain alive until their transfer event completes: ```v -mut uploaded := buffer.write_async(&queue, 0, values, []cl.Event{})! -mut dispatched := kernel.enqueue_1d_after(&queue, usize(buffer.count), 0, +mut uploaded := buffer.write_async(queue, 0, values, []cl.Event{})! +mut dispatched := kernel.enqueue_1d_after(queue, usize(buffer.count), 0, [uploaded.handle])! -mut downloaded := buffer.read_async(&queue, 0, mut result, [dispatched.handle])! +mut downloaded := buffer.read_async(queue, 0, mut result, [dispatched.handle])! downloaded.wait()! profile := downloaded.profile()! // queue must use cl.queue_profiling_enable downloaded.close()! @@ -124,6 +124,13 @@ dispatched.close()! uploaded.close()! ``` +Owned contexts, queues, buffers, images, samplers, programs, kernels, events, +and external semaphores are `@[nocopy]`, preventing accidental double release. +Constructors return owned pointers; pass them directly without adding another +`&`. When two independently closable owners are required, +call `clone_ref()`; it performs the matching OpenCL retain operation. SVM +allocations cannot be retained and therefore always have one unique owner. + For multidimensional kernels, `enqueue_nd_after()` accepts one to three global dimensions and either a matching local-size slice or an empty slice for an implementation-selected work-group size. @@ -137,14 +144,14 @@ format := cl.ImageFormat{ image_channel_order: cl.rgba image_channel_data_type: cl.unorm_int8 } -mut image := cl.new_image_2d[u32](&context, cl.mem_read_write, format, 64, 64)! +mut image := cl.new_image_2d[u32](context, cl.mem_read_write, format, 64, 64)! defer { image.close() or {} } -mut sampler := cl.new_sampler(&context, false, cl.address_clamp_to_edge, +mut sampler := cl.new_sampler(context, false, cl.address_clamp_to_edge, cl.filter_nearest)! defer { sampler.close() or {} } -image.write(&queue, pixels)! -image.set_kernel_arg(&kernel, 0)! -kernel.set_sampler_arg(1, &sampler)! +image.write(queue, pixels)! +image.set_kernel_arg(kernel, 0)! +kernel.set_sampler_arg(1, sampler)! ``` Shared virtual memory is similarly typed and capability-gated. Coarse-grained @@ -155,10 +162,10 @@ bound directly to a kernel: svm_capabilities := cl.device_svm_support(device)! if svm_capabilities & (cl.device_svm_coarse_grain_buffer | cl.device_svm_fine_grain_buffer) != 0 { - mut shared := cl.new_svm[u32](&context, cl.mem_read_write, 1024, 0)! + mut shared := cl.new_svm[u32](context, cl.mem_read_write, 1024, 0)! defer { shared.close() } - shared.write(&queue, 0, values)! - shared.set_kernel_arg(&kernel, 0)! + shared.write(queue, 0, values)! + shared.set_kernel_arg(kernel, 0)! } ``` @@ -182,15 +189,15 @@ descriptors are obtained from the exporting API; its handle-ownership rules stil ```v memory_interop := cl.load_external_memory_interop(platform, capabilities)! -mut shared := memory_interop.import_opaque_fd_buffer[f32](&context, memory_fd, +mut shared := memory_interop.import_opaque_fd_buffer[f32](context, memory_fd, element_count, cl.mem_read_write)! defer { shared.close() or {} } semaphore_interop := cl.load_external_semaphore_interop(platform, capabilities)! -mut ready := semaphore_interop.import_opaque_fd(&context, semaphore_fd)! +mut ready := semaphore_interop.import_opaque_fd(context, semaphore_fd)! defer { ready.close() or {} } -mut waited := ready.wait(&queue, [])! -mut acquired := memory_interop.acquire(&queue, [shared.handle], [waited.handle])! +mut waited := ready.wait(queue, [])! +mut acquired := memory_interop.acquire(queue, [shared.handle], [waited.handle])! defer { acquired.close() or {} } defer { waited.close() or {} } ``` @@ -207,7 +214,7 @@ ready.wait()! See [`API_DESIGN.md`](API_DESIGN.md) for the conventions shared with the companion Vulkan convenience layer. -See [`OWNERSHIP.md`](OWNERSHIP.md) for the current copy and cleanup rules. +See [`OWNERSHIP.md`](OWNERSHIP.md) for the current ownership and cleanup rules. ## Advanced example