diff --git a/LocalPreferences.toml b/LocalPreferences.toml index a2eaba4fc..b0d05202a 100644 --- a/LocalPreferences.toml +++ b/LocalPreferences.toml @@ -2,3 +2,6 @@ # which storage mode unspecified allocations should default to. # possible values: "private", "shared", "managed" #default_storage = "private" + +# when false, always use metal 3 even on a metal 4 platform +#force_metal3 = false diff --git a/lib/mtl/MTL.jl b/lib/mtl/MTL.jl index 74a6f0921..d02272c00 100644 --- a/lib/mtl/MTL.jl +++ b/lib/mtl/MTL.jl @@ -27,7 +27,7 @@ include("size.jl") include("device.jl") include("resource.jl") include("storage_type.jl") -include("compile-opts.jl") +include("compile_opts.jl") include("library.jl") include("function.jl") include("events.jl") @@ -35,11 +35,18 @@ include("fences.jl") include("heap.jl") include("buffer.jl") include("command_queue.jl") +include("command_queue4.jl") include("command_buf.jl") +include("command_buf4.jl") include("compute_pipeline.jl") +include("compute_pipeline4.jl") include("command_enc.jl") +include("command_enc4.jl") +include("command_alloc4.jl") +include("arg_table.jl") include("command_enc/blit.jl") include("command_enc/compute.jl") +include("command_enc/compute4.jl") include("binary_archive.jl") include("capture.jl") include("texture.jl") diff --git a/lib/mtl/arg_table.jl b/lib/mtl/arg_table.jl new file mode 100644 index 000000000..a8d6a465e --- /dev/null +++ b/lib/mtl/arg_table.jl @@ -0,0 +1,36 @@ +# +# command buffer +# + +export MTL4ArgumentTable, MTL4ArgumentTableDescriptor + +function MTL4ArgumentTableDescriptor() + desc = @objc [MTL4ArgumentTableDescriptor alloc]::id{MTL4ArgumentTableDescriptor} + obj = MTL4ArgumentTableDescriptor(desc) + return obj +end + +function MTL4ArgumentTable(device::MTLDevice, desc::MTL4ArgumentTableDescriptor) + err = Ref{id{NSError}}(nil) + handle = @objc [device::id{MTLDevice} newArgumentTableWithDescriptor:desc::id{MTL4ArgumentTableDescriptor} + error:err::Ptr{id{NSError}}]::id{MTL4ArgumentTable} + err[] == nil || throw(NSError(err[])) + obj = MTL4ArgumentTable(handle) + # finalizer(release, obj) + return obj +end + +# Buffer Arguments +function set_address!(cce::MTL4ArgumentTable, address, bindingIndex) + @objc [cce::id{MTL4ArgumentTable} setAddress:address::NSUInteger + atIndex:(bindingIndex-1)::NSUInteger]::Nothing +end + +function set_address!(cce::MTL4ArgumentTable, buf::MTLBuffer, bindingIndex) + @objc [cce::id{MTL4ArgumentTable} setAddress:contents(buf)::NSUInteger + atIndex:(bindingIndex-1)::NSUInteger]::Nothing +end + +function set_buffer!(cce::MTL4ArgumentTable, buf::MTLBuffer, offset, index) + @inline set_address!(cce, contents(buf)+offset, index) +end diff --git a/lib/mtl/command_alloc4.jl b/lib/mtl/command_alloc4.jl new file mode 100644 index 000000000..e206c20b3 --- /dev/null +++ b/lib/mtl/command_alloc4.jl @@ -0,0 +1,53 @@ + +export MTL4CommandAllocatorDescriptor + +function MTL4CommandAllocatorDescriptor() + handle = @objc [MTL4CommandAllocatorDescriptor new]::id{MTL4CommandAllocatorDescriptor} + obj = MTL4CommandAllocatorDescriptor(handle) + finalizer(release, obj) + return obj +end +function MTL4CommandAllocatorDescriptor(label) + desc = MTL4CommandAllocatorDescriptor() + desc.label = label + return desc +end + + + +# +# command allocator +# + +export MTL4CommandAllocator + +# @objcwrapper immutable=false MTL4CommandAllocator <: NSObject + +function MTL4CommandAllocator(device::MTLDevice) + handle = @objc [device::id{MTLDevice} newCommandAllocator]::id{MTL4CommandAllocator} + obj = MTL4CommandAllocator(handle) + finalizer(release, obj) + return obj +end + +function MTL4CommandAllocator(dev::MTLDevice, descriptor::MTL4CommandAllocatorDescriptor) + err = Ref{id{NSError}}(nil) + handle = @objc [dev::id{MTLDevice} newCommandAllocatorWithDescriptor:descriptor::id{MTL4CommandAllocatorDescriptor} + error:err::Ptr{id{NSError}}]::id{MTL4CommandAllocator} + obj = MTL4CommandAllocator(handle) + finalizer(release, obj) + return obj +end + +function MTL4CommandAllocator(dev::MTLDevice, label) + desc = MTL4CommandAllocatorDescriptor(label) + return MTL4CommandAllocator(dev, desc) +end + +function allocatedSize(alloc::MTL4CommandAllocator)::UInt64 + @objc [alloc::id{MTL4CommandAllocator} allocatedSize]::UInt64 +end + +function reset!(alloc::MTL4CommandAllocator) + @objc [alloc::id{MTL4CommandAllocator} reset]::Nothing +end diff --git a/lib/mtl/command_buf4.jl b/lib/mtl/command_buf4.jl new file mode 100644 index 000000000..f2b51fdad --- /dev/null +++ b/lib/mtl/command_buf4.jl @@ -0,0 +1,59 @@ +# +# command buffer +# + +export MTL4CommandBuffer, commit!, beginCommandBufferWithAllocator!, endCommandBuffer! + +# @objcwrapper immutable=false MTL4CommandBuffer <: NSObject + +function MTL4CommandBuffer(device::MTLDevice, label=nothing) + handle = @objc [device::id{MTLDevice} newCommandBuffer]::id{MTL4CommandBuffer} + buf = MTL4CommandBuffer(handle) + if !isnothing(label) + buf.label = label + end + return buf +end + +function MTL4CommandBuffer(f::Base.Callable, device::MTLDevice, label=nothing; queue::MTL4CommandQueue=MTL4CommandQueue(device), allocator::MTL4CommandAllocator=MTL4CommandAllocator(device)) + cmdbuf = MTL4CommandBuffer(device, label) + + commit!(f, cmdbuf, queue, allocator) +end + +function beginCommandBufferWithAllocator!(cmdbuf::MTL4CommandBuffer, allocator::MTL4CommandAllocator, options::Union{Nothing, MTL4CommandBufferOptions} = nothing) + if isnothing(options) + @objc [cmdbuf::id{MTL4CommandBuffer} beginCommandBufferWithAllocator:allocator::id{MTL4CommandAllocator}]::Nothing + else + @objc [cmdbuf::id{MTL4CommandBuffer} beginCommandBufferWithAllocator:allocator::id{MTL4CommandAllocator} + options:options::id{MTL4CommandBufferOptions}]::Nothing + end +end + +function endCommandBuffer!(cmdbuf::MTL4CommandBuffer) + @objc [cmdbuf::id{MTL4CommandBuffer} endCommandBuffer]::Nothing +end + +function commit!(cmdqueue::MTL4CommandQueue, cmdbuf::MTL4CommandBuffer) + cmdbufRef = Ref{MTL4CommandBuffer}(cmdbuf) + @objc [cmdqueue::id{MTL4CommandQueue} commit:cmdbufRef::Ref{MTL4CommandBuffer} + count:1::NSUInteger]::Nothing +end +function commit!(cmdqueue::MTL4CommandQueue, cmdbuf::MTL4CommandBuffer, options::MTL4CommitOptions) + cmdbufRef = Ref{MTL4CommandBuffer}(cmdbuf) + @objc [cmdqueue::id{MTL4CommandQueue} commit:cmdbufRef::Ref{MTL4CommandBuffer} + count:1::NSUInteger + options:options::id{MTL4CommitOptions}]::Nothing +end + +function commit!(f::Base.Callable, cmdbuf::MTL4CommandBuffer, queue::MTL4CommandQueue, allocator::MTL4CommandAllocator) + beginCommandBufferWithAllocator!(cmdbuf, allocator) + + try + ret = f(cmdbuf) + return ret + finally + endCommandBuffer!(cmdbuf) + commit!(queue, cmdbuf) + end +end diff --git a/lib/mtl/command_enc/compute4.jl b/lib/mtl/command_enc/compute4.jl new file mode 100644 index 000000000..ecc4ec1ca --- /dev/null +++ b/lib/mtl/command_enc/compute4.jl @@ -0,0 +1,91 @@ +export MTL4ComputeCommandEncoder +export set_function!, set_buffer!, set_bytes!, set_texture!, set_sampler_state! +export stages +export dispatchThreadgroups!, dispatchThreads!, endEncoding! +export use!, memoryBarrier!, append_copy!, append_fillbuffer!, append_sync! + +# @objcwrapper immutable=false MTL4ComputeCommandEncoder <: MTL4CommandEncoder + +function MTL4ComputeCommandEncoder(cmdbuf::MTL4CommandBuffer) + handle = @objc [cmdbuf::id{MTL4CommandBuffer} computeCommandEncoder]::id{MTL4ComputeCommandEncoder} + obj = MTL4ComputeCommandEncoder(handle) + # finalizer(release, obj) + return obj +end + + +function MTL4ComputeCommandEncoder(f::Base.Callable, cmdbuf::MTL4CommandBuffer, sync=false) + encoder = MTL4ComputeCommandEncoder(cmdbuf) + try + f(encoder) + finally + sync && barrierAfterStages!(encoder) + close(encoder) + end +end + +# Pipeline State +function set_function!(cce::MTL4ComputeCommandEncoder, pipeline::MTLComputePipelineState) + @objc [cce::id{MTL4ComputeCommandEncoder} setComputePipelineState:pipeline::id{MTLComputePipelineState}]::Nothing +end + +function set_argument_table!(cce::MTL4ComputeCommandEncoder, arg_table::MTL4ArgumentTable) + @objc [cce::id{MTL4ComputeCommandEncoder} setArgumentTable:arg_table::id{MTL4ArgumentTable}]::Nothing +end + +# Dispatch Commands +function dispatchThreadgroups!(cce::MTL4ComputeCommandEncoder, gridSize::MTLSize, threadGroupSize::MTLSize) + @objc [cce::id{MTL4ComputeCommandEncoder} dispatchThreadgroups:gridSize::MTLSize + threadsPerThreadgroup:threadGroupSize::MTLSize]::Nothing +end + +function dispatchThreads!(cce::MTL4ComputeCommandEncoder, threadsSize::MTLSize, threadsPerThreadgroup::MTLSize) + @objc [cce::id{MTL4ComputeCommandEncoder} dispatchThreads:threadsSize::MTLSize + threadsPerThreadgroup:threadsPerThreadgroup::MTLSize]::Nothing +end + +# Copy Operations (Blit functionality integrated into compute encoder in Metal 4) +function append_copy!(cce::MTL4ComputeCommandEncoder, dst::MTLBuffer, dstOffset::Integer, + src::MTLBuffer, srcOffset::Integer, size::Integer) + @objc [cce::id{MTL4ComputeCommandEncoder} copyFromBuffer:src::id{MTLBuffer} + sourceOffset:srcOffset::NSUInteger + toBuffer:dst::id{MTLBuffer} + destinationOffset:dstOffset::NSUInteger + size:size::NSUInteger]::Nothing +end + +# function append_copy!(cce::MTL4ComputeCommandEncoder, dst::MTLTexture, dstSlice::Integer, dstLevel::Integer, dstOrigin::MTLOrigin, +# src::MTLBuffer, srcOffset::Integer, srcBytesPerRow::Integer, srcBytesPerImage::Integer, +# size::MTLSize) +# @objc [cce::id{MTL4ComputeCommandEncoder} copyFromBuffer:src::id{MTLBuffer} +# sourceOffset:srcOffset::NSUInteger +# sourceBytesPerRow:srcBytesPerRow::NSUInteger +# sourceBytesPerImage:srcBytesPerImage::NSUInteger +# sourceSize:size::MTLSize +# toTexture:dst::id{MTLTexture} +# destinationSlice:dstSlice::NSUInteger +# destinationLevel:dstLevel::NSUInteger +# destinationOrigin:dstOrigin::MTLOrigin]::Nothing +# end + +function stages(cce::MTL4ComputeCommandEncoder) + @objc [cce::id{MTL4ComputeCommandEncoder} stages]::MTLStages +end + +# Fill Buffer +function append_fillbuffer!(cce::MTL4ComputeCommandEncoder, buffer::MTLBuffer, range::NSRange, value::UInt8) + @objc [cce::id{MTL4ComputeCommandEncoder} fillBuffer:buffer::id{MTLBuffer} + range:range::NSRange + value:value::UInt8]::Nothing +end + +function append_fillbuffer!(cce::MTL4ComputeCommandEncoder, buffer::MTLBuffer, value::UInt8, + byteSize::Integer, offset::Integer=0) + range = NSRange(offset, byteSize) + append_fillbuffer!(cce, buffer, range, value) +end + +# Convenience dispatch function for encoding +function append_current_function!(cce::MTL4ComputeCommandEncoder, gridSize::MTLSize, threadGroupSize::MTLSize) + dispatchThreadgroups!(cce, gridSize, threadGroupSize) +end diff --git a/lib/mtl/command_enc4.jl b/lib/mtl/command_enc4.jl new file mode 100644 index 000000000..d2eda756e --- /dev/null +++ b/lib/mtl/command_enc4.jl @@ -0,0 +1,35 @@ +export endEncoding!, updateFence!, waitForFence +export barrierAfterEncoderStages!, barrierAfterQueueStages!, barrierAfterStages! + +# @objcwrapper immutable=true MTL4CommandEncoder <: NSObject + +function updateFence!(encoder::MTL4CommandEncoder, fence::MTLFence, afterEncoderStages::MTLStages=MTLStageAll) + @objc [encoder::id{MTL4CommandEncoder} updateFence:fence::id{MTLFence} + afterEncoderStages:afterEncoderStages::MTLStages]::Nothing +end + +function waitForFence(encoder::MTL4CommandEncoder, fence::MTLFence, beforeEncoderStages::MTLStages=MTLStageAll) + @objc [encoder::id{MTL4CommandEncoder} waitForFence:fence::id{MTLFence} + beforeEncoderStages:afterEncoderStages::MTLStages]::Nothing +end + +function barrierAfterEncoderStages!(encoder::MTL4CommandEncoder, afterEncoderStages::MTLStages=MTLStageAll, beforeEncoderStages::MTLStages=MTLStageAll, visibilityOptions::MTL4VisibilityOptions=MTL4VisibilityOptionResourceAlias) + @objc [encoder::id{MTL4CommandEncoder} barrierAfterEncoderStages:afterEncoderStages::MTLStages + beforeEncoderStages:beforeEncoderStages::MTLStages + visibilityOptions:visibilityOptions::MTL4VisibilityOptions]::Nothing +end + +function barrierAfterQueueStages!(encoder::MTL4CommandEncoder, afterQueueStages::MTLStages=MTLStageAll, beforeStages::MTLStages=MTLStageAll, visibilityOptions::MTL4VisibilityOptions=MTL4VisibilityOptionResourceAlias) + @objc [encoder::id{MTL4CommandEncoder} barrierAfterQueueStages:afterQueueStages::MTLStages + beforeStages:beforeStages::MTLStages + visibilityOptions:visibilityOptions::MTL4VisibilityOptions]::Nothing +end + +function barrierAfterStages!(encoder::MTL4CommandEncoder, afterStages::MTLStages=MTLStageAll, beforeQueueStages::MTLStages=MTLStageAll, visibilityOptions::MTL4VisibilityOptions=MTL4VisibilityOptionResourceAlias) + @objc [encoder::id{MTL4CommandEncoder} barrierAfterStages:afterStages::MTLStages + beforeQueueStages:beforeQueueStages::MTLStages + visibilityOptions:visibilityOptions::MTL4VisibilityOptions]::Nothing +end + +endEncoding!(ce::MTL4CommandEncoder) = @objc [ce::id{MTL4CommandEncoder} endEncoding]::Nothing +Base.close(ce::MTL4CommandEncoder) = endEncoding!(ce) diff --git a/lib/mtl/command_queue4.jl b/lib/mtl/command_queue4.jl new file mode 100644 index 000000000..d736408ee --- /dev/null +++ b/lib/mtl/command_queue4.jl @@ -0,0 +1,72 @@ +export MTL4CommitOptions + +# @objcwrapper immutable=false MTL4CommitOptions <: NSObject + +function MTL4CommitOptions() + handle = @objc [MTL4CommitOptions new]::id{MTL4CommitOptions} + obj = MTL4CommitOptions(handle) + finalizer(release, obj) + return obj +end +function MTL4CommitOptions(f::Base.Callable) + options = MTL4CommitOptions() + addFeedbackHandler(f, options) + return options +end + +function _command_buffer4_callback(f) + # convert the incoming pointer, and discard any return value + function wrapper(ptr) + try + f(ptr == nil ? nothing : MTL4CommitFeedback(ptr)) + catch err + # we might be on an unmanaged thread here, so display the error + # (otherwise it may get lost, or worse, crash Julia) + @error "Command buffer callback encountered an error: " * sprint(showerror, err) + end + return + end + @objcblock(wrapper, Nothing, (id{MTL4CommitFeedback},)) +end + +function addFeedbackHandler(f::Base.Callable, options::MTL4CommitOptions) + block = _command_buffer4_callback(f) + @objc [options::id{MTL4CommitOptions} addFeedbackHandler:block::id{NSBlock}]::Nothing +end + + + +export MTL4CommandQueueDescriptor + +function MTL4CommandQueueDescriptor() + handle = @objc [MTL4CommandQueueDescriptor new]::id{MTL4CommandQueueDescriptor} + obj = MTL4CommandQueueDescriptor(handle) + finalizer(release, obj) + return obj +end +function MTL4CommandQueueDescriptor(label) + desc = MTL4CommandQueueDescriptor() + desc.label = label + return desc +end + + +export MTL4CommandQueue + +# @objcwrapper immutable=false MTL4CommandQueue <: NSObject + +function MTL4CommandQueue(dev::MTLDevice) + handle = @objc [dev::id{MTLDevice} newMTL4CommandQueue]::id{MTL4CommandQueue} + obj = MTL4CommandQueue(handle) + finalizer(release, obj) + return obj +end + +function MTL4CommandQueue(dev::MTLDevice, descriptor::MTL4CommandQueueDescriptor) + err = Ref{id{NSError}}(nil) + handle = @objc [dev::id{MTLDevice} newMTL4CommandQueueWithDescriptor:descriptor::id{MTL4CommandQueueDescriptor} + error:err::Ptr{id{NSError}}]::id{MTL4CommandQueue} + obj = MTL4CommandQueue(handle) + finalizer(release, obj) + return obj +end diff --git a/lib/mtl/compile-opts.jl b/lib/mtl/compile_opts.jl similarity index 100% rename from lib/mtl/compile-opts.jl rename to lib/mtl/compile_opts.jl diff --git a/lib/mtl/compute_pipeline4.jl b/lib/mtl/compute_pipeline4.jl new file mode 100644 index 000000000..641c5afa3 --- /dev/null +++ b/lib/mtl/compute_pipeline4.jl @@ -0,0 +1,14 @@ +# +# compute pipeline descriptor +# + +export MTL4ComputePipelineDescriptor + +# @objcwrapper immutable=false MTL4ComputePipelineDescriptor <: NSObject + +function MTL4ComputePipelineDescriptor() + handle = @objc [MTL4ComputePipelineDescriptor new]::id{MTL4ComputePipelineDescriptor} + obj = MTL4ComputePipelineDescriptor(handle) + finalizer(release, obj) + return obj +end diff --git a/lib/mtl/libmtl.jl b/lib/mtl/libmtl.jl index 196bd7c8f..5d899d5cd 100644 --- a/lib/mtl/libmtl.jl +++ b/lib/mtl/libmtl.jl @@ -3,6 +3,9 @@ using CEnum: CEnum, @cenum +const _NSRange = NSRange + + @cenum MTLTextureSwizzle::UInt8 begin MTLTextureSwizzleZero = 0x0000000000000000 MTLTextureSwizzleOne = 0x0000000000000001 @@ -91,6 +94,20 @@ end MTLBarrierScopeRenderTargets = 0x0000000000000004 end +@cenum MTLStages::UInt64 begin + MTLStageVertex = 0x0000000000000001 + MTLStageFragment = 0x0000000000000002 + MTLStageTile = 0x0000000000000004 + MTLStageObject = 0x0000000000000008 + MTLStageMesh = 0x0000000000000010 + MTLStageResourceState = 0x0000000004000000 + MTLStageDispatch = 0x0000000008000000 + MTLStageBlit = 0x0000000010000000 + MTLStageAccelerationStructure = 0x0000000020000000 + MTLStageMachineLearning = 0x0000000040000000 + MTLStageAll = 0x7fffffffffffffff +end + @objcwrapper immutable = true availability = macos(v"14.0.0") MTLArchitecture <: NSObject @objcproperties MTLArchitecture begin @@ -227,6 +244,23 @@ end MTLResourceOptionCPUCacheModeWriteCombined = 0x0000000000000001 end +@cenum MTLSparsePageSize::Int64 begin + MTLSparsePageSize16 = 101 + MTLSparsePageSize64 = 102 + MTLSparsePageSize256 = 103 +end + +@cenum MTLBufferSparseTier::Int64 begin + MTLBufferSparseTierNone = 0 + MTLBufferSparseTier1 = 1 +end + +@cenum MTLTextureSparseTier::Int64 begin + MTLTextureSparseTierNone = 0 + MTLTextureSparseTier1 = 1 + MTLTextureSparseTier2 = 2 +end + @cenum MTLHeapType::Int64 begin MTLHeapTypeAutomatic = 0 MTLHeapTypePlacement = 1 @@ -402,102 +436,7 @@ end MTLPixelFormatDepth32Float_Stencil8 = 0x0000000000000104 MTLPixelFormatX32_Stencil8 = 0x0000000000000105 MTLPixelFormatX24_Stencil8 = 0x0000000000000106 -end - -@objcwrapper immutable = true MTLBuffer <: MTLResource - -@objcproperties MTLBuffer begin - @autoproperty length::UInt64 - @autoproperty remoteStorageBuffer::id{MTLBuffer} - @autoproperty gpuAddress::UInt64 type = Ptr{Cvoid} -end - -@cenum MTLTextureType::UInt64 begin - MTLTextureType1D = 0x0000000000000000 - MTLTextureType1DArray = 0x0000000000000001 - MTLTextureType2D = 0x0000000000000002 - MTLTextureType2DArray = 0x0000000000000003 - MTLTextureType2DMultisample = 0x0000000000000004 - MTLTextureTypeCube = 0x0000000000000005 - MTLTextureTypeCubeArray = 0x0000000000000006 - MTLTextureType3D = 0x0000000000000007 - MTLTextureType2DMultisampleArray = 0x0000000000000008 - MTLTextureTypeTextureBuffer = 0x0000000000000009 -end - -@objcwrapper immutable = true MTLSharedTextureHandle <: NSObject - -@objcproperties MTLSharedTextureHandle begin - @autoproperty device::id{MTLDevice} - @autoproperty label::id{NSString} -end - -@cenum MTLTextureUsage::UInt64 begin - MTLTextureUsageUnknown = 0x0000000000000000 - MTLTextureUsageShaderRead = 0x0000000000000001 - MTLTextureUsageShaderWrite = 0x0000000000000002 - MTLTextureUsageRenderTarget = 0x0000000000000004 - MTLTextureUsagePixelFormatView = 0x0000000000000010 - MTLTextureUsageShaderAtomic = 0x0000000000000020 -end - -@cenum MTLTextureCompressionType::Int64 begin - MTLTextureCompressionTypeLossless = 0 - MTLTextureCompressionTypeLossy = 1 -end - -@objcwrapper immutable = false MTLTextureDescriptor <: NSObject - -@objcproperties MTLTextureDescriptor begin - @autoproperty textureType::MTLTextureType setter = setTextureType - @autoproperty pixelFormat::MTLPixelFormat setter = setPixelFormat - @autoproperty width::UInt64 setter = setWidth - @autoproperty height::UInt64 setter = setHeight - @autoproperty depth::UInt64 setter = setDepth - @autoproperty mipmapLevelCount::UInt64 setter = setMipmapLevelCount - @autoproperty sampleCount::UInt64 setter = setSampleCount - @autoproperty arrayLength::UInt64 setter = setArrayLength - @autoproperty resourceOptions::MTLResourceOptions setter = setResourceOptions - @autoproperty cpuCacheMode::MTLCPUCacheMode setter = setCpuCacheMode - @autoproperty storageMode::MTLStorageMode setter = setStorageMode - @autoproperty hazardTrackingMode::MTLHazardTrackingMode setter = setHazardTrackingMode - @autoproperty usage::MTLTextureUsage setter = setUsage - @autoproperty allowGPUOptimizedContents::Bool setter = setAllowGPUOptimizedContents - @autoproperty compressionType::MTLTextureCompressionType setter = setCompressionType - @autoproperty swizzle::MTLTextureSwizzleChannels setter = setSwizzle -end - -@objcwrapper immutable = false MTLTexture <: MTLResource - -@objcproperties MTLTexture begin - @autoproperty rootResource::id{MTLResource} - @autoproperty parentTexture::id{MTLTexture} - @autoproperty parentRelativeLevel::UInt64 - @autoproperty parentRelativeSlice::UInt64 - @autoproperty buffer::id{MTLBuffer} - @autoproperty bufferOffset::UInt64 - @autoproperty bufferBytesPerRow::UInt64 - @autoproperty iosurface::Ptr{Cvoid} - @autoproperty iosurfacePlane::UInt64 - @autoproperty textureType::MTLTextureType - @autoproperty pixelFormat::MTLPixelFormat - @autoproperty width::UInt64 - @autoproperty height::UInt64 - @autoproperty depth::UInt64 - @autoproperty mipmapLevelCount::UInt64 - @autoproperty sampleCount::UInt64 - @autoproperty arrayLength::UInt64 - @autoproperty usage::MTLTextureUsage - @autoproperty shareable::Bool getter = isShareable - @autoproperty framebufferOnly::Bool getter = isFramebufferOnly - @autoproperty firstMipmapInTail::UInt64 - @autoproperty tailSizeInBytes::UInt64 - @autoproperty isSparse::Bool - @autoproperty allowGPUOptimizedContents::Bool - @autoproperty compressionType::MTLTextureCompressionType - @autoproperty gpuResourceID::MTLResourceID - @autoproperty remoteStorageTexture::id{MTLTexture} - @autoproperty swizzle::MTLTextureSwizzleChannels + MTLPixelFormatUnspecialized = 0x0000000000000107 end @cenum MTLDataType::UInt64 begin @@ -596,6 +535,179 @@ end MTLDataTypeBFloat2 = 0x000000000000007a MTLDataTypeBFloat3 = 0x000000000000007b MTLDataTypeBFloat4 = 0x000000000000007c + MTLDataTypeTensor = 0x000000000000008c +end + +@cenum MTLTensorDataType::Int64 begin + MTLTensorDataTypeNone = 0 + MTLTensorDataTypeFloat32 = 3 + MTLTensorDataTypeFloat16 = 16 + MTLTensorDataTypeBFloat16 = 121 + MTLTensorDataTypeInt8 = 45 + MTLTensorDataTypeUInt8 = 49 + MTLTensorDataTypeInt16 = 37 + MTLTensorDataTypeUInt16 = 41 + MTLTensorDataTypeInt32 = 29 + MTLTensorDataTypeUInt32 = 33 +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTLTensorExtents <: NSObject + +@objcproperties MTLTensorExtents begin + @autoproperty rank::UInt64 +end + +@cenum MTLTensorError::UInt64 begin + MTLTensorErrorNone = 0x0000000000000000 + MTLTensorErrorInternalError = 0x0000000000000001 + MTLTensorErrorInvalidDescriptor = 0x0000000000000002 +end + +@cenum MTLTensorUsage::Int64 begin + MTLTensorUsageCompute = 1 + MTLTensorUsageRender = 2 + MTLTensorUsageMachineLearning = 4 +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTLTensorDescriptor <: NSObject + +@objcproperties MTLTensorDescriptor begin + @autoproperty dimensions::id{MTLTensorExtents} setter = setDimensions + @autoproperty strides::id{MTLTensorExtents} setter = setStrides + @autoproperty dataType::MTLTensorDataType setter = setDataType + @autoproperty usage::MTLTensorUsage setter = setUsage + @autoproperty resourceOptions::MTLResourceOptions setter = setResourceOptions + @autoproperty cpuCacheMode::MTLCPUCacheMode setter = setCpuCacheMode + @autoproperty storageMode::MTLStorageMode setter = setStorageMode + @autoproperty hazardTrackingMode::MTLHazardTrackingMode setter = setHazardTrackingMode +end + +@objcwrapper immutable = true MTLBuffer <: MTLResource + +@objcproperties MTLBuffer begin + @autoproperty length::UInt64 + @autoproperty remoteStorageBuffer::id{MTLBuffer} + @autoproperty gpuAddress::UInt64 type = Ptr{Cvoid} + @autoproperty sparseBufferTier::MTLBufferSparseTier availability = macos(v"26.0.0") +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTLTensor <: MTLResource + +@objcproperties MTLTensor begin + @autoproperty gpuResourceID::MTLResourceID + @autoproperty buffer::id{MTLBuffer} + @autoproperty bufferOffset::UInt64 + @autoproperty strides::id{MTLTensorExtents} + @autoproperty dimensions::id{MTLTensorExtents} + @autoproperty dataType::MTLTensorDataType + @autoproperty usage::MTLTensorUsage +end + +const MTLGPUAddress = UInt64 + +@cenum MTLTextureType::UInt64 begin + MTLTextureType1D = 0x0000000000000000 + MTLTextureType1DArray = 0x0000000000000001 + MTLTextureType2D = 0x0000000000000002 + MTLTextureType2DArray = 0x0000000000000003 + MTLTextureType2DMultisample = 0x0000000000000004 + MTLTextureTypeCube = 0x0000000000000005 + MTLTextureTypeCubeArray = 0x0000000000000006 + MTLTextureType3D = 0x0000000000000007 + MTLTextureType2DMultisampleArray = 0x0000000000000008 + MTLTextureTypeTextureBuffer = 0x0000000000000009 +end + +@objcwrapper immutable = true MTLSharedTextureHandle <: NSObject + +@objcproperties MTLSharedTextureHandle begin + @autoproperty device::id{MTLDevice} + @autoproperty label::id{NSString} +end + +@cenum MTLTextureUsage::UInt64 begin + MTLTextureUsageUnknown = 0x0000000000000000 + MTLTextureUsageShaderRead = 0x0000000000000001 + MTLTextureUsageShaderWrite = 0x0000000000000002 + MTLTextureUsageRenderTarget = 0x0000000000000004 + MTLTextureUsagePixelFormatView = 0x0000000000000010 + MTLTextureUsageShaderAtomic = 0x0000000000000020 +end + +@cenum MTLTextureCompressionType::Int64 begin + MTLTextureCompressionTypeLossless = 0 + MTLTextureCompressionTypeLossy = 1 +end + +@objcwrapper immutable = false MTLTextureDescriptor <: NSObject + +@objcproperties MTLTextureDescriptor begin + @autoproperty textureType::MTLTextureType setter = setTextureType + @autoproperty pixelFormat::MTLPixelFormat setter = setPixelFormat + @autoproperty width::UInt64 setter = setWidth + @autoproperty height::UInt64 setter = setHeight + @autoproperty depth::UInt64 setter = setDepth + @autoproperty mipmapLevelCount::UInt64 setter = setMipmapLevelCount + @autoproperty sampleCount::UInt64 setter = setSampleCount + @autoproperty arrayLength::UInt64 setter = setArrayLength + @autoproperty resourceOptions::MTLResourceOptions setter = setResourceOptions + @autoproperty cpuCacheMode::MTLCPUCacheMode setter = setCpuCacheMode + @autoproperty storageMode::MTLStorageMode setter = setStorageMode + @autoproperty hazardTrackingMode::MTLHazardTrackingMode setter = setHazardTrackingMode + @autoproperty usage::MTLTextureUsage setter = setUsage + @autoproperty allowGPUOptimizedContents::Bool setter = setAllowGPUOptimizedContents + @autoproperty compressionType::MTLTextureCompressionType setter = setCompressionType + @autoproperty swizzle::MTLTextureSwizzleChannels setter = setSwizzle + @autoproperty placementSparsePageSize::MTLSparsePageSize setter = setPlacementSparsePageSize availability = macos(v"26.0.0") +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTLTextureViewDescriptor <: NSObject + +@objcproperties MTLTextureViewDescriptor begin + @autoproperty pixelFormat::MTLPixelFormat setter = setPixelFormat + @autoproperty textureType::MTLTextureType setter = setTextureType + @autoproperty levelRange::_NSRange setter = setLevelRange + @autoproperty sliceRange::_NSRange setter = setSliceRange + @autoproperty swizzle::MTLTextureSwizzleChannels setter = setSwizzle +end + +@objcwrapper immutable = false MTLTexture <: MTLResource + +@objcproperties MTLTexture begin + @autoproperty rootResource::id{MTLResource} + @autoproperty parentTexture::id{MTLTexture} + @autoproperty parentRelativeLevel::UInt64 + @autoproperty parentRelativeSlice::UInt64 + @autoproperty buffer::id{MTLBuffer} + @autoproperty bufferOffset::UInt64 + @autoproperty bufferBytesPerRow::UInt64 + @autoproperty iosurface::Ptr{Cvoid} + @autoproperty iosurfacePlane::UInt64 + @autoproperty textureType::MTLTextureType + @autoproperty pixelFormat::MTLPixelFormat + @autoproperty width::UInt64 + @autoproperty height::UInt64 + @autoproperty depth::UInt64 + @autoproperty mipmapLevelCount::UInt64 + @autoproperty sampleCount::UInt64 + @autoproperty arrayLength::UInt64 + @autoproperty usage::MTLTextureUsage + @autoproperty shareable::Bool getter = isShareable + @autoproperty framebufferOnly::Bool getter = isFramebufferOnly + @autoproperty firstMipmapInTail::UInt64 + @autoproperty tailSizeInBytes::UInt64 + @autoproperty isSparse::Bool + @autoproperty allowGPUOptimizedContents::Bool + @autoproperty compressionType::MTLTextureCompressionType + @autoproperty gpuResourceID::MTLResourceID + @autoproperty remoteStorageTexture::id{MTLTexture} + @autoproperty swizzle::MTLTextureSwizzleChannels + @autoproperty sparseTextureTier::MTLTextureSparseTier availability = macos(v"26.0.0") +end + +@cenum MTLIndexType::UInt64 begin + MTLIndexTypeUInt16 = 0x0000000000000000 + MTLIndexTypeUInt32 = 0x0000000000000001 end @cenum MTLBindingType::Int64 begin @@ -610,6 +722,7 @@ end MTLBindingTypeInstanceAccelerationStructure = 26 MTLBindingTypeIntersectionFunctionTable = 27 MTLBindingTypeObjectPayload = 34 + MTLBindingTypeTensor = 37 end @cenum MTLArgumentType::UInt64 begin @@ -685,6 +798,15 @@ end @autoproperty isDepthTexture::Bool end +@objcwrapper immutable = true availability = macos(v"26.0.0") MTLTensorReferenceType <: MTLType + +@objcproperties MTLTensorReferenceType begin + @autoproperty tensorDataType::MTLTensorDataType + @autoproperty indexType::MTLDataType + @autoproperty dimensions::id{MTLTensorExtents} + @autoproperty access::MTLBindingAccess +end + @objcwrapper immutable = true MTLArgument <: NSObject @objcproperties MTLArgument begin @@ -750,6 +872,14 @@ end @autoproperty objectPayloadDataSize::UInt64 end +@objcwrapper immutable = true availability = macos(v"26.0.0") MTLTensorBinding <: MTLBinding + +@objcproperties MTLTensorBinding begin + @autoproperty tensorDataType::MTLTensorDataType + @autoproperty indexType::MTLDataType + @autoproperty dimensions::id{MTLTensorExtents} +end + @objcwrapper immutable = true MTLFunctionConstantValues <: NSObject @cenum MTLFunctionOptions::UInt64 begin @@ -758,6 +888,7 @@ end MTLFunctionOptionStoreFunctionInMetalPipelinesScript = 0x0000000000000002 MTLFunctionOptionStoreFunctionInMetalScript = 0x0000000000000002 MTLFunctionOptionFailOnBinaryArchiveMiss = 0x0000000000000004 + MTLFunctionOptionPipelineIndependent = 0x0000000000000008 end @objcwrapper immutable = false MTLBinaryArchive <: NSObject @@ -906,9 +1037,16 @@ end @autoproperty compileSymbolVisibility::MTLCompileSymbolVisibility setter = setCompileSymbolVisibility availability = macos(v"13.3.0") @autoproperty allowReferencingUndefinedSymbols::Bool setter = setAllowReferencingUndefinedSymbols availability = macos(v"13.3.0") @autoproperty maxTotalThreadsPerThreadgroup::UInt64 setter = setMaxTotalThreadsPerThreadgroup availability = macos(v"13.3.0") + @autoproperty requiredThreadsPerThreadgroup::MTLSize setter = setRequiredThreadsPerThreadgroup availability = macos(v"26.0.0") @autoproperty enableLogging::Bool setter = setEnableLogging availability = macos(v"15.0.0") end +@objcwrapper immutable = true availability = macos(v"26.0.0") MTLFunctionReflection <: NSObject + +@objcproperties MTLFunctionReflection begin + @autoproperty bindings::id{NSArray} type = Vector{MTLBinding} +end + @cenum MTLLibraryError::UInt64 begin MTLLibraryErrorUnsupported = 0x0000000000000001 MTLLibraryErrorInternal = 0x0000000000000002 @@ -975,11 +1113,83 @@ end MTLCounterSampleBufferErrorInternal = 2 end -@cenum MTLIOCompressionMethod::Int64 begin - MTLIOCompressionMethodZlib = 0 - MTLIOCompressionMethodLZFSE = 1 - MTLIOCompressionMethodLZ4 = 2 - MTLIOCompressionMethodLZMA = 3 +@cenum MTL4CompilerTaskStatus::Int64 begin + MTL4CompilerTaskStatusNone = 0 + MTL4CompilerTaskStatusScheduled = 1 + MTL4CompilerTaskStatusCompiling = 2 + MTL4CompilerTaskStatusFinished = 3 +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4PipelineDataSetSerializer <: NSObject + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4Compiler <: NSObject + +@objcproperties MTL4Compiler begin + @autoproperty device::id{MTLDevice} + @autoproperty label::id{NSString} + @autoproperty pipelineDataSetSerializer::id{MTL4PipelineDataSetSerializer} +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4CompilerTask <: NSObject + +@objcproperties MTL4CompilerTask begin + @autoproperty compiler::id{MTL4Compiler} + @autoproperty status::MTL4CompilerTaskStatus +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4CompilerDescriptor <: NSObject + +@objcproperties MTL4CompilerDescriptor begin + @autoproperty label::id{NSString} setter = setLabel + @autoproperty pipelineDataSetSerializer::id{MTL4PipelineDataSetSerializer} setter = setPipelineDataSetSerializer +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4Archive <: NSObject + +@objcproperties MTL4Archive begin + @autoproperty label::id{NSString} setter = setLabel +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4CompilerTaskOptions <: NSObject + +@objcproperties MTL4CompilerTaskOptions begin + @autoproperty lookupArchives::id{NSArray} type = Vector{MTL4Archive} setter = setLookupArchives +end + +struct MTL4TimestampHeapEntry + timestamp::UInt64 +end + +@cenum MTL4CounterHeapType::UInt64 begin + MTL4CounterHeapTypeInvalid = 0x0000000000000000 + MTL4CounterHeapTypeTimestamp = 0x0000000000000001 +end + +@cenum MTL4TimestampGranularity::Int64 begin + MTL4TimestampGranularityRelaxed = 0 + MTL4TimestampGranularityPrecise = 1 +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4CounterHeapDescriptor <: NSObject + +@objcproperties MTL4CounterHeapDescriptor begin + @autoproperty type::MTL4CounterHeapType setter = setType + @autoproperty count::UInt64 setter = setCount +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4CounterHeap <: NSObject + +@objcproperties MTL4CounterHeap begin + @autoproperty label::id{NSString} setter = setLabel + @autoproperty count::UInt64 + @autoproperty type::MTL4CounterHeapType +end + +@cenum MTLIOCompressionMethod::Int64 begin + MTLIOCompressionMethodZlib = 0 + MTLIOCompressionMethodLZFSE = 1 + MTLIOCompressionMethodLZ4 = 2 + MTLIOCompressionMethodLZMA = 3 MTLIOCompressionMethodLZBitmap = 4 end @@ -1037,6 +1247,7 @@ end MTLGPUFamilyMacCatalyst1 = 4001 MTLGPUFamilyMacCatalyst2 = 4002 MTLGPUFamilyMetal3 = 5001 + MTLGPUFamilyMetal4 = 5002 end @cenum MTLPipelineOption::UInt64 begin @@ -1052,12 +1263,6 @@ end MTLSparseTextureRegionAlignmentModeInward = 0x0000000000000001 end -@cenum MTLSparsePageSize::Int64 begin - MTLSparsePageSize16 = 101 - MTLSparsePageSize64 = 102 - MTLSparsePageSize256 = 103 -end - struct MTLAccelerationStructureSizes accelerationStructureSize::NSUInteger buildScratchBufferSize::NSUInteger @@ -1162,6 +1367,11 @@ end MTLStoreActionOptionCustomSamplePositions = 0x0000000000000001 end +@cenum MTLVisibilityResultType::Int64 begin + MTLVisibilityResultTypeReset = 0 + MTLVisibilityResultTypeAccumulate = 1 +end + @objcwrapper immutable = true MTLRenderPassAttachmentDescriptor <: NSObject @objcproperties MTLRenderPassAttachmentDescriptor begin @@ -1251,6 +1461,8 @@ end @autoproperty renderTargetHeight::UInt64 setter = setRenderTargetHeight @autoproperty rasterizationRateMap::id{MTLRasterizationRateMap} setter = setRasterizationRateMap @autoproperty sampleBufferAttachments::id{MTLRenderPassSampleBufferAttachmentDescriptorArray} + @autoproperty visibilityResultType::MTLVisibilityResultType setter = setVisibilityResultType availability = macos(v"26.0.0") + @autoproperty supportColorAttachmentMapping::Bool setter = setSupportColorAttachmentMapping availability = macos(v"26.0.0") end @objcwrapper immutable = true MTLBlitPassSampleBufferAttachmentDescriptor <: NSObject @@ -1385,6 +1597,11 @@ struct MTLDispatchThreadgroupsIndirectArguments threadgroupsPerGrid::NTuple{3, UInt32} end +struct MTLDispatchThreadsIndirectArguments + threadsPerGrid::NTuple{3, UInt32} + threadsPerThreadgroup::NTuple{3, UInt32} +end + struct MTLStageInRegionIndirectArguments stageInOrigin::NTuple{3, UInt32} stageInSize::NTuple{3, UInt32} @@ -1613,11 +1830,6 @@ end MTLAttributeFormatFloatRGB9E5 = 0x0000000000000037 end -@cenum MTLIndexType::UInt64 begin - MTLIndexTypeUInt16 = 0x0000000000000000 - MTLIndexTypeUInt32 = 0x0000000000000001 -end - @cenum MTLStepFunction::UInt64 begin MTLStepFunctionConstant = 0x0000000000000000 MTLStepFunctionPerVertex = 0x0000000000000001 @@ -1712,12 +1924,14 @@ end @autoproperty supportAddingBinaryFunctions::Bool setter = setSupportAddingBinaryFunctions @autoproperty maxCallStackDepth::UInt64 setter = setMaxCallStackDepth @autoproperty shaderValidation::MTLShaderValidation setter = setShaderValidation availability = macos(v"15.0.0") + @autoproperty requiredThreadsPerThreadgroup::MTLSize setter = setRequiredThreadsPerThreadgroup availability = macos(v"26.0.0") end -@objcwrapper immutable = false MTLComputePipelineState <: NSObject +@objcwrapper immutable = false MTLComputePipelineState <: MTLAllocation @objcproperties MTLComputePipelineState begin @autoproperty label::id{NSString} + @autoproperty reflection::id{MTLComputePipelineReflection} availability = macos(v"26.0.0") @autoproperty device::id{MTLDevice} @autoproperty maxTotalThreadsPerThreadgroup::UInt64 @autoproperty threadExecutionWidth::UInt64 @@ -1725,6 +1939,7 @@ end @autoproperty supportIndirectCommandBuffers::Bool @autoproperty gpuResourceID::MTLResourceID @autoproperty shaderValidation::MTLShaderValidation availability = macos(v"15.0.0") + @autoproperty requiredThreadsPerThreadgroup::MTLSize availability = macos(v"26.0.0") end @cenum MTLPrimitiveType::UInt64 begin @@ -1836,6 +2051,7 @@ end @autoproperty functionType::MTLFunctionType @autoproperty name::id{NSString} @autoproperty device::id{MTLDevice} + @autoproperty gpuResourceID::MTLResourceID availability = macos(v"26.0.0") end @objcwrapper immutable = true MTLVisibleFunctionTableDescriptor <: NSObject @@ -1870,6 +2086,7 @@ end MTLBlendFactorOneMinusSource1Color = 0x0000000000000010 MTLBlendFactorSource1Alpha = 0x0000000000000011 MTLBlendFactorOneMinusSource1Alpha = 0x0000000000000012 + MTLBlendFactorUnspecialized = 0x0000000000000013 end @cenum MTLBlendOperation::UInt64 begin @@ -1878,6 +2095,7 @@ end MTLBlendOperationReverseSubtract = 0x0000000000000002 MTLBlendOperationMin = 0x0000000000000003 MTLBlendOperationMax = 0x0000000000000004 + MTLBlendOperationUnspecialized = 0x0000000000000005 end @cenum MTLColorWriteMask::UInt64 begin @@ -1887,6 +2105,7 @@ end MTLColorWriteMaskBlue = 0x0000000000000002 MTLColorWriteMaskAlpha = 0x0000000000000001 MTLColorWriteMaskAll = 0x000000000000000f + MTLColorWriteMaskUnspecialized = 0x0000000000000010 end @cenum MTLPrimitiveTopologyClass::UInt64 begin @@ -1934,6 +2153,8 @@ end @autoproperty writeMask::MTLColorWriteMask setter = setWriteMask end +@objcwrapper immutable = true availability = macos(v"26.0.0") MTLLogicalToPhysicalColorAttachmentMap <: NSObject + @objcwrapper immutable = true MTLRenderPipelineReflection <: NSObject @objcproperties MTLRenderPipelineReflection begin @@ -1996,11 +2217,12 @@ end @autoproperty tileAdditionalBinaryFunctions::id{NSArray} type = Vector{MTLFunction} setter = setTileAdditionalBinaryFunctions end -@objcwrapper immutable = true MTLRenderPipelineState <: NSObject +@objcwrapper immutable = true MTLRenderPipelineState <: MTLAllocation @objcproperties MTLRenderPipelineState begin @autoproperty label::id{NSString} @autoproperty device::id{MTLDevice} + @autoproperty reflection::id{MTLRenderPipelineReflection} availability = macos(v"26.0.0") @autoproperty maxTotalThreadsPerThreadgroup::UInt64 @autoproperty threadgroupSizeMatchesTileSize::Bool @autoproperty imageblockSampleLength::UInt64 @@ -2012,6 +2234,9 @@ end @autoproperty maxTotalThreadgroupsPerMeshGrid::UInt64 @autoproperty gpuResourceID::MTLResourceID @autoproperty shaderValidation::MTLShaderValidation availability = macos(v"15.0.0") + @autoproperty requiredThreadsPerTileThreadgroup::MTLSize availability = macos(v"26.0.0") + @autoproperty requiredThreadsPerObjectThreadgroup::MTLSize availability = macos(v"26.0.0") + @autoproperty requiredThreadsPerMeshThreadgroup::MTLSize availability = macos(v"26.0.0") end @objcwrapper immutable = true MTLTileRenderPipelineColorAttachmentDescriptor <: NSObject @@ -2038,6 +2263,7 @@ end @autoproperty supportAddingBinaryFunctions::Bool setter = setSupportAddingBinaryFunctions @autoproperty maxCallStackDepth::UInt64 setter = setMaxCallStackDepth @autoproperty shaderValidation::MTLShaderValidation setter = setShaderValidation availability = macos(v"15.0.0") + @autoproperty requiredThreadsPerThreadgroup::MTLSize setter = setRequiredThreadsPerThreadgroup availability = macos(v"26.0.0") end @objcwrapper immutable = true MTLMeshRenderPipelineDescriptor <: NSObject @@ -2070,6 +2296,8 @@ end @autoproperty meshLinkedFunctions::id{MTLLinkedFunctions} setter = setMeshLinkedFunctions availability = macos(v"14.0.0") @autoproperty fragmentLinkedFunctions::id{MTLLinkedFunctions} setter = setFragmentLinkedFunctions availability = macos(v"14.0.0") @autoproperty shaderValidation::MTLShaderValidation setter = setShaderValidation availability = macos(v"15.0.0") + @autoproperty requiredThreadsPerObjectThreadgroup::MTLSize setter = setRequiredThreadsPerObjectThreadgroup availability = macos(v"26.0.0") + @autoproperty requiredThreadsPerMeshThreadgroup::MTLSize setter = setRequiredThreadsPerMeshThreadgroup availability = macos(v"26.0.0") end @objcwrapper immutable = true MTLParallelRenderCommandEncoder <: MTLCommandEncoder @@ -2128,6 +2356,15 @@ end @autoproperty gpuResourceID::MTLResourceID end +struct MTL4BufferRange + bufferAddress::MTLGPUAddress + length::UInt64 +end + +function MTL4BufferRangeMake(bufferAddress, length) + return @ccall (Symbol("/System/Library/Frameworks/Metal.framework/Resources/BridgeSupport/Metal.dylib")).MTL4BufferRangeMake(bufferAddress::MTLGPUAddress, length::UInt64)::MTL4BufferRange +end + struct _MTLPackedFloat3 data::NTuple{12, UInt8} end @@ -2251,11 +2488,18 @@ function Base.propertynames(x::MTLComponentTransform, private::Bool = false) ) end +@cenum MTLAccelerationStructureRefitOptions::UInt64 begin + MTLAccelerationStructureRefitOptionVertexData = 0x0000000000000001 + MTLAccelerationStructureRefitOptionPerPrimitiveData = 0x0000000000000002 +end + @cenum MTLAccelerationStructureUsage::UInt64 begin MTLAccelerationStructureUsageNone = 0x0000000000000000 MTLAccelerationStructureUsageRefit = 0x0000000000000001 MTLAccelerationStructureUsagePreferFastBuild = 0x0000000000000002 MTLAccelerationStructureUsageExtendedLimits = 0x0000000000000004 + MTLAccelerationStructureUsagePreferFastIntersection = 0x0000000000000010 + MTLAccelerationStructureUsageMinimizeMemory = 0x0000000000000020 end @cenum MTLAccelerationStructureInstanceOptions::UInt32 begin @@ -2542,6 +2786,7 @@ end @autoproperty hazardTrackingMode::MTLHazardTrackingMode setter = setHazardTrackingMode @autoproperty resourceOptions::MTLResourceOptions setter = setResourceOptions @autoproperty type::MTLHeapType setter = setType + @autoproperty maxCompatiblePlacementSparsePageSize::MTLSparsePageSize setter = setMaxCompatiblePlacementSparsePageSize availability = macos(v"26.0.0") end @objcwrapper immutable = true MTLArgumentEncoder <: NSObject @@ -2553,6 +2798,132 @@ end @autoproperty alignment::UInt64 end +@cenum MTL4VisibilityOptions::Int64 begin + MTL4VisibilityOptionNone = 0 + MTL4VisibilityOptionDevice = 1 + MTL4VisibilityOptionResourceAlias = 2 +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4CommandBuffer <: NSObject + +@objcproperties MTL4CommandBuffer begin + @autoproperty device::id{MTLDevice} + @autoproperty label::id{NSString} setter = setLabel +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4CommandEncoder <: NSObject + +@objcproperties MTL4CommandEncoder begin + @autoproperty label::id{NSString} setter = setLabel + @autoproperty commandBuffer::id{MTL4CommandBuffer} +end + +@cenum MTL4RenderEncoderOptions::UInt64 begin + MTL4RenderEncoderOptionNone = 0x0000000000000000 + MTL4RenderEncoderOptionSuspending = 0x0000000000000001 + MTL4RenderEncoderOptionResuming = 0x0000000000000002 +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4RenderCommandEncoder <: MTL4CommandEncoder + +@objcproperties MTL4RenderCommandEncoder begin + @autoproperty tileWidth::UInt64 + @autoproperty tileHeight::UInt64 +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4CommandBufferOptions <: NSObject + +@objcproperties MTL4CommandBufferOptions begin + @autoproperty logState::id{MTLLogState} setter = setLogState +end + +@objcwrapper immutable = false MTLEvent <: NSObject + +@objcproperties MTLEvent begin + @autoproperty device::id{MTLDevice} + @autoproperty label::id{NSString} setter = setLabel +end + +@objcwrapper immutable = true MTLSharedEventListener <: NSObject + +@objcproperties MTLSharedEventListener begin + @autoproperty dispatchQueue::id{dispatch_queue_t} +end + +@objcwrapper immutable = false MTLSharedEvent <: MTLEvent + +@objcproperties MTLSharedEvent begin + @autoproperty signaledValue::UInt64 setter = setSignaledValue +end + +@objcwrapper immutable = true MTLSharedEventHandle <: NSObject + +@objcproperties MTLSharedEventHandle begin + @autoproperty label::id{NSString} +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4CommitFeedback <: NSObject + +@objcproperties MTL4CommitFeedback begin + @autoproperty error::id{NSError} + @autoproperty GPUStartTime::Cdouble + @autoproperty GPUEndTime::Cdouble +end + +@cenum MTL4CommandQueueError::Int64 begin + MTL4CommandQueueErrorNone = 0 + MTL4CommandQueueErrorTimeout = 1 + MTL4CommandQueueErrorNotPermitted = 2 + MTL4CommandQueueErrorOutOfMemory = 3 + MTL4CommandQueueErrorDeviceRemoved = 4 + MTL4CommandQueueErrorAccessRevoked = 5 + MTL4CommandQueueErrorInternal = 6 +end + +@objcwrapper immutable = false availability = macos(v"26.0.0") MTL4CommitOptions <: NSObject + +@objcwrapper immutable = false availability = macos(v"26.0.0") MTL4CommandQueueDescriptor <: NSObject + +@objcproperties MTL4CommandQueueDescriptor begin + @autoproperty label::id{NSString} setter = setLabel + @autoproperty feedbackQueue::id{dispatch_queue_t} setter = setFeedbackQueue +end + +struct MTL4UpdateSparseTextureMappingOperation + mode::MTLSparseTextureMappingMode + textureRegion::MTLRegion + textureLevel::NSUInteger + textureSlice::NSUInteger + heapOffset::NSUInteger +end + +struct MTL4CopySparseTextureMappingOperation + sourceRegion::MTLRegion + sourceLevel::NSUInteger + sourceSlice::NSUInteger + destinationOrigin::MTLOrigin + destinationLevel::NSUInteger + destinationSlice::NSUInteger +end + +struct MTL4UpdateSparseBufferMappingOperation + mode::MTLSparseTextureMappingMode + bufferRange::NSRange + heapOffset::NSUInteger +end + +struct MTL4CopySparseBufferMappingOperation + sourceRange::NSRange + destinationOffset::NSUInteger +end + +@objcwrapper immutable = false availability = macos(v"26.0.0") MTL4CommandQueue <: NSObject + +@objcproperties MTL4CommandQueue begin + @autoproperty device::id{MTLDevice} + @autoproperty label::id{NSString} +end + @cenum MTLCaptureError::Int64 begin MTLCaptureErrorNotSupported = 1 MTLCaptureErrorAlreadyCapturing = 2 @@ -2578,6 +2949,7 @@ end @autoproperty label::id{NSString} setter = setLabel @autoproperty device::id{MTLDevice} @autoproperty commandQueue::id{MTLCommandQueue} + @autoproperty mtl4CommandQueue::id{MTL4CommandQueue} end @objcwrapper immutable = true MTLCaptureManager <: NSObject @@ -2626,6 +2998,7 @@ end @autoproperty maxObjectThreadgroupMemoryBindCount::UInt64 setter = setMaxObjectThreadgroupMemoryBindCount availability = macos(v"14.0.0") @autoproperty supportRayTracing::Bool setter = setSupportRayTracing @autoproperty supportDynamicAttributeStride::Bool setter = setSupportDynamicAttributeStride availability = macos(v"14.0.0") + @autoproperty supportColorAttachmentMapping::Bool setter = setSupportColorAttachmentMapping availability = macos(v"26.0.0") end @objcwrapper immutable = true MTLIndirectCommandBuffer <: MTLResource @@ -2635,31 +3008,6 @@ end @autoproperty gpuResourceID::MTLResourceID end -@objcwrapper immutable = false MTLEvent <: NSObject - -@objcproperties MTLEvent begin - @autoproperty device::id{MTLDevice} - @autoproperty label::id{NSString} setter = setLabel -end - -@objcwrapper immutable = true MTLSharedEventListener <: NSObject - -@objcproperties MTLSharedEventListener begin - @autoproperty dispatchQueue::id{dispatch_queue_t} -end - -@objcwrapper immutable = false MTLSharedEvent <: MTLEvent - -@objcproperties MTLSharedEvent begin - @autoproperty signaledValue::UInt64 setter = setSignaledValue -end - -@objcwrapper immutable = true MTLSharedEventHandle <: NSObject - -@objcproperties MTLSharedEventHandle begin - @autoproperty label::id{NSString} -end - @cenum MTLFunctionLogType::UInt64 begin MTLFunctionLogTypeValidation = 0x0000000000000000 end @@ -2673,11 +3021,6 @@ end @autoproperty column::UInt64 end -@cenum MTLAccelerationStructureRefitOptions::UInt64 begin - MTLAccelerationStructureRefitOptionVertexData = 0x0000000000000001 - MTLAccelerationStructureRefitOptionPerPrimitiveData = 0x0000000000000002 -end - @objcwrapper immutable = true MTLAccelerationStructureCommandEncoder <: MTLCommandEncoder @objcwrapper immutable = true MTLAccelerationStructurePassSampleBufferAttachmentDescriptor <: NSObject @@ -2764,6 +3107,12 @@ end @autoproperty url::id{NSURL} setter = setUrl end +struct MTLIntersectionFunctionBufferArguments + intersectionFunctionBuffer::UInt64 + intersectionFunctionBufferSize::UInt64 + intersectionFunctionStride::UInt64 +end + @cenum MTLIntersectionFunctionSignature::UInt64 begin MTLIntersectionFunctionSignatureNone = 0x0000000000000000 MTLIntersectionFunctionSignatureInstancing = 0x0000000000000001 @@ -2774,6 +3123,8 @@ end MTLIntersectionFunctionSignatureExtendedLimits = 0x0000000000000020 MTLIntersectionFunctionSignatureMaxLevels = 0x0000000000000040 MTLIntersectionFunctionSignatureCurveData = 0x0000000000000080 + MTLIntersectionFunctionSignatureIntersectionFunctionBuffer = 0x0000000000000100 + MTLIntersectionFunctionSignatureUserData = 0x0000000000000200 end @objcwrapper immutable = true MTLIntersectionFunctionTableDescriptor <: NSObject @@ -2933,6 +3284,473 @@ end @autoproperty allocationCount::UInt64 end +@objcwrapper immutable = true availability = macos(v"26.0.0") MTLResourceViewPoolDescriptor <: NSObject + +@objcproperties MTLResourceViewPoolDescriptor begin + @autoproperty resourceViewCount::UInt64 setter = setResourceViewCount + @autoproperty label::id{NSString} setter = setLabel +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTLResourceViewPool <: NSObject + +@objcproperties MTLResourceViewPool begin + @autoproperty baseResourceID::MTLResourceID + @autoproperty resourceViewCount::UInt64 + @autoproperty device::id{MTLDevice} + @autoproperty label::id{NSString} +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTLTextureViewPool <: MTLResourceViewPool + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4ArgumentTableDescriptor <: NSObject + +@objcproperties MTL4ArgumentTableDescriptor begin + @autoproperty maxBufferBindCount::UInt64 setter = setMaxBufferBindCount + @autoproperty maxTextureBindCount::UInt64 setter = setMaxTextureBindCount + @autoproperty maxSamplerStateBindCount::UInt64 setter = setMaxSamplerStateBindCount + @autoproperty initializeBindings::Bool setter = setInitializeBindings + @autoproperty supportAttributeStrides::Bool setter = setSupportAttributeStrides + @autoproperty label::id{NSString} setter = setLabel +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4ArgumentTable <: NSObject + +@objcproperties MTL4ArgumentTable begin + @autoproperty device::id{MTLDevice} + @autoproperty label::id{NSString} +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4BinaryFunction <: NSObject + +@objcproperties MTL4BinaryFunction begin + @autoproperty name::id{NSString} + @autoproperty functionType::MTLFunctionType +end + +@objcwrapper immutable = false availability = macos(v"26.0.0") MTL4CommandAllocatorDescriptor <: NSObject + +@objcproperties MTL4CommandAllocatorDescriptor begin + @autoproperty label::id{NSString} setter = setLabel +end + +@objcwrapper immutable = false availability = macos(v"26.0.0") MTL4CommandAllocator <: NSObject + +@objcproperties MTL4CommandAllocator begin + @autoproperty device::id{MTLDevice} + @autoproperty label::id{NSString} +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4RenderPassDescriptor <: NSObject + +@objcproperties MTL4RenderPassDescriptor begin + @autoproperty colorAttachments::id{MTLRenderPassColorAttachmentDescriptorArray} + @autoproperty depthAttachment::id{MTLRenderPassDepthAttachmentDescriptor} setter = setDepthAttachment + @autoproperty stencilAttachment::id{MTLRenderPassStencilAttachmentDescriptor} setter = setStencilAttachment + @autoproperty renderTargetArrayLength::UInt64 setter = setRenderTargetArrayLength + @autoproperty imageblockSampleLength::UInt64 setter = setImageblockSampleLength + @autoproperty threadgroupMemoryLength::UInt64 setter = setThreadgroupMemoryLength + @autoproperty tileWidth::UInt64 setter = setTileWidth + @autoproperty tileHeight::UInt64 setter = setTileHeight + @autoproperty defaultRasterSampleCount::UInt64 setter = setDefaultRasterSampleCount + @autoproperty renderTargetWidth::UInt64 setter = setRenderTargetWidth + @autoproperty renderTargetHeight::UInt64 setter = setRenderTargetHeight + @autoproperty rasterizationRateMap::id{MTLRasterizationRateMap} setter = setRasterizationRateMap + @autoproperty visibilityResultBuffer::id{MTLBuffer} setter = setVisibilityResultBuffer + @autoproperty visibilityResultType::MTLVisibilityResultType setter = setVisibilityResultType + @autoproperty supportColorAttachmentMapping::Bool setter = setSupportColorAttachmentMapping +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4AccelerationStructureDescriptor <: MTLAccelerationStructureDescriptor + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4AccelerationStructureGeometryDescriptor <: NSObject + +@objcproperties MTL4AccelerationStructureGeometryDescriptor begin + @autoproperty intersectionFunctionTableOffset::UInt64 setter = setIntersectionFunctionTableOffset + @autoproperty opaque::Bool setter = setOpaque + @autoproperty allowDuplicateIntersectionFunctionInvocation::Bool setter = setAllowDuplicateIntersectionFunctionInvocation + @autoproperty label::id{NSString} setter = setLabel + @autoproperty primitiveDataBuffer::MTL4BufferRange setter = setPrimitiveDataBuffer + @autoproperty primitiveDataStride::UInt64 setter = setPrimitiveDataStride + @autoproperty primitiveDataElementSize::UInt64 setter = setPrimitiveDataElementSize +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4PrimitiveAccelerationStructureDescriptor <: MTL4AccelerationStructureDescriptor + +@objcproperties MTL4PrimitiveAccelerationStructureDescriptor begin + @autoproperty geometryDescriptors::id{NSArray} type = Vector{MTL4AccelerationStructureGeometryDescriptor} setter = setGeometryDescriptors + @autoproperty motionStartBorderMode::MTLMotionBorderMode setter = setMotionStartBorderMode + @autoproperty motionEndBorderMode::MTLMotionBorderMode setter = setMotionEndBorderMode + @autoproperty motionStartTime::Cfloat setter = setMotionStartTime + @autoproperty motionEndTime::Cfloat setter = setMotionEndTime + @autoproperty motionKeyframeCount::UInt64 setter = setMotionKeyframeCount +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4AccelerationStructureTriangleGeometryDescriptor <: MTL4AccelerationStructureGeometryDescriptor + +@objcproperties MTL4AccelerationStructureTriangleGeometryDescriptor begin + @autoproperty vertexBuffer::MTL4BufferRange setter = setVertexBuffer + @autoproperty vertexFormat::MTLAttributeFormat setter = setVertexFormat + @autoproperty vertexStride::UInt64 setter = setVertexStride + @autoproperty indexBuffer::MTL4BufferRange setter = setIndexBuffer + @autoproperty indexType::MTLIndexType setter = setIndexType + @autoproperty triangleCount::UInt64 setter = setTriangleCount + @autoproperty transformationMatrixBuffer::MTL4BufferRange setter = setTransformationMatrixBuffer + @autoproperty transformationMatrixLayout::MTLMatrixLayout setter = setTransformationMatrixLayout +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4AccelerationStructureBoundingBoxGeometryDescriptor <: MTL4AccelerationStructureGeometryDescriptor + +@objcproperties MTL4AccelerationStructureBoundingBoxGeometryDescriptor begin + @autoproperty boundingBoxBuffer::MTL4BufferRange setter = setBoundingBoxBuffer + @autoproperty boundingBoxStride::UInt64 setter = setBoundingBoxStride + @autoproperty boundingBoxCount::UInt64 setter = setBoundingBoxCount +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4AccelerationStructureMotionTriangleGeometryDescriptor <: MTL4AccelerationStructureGeometryDescriptor + +@objcproperties MTL4AccelerationStructureMotionTriangleGeometryDescriptor begin + @autoproperty vertexBuffers::MTL4BufferRange setter = setVertexBuffers + @autoproperty vertexFormat::MTLAttributeFormat setter = setVertexFormat + @autoproperty vertexStride::UInt64 setter = setVertexStride + @autoproperty indexBuffer::MTL4BufferRange setter = setIndexBuffer + @autoproperty indexType::MTLIndexType setter = setIndexType + @autoproperty triangleCount::UInt64 setter = setTriangleCount + @autoproperty transformationMatrixBuffer::MTL4BufferRange setter = setTransformationMatrixBuffer + @autoproperty transformationMatrixLayout::MTLMatrixLayout setter = setTransformationMatrixLayout +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4AccelerationStructureMotionBoundingBoxGeometryDescriptor <: MTL4AccelerationStructureGeometryDescriptor + +@objcproperties MTL4AccelerationStructureMotionBoundingBoxGeometryDescriptor begin + @autoproperty boundingBoxBuffers::MTL4BufferRange setter = setBoundingBoxBuffers + @autoproperty boundingBoxStride::UInt64 setter = setBoundingBoxStride + @autoproperty boundingBoxCount::UInt64 setter = setBoundingBoxCount +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4AccelerationStructureCurveGeometryDescriptor <: MTL4AccelerationStructureGeometryDescriptor + +@objcproperties MTL4AccelerationStructureCurveGeometryDescriptor begin + @autoproperty controlPointBuffer::MTL4BufferRange setter = setControlPointBuffer + @autoproperty controlPointCount::UInt64 setter = setControlPointCount + @autoproperty controlPointStride::UInt64 setter = setControlPointStride + @autoproperty controlPointFormat::MTLAttributeFormat setter = setControlPointFormat + @autoproperty radiusBuffer::MTL4BufferRange setter = setRadiusBuffer + @autoproperty radiusFormat::MTLAttributeFormat setter = setRadiusFormat + @autoproperty radiusStride::UInt64 setter = setRadiusStride + @autoproperty indexBuffer::MTL4BufferRange setter = setIndexBuffer + @autoproperty indexType::MTLIndexType setter = setIndexType + @autoproperty segmentCount::UInt64 setter = setSegmentCount + @autoproperty segmentControlPointCount::UInt64 setter = setSegmentControlPointCount + @autoproperty curveType::MTLCurveType setter = setCurveType + @autoproperty curveBasis::MTLCurveBasis setter = setCurveBasis + @autoproperty curveEndCaps::MTLCurveEndCaps setter = setCurveEndCaps +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4AccelerationStructureMotionCurveGeometryDescriptor <: MTL4AccelerationStructureGeometryDescriptor + +@objcproperties MTL4AccelerationStructureMotionCurveGeometryDescriptor begin + @autoproperty controlPointBuffers::MTL4BufferRange setter = setControlPointBuffers + @autoproperty controlPointCount::UInt64 setter = setControlPointCount + @autoproperty controlPointStride::UInt64 setter = setControlPointStride + @autoproperty controlPointFormat::MTLAttributeFormat setter = setControlPointFormat + @autoproperty radiusBuffers::MTL4BufferRange setter = setRadiusBuffers + @autoproperty radiusFormat::MTLAttributeFormat setter = setRadiusFormat + @autoproperty radiusStride::UInt64 setter = setRadiusStride + @autoproperty indexBuffer::MTL4BufferRange setter = setIndexBuffer + @autoproperty indexType::MTLIndexType setter = setIndexType + @autoproperty segmentCount::UInt64 setter = setSegmentCount + @autoproperty segmentControlPointCount::UInt64 setter = setSegmentControlPointCount + @autoproperty curveType::MTLCurveType setter = setCurveType + @autoproperty curveBasis::MTLCurveBasis setter = setCurveBasis + @autoproperty curveEndCaps::MTLCurveEndCaps setter = setCurveEndCaps +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4InstanceAccelerationStructureDescriptor <: MTL4AccelerationStructureDescriptor + +@objcproperties MTL4InstanceAccelerationStructureDescriptor begin + @autoproperty instanceDescriptorBuffer::MTL4BufferRange setter = setInstanceDescriptorBuffer + @autoproperty instanceDescriptorStride::UInt64 setter = setInstanceDescriptorStride + @autoproperty instanceCount::UInt64 setter = setInstanceCount + @autoproperty instanceDescriptorType::MTLAccelerationStructureInstanceDescriptorType setter = setInstanceDescriptorType + @autoproperty motionTransformBuffer::MTL4BufferRange setter = setMotionTransformBuffer + @autoproperty motionTransformCount::UInt64 setter = setMotionTransformCount + @autoproperty instanceTransformationMatrixLayout::MTLMatrixLayout setter = setInstanceTransformationMatrixLayout + @autoproperty motionTransformType::MTLTransformType setter = setMotionTransformType + @autoproperty motionTransformStride::UInt64 setter = setMotionTransformStride +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4IndirectInstanceAccelerationStructureDescriptor <: MTL4AccelerationStructureDescriptor + +@objcproperties MTL4IndirectInstanceAccelerationStructureDescriptor begin + @autoproperty instanceDescriptorBuffer::MTL4BufferRange setter = setInstanceDescriptorBuffer + @autoproperty instanceDescriptorStride::UInt64 setter = setInstanceDescriptorStride + @autoproperty maxInstanceCount::UInt64 setter = setMaxInstanceCount + @autoproperty instanceCountBuffer::MTL4BufferRange setter = setInstanceCountBuffer + @autoproperty instanceDescriptorType::MTLAccelerationStructureInstanceDescriptorType setter = setInstanceDescriptorType + @autoproperty motionTransformBuffer::MTL4BufferRange setter = setMotionTransformBuffer + @autoproperty maxMotionTransformCount::UInt64 setter = setMaxMotionTransformCount + @autoproperty motionTransformCountBuffer::MTL4BufferRange setter = setMotionTransformCountBuffer + @autoproperty instanceTransformationMatrixLayout::MTLMatrixLayout setter = setInstanceTransformationMatrixLayout + @autoproperty motionTransformType::MTLTransformType setter = setMotionTransformType + @autoproperty motionTransformStride::UInt64 setter = setMotionTransformStride +end + +@objcwrapper immutable = false availability = macos(v"26.0.0") MTL4ComputeCommandEncoder <: MTL4CommandEncoder + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4MachineLearningCommandEncoder <: MTL4CommandEncoder + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4LibraryDescriptor <: NSObject + +@objcproperties MTL4LibraryDescriptor begin + @autoproperty source::id{NSString} setter = setSource + @autoproperty options::id{MTLCompileOptions} setter = setOptions + @autoproperty name::id{NSString} setter = setName +end + +@objcwrapper immutable = false availability = macos(v"26.0.0") MTL4FunctionDescriptor <: NSObject + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4LibraryFunctionDescriptor <: MTL4FunctionDescriptor + +@objcproperties MTL4LibraryFunctionDescriptor begin + @autoproperty name::id{NSString} setter = setName + @autoproperty library::id{MTLLibrary} setter = setLibrary +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4SpecializedFunctionDescriptor <: MTL4FunctionDescriptor + +@objcproperties MTL4SpecializedFunctionDescriptor begin + @autoproperty functionDescriptor::id{MTL4FunctionDescriptor} setter = setFunctionDescriptor + @autoproperty specializedName::id{NSString} setter = setSpecializedName + @autoproperty constantValues::id{MTLFunctionConstantValues} setter = setConstantValues +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4StitchedFunctionDescriptor <: MTL4FunctionDescriptor + +@objcproperties MTL4StitchedFunctionDescriptor begin + @autoproperty functionGraph::id{MTLFunctionStitchingGraph} setter = setFunctionGraph + @autoproperty functionDescriptors::id{NSArray} type = Vector{MTL4FunctionDescriptor} setter = setFunctionDescriptors +end + +@cenum MTL4ShaderReflection::UInt64 begin + MTL4ShaderReflectionNone = 0x0000000000000000 + MTL4ShaderReflectionBindingInfo = 0x0000000000000001 + MTL4ShaderReflectionBufferTypeInfo = 0x0000000000000002 +end + +@cenum MTL4AlphaToOneState::Int64 begin + MTL4AlphaToOneStateDisabled = 0 + MTL4AlphaToOneStateEnabled = 1 +end + +@cenum MTL4AlphaToCoverageState::Int64 begin + MTL4AlphaToCoverageStateDisabled = 0 + MTL4AlphaToCoverageStateEnabled = 1 +end + +@cenum MTL4BlendState::Int64 begin + MTL4BlendStateDisabled = 0 + MTL4BlendStateEnabled = 1 + MTL4BlendStateUnspecialized = 2 +end + +@cenum MTL4IndirectCommandBufferSupportState::Int64 begin + MTL4IndirectCommandBufferSupportStateDisabled = 0 + MTL4IndirectCommandBufferSupportStateEnabled = 1 +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4PipelineOptions <: NSObject + +@objcproperties MTL4PipelineOptions begin + @autoproperty shaderValidation::MTLShaderValidation setter = setShaderValidation + @autoproperty shaderReflection::MTL4ShaderReflection setter = setShaderReflection +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4PipelineDescriptor <: NSObject + +@objcproperties MTL4PipelineDescriptor begin + @autoproperty label::id{NSString} setter = setLabel + @autoproperty options::id{MTL4PipelineOptions} setter = setOptions +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4StaticLinkingDescriptor <: NSObject + +@objcproperties MTL4StaticLinkingDescriptor begin + @autoproperty functionDescriptors::id{NSArray} type = Vector{MTL4FunctionDescriptor} setter = setFunctionDescriptors + @autoproperty privateFunctionDescriptors::id{NSArray} type = Vector{MTL4FunctionDescriptor} setter = setPrivateFunctionDescriptors + @autoproperty groups::id{NSDictionary} setter = setGroups +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4ComputePipelineDescriptor <: MTL4PipelineDescriptor + +@objcproperties MTL4ComputePipelineDescriptor begin + @autoproperty computeFunctionDescriptor::id{MTL4FunctionDescriptor} setter = setComputeFunctionDescriptor + @autoproperty threadGroupSizeIsMultipleOfThreadExecutionWidth::Bool setter = setThreadGroupSizeIsMultipleOfThreadExecutionWidth + @autoproperty maxTotalThreadsPerThreadgroup::UInt64 setter = setMaxTotalThreadsPerThreadgroup + @autoproperty requiredThreadsPerThreadgroup::MTLSize setter = setRequiredThreadsPerThreadgroup + @autoproperty supportBinaryLinking::Bool setter = setSupportBinaryLinking + @autoproperty staticLinkingDescriptor::id{MTL4StaticLinkingDescriptor} setter = setStaticLinkingDescriptor + @autoproperty supportIndirectCommandBuffers::MTL4IndirectCommandBufferSupportState setter = setSupportIndirectCommandBuffers +end + +@cenum MTL4LogicalToPhysicalColorAttachmentMappingState::Int64 begin + MTL4LogicalToPhysicalColorAttachmentMappingStateIdentity = 0 + MTL4LogicalToPhysicalColorAttachmentMappingStateInherited = 1 +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4RenderPipelineColorAttachmentDescriptor <: NSObject + +@objcproperties MTL4RenderPipelineColorAttachmentDescriptor begin + @autoproperty pixelFormat::MTLPixelFormat setter = setPixelFormat + @autoproperty blendingState::MTL4BlendState setter = setBlendingState + @autoproperty sourceRGBBlendFactor::MTLBlendFactor setter = setSourceRGBBlendFactor + @autoproperty destinationRGBBlendFactor::MTLBlendFactor setter = setDestinationRGBBlendFactor + @autoproperty rgbBlendOperation::MTLBlendOperation setter = setRgbBlendOperation + @autoproperty sourceAlphaBlendFactor::MTLBlendFactor setter = setSourceAlphaBlendFactor + @autoproperty destinationAlphaBlendFactor::MTLBlendFactor setter = setDestinationAlphaBlendFactor + @autoproperty alphaBlendOperation::MTLBlendOperation setter = setAlphaBlendOperation + @autoproperty writeMask::MTLColorWriteMask setter = setWriteMask +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4RenderPipelineColorAttachmentDescriptorArray <: NSObject + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4RenderPipelineBinaryFunctionsDescriptor <: NSObject + +@objcproperties MTL4RenderPipelineBinaryFunctionsDescriptor begin + @autoproperty vertexAdditionalBinaryFunctions::id{NSArray} type = Vector{MTL4BinaryFunction} setter = setVertexAdditionalBinaryFunctions + @autoproperty fragmentAdditionalBinaryFunctions::id{NSArray} type = Vector{MTL4BinaryFunction} setter = setFragmentAdditionalBinaryFunctions + @autoproperty tileAdditionalBinaryFunctions::id{NSArray} type = Vector{MTL4BinaryFunction} setter = setTileAdditionalBinaryFunctions + @autoproperty objectAdditionalBinaryFunctions::id{NSArray} type = Vector{MTL4BinaryFunction} setter = setObjectAdditionalBinaryFunctions + @autoproperty meshAdditionalBinaryFunctions::id{NSArray} type = Vector{MTL4BinaryFunction} setter = setMeshAdditionalBinaryFunctions +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4RenderPipelineDescriptor <: MTL4PipelineDescriptor + +@objcproperties MTL4RenderPipelineDescriptor begin + @autoproperty vertexFunctionDescriptor::id{MTL4FunctionDescriptor} setter = setVertexFunctionDescriptor + @autoproperty fragmentFunctionDescriptor::id{MTL4FunctionDescriptor} setter = setFragmentFunctionDescriptor + @autoproperty vertexDescriptor::id{MTLVertexDescriptor} setter = setVertexDescriptor + @autoproperty rasterSampleCount::UInt64 setter = setRasterSampleCount + @autoproperty alphaToCoverageState::MTL4AlphaToCoverageState setter = setAlphaToCoverageState + @autoproperty alphaToOneState::MTL4AlphaToOneState setter = setAlphaToOneState + @autoproperty rasterizationEnabled::Bool getter = isRasterizationEnabled setter = setRasterizationEnabled + @autoproperty maxVertexAmplificationCount::UInt64 setter = setMaxVertexAmplificationCount + @autoproperty colorAttachments::id{MTL4RenderPipelineColorAttachmentDescriptorArray} + @autoproperty inputPrimitiveTopology::MTLPrimitiveTopologyClass setter = setInputPrimitiveTopology + @autoproperty vertexStaticLinkingDescriptor::id{MTL4StaticLinkingDescriptor} setter = setVertexStaticLinkingDescriptor + @autoproperty fragmentStaticLinkingDescriptor::id{MTL4StaticLinkingDescriptor} setter = setFragmentStaticLinkingDescriptor + @autoproperty supportVertexBinaryLinking::Bool setter = setSupportVertexBinaryLinking + @autoproperty supportFragmentBinaryLinking::Bool setter = setSupportFragmentBinaryLinking + @autoproperty colorAttachmentMappingState::MTL4LogicalToPhysicalColorAttachmentMappingState setter = setColorAttachmentMappingState + @autoproperty supportIndirectCommandBuffers::MTL4IndirectCommandBufferSupportState setter = setSupportIndirectCommandBuffers +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4MachineLearningPipelineDescriptor <: MTL4PipelineDescriptor + +@objcproperties MTL4MachineLearningPipelineDescriptor begin + @autoproperty label::id{NSString} setter = setLabel + @autoproperty machineLearningFunctionDescriptor::id{MTL4FunctionDescriptor} setter = setMachineLearningFunctionDescriptor +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4MachineLearningPipelineReflection <: NSObject + +@objcproperties MTL4MachineLearningPipelineReflection begin + @autoproperty bindings::id{NSArray} type = Vector{MTLBinding} +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4MachineLearningPipelineState <: MTLAllocation + +@objcproperties MTL4MachineLearningPipelineState begin + @autoproperty label::id{NSString} + @autoproperty device::id{MTLDevice} + @autoproperty reflection::id{MTL4MachineLearningPipelineReflection} + @autoproperty intermediatesHeapSize::UInt64 +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4TileRenderPipelineDescriptor <: MTL4PipelineDescriptor + +@objcproperties MTL4TileRenderPipelineDescriptor begin + @autoproperty tileFunctionDescriptor::id{MTL4FunctionDescriptor} setter = setTileFunctionDescriptor + @autoproperty rasterSampleCount::UInt64 setter = setRasterSampleCount + @autoproperty colorAttachments::id{MTLTileRenderPipelineColorAttachmentDescriptorArray} + @autoproperty threadgroupSizeMatchesTileSize::Bool setter = setThreadgroupSizeMatchesTileSize + @autoproperty maxTotalThreadsPerThreadgroup::UInt64 setter = setMaxTotalThreadsPerThreadgroup + @autoproperty requiredThreadsPerThreadgroup::MTLSize setter = setRequiredThreadsPerThreadgroup + @autoproperty staticLinkingDescriptor::id{MTL4StaticLinkingDescriptor} setter = setStaticLinkingDescriptor + @autoproperty supportBinaryLinking::Bool setter = setSupportBinaryLinking +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4MeshRenderPipelineDescriptor <: MTL4PipelineDescriptor + +@objcproperties MTL4MeshRenderPipelineDescriptor begin + @autoproperty objectFunctionDescriptor::id{MTL4FunctionDescriptor} setter = setObjectFunctionDescriptor + @autoproperty meshFunctionDescriptor::id{MTL4FunctionDescriptor} setter = setMeshFunctionDescriptor + @autoproperty fragmentFunctionDescriptor::id{MTL4FunctionDescriptor} setter = setFragmentFunctionDescriptor + @autoproperty maxTotalThreadsPerObjectThreadgroup::UInt64 setter = setMaxTotalThreadsPerObjectThreadgroup + @autoproperty maxTotalThreadsPerMeshThreadgroup::UInt64 setter = setMaxTotalThreadsPerMeshThreadgroup + @autoproperty requiredThreadsPerObjectThreadgroup::MTLSize setter = setRequiredThreadsPerObjectThreadgroup + @autoproperty requiredThreadsPerMeshThreadgroup::MTLSize setter = setRequiredThreadsPerMeshThreadgroup + @autoproperty objectThreadgroupSizeIsMultipleOfThreadExecutionWidth::Bool setter = setObjectThreadgroupSizeIsMultipleOfThreadExecutionWidth + @autoproperty meshThreadgroupSizeIsMultipleOfThreadExecutionWidth::Bool setter = setMeshThreadgroupSizeIsMultipleOfThreadExecutionWidth + @autoproperty payloadMemoryLength::UInt64 setter = setPayloadMemoryLength + @autoproperty maxTotalThreadgroupsPerMeshGrid::UInt64 setter = setMaxTotalThreadgroupsPerMeshGrid + @autoproperty rasterSampleCount::UInt64 setter = setRasterSampleCount + @autoproperty alphaToCoverageState::MTL4AlphaToCoverageState setter = setAlphaToCoverageState + @autoproperty alphaToOneState::MTL4AlphaToOneState setter = setAlphaToOneState + @autoproperty rasterizationEnabled::Bool getter = isRasterizationEnabled setter = setRasterizationEnabled + @autoproperty maxVertexAmplificationCount::UInt64 setter = setMaxVertexAmplificationCount + @autoproperty colorAttachments::id{MTL4RenderPipelineColorAttachmentDescriptorArray} + @autoproperty objectStaticLinkingDescriptor::id{MTL4StaticLinkingDescriptor} setter = setObjectStaticLinkingDescriptor + @autoproperty meshStaticLinkingDescriptor::id{MTL4StaticLinkingDescriptor} setter = setMeshStaticLinkingDescriptor + @autoproperty fragmentStaticLinkingDescriptor::id{MTL4StaticLinkingDescriptor} setter = setFragmentStaticLinkingDescriptor + @autoproperty supportObjectBinaryLinking::Bool setter = setSupportObjectBinaryLinking + @autoproperty supportMeshBinaryLinking::Bool setter = setSupportMeshBinaryLinking + @autoproperty supportFragmentBinaryLinking::Bool setter = setSupportFragmentBinaryLinking + @autoproperty colorAttachmentMappingState::MTL4LogicalToPhysicalColorAttachmentMappingState setter = setColorAttachmentMappingState + @autoproperty supportIndirectCommandBuffers::MTL4IndirectCommandBufferSupportState setter = setSupportIndirectCommandBuffers +end + +@cenum MTL4PipelineDataSetSerializerConfiguration::Int64 begin + MTL4PipelineDataSetSerializerConfigurationCaptureDescriptors = 1 + MTL4PipelineDataSetSerializerConfigurationCaptureBinaries = 2 +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4PipelineDataSetSerializerDescriptor <: NSObject + +@objcproperties MTL4PipelineDataSetSerializerDescriptor begin + @autoproperty configuration::MTL4PipelineDataSetSerializerConfiguration setter = setConfiguration +end + +@cenum MTL4BinaryFunctionOptions::UInt64 begin + MTL4BinaryFunctionOptionNone = 0x0000000000000000 + MTL4BinaryFunctionOptionPipelineIndependent = 0x0000000000000002 +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4BinaryFunctionDescriptor <: NSObject + +@objcproperties MTL4BinaryFunctionDescriptor begin + @autoproperty name::id{NSString} setter = setName + @autoproperty functionDescriptor::id{MTL4FunctionDescriptor} setter = setFunctionDescriptor + @autoproperty options::MTL4BinaryFunctionOptions setter = setOptions +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4PipelineStageDynamicLinkingDescriptor <: NSObject + +@objcproperties MTL4PipelineStageDynamicLinkingDescriptor begin + @autoproperty maxCallStackDepth::UInt64 setter = setMaxCallStackDepth + @autoproperty binaryLinkedFunctions::id{NSArray} type = Vector{MTL4BinaryFunction} setter = setBinaryLinkedFunctions + @autoproperty preloadedLibraries::id{NSArray} type = Vector{MTLDynamicLibrary} setter = setPreloadedLibraries +end + +@objcwrapper immutable = true availability = macos(v"26.0.0") MTL4RenderPipelineDynamicLinkingDescriptor <: NSObject + +@objcproperties MTL4RenderPipelineDynamicLinkingDescriptor begin + @autoproperty vertexLinkingDescriptor::id{MTL4PipelineStageDynamicLinkingDescriptor} + @autoproperty fragmentLinkingDescriptor::id{MTL4PipelineStageDynamicLinkingDescriptor} + @autoproperty tileLinkingDescriptor::id{MTL4PipelineStageDynamicLinkingDescriptor} + @autoproperty objectLinkingDescriptor::id{MTL4PipelineStageDynamicLinkingDescriptor} + @autoproperty meshLinkingDescriptor::id{MTL4PipelineStageDynamicLinkingDescriptor} +end + mutable struct MTLSharedEventHandlePrivate end mutable struct MTLSharedTextureHandlePrivate end diff --git a/res/wrap/libmtl.toml b/res/wrap/libmtl.toml index 9f1fcc331..875365d96 100644 --- a/res/wrap/libmtl.toml +++ b/res/wrap/libmtl.toml @@ -1,6 +1,7 @@ [general] library_name = "Symbol(\"/System/Library/Frameworks/Metal.framework/Resources/BridgeSupport/Metal.dylib\")" output_file_path = "../../lib/mtl/libmtl.jl" +prologue_file_path = "libmtl_prologue.jl" minimum_macos_supported = "13" @@ -50,6 +51,15 @@ immutable=false [api.MTLCommandQueue] immutable=false +[api.MTL4CommandQueueDescriptor] +immutable=false + +[api.MTL4CommandQueue] +immutable=false + +[api.MTL4CommitOptions] +immutable=false + [api.MTLCompileOptions] immutable=false [api.MTLCompileOptions.proptype] @@ -58,6 +68,15 @@ immutable=false [api.MTLComputeCommandEncoder] immutable=false +[api.MTL4ComputeCommandEncoder] +immutable=false + +[api.MTL4CommandAllocatorDescriptor] +immutable=false + +[api.MTL4CommandAllocator] +immutable=false + [api.MTLComputePipelineDescriptor] immutable=false @@ -76,6 +95,9 @@ immutable=false [api.MTLFunctionDescriptor] immutable=false +[api.MTL4FunctionDescriptor] +immutable=false + [api.MTLHeap] immutable=false diff --git a/res/wrap/libmtl_prologue.jl b/res/wrap/libmtl_prologue.jl new file mode 100644 index 000000000..fbeeb15c9 --- /dev/null +++ b/res/wrap/libmtl_prologue.jl @@ -0,0 +1 @@ +const _NSRange = NSRange diff --git a/src/compiler/execution.jl b/src/compiler/execution.jl index cfc31e4c0..80c5ff669 100644 --- a/src/compiler/execution.jl +++ b/src/compiler/execution.jl @@ -100,16 +100,17 @@ end ## argument conversion -struct Adaptor +struct Adaptor{T <: Union{Nothing,MTLComputeCommandEncoder,MTL4ArgumentTable}} # the current command encoder, if any. - cce::Union{Nothing,MTLComputeCommandEncoder} + cce::T end # convert Metal buffers to their GPU address +function Adapt.adapt_storage(to::Adaptor{<:MTLComputeCommandEncoder}, buf::MTLBuffer) + MTL.use!(to.cce, buf, MTL.ReadWriteUsage) + reinterpret(Core.LLVMPtr{Nothing,AS.Device}, buf.gpuAddress) +end function Adapt.adapt_storage(to::Adaptor, buf::MTLBuffer) - if to.cce !== nothing - MTL.use!(to.cce, buf, MTL.ReadWriteUsage) - end reinterpret(Core.LLVMPtr{Nothing,AS.Device}, buf.gpuAddress) end function Adapt.adapt_storage(to::Adaptor, ptr::MtlPtr{T}) where {T} @@ -264,7 +265,9 @@ end end @autoreleasepool function (kernel::HostKernel)(args...; groups=1, threads=1, - queue=global_queue(device())) + queue=use_metal4() ? global_queue4(device()) : global_queue(device())) + use_mtl4 = queue isa MTL4CommandQueue + groups = MTLSize(groups) threads = MTLSize(threads) (groups.width>0 && groups.height>0 && groups.depth>0) || @@ -275,16 +278,37 @@ end (threads.width * threads.height * threads.depth) > kernel.pipeline.maxTotalThreadsPerThreadgroup && throw(ArgumentError("Number of threads in group ($(threads.width * threads.height * threads.depth)) should not exceed $(kernel.pipeline.maxTotalThreadsPerThreadgroup)")) - cmdbuf = MTLCommandBuffer(queue) - cmdbuf.label = "MTLCommandBuffer($(nameof(kernel.f)))" - cce = MTLComputeCommandEncoder(cmdbuf) + if use_mtl4 + allocator = MTL4CommandAllocator(device()) + cmdbuf = MTL4CommandBuffer(device()) + cmdbuf.label = "MTL4CommandBuffer($(nameof(kernel.f)))" + + beginCommandBufferWithAllocator!(cmdbuf, allocator) + cce = MTL4ComputeCommandEncoder(cmdbuf) + else + cmdbuf = MTLCommandBuffer(queue) + cmdbuf.label = "MTLCommandBuffer($(nameof(kernel.f)))" + cce = MTLComputeCommandEncoder(cmdbuf) + end + argument_buffers = try MTL.set_function!(cce, kernel.pipeline) - bufs = encode_arguments!(cce, kernel, kernel.f, args...) + if use_mtl4 + argtabdesc = MTL.MTL4ArgumentTableDescriptor() + argtabdesc.maxBufferBindCount = min(31, length(args) + 1) + argtab = MTL.MTL4ArgumentTable(device(), argtabdesc) + bufs = encode_arguments!(argtab, kernel, kernel.f, args...) + + MTL.set_argument_table!(cce, argtab) + else + bufs = encode_arguments!(cce, kernel, kernel.f, args...) + end MTL.append_current_function!(cce, groups, threads) bufs finally + use_mtl4 && barrierAfterStages!(cce) close(cce) + use_mtl4 && endCommandBuffer!(cmdbuf) end # the command buffer retains resources that are explicitly encoded (i.e. direct buffer @@ -295,20 +319,38 @@ end # kernel has actually completed. # # TODO: is there a way to bind additional resources to the command buffer? - roots = [kernel.f, args] - MTL.on_completed(cmdbuf) do buf - empty!(roots) - foreach(free, argument_buffers) - - # Check for errors - # XXX: we cannot do this nicely, e.g. throwing an `error` or reporting with `@error` - # because we're not allowed to switch tasks from this contexts. - if buf.status == MTL.MTLCommandBufferStatusError - Core.println("ERROR: Failed to submit command buffer: $(buf.error.localizedDescription)") + if use_mtl4 + options = MTL.MTL4CommitOptions() do feedback + # TODO: RESOURCE MANAGEMENT STUFF + + # Check for errors + # XXX: we cannot do this nicely, e.g. throwing an `error` or reporting with `@error` + # because we're not allowed to switch tasks from this contexts. + + if !isnothing(feedback.error) + Core.println("ERROR: Failed to submit command buffer: $(feedback.error.localizedDescription)") + end end + else + roots = [kernel.f, args] + MTL.on_completed(cmdbuf) do buf + empty!(roots) + foreach(free, argument_buffers) + + # Check for errors + # XXX: we cannot do this nicely, e.g. throwing an `error` or reporting with `@error` + # because we're not allowed to switch tasks from this contexts. + if buf.status == MTL.MTLCommandBufferStatusError + Core.println("ERROR: Failed to submit command buffer: $(buf.error.localizedDescription)") + end + end + end + if use_mtl4 + commit!(queue, cmdbuf, options) + else + commit!(cmdbuf) end - commit!(cmdbuf) end ## Intra-warp Helpers diff --git a/src/initialization.jl b/src/initialization.jl index fe6fca31f..683cd3f94 100644 --- a/src/initialization.jl +++ b/src/initialization.jl @@ -8,11 +8,17 @@ return false end end + + const use_metal4 = OncePerProcess{Bool}() do + dev = device() + force_metal3 = @load_preference("force_metal3", false) + return functional() && !force_metal3 && supports_family(dev, MTL.MTLGPUFamilyMetal4) + end else # Becomes `nothing` once it has been determined that the device is on macOS const _functional = Ref{Union{Nothing,Bool}}(false) - function functional() + function functional()::Bool if isnothing(_functional[]) dev = device() @@ -22,6 +28,16 @@ else end _functional[] end + + const _use_metal4 = Ref{Union{Nothing,Bool}}(nothing) + function use_metal4()::Bool + if isnothing(_use_metal4[]) + dev = device() + force_metal3 = @load_preference("force_metal3", false) + _use_metal4[] = functional() && !force_metal3 && supports_family(dev, MTL.MTLGPUFamilyMetal4) + end + _use_metal4[] + end end function __init__() diff --git a/src/memory.jl b/src/memory.jl index 1268f1eb0..5dcb7ed3b 100644 --- a/src/memory.jl +++ b/src/memory.jl @@ -31,7 +31,7 @@ Base.convert(::Type{Ptr{T}}, ptr::MtlPtr) where {T} = # CPU -> GPU function Base.unsafe_copyto!(dev::MTLDevice, dst::MtlPtr{T}, src::Ptr{T}, N::Integer; - queue::MTLCommandQueue=global_queue(dev), async::Bool=false) where T + queue=use_metal4() ? global_queue4(dev) : global_queue(dev), async::Bool=false) where T storage_type = dst.buffer.storageMode if storage_type == MTL.MTLStorageModePrivate # stage through a shared buffer @@ -53,7 +53,7 @@ end # GPU -> CPU function Base.unsafe_copyto!(dev::MTLDevice, dst::Ptr{T}, src::MtlPtr{T}, N::Integer; - queue::MTLCommandQueue=global_queue(dev), async::Bool=false) where T + queue=use_metal4() ? global_queue4(dev) : global_queue(dev), async::Bool=false) where T storage_type = src.buffer.storageMode if storage_type == MTL.MTLStorageModePrivate # stage through a shared buffer @@ -71,9 +71,9 @@ function Base.unsafe_copyto!(dev::MTLDevice, dst::Ptr{T}, src::MtlPtr{T}, N::Int unsafe_copyto!(dst, convert(Ptr{T}, tmp_buf), N) end free(tmp_buf) - elseif storage_type == MTL.MTLStorageModeShared + elseif storage_type == MTL.MTLStorageModeShared unsafe_copyto!(dst, convert(Ptr{T}, src), N) - elseif storage_type == MTL.MTLStorageModeManaged + elseif storage_type == MTL.MTLStorageModeManaged cmdbuf = MTLCommandBuffer(queue) do cmdbuf MTLBlitCommandEncoder(cmdbuf) do enc append_sync!(enc, src.buffer) @@ -88,32 +88,49 @@ end # GPU -> GPU @autoreleasepool function Base.unsafe_copyto!(dev::MTLDevice, dst::MtlPtr{T}, src::MtlPtr{T}, N::Integer; - queue::MTLCommandQueue=global_queue(dev), + queue=use_metal4() ? global_queue4(dev) : global_queue(dev), async::Bool=false) where T if N > 0 - cmdbuf = MTLCommandBuffer(queue) - MTLBlitCommandEncoder(cmdbuf) do enc - append_copy!(enc, dst.buffer, dst.offset, src.buffer, src.offset, N * sizeof(T)) + if queue isa MTL4CommandQueue + @info "MTL4" + cmdbuf = MTL4CommandBuffer(dev; queue) do cmdbuf + MTL4ComputeCommandEncoder(cmdbuf, !async) do enc + append_copy!(enc, dst.buffer, dst.offset, src.buffer, src.offset, N * sizeof(T)) + end + end + else + @info "MTL3" + cmdbuf = MTLCommandBuffer(queue) + MTLBlitCommandEncoder(cmdbuf) do enc + append_copy!(enc, dst.buffer, dst.offset, src.buffer, src.offset, N * sizeof(T)) + end + commit!(cmdbuf) + async || wait_completed(cmdbuf) end - commit!(cmdbuf) - async || wait_completed(cmdbuf) end return dst end @autoreleasepool function unsafe_fill!(dev::MTLDevice, dst::MtlPtr{T}, value::Union{UInt8,Int8}, N::Integer; - queue::MTLCommandQueue=global_queue(dev), + queue=use_metal4() ? global_queue4(dev) : global_queue(dev), async::Bool=false) where T if N > 0 - cmdbuf = MTLCommandBuffer(queue) - MTLBlitCommandEncoder(cmdbuf) do enc - append_fillbuffer!(enc, dst.buffer, value, N * sizeof(T), dst.offset) + if queue isa MTL4CommandQueue + @info "MTL4" + cmdbuf = MTL4CommandBuffer(dev; queue) do cmdbuf + MTL4ComputeCommandEncoder(cmdbuf, !async) do enc + append_fillbuffer!(enc, dst.buffer, value, N * sizeof(T), dst.offset) + end + end + else + cmdbuf = MTLCommandBuffer(queue) + MTLBlitCommandEncoder(cmdbuf) do enc + append_fillbuffer!(enc, dst.buffer, value, N * sizeof(T), dst.offset) + end + commit!(cmdbuf) + async || wait_completed(cmdbuf) end - commit!(cmdbuf) - async || wait_completed(cmdbuf) end return dst end - -# TODO: Implement generic fill since mtBlitCommandEncoderFillBuffer is limiting diff --git a/src/state.jl b/src/state.jl index 3a0512e52..9c064dc91 100644 --- a/src/state.jl +++ b/src/state.jl @@ -36,11 +36,12 @@ Sets the Metal GPU device associated with the current Julia task. device!(dev::MTLDevice) = task_local_storage(:MTLDevice, dev) const global_queues = WeakKeyDict{MTLCommandQueue,Nothing}() +const global_queues4 = WeakKeyDict{MTL4CommandQueue,Tuple{MTL4CommandBuffer, MTL4CommandAllocator}}() """ global_queue(dev::MTLDevice)::MTLCommandQueue -Return the Metal command queue associated with the current Julia thread. +Return the Metal 3 command queue associated with the current Julia thread. """ function global_queue(dev::MTLDevice) get!(task_local_storage(), (:MTLCommandQueue, dev)) do @@ -55,6 +56,27 @@ function global_queue(dev::MTLDevice) end::MTLCommandQueue end +""" + global_queue4(dev::MTLDevice)::MTL4CommandQueue + +Return the Metal 4 command queue associated with the current Julia thread. +""" +function global_queue4(dev::MTLDevice) + get!(task_local_storage(), (:MTL4CommandQueue, dev)) do + @autoreleasepool begin + desc = MTL4CommandQueueDescriptor("global_queue4($(current_task()))") + + # NOTE: MTL4CommandQueue itself is manually reference-counted, + # the release pool is for resources used during its construction. + queue = MTL4CommandQueue(dev, desc) + + + global_queues4[queue] = (MTL4CommandBuffer(dev, "sync_buffer($(current_task()))"), MTL4CommandAllocator(dev, "sync_allocater($(current_task()))")) + queue + end + end::MTL4CommandQueue +end + # TODO: Increase performance (currently ~15us) """ synchronize(queue) @@ -65,10 +87,37 @@ Create a new MTLCommandBuffer from the global command queue, commit it to the qu and simply wait for it to be completed. Since command buffers *should* execute in a First-In-First-Out manner, this synchronizes the GPU. """ -@autoreleasepool function synchronize(queue::MTLCommandQueue=global_queue(device())) +@autoreleasepool function synchronize(queue::MTLCommandQueue) cmdbuf = MTLCommandBuffer(queue) commit!(cmdbuf) wait_completed(cmdbuf) + return +end +@autoreleasepool function synchronize(queue::MTL4CommandQueue) + cmdbuf, allocator = get(global_queues4, queue) do + dev = queue.device + MTL4CommandBuffer(dev), MTL4CommandAllocator(dev) + end + + cmdbuf = commit!(cmdbuf, queue, allocator) do cmdbuf + encoder = MTL4ComputeCommandEncoder(cmdbuf) + MTL.barrierAfterQueueStages!(encoder) + close(encoder) + end + return +end +function synchronize() + dev = device() + tlskeys = keys(task_local_storage()) + # hasmtl3key = (:MTLCommandQueue, dev) in tlskeys + # hasmtl4key = use_metal4() && (:MTL4CommandQueue, dev) in tlskeys + if (:MTLCommandQueue, dev) in tlskeys + synchronize(global_queue(dev)) + end + if use_metal4() && (:MTL4CommandQueue, dev) in tlskeys + synchronize(global_queue4(dev)) + end + return end """ @@ -80,4 +129,7 @@ function device_synchronize() for queue in keys(global_queues) synchronize(queue) end + for queue in keys(global_queues4) + synchronize(queue) + end end diff --git a/src/utilities.jl b/src/utilities.jl index a6ee9c042..adec4bf43 100644 --- a/src/utilities.jl +++ b/src/utilities.jl @@ -42,6 +42,7 @@ function versioninfo(io::IO=stdout) prefs = [ "default_storage" => load_preference(Metal, "default_storage"), + "force_metal3" => load_preference(Metal, "force_metal3"), ] if any(x->!isnothing(x[2]), prefs) println(io, "Preferences:")