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 _______________________________________________ Bf-blender-cvs mailing list Bf-blender-cvs@blender.org List details, subscription details or unsubscribe: https://lists.blender.org/mailman/listinfo/bf-blender-cvs