Cycles: support printf in Metal 15.0

For metal version after 3.2 it's possible to log debugging messages, it
works similar to `printf()`, except for a few differences:
- `%s` is not supported,
- `double` doesn't exist, so no casting to double for `%f`,
- no `\n` needed at the end of the format string.

To see the print in the console, environment variables `MTL_LOG_LEVEL`
should be set to `MTLLogLevelDebug`, and `MTL_LOG_TO_STDERR` should be
set to `1`. See
https://developer.apple.com/documentation/metal/logging-shader-debug-messages

Right now `printf()`, `print_float()`, `print_float2()`,
`print_float3()` and `print_float4()` are supported.

Thanks to @fclem for finding this out.

Pull Request: https://projects.blender.org/blender/blender/pulls/146585

intern/cycles/device/metal/device_impl.mm

@@ -479,6 +479,16 @@ void MetalDevice::compile_and_load(const int device_id, MetalPipelineType pso_ty
       options.languageVersion = MTLLanguageVersion3_1;
     }
 #  endif
+#  if defined(MAC_OS_VERSION_15_0)
+    if (@available(macos 15.0, *)) {
+      options.languageVersion = MTLLanguageVersion3_2;
+      if (const char *loglevel = getenv("MTL_LOG_LEVEL")) {
+        if (strcmp(loglevel, "MTLLogLevelDebug") == 0) {
+          options.enableLogging = true;
+        }
+      }
+    }
+#  endif
 
     if (getenv("CYCLES_METAL_PROFILING") || getenv("CYCLES_METAL_DEBUG")) {
       path_write_text(path_cache_get(string_printf("%s.metal", kernel_type_as_string(pso_type))),

intern/cycles/util/defines.h

@@ -114,3 +114,18 @@
 
 #define CONCAT_HELPER(a, ...) a##__VA_ARGS__
 #define CONCAT(a, ...) CONCAT_HELPER(a, __VA_ARGS__)
+
+#if (defined __KERNEL_METAL__) && (__METAL_VERSION__ >= 320)
+#  define __METAL_PRINTF__
+#endif
+
+/* Metal's logging works very similar to `printf()`, except for a few differences:
+ * - %s is not supported,
+ * - double doesn't exist, so no casting to double for %f,
+ * - no `\n` needed at the end of the format string.
+ * NOTE: To see the print in the console, environment variables `MTL_LOG_LEVEL` should be set to
+ * `MTLLogLevelDebug`, and `MTL_LOG_TO_STDERR` should be set to `1`.
+ * See https://developer.apple.com/documentation/metal/logging-shader-debug-messages */
+#  ifdef __METAL_PRINTF__
+#    define printf(...) metal::os_log_default.log_debug(__VA_ARGS__)
+#  endif

intern/cycles/util/types_base.h

@@ -76,12 +76,16 @@ CCL_NAMESPACE_END
 #  define __KERNEL_PRINTF__
 #endif
 
+#if defined __METAL_PRINTF__
+#  define print_float(label, a) metal::os_log_default.log_debug(label ": %.8f", a)
+#else
 ccl_device_inline void print_float(const ccl_private char *label, const float a)
 {
-#ifdef __KERNEL_PRINTF__
+#  ifdef __KERNEL_PRINTF__
   printf("%s: %.8f\n", label, (double)a);
-#endif
+#  endif
 }
+#endif
 
 /* Most GPU APIs matching native vector types, so we only need to implement them for
  * CPU and oneAPI. */

intern/cycles/util/types_float2.h

@@ -57,11 +57,15 @@ ccl_device_inline int2 make_int2(const float2 f)
   return make_int2((int)f.x, (int)f.y);
 }
 
+#if defined __METAL_PRINTF__
+#  define print_float2(label, a) metal::os_log_default.log_debug(label ": %.8f %.8f", a.x, a.y)
+#else
 ccl_device_inline void print_float2(const ccl_private char *label, const float2 a)
 {
-#ifdef __KERNEL_PRINTF__
+#  ifdef __KERNEL_PRINTF__
   printf("%s: %.8f %.8f\n", label, (double)a.x, (double)a.y);
-#endif
+#  endif
 }
+#endif
 
 CCL_NAMESPACE_END

intern/cycles/util/types_float3.h

@@ -122,15 +122,20 @@ ccl_device_inline float3 make_float3(const float3 a)
   return a;
 }
 
+#if defined __METAL_PRINTF__
+#  define print_float3(label, a) \
+    metal::os_log_default.log_debug(label ":  %.8f %.8f %.8f", a.x, a.y, a.z)
+#else
 ccl_device_inline void print_float3(const ccl_private char *label, const float3 a)
 {
-#ifdef __KERNEL_PRINTF__
+#  ifdef __KERNEL_PRINTF__
   printf("%s: %.8f %.8f %.8f\n", label, (double)a.x, (double)a.y, (double)a.z);
-#else
+#  else
   (void)label;
   (void)a;
-#endif
+#  endif
 }
+#endif
 
 ccl_device_inline float2 make_float2(const float3 a)
 {

intern/cycles/util/types_float4.h

@@ -120,11 +120,16 @@ ccl_device_inline int4 make_int4(const float4 f)
 #endif
 }
 
+#if defined __METAL_PRINTF__
+#  define print_float4(label, a) \
+    metal::os_log_default.log_debug(label ": %.8f %.8f %.8f %.8f", a.x, a.y, a.z, a.w)
+#else
 ccl_device_inline void print_float4(const ccl_private char *label, const float4 a)
 {
-#ifdef __KERNEL_PRINTF__
+#  ifdef __KERNEL_PRINTF__
   printf("%s: %.8f %.8f %.8f %.8f\n", label, (double)a.x, (double)a.y, (double)a.z, (double)a.w);
-#endif
+#  endif
 }
+#endif
 
 CCL_NAMESPACE_END