From 04166ea0ea7dad9c4642ddd86fd005a0c16257f6 Mon Sep 17 00:00:00 2001 From: Weizhen Huang Date: Wed, 1 Oct 2025 14:25:44 +0200 Subject: [PATCH] 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 | 10 ++++++++++ intern/cycles/util/defines.h | 15 +++++++++++++++ intern/cycles/util/types_base.h | 8 ++++++-- intern/cycles/util/types_float2.h | 8 ++++++-- intern/cycles/util/types_float3.h | 11 ++++++++--- intern/cycles/util/types_float4.h | 9 +++++++-- 6 files changed, 52 insertions(+), 9 deletions(-) diff --git a/intern/cycles/device/metal/device_impl.mm b/intern/cycles/device/metal/device_impl.mm index 2e53a6ac47a..d8836cdd256 100644 --- a/intern/cycles/device/metal/device_impl.mm +++ b/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))), diff --git a/intern/cycles/util/defines.h b/intern/cycles/util/defines.h index d2464242ce5..9a161f28371 100644 --- a/intern/cycles/util/defines.h +++ b/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 diff --git a/intern/cycles/util/types_base.h b/intern/cycles/util/types_base.h index cbc50971c58..235455a283f 100644 --- a/intern/cycles/util/types_base.h +++ b/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. */ diff --git a/intern/cycles/util/types_float2.h b/intern/cycles/util/types_float2.h index 82c89c29af5..80dad970ec6 100644 --- a/intern/cycles/util/types_float2.h +++ b/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 diff --git a/intern/cycles/util/types_float3.h b/intern/cycles/util/types_float3.h index 44a4c57f78d..157a4756eaf 100644 --- a/intern/cycles/util/types_float3.h +++ b/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) { diff --git a/intern/cycles/util/types_float4.h b/intern/cycles/util/types_float4.h index d8980306749..22ae3d9c7ea 100644 --- a/intern/cycles/util/types_float4.h +++ b/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