@mtlprintf

docs/src/api/kernel.md

@@ -54,3 +54,12 @@ MemoryFlags
 threadgroup_barrier
 simdgroup_barrier
 ```
+
+## Printing
+
+```@docs
+@mtlprintf
+@mtlprint
+@mtlprintln
+@mtlshow
+```

docs/src/usage/kernel.md

@@ -90,6 +90,34 @@ Additional notes:
 - Kernels must always return nothing
 - Kernels are asynchronous. To synchronize, use the `Metal.@sync` macro.
 
+## Printing
+
+When debugging, it's not uncommon to want to print some values. This is achieved with `@mtlprintf`:
+
+```julia
+function gpu_add2_print!(y, x)
+    index = thread_position_in_grid_1d()
+    @mtlprintf("thread %d", index)
+    @inbounds y[i] += x[i]
+    return nothing
+end
+
+A = Metal.ones(Float32, 8);
+B = Metal.rand(Float32, 8);
+
+@metal threads=length(A) gpu_add2_print!(A, B)
+```
+
+`@mtlprintf` is supported on macOS 15 and later. `@mtlprintf` support most of the format specifiers that `printf`
+supports in C with the following exceptions:
+ - `%n` and `%s` conversion specifiers are not supported
+ - Default argument promotion applies to arguments of half type which promote to the `double` type
+ - The format string must be a string literal
+
+Metal places output from `@mtlprintf` into a log buffer. The system only removes the messages from the log buffer when the command buffer completes. When the log buffer becomes full, the system drops all subsequent messages.
+
+See also: `@mtlprint`, `@mtlprintln` and `@mtlshow`
+
 ## Other Helpful Links
 
 [Metal Shading Language Specification](https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf)

lib/mtl/MTL.jl

@@ -34,6 +34,7 @@ include("events.jl")
 include("fences.jl")
 include("heap.jl")
 include("buffer.jl")
+include("log_state.jl")
 include("residency_set.jl")
 include("command_queue.jl")
 include("command_buf.jl")

lib/mtl/command_queue.jl

@@ -1,3 +1,24 @@
+
+export MTLCommandQueueDescriptor
+
+# @objcwrapper immutable=false MTLCommandQueueDescriptor <: NSObject
+
+function MTLCommandQueueDescriptor()
+    handle = @objc [MTLCommandQueueDescriptor alloc]::id{MTLCommandQueueDescriptor}
+    obj = MTLCommandQueueDescriptor(handle)
+    finalizer(release, obj)
+    @objc [obj::id{MTLCommandQueueDescriptor} init]::id{MTLCommandQueueDescriptor}
+    return obj
+end
+
+function MTLCommandQueue(dev::MTLDevice, descriptor::MTLCommandQueueDescriptor)
+    handle = @objc [dev::id{MTLDevice} newCommandQueueWithDescriptor:descriptor::id{MTLCommandQueueDescriptor}]::id{MTLCommandQueue}
+    obj = MTLCommandQueue(handle)
+    finalizer(release, obj)
+    return obj
+end
+
+
 export MTLCommandQueue
 
 # @objcwrapper immutable=false MTLCommandQueue <: NSObject
@@ -8,3 +29,4 @@ function MTLCommandQueue(dev::MTLDevice)
     finalizer(release, obj)
     return obj
 end
+

lib/mtl/libmtl.jl

@@ -1618,7 +1618,7 @@ end
     @autoproperty dispatchType::MTLDispatchType
 end
 
-@objcwrapper immutable = true availability = macos(v"15.0.0") MTLCommandQueueDescriptor <: NSObject
+@objcwrapper immutable = false availability = macos(v"15.0.0") MTLCommandQueueDescriptor <: NSObject
 
 @objcproperties MTLCommandQueueDescriptor begin
     @autoproperty maxCommandBufferCount::UInt64 setter = setMaxCommandBufferCount
@@ -3101,7 +3101,7 @@ end
     MTLLogLevelFault = 5
 end
 
-@objcwrapper immutable = true availability = macos(v"15.0.0") MTLLogStateDescriptor <: NSObject
+@objcwrapper immutable = false availability = macos(v"15.0.0") MTLLogStateDescriptor <: NSObject
 
 @objcproperties MTLLogStateDescriptor begin
     @autoproperty level::MTLLogLevel setter = setLevel

lib/mtl/log_state.jl

@@ -0,0 +1,26 @@
+export MTLLogLevel
+
+export MTLLogStateDescriptor
+
+# @objcwrapper immutable = false MTLLogStateDescriptor <: NSObject
+
+function MTLLogStateDescriptor()
+    handle = @objc [MTLLogStateDescriptor alloc]::id{MTLLogStateDescriptor}
+    obj = MTLLogStateDescriptor(handle)
+    finalizer(release, obj)
+    @objc [obj::id{MTLLogStateDescriptor} init]::id{MTLLogStateDescriptor}
+    return obj
+end
+
+
+export MTLLogState
+
+# @objcwrapper immutable = true MTLLogState <: NSObject
+
+function MTLLogState(dev::MTLDevice, descriptor::MTLLogStateDescriptor)
+    err = Ref{id{NSError}}(nil)
+    handle = @objc [dev::id{MTLDevice} newLogStateWithDescriptor:descriptor::id{MTLLogStateDescriptor}
+        error:err::Ptr{id{NSError}}]::id{MTLLogState}
+    err[] == nil || throw(NSError(err[]))
+    MTLLogState(handle)
+end

res/wrap/libmtl.toml

@@ -50,6 +50,9 @@ immutable=false
 [api.MTLCommandQueue]
 immutable=false
 
+[api.MTLCommandQueueDescriptor]
+immutable=false
+
 [api.MTL4CommandQueueDescriptor]
 immutable=false
 
@@ -106,6 +109,9 @@ immutable=false
 [api.MTLLibrary]
 immutable=false
 
+[api.MTLLogStateDescriptor]
+immutable=false
+
 [api.MTLResidencySet]
 immutable=false
 

src/Metal.jl

@@ -37,6 +37,7 @@ include("device/intrinsics/synchronization.jl")
 include("device/intrinsics/memory.jl")
 include("device/intrinsics/simd.jl")
 include("device/intrinsics/atomics.jl")
+include("device/intrinsics/output.jl")
 include("device/random.jl")
 include("device/quirks.jl")
 

src/MetalKernels.jl

@@ -195,7 +195,7 @@ end
 end
 
 @device_override @inline function KA.__print(args...)
-    # TODO
+    Metal._mtlprint(args...)
 end
 ## COV_EXCL_STOP
 

src/compiler/compilation.jl

@@ -169,9 +169,9 @@ function compile(@nospecialize(job::CompilerJob))
 
     @signpost_interval log=log_compiler() "Generate LLVM IR" begin
         # TODO: on 1.9, this actually creates a context. cache those.
-        ir, entry = JuliaContext() do ctx
+        ir, entry, loggingEnabled = JuliaContext() do ctx
             mod, meta = GPUCompiler.compile(:llvm, job)
-            string(mod), LLVM.name(meta.entry)
+            string(mod), LLVM.name(meta.entry), haskey(functions(mod), "air.os_log")
         end
     end
 
@@ -237,7 +237,7 @@ function compile(@nospecialize(job::CompilerJob))
         end
     end
 
-    return (; ir, air, metallib, entry)
+    return (; ir, air, metallib, entry, loggingEnabled)
 end
 
 # link into an executable kernel
@@ -275,5 +275,5 @@ end
         end
     end
 
-    pipeline_state
+    pipeline_state, compiled.loggingEnabled
 end

src/compiler/execution.jl

@@ -161,6 +161,7 @@ mtlconvert(arg, cce=nothing) = adapt(Adaptor(cce), arg)
 struct HostKernel{F,TT}
     f::F
     pipeline::MTLComputePipelineState
+    loggingEnabled::Bool
 end
 
 const mtlfunction_lock = ReentrantLock()
@@ -186,15 +187,15 @@ function mtlfunction(f::F, tt::TT=Tuple{}; name=nothing, kwargs...) where {F,TT}
         cache = compiler_cache(dev)
         source = methodinstance(F, tt)
         config = compiler_config(dev; name, kwargs...)::MetalCompilerConfig
-        pipeline = GPUCompiler.cached_compilation(cache, source, config, compile, link)
+        pipeline, loggingEnabled = GPUCompiler.cached_compilation(cache, source, config, compile, link)
 
         # create a callable object that captures the function instance. we don't need to think
         # about world age here, as GPUCompiler already does and will return a different object
         h = hash(pipeline, hash(f, hash(tt)))
         kernel = get(_kernel_instances, h, nothing)
         if kernel === nothing
             # create the kernel state object
-            kernel = HostKernel{F,tt}(f, pipeline)
+            kernel = HostKernel{F, tt}(f, pipeline, loggingEnabled)
             _kernel_instances[h] = kernel
         end
         return kernel::HostKernel{F,tt}
@@ -277,7 +278,35 @@ end
 
     kernel_state = KernelState(Random.rand(UInt32))
 
-    cmdbuf = MTLCommandBuffer(queue)
+    cmdbuf = if kernel.loggingEnabled
+        # TODO: make this a dynamic error, i.e., from the kernel (JuliaGPU/Metal.jl#433)
+        @static if !is_macos(v"15.0.0")
+            error("Logging is only supported on macOS 15 or higher")
+        end
+
+        if MTLCaptureManager().isCapturing
+            error("Logging is not supported while GPU frame capturing")
+        end
+
+        log_state_descriptor = MTLLogStateDescriptor()
+        log_state_descriptor.level = MTL.MTLLogLevelDebug
+        log_state = MTLLogState(queue.device, log_state_descriptor)
+
+        function log_handler(subSystem, category, logLevel, message)
+            Core.print(String(NSString(message)))
+            return nothing
+        end
+
+        block = @objcblock(log_handler, Nothing, (id{NSString}, id{NSString}, NSInteger, id{NSString}))
+        @objc [log_state::id{MTLLogState} addLogHandler:block::id{NSBlock}]::Nothing
+
+        cmdbuf_descriptor = MTLCommandBufferDescriptor()
+        cmdbuf_descriptor.logState = log_state
+        MTLCommandBuffer(queue, cmdbuf_descriptor)
+    else
+        MTLCommandBuffer(queue)
+    end
+
     cmdbuf.label = "MTLCommandBuffer($(nameof(kernel.f)))"
     cce = MTLComputeCommandEncoder(cmdbuf)
     argument_buffers = try