diff options
author | Brecht Van Lommel <brecht@blender.org> | 2022-08-18 21:31:34 +0300 |
---|---|---|
committer | Brecht Van Lommel <brecht@blender.org> | 2022-08-18 21:31:34 +0300 |
commit | 9961aae1e600f2154a1e14f2ae9ab21936bf781a (patch) | |
tree | 0600e98b5566e1f16f905644c3b80d41e7bf89fd /intern/cycles | |
parent | d2255aa4ed6d6b3fc3a42871a649682e357a305e (diff) | |
parent | 6a4f4810f38b2efc49d55dad6960f610f166773f (diff) |
Merge branch 'blender-v3.3-release'
Diffstat (limited to 'intern/cycles')
-rw-r--r-- | intern/cycles/integrator/path_trace_work_gpu.cpp | 14 | ||||
-rw-r--r-- | intern/cycles/kernel/device/metal/compat.h | 9 |
2 files changed, 14 insertions, 9 deletions
diff --git a/intern/cycles/integrator/path_trace_work_gpu.cpp b/intern/cycles/integrator/path_trace_work_gpu.cpp index fa313f6460a..ee250a6916b 100644 --- a/intern/cycles/integrator/path_trace_work_gpu.cpp +++ b/intern/cycles/integrator/path_trace_work_gpu.cpp @@ -204,22 +204,26 @@ void PathTraceWorkGPU::alloc_integrator_sorting() integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE] = (int *)integrator_shader_sort_counter_.device_pointer; - if (device_scene_->data.kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) { + integrator_shader_sort_prefix_sum_.alloc(sort_buckets); + integrator_shader_sort_prefix_sum_.zero_to_device(); + } + + if (device_scene_->data.kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) { + if (integrator_shader_raytrace_sort_counter_.size() < sort_buckets) { integrator_shader_raytrace_sort_counter_.alloc(sort_buckets); integrator_shader_raytrace_sort_counter_.zero_to_device(); integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE] = (int *)integrator_shader_raytrace_sort_counter_.device_pointer; } + } - if (device_scene_->data.kernel_features & KERNEL_FEATURE_MNEE) { + if (device_scene_->data.kernel_features & KERNEL_FEATURE_MNEE) { + if (integrator_shader_mnee_sort_counter_.size() < sort_buckets) { integrator_shader_mnee_sort_counter_.alloc(sort_buckets); integrator_shader_mnee_sort_counter_.zero_to_device(); integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE] = (int *)integrator_shader_mnee_sort_counter_.device_pointer; } - - integrator_shader_sort_prefix_sum_.alloc(sort_buckets); - integrator_shader_sort_prefix_sum_.zero_to_device(); } } diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h index b20cfca9a9c..a04261011f0 100644 --- a/intern/cycles/kernel/device/metal/compat.h +++ b/intern/cycles/kernel/device/metal/compat.h @@ -29,11 +29,12 @@ using namespace metal::raytracing; /* Qualifiers */ -#if defined(__KERNEL_METAL_APPLE__) +/* Inline everything for Apple GPUs. This gives ~1.1x speedup and 10% spill + * reduction for integator_shade_surface. However it comes at the cost of + * longer compile times (~4.5 minutes on M1 Max) and is disabled for that + * reason, until there is a user option to manually enable it. */ -/* Inline everything for Apple GPUs. - * This gives ~1.1x speedup and 10% spill reduction for integator_shade_surface - * at the cost of longer compile times (~4.5 minutes on M1 Max). */ +#if 0 // defined(__KERNEL_METAL_APPLE__) # define ccl_device __attribute__((always_inline)) # define ccl_device_inline __attribute__((always_inline)) |