diff options
| author | Brecht Van Lommel <brecht@blender.org> | 2022-04-28 01:46:14 +0300 |
|---|---|---|
| committer | Brecht Van Lommel <brecht@blender.org> | 2022-04-28 01:46:43 +0300 |
| commit | 52a5f68562680c0ccd6d4e525098bb5e2af7d0bd (patch) | |
| tree | e370e4a7fca575aad8ff8e9f1b975b4fd98a9c04 /intern/cycles/device/metal/queue.mm | |
| parent | 3558f565f1e8a8e5dc49067cc0500cbf993af69e (diff) | |
Revert "Cycles: Enable inlining on Apple Silicon for 1.1x speedup"
This reverts commit b82de02e7ce857e20b842a074c0068b146a9fd79. It is causing
crashes in various regression tests.
Ref D14763
Diffstat (limited to 'intern/cycles/device/metal/queue.mm')
| -rw-r--r-- | intern/cycles/device/metal/queue.mm | 9 |
1 files changed, 4 insertions, 5 deletions
diff --git a/intern/cycles/device/metal/queue.mm b/intern/cycles/device/metal/queue.mm index 2079aa65499..1686ab95ffa 100644 --- a/intern/cycles/device/metal/queue.mm +++ b/intern/cycles/device/metal/queue.mm @@ -108,6 +108,9 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel, VLOG(3) << "Metal queue launch " << device_kernel_as_string(kernel) << ", work_size " << work_size; + const MetalDeviceKernel &metal_kernel = metal_device->kernels.get(kernel); + const MetalKernelPipeline &metal_kernel_pso = metal_kernel.get_pso(); + id<MTLComputeCommandEncoder> mtlComputeCommandEncoder = get_compute_encoder(kernel); /* Determine size requirement for argument buffer. */ @@ -209,8 +212,6 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel, } bytes_written = globals_offsets + sizeof(KernelParamsMetal); - const MetalKernelPipeline &metal_kernel_pso = metal_device->get_best_pipeline(kernel); - /* Encode ancillaries */ [metal_device->mtlAncillaryArgEncoder setArgumentBuffer:arg_buffer offset:metal_offsets]; [metal_device->mtlAncillaryArgEncoder setBuffer:metal_device->texture_bindings_2d @@ -283,7 +284,7 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel, [mtlComputeCommandEncoder setComputePipelineState:metal_kernel_pso.pipeline]; /* Compute kernel launch parameters. */ - const int num_threads_per_block = metal_kernel_pso.num_threads_per_block; + const int num_threads_per_block = metal_kernel.get_num_threads_per_block(); int shared_mem_bytes = 0; @@ -546,8 +547,6 @@ id<MTLComputeCommandEncoder> MetalDeviceQueue::get_compute_encoder(DeviceKernel computeCommandEncoderWithDispatchType:concurrent ? MTLDispatchTypeConcurrent : MTLDispatchTypeSerial]; - [mtlComputeEncoder setLabel:@(device_kernel_as_string(kernel))]; - /* declare usage of MTLBuffers etc */ prepare_resources(kernel); } |