[Bf-blender-cvs] [021c8c7cd0c] blender-v3.3-release: Cycles: Tweak inlining policy on Metal

[Michael Jones]

Commit: 021c8c7cd0c7472eb182d72c11d7201faa13c1f2
Author: Michael Jones
Date:   Tue Sep 27 17:01:17 2022 +0100
Branches: blender-v3.3-release
https://developer.blender.org/rB021c8c7cd0c7472eb182d72c11d7201faa13c1f2

Cycles: Tweak inlining policy on Metal

This patch optimises the Metal inlining policy. It gives a small speedup
(2-3% on M1 Max) with no notable compilation slowdown vs what is already
in master. Previously noted compilation slowdowns (as reported in T100102)
were caused by forcing inlining for `ccl_device`, but we get better
rendering perf by relying on compiler heuristics in these cases.

Backported to 3.3 because this also fixes a test failure.

Differential Revision: https://developer.blender.org/D16081

===================================================================

M	intern/cycles/kernel/device/metal/compat.h

===================================================================

diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h
index b86d1f64307..c321f4451f6 100644
--- a/intern/cycles/kernel/device/metal/compat.h
+++ b/intern/cycles/kernel/device/metal/compat.h
@@ -29,28 +29,13 @@ using namespace metal::raytracing;
 
 /* Qualifiers */
 
-/* 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. */
-
-#if 0  // defined(__KERNEL_METAL_APPLE__)
-
-#  define ccl_device __attribute__((always_inline))
-#  define ccl_device_inline __attribute__((always_inline))
-#  define ccl_device_forceinline __attribute__((always_inline))
-#  define ccl_device_noinline __attribute__((always_inline))
-
+#define ccl_device
+#define ccl_device_inline ccl_device __attribute__((always_inline))
+#define ccl_device_forceinline ccl_device __attribute__((always_inline))
+#if defined(__KERNEL_METAL_APPLE__)
+#  define ccl_device_noinline ccl_device
 #else
-
-#  define ccl_device
-#  define ccl_device_inline ccl_device
-#  define ccl_device_forceinline ccl_device
-#  if defined(__KERNEL_METAL_APPLE__)
-#    define ccl_device_noinline ccl_device
-#  else
-#    define ccl_device_noinline ccl_device __attribute__((noinline))
-#  endif
+#  define ccl_device_noinline ccl_device __attribute__((noinline))
 #endif
 
 #define ccl_device_noinline_cpu ccl_device