diff options
author | Michael Jones <michael_p_jones@apple.com> | 2021-12-07 18:11:35 +0300 |
---|---|---|
committer | Michael Jones <michael_p_jones@apple.com> | 2021-12-07 18:52:21 +0300 |
commit | 9558fa5196033390111a2348caa66ab18b8a4f89 (patch) | |
tree | acc3ed446f709390abfef5f97f82c1ed9abe0100 /intern/cycles/device/metal/queue.h | |
parent | 565b33c0ad31966b860123837d2c4b5a8cbedad2 (diff) |
Cycles: Metal host-side code
This patch adds the Metal host-side code:
- Add all core host-side Metal backend files (device_impl, queue, etc)
- Add MetalRT BVH setup files
- Integrate with Cycles device enumeration code
- Revive `path_source_replace_includes` in util/path (required for MSL compilation)
This patch also includes a couple of small kernel-side fixes:
- Add an implementation of `lgammaf` for Metal [Nemes, Gergő (2010), "New asymptotic expansion for the Gamma function", Archiv der Mathematik](https://users.renyi.hu/~gergonemes/)
- include "work_stealing.h" inside the Metal context class because it accesses state now
Ref T92212
Reviewed By: brecht
Maniphest Tasks: T92212
Differential Revision: https://developer.blender.org/D13423
Diffstat (limited to 'intern/cycles/device/metal/queue.h')
-rw-r--r-- | intern/cycles/device/metal/queue.h | 97 |
1 files changed, 97 insertions, 0 deletions
diff --git a/intern/cycles/device/metal/queue.h b/intern/cycles/device/metal/queue.h new file mode 100644 index 00000000000..7aafcb2efe4 --- /dev/null +++ b/intern/cycles/device/metal/queue.h @@ -0,0 +1,97 @@ +/* + * Copyright 2021 Blender Foundation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#ifdef WITH_METAL + +# include "device/kernel.h" +# include "device/memory.h" +# include "device/queue.h" + +# include "device/metal/util.h" +# include "kernel/device/metal/globals.h" + +# define metal_printf VLOG(4) << string_printf + +CCL_NAMESPACE_BEGIN + +class MetalDevice; + +/* Base class for Metal queues. */ +class MetalDeviceQueue : public DeviceQueue { + public: + MetalDeviceQueue(MetalDevice *device); + ~MetalDeviceQueue(); + + virtual int num_concurrent_states(const size_t) const override; + virtual int num_concurrent_busy_states() const override; + + virtual void init_execution() override; + + virtual bool enqueue(DeviceKernel kernel, + const int work_size, + DeviceKernelArguments const &args) override; + + virtual bool synchronize() override; + + virtual void zero_to_device(device_memory &mem) override; + virtual void copy_to_device(device_memory &mem) override; + virtual void copy_from_device(device_memory &mem) override; + + virtual bool kernel_available(DeviceKernel kernel) const override; + + protected: + void prepare_resources(DeviceKernel kernel); + + id<MTLComputeCommandEncoder> get_compute_encoder(DeviceKernel kernel); + id<MTLBlitCommandEncoder> get_blit_encoder(); + + MetalDevice *metal_device; + MetalBufferPool temp_buffer_pool; + + API_AVAILABLE(macos(11.0), ios(14.0)) + MTLCommandBufferDescriptor *command_buffer_desc = nullptr; + id<MTLDevice> mtlDevice = nil; + id<MTLCommandQueue> mtlCommandQueue = nil; + id<MTLCommandBuffer> mtlCommandBuffer = nil; + id<MTLComputeCommandEncoder> mtlComputeEncoder = nil; + id<MTLBlitCommandEncoder> mtlBlitEncoder = nil; + id<MTLSharedEvent> shared_event = nil; + MTLSharedEventListener *shared_event_listener = nil; + + dispatch_queue_t event_queue; + dispatch_semaphore_t wait_semaphore; + + struct CopyBack { + void *host_pointer; + void *gpu_mem; + uint64_t size; + }; + std::vector<CopyBack> copy_back_mem; + + uint64_t shared_event_id; + uint64_t command_buffers_submitted = 0; + uint64_t command_buffers_completed = 0; + Stats &stats; + + void close_compute_encoder(); + void close_blit_encoder(); +}; + +CCL_NAMESPACE_END + +#endif /* WITH_METAL */ |