ARMSX2/pcsx2/GS/Renderers/Metal/GSDeviceMTL.mm

// SPDX-FileCopyrightText: 2002-2026 PCSX2 Dev Team
// SPDX-License-Identifier: GPL-3.0+

#include "Host.h"
#include "GS/GSGL.h"
#include "GS/Renderers/Metal/GSMetalCPPAccessible.h"
#include "GS/Renderers/Metal/GSDeviceMTL.h"
#include "GS/Renderers/Metal/GSTextureMTL.h"
#include "GS/GSPerfMon.h"
#include "GS/GSShaderCompileIndicator.h"

#include "common/Console.h"
#include "common/HostSys.h"

#include "cpuinfo.h"
#include "imgui.h"

#ifdef __APPLE__
#include "GSMTLSharedHeader.h"

static constexpr simd::float2 ToSimd(const GSVector2& vec)
{
	return simd::make_float2(vec.x, vec.y);
}

GSDevice* MakeGSDeviceMTL()
{
	return new GSDeviceMTL();
}

static GSAdapterInfo GetMetalAdapterInfo(id<MTLDevice> dev)
{
	GSAdapterInfo ai;
	ai.name = [[dev name] UTF8String];
	ai.max_texture_size = GSMTLDevice::GetMaxTextureSize(dev);
	ai.max_upscale_multiplier = GSGetMaxUpscaleMultiplier(ai.max_texture_size);
	return ai;
}

std::vector<GSAdapterInfo> GetMetalAdapterList()
{ @autoreleasepool {
	std::vector<GSAdapterInfo> list;
	// MTLCopyAllDevices only exists on iOS 18. We ship down to 17, where there's a
	// single GPU and nothing to enumerate anyway.
	if (@available(macOS 10.11, iOS 18.0, *))
	{
		auto devs = MRCTransfer(MTLCopyAllDevices());
		for (id<MTLDevice> dev in devs.Get())
			list.push_back(GetMetalAdapterInfo(dev));
	}
	else
	{
		auto dev = MRCTransfer(MTLCreateSystemDefaultDevice());
		if (dev.Get())
			list.push_back(GetMetalAdapterInfo(dev.Get()));
	}
	return list;
}}

bool GSDeviceMTL::UsageTracker::PrepareForAllocation(u64 last_draw, size_t amt)
{
	auto removeme = std::find_if(m_usage.begin(), m_usage.end(), [last_draw](UsageEntry usage){ return usage.drawno > last_draw; });
	if (removeme != m_usage.begin())
		m_usage.erase(m_usage.begin(), removeme);

	bool still_in_use = false;
	bool needs_wrap = m_pos + amt > m_size;
	if (!m_usage.empty())
	{
		size_t used = m_usage.front().pos;
		if (needs_wrap)
			still_in_use = used >= m_pos || used < amt;
		else
			still_in_use = used >= m_pos && used < m_pos + amt;
	}
	if (needs_wrap)
		m_pos = 0;

	return still_in_use || amt > m_size;
}

size_t GSDeviceMTL::UsageTracker::Allocate(u64 current_draw, size_t amt)
{
	if (m_usage.empty() || m_usage.back().drawno != current_draw)
		m_usage.push_back({current_draw, m_pos});
	size_t ret = m_pos;
	m_pos += amt;
	return ret;
}

void GSDeviceMTL::UsageTracker::Reset(size_t new_size)
{
	m_usage.clear();
	m_size = new_size;
	m_pos = 0;
}

GSDeviceMTL::GSDeviceMTL()
	: m_backref(std::make_shared<std::pair<std::mutex, GSDeviceMTL*>>())
	, m_dev(nil)
{
	m_backref->second = this;
	m_resource_options_shared_wc = MTLResourceStorageModeShared | MTLResourceCPUCacheModeWriteCombined;
#ifdef _M_X86
	// WC memory doesn't work properly on AMD hackintoshes, and ends up being horribly slow even when only writing
	if (cpuinfo_get_core(0)->vendor == cpuinfo_vendor_amd)
		m_resource_options_shared_wc = MTLResourceStorageModeShared;
#endif
}

GSDeviceMTL::~GSDeviceMTL()
{
	// m_ds_as_rt_texture is owned if the device has memoryless textures
	if (m_dev.features.memoryless_textures)
		[m_ds_as_rt_texture release];
	else if (m_ds_as_rt_gstexture)
		delete m_ds_as_rt_gstexture;
}

GSDeviceMTL::Map GSDeviceMTL::Allocate(UploadBuffer& buffer, size_t amt)
{
	amt = (amt + 31) & ~31ull;
	u64 last_draw = m_last_finished_draw.load(std::memory_order_acquire);
	bool needs_new = buffer.usage.PrepareForAllocation(last_draw, amt);
	if (needs_new) [[unlikely]]
	{
		// Orphan buffer
		size_t newsize = std::max<size_t>(buffer.usage.Size() * 2, 4096);
		while (newsize < amt)
			newsize *= 2;
		MTLResourceOptions options = m_resource_options_shared_wc;
		buffer.mtlbuffer = MRCTransfer([m_dev.dev newBufferWithLength:newsize options:options]);
		pxAssertRel(buffer.mtlbuffer, "Failed to allocate MTLBuffer (out of memory?)");
		buffer.buffer = [buffer.mtlbuffer contents];
		buffer.usage.Reset(newsize);
	}

	size_t pos = buffer.usage.Allocate(m_current_draw, amt);

	Map ret = {buffer.mtlbuffer, pos, reinterpret_cast<char*>(buffer.buffer) + pos};
	pxAssertMsg(pos <= buffer.usage.Size(), "Previous code should have guaranteed there was enough space");
	return ret;
}

/// Allocate space in the given buffer for use with the given render command encoder
GSDeviceMTL::Map GSDeviceMTL::Allocate(BufferPair& buffer, size_t amt)
{
	amt = (amt + 31) & ~31ull;
	u64 last_draw = m_last_finished_draw.load(std::memory_order_acquire);
	size_t base_pos = buffer.usage.Pos();
	bool needs_new = buffer.usage.PrepareForAllocation(last_draw, amt);
	bool needs_upload = needs_new || buffer.usage.Pos() == 0;
	if (!m_dev.features.unified_memory && needs_upload)
	{
		if (base_pos != buffer.last_upload)
		{
			id<MTLBlitCommandEncoder> enc = GetVertexUploadEncoder();
			[enc copyFromBuffer:buffer.cpubuffer
			       sourceOffset:buffer.last_upload
			           toBuffer:buffer.gpubuffer
			  destinationOffset:buffer.last_upload
			               size:base_pos - buffer.last_upload];
		}
		buffer.last_upload = 0;
	}
	if (needs_new) [[unlikely]]
	{
		// Orphan buffer
		size_t newsize = std::max<size_t>(buffer.usage.Size() * 2, 4096);
		while (newsize < amt)
			newsize *= 2;
		MTLResourceOptions options = m_resource_options_shared_wc;
		buffer.cpubuffer = MRCTransfer([m_dev.dev newBufferWithLength:newsize options:options]);
		pxAssertRel(buffer.cpubuffer, "Failed to allocate MTLBuffer (out of memory?)");
		buffer.buffer = [buffer.cpubuffer contents];
		buffer.usage.Reset(newsize);
		if (!m_dev.features.unified_memory)
		{
			options = MTLResourceStorageModePrivate | MTLResourceHazardTrackingModeUntracked;
			buffer.gpubuffer = MRCTransfer([m_dev.dev newBufferWithLength:newsize options:options]);
			pxAssertRel(buffer.gpubuffer, "Failed to allocate MTLBuffer (out of memory?)");
		}
	}

	size_t pos = buffer.usage.Allocate(m_current_draw, amt);
	Map ret = {nil, pos, reinterpret_cast<char*>(buffer.buffer) + pos};
	ret.gpu_buffer = m_dev.features.unified_memory ? buffer.cpubuffer : buffer.gpubuffer;
	pxAssertMsg(pos <= buffer.usage.Size(), "Previous code should have guaranteed there was enough space");
	return ret;
}

void GSDeviceMTL::Sync(BufferPair& buffer)
{
	if (m_dev.features.unified_memory || buffer.usage.Pos() == buffer.last_upload)
		return;

	id<MTLBlitCommandEncoder> enc = GetVertexUploadEncoder();
	[enc copyFromBuffer:buffer.cpubuffer
	       sourceOffset:buffer.last_upload
	           toBuffer:buffer.gpubuffer
	  destinationOffset:buffer.last_upload
	               size:buffer.usage.Pos() - buffer.last_upload];
	[enc updateFence:m_draw_sync_fence];
	buffer.last_upload = buffer.usage.Pos();
}

id<MTLBlitCommandEncoder> GSDeviceMTL::GetTextureUploadEncoder()
{
	if (!m_texture_upload_cmdbuf)
	{
		m_texture_upload_cmdbuf = MRCRetain([m_queue commandBuffer]);
		m_texture_upload_encoder = MRCRetain([m_texture_upload_cmdbuf blitCommandEncoder]);
		pxAssertRel(m_texture_upload_encoder, "Failed to create texture upload encoder!");
		[m_texture_upload_cmdbuf setLabel:@"Texture Upload"];
	}
	return m_texture_upload_encoder;
}

id<MTLBlitCommandEncoder> GSDeviceMTL::GetLateTextureUploadEncoder()
{
	if (!m_late_texture_upload_encoder)
	{
		EndRenderPass();
		m_late_texture_upload_encoder = MRCRetain([GetRenderCmdBuf() blitCommandEncoder]);
		pxAssertRel(m_late_texture_upload_encoder, "Failed to create late texture upload encoder!");
		[m_late_texture_upload_encoder setLabel:@"Late Texture Upload"];
		if (!m_dev.features.unified_memory)
			[m_late_texture_upload_encoder waitForFence:m_draw_sync_fence];
	}
	return m_late_texture_upload_encoder;
}

id<MTLBlitCommandEncoder> GSDeviceMTL::GetVertexUploadEncoder()
{
	if (!m_vertex_upload_cmdbuf)
	{
		m_vertex_upload_cmdbuf = MRCRetain([m_queue commandBuffer]);
		m_vertex_upload_encoder = MRCRetain([m_vertex_upload_cmdbuf blitCommandEncoder]);
		pxAssertRel(m_vertex_upload_encoder, "Failed to create vertex upload encoder!");
		[m_vertex_upload_cmdbuf setLabel:@"Vertex Upload"];
	}
	return m_vertex_upload_encoder;
}

/// Get the draw command buffer, creating a new one if it doesn't exist
id<MTLCommandBuffer> GSDeviceMTL::GetRenderCmdBuf()
{
	if (!m_current_render_cmdbuf)
	{
		m_encoders_in_current_cmdbuf = 0;
		m_current_render_cmdbuf = MRCRetain([m_queue commandBuffer]);
		pxAssertRel(m_current_render_cmdbuf, "Failed to create draw command buffer!");
		[m_current_render_cmdbuf setLabel:@"Draw"];
	}
	return m_current_render_cmdbuf;
}

id<MTLCommandBuffer> GSDeviceMTL::GetRenderCmdBufWithoutCreate()
{
	return m_current_render_cmdbuf;
}

id<MTLFence> GSDeviceMTL::GetSpinFence()
{
	return m_spin_timer ? m_spin_fence : nil;
}

id<MTLTexture> GSDeviceMTL::GetRT1DepthTexture(GSTextureMTL* depth)
{
	if (m_dev.features.framebuffer_fetch)
		return m_ds_as_rt_texture;
	else
		return static_cast<GSTextureMTL*>(m_ds_as_rt)->GetTexture();
}

void GSDeviceMTL::DrawCommandBufferFinished(u64 draw, id<MTLCommandBuffer> buffer)
{
	// We can do the update non-atomically because we only ever update under the lock
	u64 newval = std::max(draw, m_last_finished_draw.load(std::memory_order_relaxed));
	m_last_finished_draw.store(newval, std::memory_order_release);
	AccumulateCommandBufferTime(buffer);
}

void GSDeviceMTL::FlushEncoders()
{
	bool needs_submit = m_current_render_cmdbuf;
	if (needs_submit)
	{
		EndRenderPass();
		Sync(m_vertex_upload_buf);
	}
	if (m_dev.features.unified_memory)
	{
		pxAssertMsg(!m_vertex_upload_cmdbuf, "Should never be used!");
	}
	else if (m_vertex_upload_cmdbuf)
	{
		[m_vertex_upload_encoder endEncoding];
		[m_vertex_upload_cmdbuf commit];
		m_vertex_upload_encoder = nil;
		m_vertex_upload_cmdbuf = nil;
	}
	if (m_texture_upload_cmdbuf)
	{
		[m_texture_upload_encoder endEncoding];
		[m_texture_upload_cmdbuf commit];
		m_texture_upload_encoder = nil;
		m_texture_upload_cmdbuf = nil;
	}
	if (!needs_submit)
		return;
	if (m_late_texture_upload_encoder)
	{
		[m_late_texture_upload_encoder endEncoding];
		m_late_texture_upload_encoder = nil;
	}
	u32 spin_cycles = 0;
	constexpr double s_to_ns = 1000000000;
	if (m_spin_timer)
	{
		u32 spin_id;
		{
			std::lock_guard<std::mutex> guard(m_backref->first);
			auto draw = m_spin_manager.DrawSubmitted(m_encoders_in_current_cmdbuf);
			u32 constant_offset = 200000 * m_spin_manager.SpinsPerUnitTime(); // 200µs
			u32 minimum_spin = 2 * constant_offset; // 400µs (200µs after subtracting constant_offset)
			u32 maximum_spin = std::max<u32>(1024, 16000000 * m_spin_manager.SpinsPerUnitTime()); // 16ms
			if (draw.recommended_spin > minimum_spin)
				spin_cycles = std::min(draw.recommended_spin - constant_offset, maximum_spin);
			spin_id = draw.id;
		}
		[m_current_render_cmdbuf addCompletedHandler:[backref = m_backref, draw = m_current_draw, spin_id](id<MTLCommandBuffer> buf)
		{
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wunguarded-availability"
			// Starting from kernelStartTime includes time the command buffer spent waiting to execute
			// This is useful for avoiding issues on GPUs without async compute (Intel) where spinning
			// delays the next command buffer start, which then makes the spin manager think it should spin more
			// (If a command buffer contains multiple encoders, the GPU will start before the kernel finishes,
			//  so we choose kernelStartTime over kernelEndTime)
			u64 begin = [buf kernelStartTime] * s_to_ns;
			u64 end = [buf GPUEndTime] * s_to_ns;
#pragma clang diagnostic pop
			std::lock_guard<std::mutex> guard(backref->first);
			if (GSDeviceMTL* dev = backref->second)
			{
				dev->DrawCommandBufferFinished(draw, buf);
				dev->m_spin_manager.DrawCompleted(spin_id, static_cast<u32>(begin), static_cast<u32>(end));
			}
		}];
	}
	else
	{
		[m_current_render_cmdbuf addCompletedHandler:[backref = m_backref, draw = m_current_draw](id<MTLCommandBuffer> buf)
		{
			std::lock_guard<std::mutex> guard(backref->first);
			if (GSDeviceMTL* dev = backref->second)
				dev->DrawCommandBufferFinished(draw, buf);
		}];
	}
	[m_current_render_cmdbuf commit];
	m_current_render_cmdbuf = nil;
	m_current_draw++;
	if (spin_cycles)
	{
		id<MTLCommandBuffer> spinCmdBuf = [m_queue commandBuffer];
		[spinCmdBuf setLabel:@"Spin"];
		id<MTLComputeCommandEncoder> spinCmdEncoder = [spinCmdBuf computeCommandEncoder];
		[spinCmdEncoder setLabel:@"Spin"];
		[spinCmdEncoder waitForFence:m_spin_fence];
		[spinCmdEncoder setComputePipelineState:m_spin_pipeline];
		[spinCmdEncoder setBytes:&spin_cycles length:sizeof(spin_cycles) atIndex:0];
		[spinCmdEncoder setBuffer:m_spin_buffer offset:0 atIndex:1];
		[spinCmdEncoder dispatchThreadgroups:MTLSizeMake(1, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
		[spinCmdEncoder endEncoding];
		[spinCmdBuf addCompletedHandler:[backref = m_backref, spin_cycles](id<MTLCommandBuffer> buf)
		{
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wunguarded-availability"
			u64 begin = [buf GPUStartTime] * s_to_ns;
			u64 end = [buf GPUEndTime] * s_to_ns;
#pragma clang diagnostic pop
			std::lock_guard<std::mutex> guard(backref->first);
			if (GSDeviceMTL* dev = backref->second)
				dev->m_spin_manager.SpinCompleted(spin_cycles, static_cast<u32>(begin), static_cast<u32>(end));
		}];
		[spinCmdBuf commit];
	}
}

void GSDeviceMTL::FlushEncodersForReadback()
{
	FlushEncoders();
	if (@available(macOS 10.15, iOS 10.3, *))
	{
		if (GSConfig.HWSpinGPUForReadbacks)
		{
			m_spin_manager.ReadbackRequested();
			m_spin_timer = 30;
		}
	}
}

void GSDeviceMTL::EndRenderPass()
{
	if (m_current_render.encoder)
	{
		EndDebugGroup(m_current_render.encoder);
		g_perfmon.Put(GSPerfMon::RenderPasses, 1);
		if (m_spin_timer)
			[m_current_render.encoder updateFence:m_spin_fence afterStages:MTLRenderStageFragment];
		[m_current_render.encoder endEncoding];
		m_current_render.encoder = nil;
		memset(&m_current_render, 0, offsetof(MainRenderEncoder, depth_sel));
		m_current_render.depth_sel = DepthStencilSelector::NoDepth();
	}
	// The late-upload blit encoder also lives on the render command buffer; any code
	// that ends the render pass and then opens a new encoder on that buffer (e.g.
	// DoCopyRect's blit encoder during OI_BlitFMV) must leave it closed, or Metal aborts
	// with "A command encoder is already encoding to this command buffer".
	if (m_late_texture_upload_encoder)
	{
		[m_late_texture_upload_encoder endEncoding];
		m_late_texture_upload_encoder = nil;
	}
}

static GSVector4 GetRTLoadInfo(GSTextureMTL* tex, MTLLoadAction* load_action)
{
	if (tex)
	{
		if (tex->GetState() == GSTexture::State::Invalidated)
		{
			*load_action = MTLLoadActionDontCare;
		}
		else if (tex->GetState() == GSTexture::State::Cleared && *load_action != MTLLoadActionDontCare)
		{
			*load_action = MTLLoadActionClear;
			return tex->GetClearForFormat();
		}
	}
	return {};
};

void GSDeviceMTL::PrepareBeginRenderPass()
{
	m_encoders_in_current_cmdbuf++;

	if (m_late_texture_upload_encoder)
	{
		[m_late_texture_upload_encoder endEncoding];
		m_late_texture_upload_encoder = nullptr;
	}

	EndRenderPass();
}

void GSDeviceMTL::BeginRenderPass(NSString* name, GSTexture* color, MTLLoadAction color_load,
	GSTexture* depth, MTLLoadAction depth_load, GSTexture* stencil, MTLLoadAction stencil_load, bool rt1)
{
	GSTextureMTL* mc = static_cast<GSTextureMTL*>(color);
	GSTextureMTL* md = static_cast<GSTextureMTL*>(depth);
	GSTextureMTL* ms = static_cast<GSTextureMTL*>(stencil);
	bool needs_new = color   != m_current_render.color_target
	              || depth   != m_current_render.depth_target
	              || stencil != m_current_render.stencil_target
	              || rt1     != m_current_render.has.rt1_depth;

	// Depth and stencil might be the same, so do all invalidation checks before resetting invalidation
	GSVector4 color_clear = GetRTLoadInfo(mc, &color_load);
	GSVector4 depth_clear = GetRTLoadInfo(md, &depth_load);

	// Stencil and depth are one texture, stencil clears aren't supported
	if (ms && ms->GetState() == GSTexture::State::Invalidated)
		stencil_load = MTLLoadActionDontCare;
	needs_new |= mc && color_load == MTLLoadActionClear;
	needs_new |= md && depth_load == MTLLoadActionClear;

	// Reset texture state
	if (mc) mc->SetState(GSTexture::State::Dirty);
	if (md) md->SetState(GSTexture::State::Dirty);
	if (ms) ms->SetState(GSTexture::State::Dirty);

	if (!needs_new)
	{
		if (m_current_render.name != (__bridge void*)name)
		{
			m_current_render.name = (__bridge void*)name;
			[m_current_render.encoder setLabel:name];
		}
		return;
	}

	m_encoders_in_current_cmdbuf++;

	// Note: any open late-upload encoder is closed by EndRenderPass() below.

	int idx = 0;
	if (mc) idx |= 1;
	if (md) idx |= 2;
	if (ms) idx |= 4;
	if (rt1) idx |= 8;

	MTLRenderPassDescriptor* desc = m_render_pass_desc[idx];
	if (mc)
	{
		mc->m_last_write = m_current_draw;
		desc.colorAttachments[0].texture = mc->GetTexture();
		if (color_load == MTLLoadActionClear)
			desc.colorAttachments[0].clearColor = MTLClearColorMake(color_clear.r, color_clear.g, color_clear.b, color_clear.a);
		desc.colorAttachments[0].loadAction = color_load;
	}
	if (md)
	{
		md->m_last_write = m_current_draw;
		desc.depthAttachment.texture = md->GetTexture();
		if (depth_load == MTLLoadActionClear)
			desc.depthAttachment.clearDepth = depth_clear.x;
		desc.depthAttachment.loadAction = depth_load;
	}
	if (ms)
	{
		ms->m_last_write = m_current_draw;
		desc.stencilAttachment.texture = ms->GetTexture();
		pxAssert(stencil_load != MTLLoadActionClear);
		desc.stencilAttachment.loadAction = stencil_load;
	}
	if (rt1)
	{
		pxAssert(md);
		MTLRenderPassColorAttachmentDescriptor* color1 = desc.colorAttachments[1];
		color1.texture = GetRT1DepthTexture(md);
		if (m_features.framebuffer_fetch)
			color1.clearColor = MTLClearColorMake(depth_load == MTLLoadActionClear ? depth_clear.x : -1, 0, 0, 0);
	}

	PrepareBeginRenderPass();
	m_current_render.encoder = MRCRetain([GetRenderCmdBuf() renderCommandEncoderWithDescriptor:desc]);
	m_current_render.name = (__bridge void*)name;
	[m_current_render.encoder setLabel:name];
	if (!m_dev.features.unified_memory)
		[m_current_render.encoder waitForFence:m_draw_sync_fence
		                          beforeStages:MTLRenderStageVertex];
	m_current_render.color_target = color;
	m_current_render.depth_target = depth;
	m_current_render.stencil_target = stencil;
	m_current_render.has.rt1_depth |= rt1;
	pxAssertRel(m_current_render.encoder, "Failed to create render encoder!");
}

void GSDeviceMTL::BeginFullROV(NSString* name, uint32_t width, uint32_t height)
{
	bool needs_new = !m_current_render.is_full_rov()
	              || width  > m_current_render.full_rov_size.w
	              || height > m_current_render.full_rov_size.h;

	if (!needs_new)
	{
		if (m_current_render.name != (__bridge void*)name)
		{
			m_current_render.name = (__bridge void*)name;
			[m_current_render.encoder setLabel:name];
		}
		return;
	}

	if (m_current_render.is_full_rov())
	{
		// Render was too small.  Don't let it shrink in one dimension while growing in the other.
		width  = std::max(width,  m_current_render.full_rov_size.w);
		height = std::max(height, m_current_render.full_rov_size.h);
	}

	m_current_render.full_rov_size.w = width;
	m_current_render.full_rov_size.h = height;

	MTLRenderPassDescriptor* desc = m_full_rov_render_pass_desc;
	[desc setRenderTargetWidth:width];
	[desc setRenderTargetHeight:height];

	if (m_dev.features.rov_requires_rt && (width > m_rov_dummy_texture_size.w || height > m_rov_dummy_texture_size.h))
	{
		m_rov_dummy_texture_size.w = std::max(m_rov_dummy_texture_size.w, width);
		m_rov_dummy_texture_size.h = std::max(m_rov_dummy_texture_size.h, height);
		MTLTextureDescriptor* tdesc = [MTLTextureDescriptor
			texture2DDescriptorWithPixelFormat:MTLPixelFormatR8Unorm
			                             width:m_rov_dummy_texture_size.w
			                            height:m_rov_dummy_texture_size.h
			                         mipmapped:NO];
		[tdesc setUsage:MTLTextureUsageRenderTarget];
		[tdesc setStorageMode:MTLStorageModePrivate];
		id<MTLTexture> tex = m_rov_dummy_texture = MRCTransfer([m_dev.dev newTextureWithDescriptor:tdesc]);
		[tex setLabel:@"ROV Dummy Texture"];
		[[[desc colorAttachments] objectAtIndexedSubscript:0] setTexture:tex];
	}

	PrepareBeginRenderPass();
	m_current_render.encoder = MRCRetain([GetRenderCmdBuf() renderCommandEncoderWithDescriptor:desc]);
	m_current_render.name = (__bridge void*)name;
	[m_current_render.encoder setLabel:name];
	if (!m_dev.features.unified_memory)
		[m_current_render.encoder waitForFence:m_draw_sync_fence
		                          beforeStages:MTLRenderStageVertex];
}

void GSDeviceMTL::FrameCompleted()
{
	if (m_spin_timer)
		m_spin_timer--;
	m_spin_manager.NextFrame();
}

static constexpr MTLPixelFormat ConvertPixelFormat(GSTexture::Format format)
{
	switch (format)
	{
		case GSTexture::Format::DepthColor:   return MTLPixelFormatR32Float;
		case GSTexture::Format::PrimID:       return MTLPixelFormatR32Float;
		case GSTexture::Format::UInt32:       return MTLPixelFormatR32Uint;
		case GSTexture::Format::UInt16:       return MTLPixelFormatR16Uint;
		case GSTexture::Format::UNorm8:       return MTLPixelFormatA8Unorm;
		case GSTexture::Format::Color:        return MTLPixelFormatRGBA8Unorm;
		case GSTexture::Format::ColorHQ:      return MTLPixelFormatRGB10A2Unorm;
		case GSTexture::Format::ColorHDR:     return MTLPixelFormatRGBA16Float;
		case GSTexture::Format::ColorClip:    return MTLPixelFormatRGBA16Unorm;
		case GSTexture::Format::DepthStencil: return MTLPixelFormatDepth32Float_Stencil8;
		case GSTexture::Format::Invalid:      return MTLPixelFormatInvalid;
		case GSTexture::Format::BC1:          return MTLPixelFormatBC1_RGBA;
		case GSTexture::Format::BC2:          return MTLPixelFormatBC2_RGBA;
		case GSTexture::Format::BC3:          return MTLPixelFormatBC3_RGBA;
		case GSTexture::Format::BC7:          return MTLPixelFormatBC7_RGBAUnorm;
	}
}

GSTexture* GSDeviceMTL::CreateSurface(GSTexture::Usage usage, int width, int height, int levels, GSTexture::Format format)
{ @autoreleasepool {
	pxAssert(GSTexture::ValidateUsageAndFormat(usage, format));

	MTLPixelFormat fmt = ConvertPixelFormat(format);
	pxAssertRel(format != GSTexture::Format::Invalid, "Can't create surface of this format!");

	MTLTextureDescriptor* desc = [MTLTextureDescriptor
		texture2DDescriptorWithPixelFormat:fmt
		                             width:width
		                            height:height
		                         mipmapped:levels > 1];

	if (levels > 1)
		[desc setMipmapLevelCount:levels];

	[desc setStorageMode:MTLStorageModePrivate];

	MTLTextureUsage mtl_usage = MTLTextureUsageShaderRead;
	bool needs_rov_tex = false;

	if (GSTexture::IsRenderTargetOrDepthStencil(usage))
		mtl_usage |= MTLTextureUsageRenderTarget;

	if ((usage & GSTexture::FeedbackTarget) == GSTexture::FeedbackTarget)
	{
		if (m_dev.features.slow_color_compression)
			mtl_usage |= MTLTextureUsagePixelFormatView; // Force color compression off by including PixelFormatView
	}

	if (usage == GSTexture::ShaderWriteTarget && format == GSTexture::Format::Color && m_dev.features.rov_requires_r32)
	{
		// Need to make an R32 view of the texture for ROV
		mtl_usage |= MTLTextureUsagePixelFormatView;
		needs_rov_tex = true;
	}

	if (GSTexture::IsShaderWrite(usage))
		mtl_usage |= MTLTextureUsageShaderWrite;

	[desc setUsage:mtl_usage];

	MRCOwned<id<MTLTexture>> tex = MRCTransfer([m_dev.dev newTextureWithDescriptor:desc]);
	if (tex)
	{
		MRCOwned<id<MTLTexture>> rov_tex = needs_rov_tex ? MRCTransfer([tex newTextureViewWithPixelFormat:MTLPixelFormatR32Uint]) : nil;
		GSTextureMTL* t = new GSTextureMTL(this, tex, rov_tex, usage, format);
		if (GSTexture::IsRenderTarget(usage))
		{
			ClearRenderTarget(t, 0);
		}
		else if (GSTexture::IsDepthStencil(usage))
		{
			ClearDepth(t, 0.0f);
		}
		return t;
	}
	else
	{
		return nullptr;
	}
}}

void GSDeviceMTL::DoMerge(GSTexture* sTex[3], GSVector4* sRect, GSTexture* dTex, GSVector4* dRect, const GSRegPMODE& PMODE, const GSRegEXTBUF& EXTBUF, u32 c, const Filter filter)
{ @autoreleasepool {
	id<MTLCommandBuffer> cmdbuf = GetRenderCmdBuf();
	GSScopedDebugGroupMTL dbg(cmdbuf, @"DoMerge");

	GSVector4 full_r(0.0f, 0.0f, 1.0f, 1.0f);
	bool feedback_write_2 = PMODE.EN2 && sTex[2] != nullptr && EXTBUF.FBIN == 1;
	bool feedback_write_1 = PMODE.EN1 && sTex[2] != nullptr && EXTBUF.FBIN == 0;
	bool feedback_write_2_but_blend_bg = feedback_write_2 && PMODE.SLBG == 1;

	ClearRenderTarget(dTex, c);

	const GSVector4 unorm_c = GSVector4::unorm8(c);
	vector_float4 cb_c = { unorm_c.r, unorm_c.g, unorm_c.b, unorm_c.a };
	GSMTLConvertPSUniform cb_yuv = {};
	cb_yuv.emoda = EXTBUF.EMODA;
	cb_yuv.emodc = EXTBUF.EMODC;

	if (sTex[1] && (PMODE.SLBG == 0 || feedback_write_2_but_blend_bg))
	{
		// 2nd output is enabled and selected. Copy it to destination so we can blend it with 1st output
		// Note: value outside of dRect must contains the background color (c)
		StretchRect(sTex[1], sRect[1], dTex, dRect[1], ShaderConvert::COPY, filter);
	}

	// Save 2nd output
	if (feedback_write_2) // FIXME I'm not sure dRect[1] is always correct
		DoStretchRect(dTex, full_r, sTex[2], dRect[1], GetConvertPipeline(ShaderConvert::YUV), filter, LoadAction::DontCareIfFull, &cb_yuv, sizeof(cb_yuv));

	if (feedback_write_2_but_blend_bg)
		ClearRenderTarget(dTex, c);

	if (sTex[0])
	{
		int idx = (PMODE.AMOD << 1) | PMODE.MMOD;
		id<MTLRenderPipelineState> pipeline = m_merge_pipeline[idx];

		// 1st output is enabled. It must be blended
		if (PMODE.MMOD == 1)
		{
			// Blend with a constant alpha
			DoStretchRect(sTex[0], sRect[0], dTex, dRect[0], pipeline, filter, LoadAction::Load, &cb_c, sizeof(cb_c));
		}
		else
		{
			// Blend with 2 * input alpha
			DoStretchRect(sTex[0], sRect[0], dTex, dRect[0], pipeline, filter, LoadAction::Load, nullptr, 0);
		}
	}

	if (feedback_write_1) // FIXME I'm not sure dRect[0] is always correct
		StretchRect(dTex, full_r, sTex[2], dRect[0], ShaderConvert::YUV, filter);
}}

void GSDeviceMTL::DoInterlace(GSTexture* sTex, const GSVector4& sRect, GSTexture* dTex, const GSVector4& dRect, ShaderInterlace shader, Filter filter, const InterlaceConstantBuffer& cb)
{ @autoreleasepool {
	id<MTLCommandBuffer> cmdbuf = GetRenderCmdBuf();
	GSScopedDebugGroupMTL dbg(cmdbuf, @"DoInterlace");

	const bool can_discard = shader == ShaderInterlace::WEAVE || shader == ShaderInterlace::MAD_BUFFER;
	DoStretchRect(sTex, sRect, dTex, dRect, m_interlace_pipeline[static_cast<int>(shader)], filter,
		!can_discard ? LoadAction::DontCareIfFull : LoadAction::Load, &cb, sizeof(cb));
}}

void GSDeviceMTL::DoFXAA(GSTexture* sTex, GSTexture* dTex)
{
	BeginRenderPass(@"FXAA", dTex, MTLLoadActionDontCare, nullptr, MTLLoadActionDontCare);
	RenderCopy(sTex, m_fxaa_pipeline, GSVector4i(0, 0, dTex->GetSize().x, dTex->GetSize().y));
}

void GSDeviceMTL::DoShadeBoost(GSTexture* sTex, GSTexture* dTex, const float params[4])
{
	BeginRenderPass(@"ShadeBoost", dTex, MTLLoadActionDontCare, nullptr, MTLLoadActionDontCare);
	[m_current_render.encoder setFragmentBytes:params
	                                    length:sizeof(float) * 4
	                                   atIndex:GSMTLBufferIndexUniforms];
	RenderCopy(sTex, m_shadeboost_pipeline, GSVector4i(0, 0, dTex->GetSize().x, dTex->GetSize().y));
}

bool GSDeviceMTL::DoCAS(GSTexture* sTex, GSTexture* dTex, bool sharpen_only, const std::array<u32, NUM_CAS_CONSTANTS>& constants)
{ @autoreleasepool {
	g_perfmon.Put(GSPerfMon::TextureCopies, 1);

	static constexpr int threadGroupWorkRegionDim = 16;
	const int dispatchX = (dTex->GetWidth() + (threadGroupWorkRegionDim - 1)) / threadGroupWorkRegionDim;
	const int dispatchY = (dTex->GetHeight() + (threadGroupWorkRegionDim - 1)) / threadGroupWorkRegionDim;
	static_assert(sizeof(constants) == sizeof(GSMTLCASPSUniform));

	EndRenderPass();
	id<MTLComputeCommandEncoder> enc = [GetRenderCmdBuf() computeCommandEncoder];
	[enc setLabel:@"CAS"];
	[enc setComputePipelineState:m_cas_pipeline[sharpen_only]];
	[enc setTexture:static_cast<GSTextureMTL*>(sTex)->GetTexture() atIndex:0];
	[enc setTexture:static_cast<GSTextureMTL*>(dTex)->GetTexture() atIndex:1];
	[enc setBytes:&constants length:sizeof(constants) atIndex:GSMTLBufferIndexUniforms];
	[enc dispatchThreadgroups:MTLSizeMake(dispatchX, dispatchY, 1)
	    threadsPerThreadgroup:MTLSizeMake(64, 1, 1)];
	[enc endEncoding];
	return true;
}}

#if PCSX2_HAS_METALFX
bool GSDeviceMTL::EnsureMetalFXSpatial(GSTexture* sTex, GSTexture* dTex)
{ @autoreleasepool {
	id<MTLTexture> src = static_cast<GSTextureMTL*>(sTex)->GetTexture();
	id<MTLTexture> dst = static_cast<GSTextureMTL*>(dTex)->GetTexture();
	const int in_w = sTex->GetWidth(), in_h = sTex->GetHeight();
	const int out_w = dTex->GetWidth(), out_h = dTex->GetHeight();
	const MTLPixelFormat in_fmt = [src pixelFormat];
	const MTLPixelFormat out_fmt = [dst pixelFormat];

	// Reuse the cached scaler if the size/format key is unchanged (creation is expensive).
	if (m_mfx_spatial && m_mfx_in_w == in_w && m_mfx_in_h == in_h &&
	    m_mfx_out_w == out_w && m_mfx_out_h == out_h &&
	    m_mfx_in_fmt == in_fmt && m_mfx_out_fmt == out_fmt)
	{
		return true;
	}

	MTLFXSpatialScalerDescriptor* desc = [[MTLFXSpatialScalerDescriptor alloc] init];
	desc.inputWidth = in_w;
	desc.inputHeight = in_h;
	desc.outputWidth = out_w;
	desc.outputHeight = out_h;
	desc.colorTextureFormat = in_fmt;
	desc.outputTextureFormat = out_fmt;
	// GS display output is already tonemapped/sRGB-ish, so treat it as perceptual.
	desc.colorProcessingMode = MTLFXSpatialScalerColorProcessingModePerceptual;

	m_mfx_spatial = MRCTransfer([desc newSpatialScalerWithDevice:m_dev.dev]);
	[desc release];
	if (!m_mfx_spatial)
	{
		Console.Error("MetalFX: Failed to create spatial scaler.");
		return false;
	}

	m_mfx_in_w = in_w; m_mfx_in_h = in_h;
	m_mfx_out_w = out_w; m_mfx_out_h = out_h;
	m_mfx_in_fmt = in_fmt; m_mfx_out_fmt = out_fmt;
	return true;
}}
#else
bool GSDeviceMTL::EnsureMetalFXSpatial(GSTexture* sTex, GSTexture* dTex)
{
	// Statically compiled out on the iOS Simulator (PCSX2_HAS_METALFX=0).
	(void)sTex; (void)dTex;
	return false;
}
#endif

bool GSDeviceMTL::DoMetalFXSpatial(GSTexture* sTex, GSTexture* dTex)
{
#if PCSX2_HAS_METALFX
	@autoreleasepool {
		if (@available(macOS 13.0, iOS 16.0, *))
		{
			if (!EnsureMetalFXSpatial(sTex, dTex))
				return false;

			g_perfmon.Put(GSPerfMon::TextureCopies, 1);
			EndRenderPass(); // MetalFX manages its own encoder; must not be inside one.
			[m_mfx_spatial setColorTexture:static_cast<GSTextureMTL*>(sTex)->GetTexture()];
			[m_mfx_spatial setOutputTexture:static_cast<GSTextureMTL*>(dTex)->GetTexture()];
			[m_mfx_spatial encodeToCommandBuffer:GetRenderCmdBuf()];
			return true;
		}
		return false;
	}
#else
	// Statically compiled out on the iOS Simulator (PCSX2_HAS_METALFX=0).
	(void)sTex; (void)dTex;
	return false;
#endif
}

MRCOwned<id<MTLFunction>> GSDeviceMTL::LoadShader(NSString* name)
{
	NSError* err = nil;
	MRCOwned<id<MTLFunction>> fn = MRCTransfer([m_dev.shaders newFunctionWithName:name constantValues:m_fn_constants error:&err]);
	if (err) [[unlikely]]
	{
		NSString* msg = [NSString stringWithFormat:@"Failed to load shader %@: %@", name, [err localizedDescription]];
		Console.Error("%s", [msg UTF8String]);
		pxFailRel([msg UTF8String]);
	}
	return fn;
}

MRCOwned<id<MTLRenderPipelineState>> GSDeviceMTL::MakePipeline(MTLRenderPipelineDescriptor* desc, id<MTLFunction> vertex, id<MTLFunction> fragment, NSString* name)
{
	const GSShaderCompileIndicator::CompileTimer compile_timer;
	[desc setLabel:name];
	[desc setVertexFunction:vertex];
	[desc setFragmentFunction:fragment];
	NSError* err;
	MRCOwned<id<MTLRenderPipelineState>> res = MRCTransfer([m_dev.dev newRenderPipelineStateWithDescriptor:desc error:&err]);
	if (err) [[unlikely]]
	{
		NSString* msg = [NSString stringWithFormat:@"Failed to create pipeline %@: %@", name, [err localizedDescription]];
		Console.Error("%s", [msg UTF8String]);
		pxFailRel([msg UTF8String]);
	}
	return res;
}

MRCOwned<id<MTLComputePipelineState>> GSDeviceMTL::MakeComputePipeline(id<MTLFunction> compute, NSString* name)
{
	const GSShaderCompileIndicator::CompileTimer compile_timer;
	MRCOwned<MTLComputePipelineDescriptor*> desc = MRCTransfer([MTLComputePipelineDescriptor new]);
	[desc setLabel:name];
	[desc setComputeFunction:compute];
	NSError* err;
	MRCOwned<id<MTLComputePipelineState>> res = MRCTransfer([m_dev.dev
		newComputePipelineStateWithDescriptor:desc
		                              options:0
		                           reflection:nil
		                                error:&err]);
	if (err) [[unlikely]]
	{
		NSString* msg = [NSString stringWithFormat:@"Failed to create pipeline %@: %@", name, [err localizedDescription]];
		Console.Error("%s", [msg UTF8String]);
		pxFailRel([msg UTF8String]);
	}
	return res;
}

static void applyAttribute(MTLVertexDescriptor* desc, NSUInteger idx, MTLVertexFormat fmt, NSUInteger offset, NSUInteger buffer_index)
{
	MTLVertexAttributeDescriptor* attrs = desc.attributes[idx];
	attrs.format = fmt;
	attrs.offset = offset;
	attrs.bufferIndex = buffer_index;
}

static void setFnConstantB(MTLFunctionConstantValues* fc, bool value, GSMTLFnConstants constant)
{
	[fc setConstantValue:&value type:MTLDataTypeBool atIndex:constant];
}

static void setFnConstantI(MTLFunctionConstantValues* fc, unsigned int value, GSMTLFnConstants constant)
{
	[fc setConstantValue:&value type:MTLDataTypeUInt atIndex:constant];
}

template <typename T, std::enable_if_t<std::is_enum_v<T>, bool> = true>
static void setFnConstantI(MTLFunctionConstantValues* fc, T value, GSMTLFnConstants constant)
{
	setFnConstantI(fc, static_cast<std::underlying_type_t<T>>(value), constant);
}

template <typename Fn>
static void OnMainThread(Fn&& fn)
{
	if ([NSThread isMainThread])
		fn();
	else
		dispatch_sync(dispatch_get_main_queue(), fn);
}

RenderAPI GSDeviceMTL::GetRenderAPI() const
{
	return RenderAPI::Metal;
}

bool GSDeviceMTL::HasSurface()  const { return static_cast<bool>(m_layer);}

void GSDeviceMTL::AttachSurfaceOnMainThread()
{
	pxAssert([NSThread isMainThread]);
	m_view = MRCRetain((__bridge GSMTLView*)m_window_info.window_handle);
#if PCSX2_MTL_USES_UIVIEW
	// On iOS the UIView already owns its backing CAMetalLayer (created by the host
	// view). Reuse that existing layer instead of allocating a detached one that
	// would never be attached to the view hierarchy.
	CALayer* view_layer = [m_view layer];
	if (![view_layer isKindOfClass:[CAMetalLayer class]])
	{
		return;
	}

	m_layer = MRCRetain((CAMetalLayer*)view_layer);
	[m_layer setDrawableSize:CGSizeMake(m_window_info.surface_width, m_window_info.surface_height)];
	[m_layer setDevice:m_dev.dev];
#else
	m_layer = MRCRetain([CAMetalLayer layer]);
	[m_layer setDrawableSize:CGSizeMake(m_window_info.surface_width, m_window_info.surface_height)];
	[m_layer setDevice:m_dev.dev];
	[m_view setWantsLayer:YES];
	[m_view setLayer:m_layer];
#endif
}

void GSDeviceMTL::DetachSurfaceOnMainThread()
{
	pxAssert([NSThread isMainThread]);
#if !PCSX2_MTL_USES_UIVIEW
	[m_view setLayer:nullptr];
	[m_view setWantsLayer:NO];
#endif
	m_view = nullptr;
	m_layer = nullptr;
}

// Metal is fun and won't let you use newBufferWithBytes for private buffers
static MRCOwned<id<MTLBuffer>> CreatePrivateBufferWithContent(
	id<MTLDevice> dev, id<MTLCommandBuffer> cb,
	MTLResourceOptions options, NSUInteger length,
	std::function<void(void*)> fill)
{
	MRCOwned<id<MTLBuffer>> tmp = MRCTransfer([dev newBufferWithLength:length options:MTLResourceStorageModeShared]);
	MRCOwned<id<MTLBuffer>> actual = MRCTransfer([dev newBufferWithLength:length options:options|MTLResourceStorageModePrivate]);
	fill([tmp contents]);
	id<MTLBlitCommandEncoder> blit = [cb blitCommandEncoder];
	[blit copyFromBuffer:tmp sourceOffset:0 toBuffer:actual destinationOffset:0 size:length];
	[blit endEncoding];
	return actual;
}

static MRCOwned<id<MTLSamplerState>> CreateSampler(id<MTLDevice> dev, GSHWDrawConfig::SamplerSelector sel)
{
	MRCOwned<MTLSamplerDescriptor*> sdesc = MRCTransfer([MTLSamplerDescriptor new]);
	const char* minname = sel.biln ? "Ln" : "Pt";
	const char* magname = minname;
	[sdesc setMinFilter:sel.biln ? MTLSamplerMinMagFilterLinear : MTLSamplerMinMagFilterNearest];
	[sdesc setMagFilter:sel.biln ? MTLSamplerMinMagFilterLinear : MTLSamplerMinMagFilterNearest];
	switch (static_cast<GS_MIN_FILTER>(sel.triln))
	{
		case GS_MIN_FILTER::Nearest:
		case GS_MIN_FILTER::Linear:
			[sdesc setMipFilter:MTLSamplerMipFilterNotMipmapped];
			break;
		case GS_MIN_FILTER::Nearest_Mipmap_Nearest:
			minname = "PtPt";
			[sdesc setMinFilter:MTLSamplerMinMagFilterNearest];
			[sdesc setMipFilter:MTLSamplerMipFilterNearest];
			break;
		case GS_MIN_FILTER::Nearest_Mipmap_Linear:
			minname = "PtLn";
			[sdesc setMinFilter:MTLSamplerMinMagFilterNearest];
			[sdesc setMipFilter:MTLSamplerMipFilterLinear];
			break;
		case GS_MIN_FILTER::Linear_Mipmap_Nearest:
			minname = "LnPt";
			[sdesc setMinFilter:MTLSamplerMinMagFilterLinear];
			[sdesc setMipFilter:MTLSamplerMipFilterNearest];
			break;
		case GS_MIN_FILTER::Linear_Mipmap_Linear:
			minname = "LnLn";
			[sdesc setMinFilter:MTLSamplerMinMagFilterLinear];
			[sdesc setMipFilter:MTLSamplerMipFilterLinear];
			break;
	}

	const char* taudesc = sel.tau ? "Repeat" : "Clamp";
	const char* tavdesc = sel.tav == sel.tau ? "" : sel.tav ? "Repeat" : "Clamp";
	[sdesc setSAddressMode:sel.tau ? MTLSamplerAddressModeRepeat : MTLSamplerAddressModeClampToEdge];
	[sdesc setTAddressMode:sel.tav ? MTLSamplerAddressModeRepeat : MTLSamplerAddressModeClampToEdge];
	[sdesc setRAddressMode:MTLSamplerAddressModeClampToEdge];

	[sdesc setMaxAnisotropy:1];
	bool clampLOD = sel.lodclamp || !sel.UseMipmapFiltering();
	const char* clampdesc = clampLOD ? " LODClamp" : "";
	[sdesc setLodMaxClamp:clampLOD ? 0.25f : FLT_MAX];

	[sdesc setLabel:[NSString stringWithFormat:@"%s%s %s%s%s", taudesc, tavdesc, magname, minname, clampdesc]];
	MRCOwned<id<MTLSamplerState>> ret = MRCTransfer([dev newSamplerStateWithDescriptor:sdesc]);
	pxAssertRel(ret, "Failed to create sampler!");
	return ret;
}

static bool getDepthFeedback(const GSMTLDevice& dev, bool fbfetch)
{
	switch (GSConfig.DepthFeedbackMode)
	{
		case GSDepthFeedbackMode::Auto:
			return dev.features.depth_feedback;
		case GSDepthFeedbackMode::Depth:
			// Depth feedback + FBFetch not supported
			return !fbfetch;
		default:
			return false;
	}
}

// Some shaders are only used by methods on MTLDevice, which currently use separately-compiled shaders
static bool ConvertShaderNotNeeded(ShaderConvert shader)
{
	switch (shader)
	{
		case ShaderConvert::DATM_0:
		case ShaderConvert::DATM_1:
		case ShaderConvert::DATM_0_RTA_CORRECTION:
		case ShaderConvert::DATM_1_RTA_CORRECTION:
		case ShaderConvert::CLUT_4:
		case ShaderConvert::CLUT_8:
		case ShaderConvert::COLCLIP_INIT:
		case ShaderConvert::COLCLIP_RESOLVE:
			return true;
		default:
			return false;
	}
}

bool GSDeviceMTL::Create(GSVSyncMode vsync_mode, bool allow_present_throttle)
{ @autoreleasepool {
	if (!GSDevice::Create(vsync_mode, allow_present_throttle))
		return false;

	NSString* ns_adapter_name = [NSString stringWithUTF8String:GSConfig.Adapter.c_str()];
	// No device enumeration below iOS 18 — fall through to the default device.
	if (@available(macOS 10.11, iOS 18.0, *))
	{
		auto devs = MRCTransfer(MTLCopyAllDevices());
		for (id<MTLDevice> dev in devs.Get())
		{
			if ([[dev name] isEqualToString:ns_adapter_name])
				m_dev = GSMTLDevice(MRCRetain(dev));
		}
	}
	if (!m_dev.dev)
	{
		if (GSConfig.Adapter == GetDefaultAdapter())
			Console.WriteLn("Metal: Using default adapter");
		else if (!GSConfig.Adapter.empty())
			Console.Warning("Metal: Couldn't find adapter %s, using default", GSConfig.Adapter.c_str());
		m_dev = GSMTLDevice(MRCTransfer(MTLCreateSystemDefaultDevice()));
		if (!m_dev.dev)
			Host::ReportErrorAsync(TRANSLATE_SV("GSDeviceMTL", "No Metal Devices Available"), TRANSLATE_SV("GSDeviceMTL", "No Metal-supporting GPUs were found.  PCSX2 requires a Metal GPU (available on all Macs from 2012 onwards)."));
	}

	m_name = [[m_dev.dev name] UTF8String];
	m_queue = MRCTransfer([m_dev.dev newCommandQueue]);

	m_pass_desc = MRCTransfer([MTLRenderPassDescriptor new]);
	[m_pass_desc colorAttachments][0].loadAction = MTLLoadActionClear;
	[m_pass_desc colorAttachments][0].clearColor = MTLClearColorMake(0, 0, 0, 0);
	[m_pass_desc colorAttachments][0].storeAction = MTLStoreActionStore;

	if (char* env = getenv("MTL_USE_PRESENT_DRAWABLE"))
		m_use_present_drawable = static_cast<UsePresentDrawable>(atoi(env));
	else if (@available(macOS 13.0, *))
		m_use_present_drawable = UsePresentDrawable::Always;
	else // Before Ventura, presentDrawable acts like vsync is on when windowed
		m_use_present_drawable = UsePresentDrawable::IfVsync;

	m_capture_start_frame = 0;
	if (char* env = getenv("MTL_CAPTURE"))
	{
		m_capture_start_frame = atoi(env);
	}
	if (m_capture_start_frame)
	{
		Console.WriteLn("Metal will capture frame %u", m_capture_start_frame);
	}

	if (m_dev.IsOk() && m_queue)
	{
		// This is a little less than ideal, pinging back and forward between threads, but we don't really
		// have any other option, because Qt uses a blocking queued connection for window acquire.
		if (!AcquireWindow(true))
			return false;

		OnMainThread([this]
		{
			AttachSurfaceOnMainThread();
		});

		// Metal does not support mailbox.
		m_vsync_mode = (m_vsync_mode == GSVSyncMode::Mailbox) ? GSVSyncMode::FIFO : m_vsync_mode;
#if !TARGET_OS_IPHONE
		[m_layer setDisplaySyncEnabled:m_vsync_mode == GSVSyncMode::FIFO];
#endif
	}
	else
	{
		return false;
	}

	MTLPixelFormat layer_px_fmt = [m_layer pixelFormat];

	m_features.broken_point_sampler = false;
	m_features.vs_expand = !GSConfig.DisableVertexShaderExpand;
	m_features.primitive_id = m_dev.features.primid;
	// Apple GPUs cannot synchronize fragment writes to fragment reads within a
	// render pass. Let the GS renderer choose its non-barrier feedback fallback.
#if TARGET_OS_IPHONE
	m_features.texture_barrier = false;
#else
	m_features.texture_barrier = true;
#endif
	m_features.multidraw_fb_copy = false;
	m_features.provoking_vertex_last = false;
	m_features.point_expand = true;
	m_features.line_expand = false;
	m_features.prefer_new_textures = true;
	m_features.dxt_textures = true;
	m_features.bptc_textures = true;
	m_features.framebuffer_fetch = m_dev.features.framebuffer_fetch && !GSConfig.DisableFramebufferFetch;
	m_features.stencil_buffer = true;
	m_features.cas_sharpening = true;
	m_features.test_and_sample_depth = true;
	m_features.depth_feedback = getDepthFeedback(m_dev, m_features.framebuffer_fetch);
	m_features.aa1 = GSConfig.HWAA1 && m_features.vs_expand;
	// Apple GPUs miscompare depth written from the shader (the PS2 32-bit Z floor) against
	// the fixed-function interpolation a later read-only pass tests with, so a GEQUAL retest
	// of the same geometry drops out along shared triangle edges and the layer underneath
	// shows through as pinpoints -- God of War II's Athena statue, and dark walls in Black.
	// Skipping the floor also drops [[depth(less)]] output, restoring early-ZS on a TBDR.
	// See the matching gate in GSDeviceVK::CheckFeatures for the measurements.
	m_features.no_ps2_z_quantization = GSConfig.DisablePS2DepthQuantization || m_dev.features.apple_gpu;
	// MetalFX spatial upscaler: macOS 13+ / iOS 16+ device. The supportsDevice:
	// probe returns NO on devices whose GPU lacks the hardware. On the iOS Simulator
	// the MetalFX framework is absent at compile time (PCSX2_HAS_METALFX=0), so the
	// feature is statically disabled here -- m_features.metalfx_spatial keeps its
	// default false value and the upscaler UI will report unavailable.
#if PCSX2_HAS_METALFX
	if (@available(macOS 13.0, iOS 16.0, *))
		m_features.metalfx_spatial = [MTLFXSpatialScalerDescriptor supportsDevice:m_dev.dev];
#endif
	m_features.rov = m_dev.features.rov && !m_features.framebuffer_fetch;
	m_max_texture_size = m_dev.features.max_texsize;

	// Init metal stuff
	m_fn_constants = MRCTransfer([MTLFunctionConstantValues new]);
	setFnConstantB(m_fn_constants, m_features.framebuffer_fetch,    GSMTLConstantIndex_FRAMEBUFFER_FETCH);
	setFnConstantB(m_fn_constants, m_features.depth_feedback,       GSMTLConstantIndex_DEPTH_FEEDBACK);
	setFnConstantB(m_fn_constants, m_dev.features.rov_requires_r32, GSMTLConstantIndex_ROV_NEEDS_R32);

	m_draw_sync_fence = MRCTransfer([m_dev.dev newFence]);
	[m_draw_sync_fence setLabel:@"Draw Sync Fence"];
	m_spin_fence = MRCTransfer([m_dev.dev newFence]);
	[m_spin_fence setLabel:@"Spin Fence"];
	constexpr MTLResourceOptions spin_opts = MTLResourceStorageModePrivate | MTLResourceHazardTrackingModeUntracked;
	m_spin_buffer = MRCTransfer([m_dev.dev newBufferWithLength:4 options:spin_opts]);
	[m_spin_buffer setLabel:@"Spin Buffer"];
	id<MTLCommandBuffer> initCommands = [m_queue commandBuffer];
	id<MTLBlitCommandEncoder> clearSpinBuffer = [initCommands blitCommandEncoder];
	[clearSpinBuffer fillBuffer:m_spin_buffer range:NSMakeRange(0, 4) value:0];
	[clearSpinBuffer updateFence:m_spin_fence];
	[clearSpinBuffer endEncoding];
	m_spin_pipeline = MakeComputePipeline(LoadShader(@"waste_time"), @"waste_time");

	for (int sharpen_only = 0; sharpen_only < 2; sharpen_only++)
	{
		setFnConstantB(m_fn_constants, sharpen_only, GSMTLConstantIndex_CAS_SHARPEN_ONLY);
		NSString* shader = m_dev.features.has_fast_half ? @"CASHalf" : @"CASFloat";
		m_cas_pipeline[sharpen_only] = MakeComputePipeline(LoadShader(shader), sharpen_only ? @"CAS Sharpen" : @"CAS Upscale");
	}

	m_expand_index_buffer = CreatePrivateBufferWithContent(m_dev.dev, initCommands, MTLResourceHazardTrackingModeUntracked, EXPAND_BUFFER_SIZE, GenerateExpansionIndexBuffer);
	[m_expand_index_buffer setLabel:@"Point/Sprite Expand Indices"];

	m_hw_vertex = MRCTransfer([MTLVertexDescriptor new]);
	[[[m_hw_vertex layouts] objectAtIndexedSubscript:GSMTLBufferIndexHWVertices] setStride:sizeof(GSVertex)];
	applyAttribute(m_hw_vertex, GSMTLAttributeIndexST, MTLVertexFormatFloat2,           offsetof(GSVertex, ST),      GSMTLBufferIndexHWVertices);
	applyAttribute(m_hw_vertex, GSMTLAttributeIndexC,  MTLVertexFormatUChar4,           offsetof(GSVertex, RGBAQ.R), GSMTLBufferIndexHWVertices);
	applyAttribute(m_hw_vertex, GSMTLAttributeIndexQ,  MTLVertexFormatFloat,            offsetof(GSVertex, RGBAQ.Q), GSMTLBufferIndexHWVertices);
	applyAttribute(m_hw_vertex, GSMTLAttributeIndexXY, MTLVertexFormatUShort2,          offsetof(GSVertex, XYZ.X),   GSMTLBufferIndexHWVertices);
	applyAttribute(m_hw_vertex, GSMTLAttributeIndexZ,  MTLVertexFormatUInt,             offsetof(GSVertex, XYZ.Z),   GSMTLBufferIndexHWVertices);
	applyAttribute(m_hw_vertex, GSMTLAttributeIndexUV, MTLVertexFormatUShort2,          offsetof(GSVertex, UV),      GSMTLBufferIndexHWVertices);
	applyAttribute(m_hw_vertex, GSMTLAttributeIndexF,  MTLVertexFormatUChar4Normalized, offsetof(GSVertex, FOG),     GSMTLBufferIndexHWVertices);

	for (size_t i = 0; i < std::size(m_render_pass_desc); i++)
	{
		const bool depth   = i & 2;
		const bool stencil = i & 4;
		const bool rt1     = i & 8;
		if (rt1 && m_features.depth_feedback)
			continue;
		if (rt1 && !depth)
			continue;
		if (stencil && m_features.framebuffer_fetch)
			continue;
		auto desc = MRCTransfer([MTLRenderPassDescriptor new]);
		[[desc   depthAttachment] setStoreAction:MTLStoreActionStore];
		[[desc stencilAttachment] setStoreAction:MTLStoreActionStore];
		if (rt1)
		{
			MTLRenderPassColorAttachmentDescriptor* color1 = [[desc colorAttachments] objectAtIndexedSubscript:1];
			[color1 setStoreAction:MTLStoreActionDontCare];
			[color1 setLoadAction:m_features.framebuffer_fetch ? MTLLoadActionClear : MTLLoadActionLoad];
		}
		m_render_pass_desc[i] = desc;
	}
	if (m_features.rov)
	{
		m_full_rov_render_pass_desc = MRCTransfer([MTLRenderPassDescriptor new]);
		[m_full_rov_render_pass_desc setDefaultRasterSampleCount:1];
		if (m_dev.features.rov_requires_rt)
		{
			MTLRenderPassColorAttachmentDescriptor* color0 = [[m_full_rov_render_pass_desc colorAttachments] objectAtIndexedSubscript:0];
			[color0 setLoadAction:MTLLoadActionDontCare];
			[color0 setStoreAction:MTLStoreActionDontCare];
		}
	}

	// Init samplers
	m_sampler_hw[SamplerSelector::Linear().key] = CreateSampler(m_dev.dev, SamplerSelector::Linear());
	m_sampler_hw[SamplerSelector::Point().key] = CreateSampler(m_dev.dev, SamplerSelector::Point());

	// Init depth stencil states
	MTLDepthStencilDescriptor* dssdesc = [[MTLDepthStencilDescriptor new] autorelease];
	MTLStencilDescriptor* stencildesc = [[MTLStencilDescriptor new] autorelease];
	stencildesc.stencilCompareFunction = MTLCompareFunctionAlways;
	stencildesc.depthFailureOperation = MTLStencilOperationKeep;
	stencildesc.stencilFailureOperation = MTLStencilOperationKeep;
	stencildesc.depthStencilPassOperation = MTLStencilOperationReplace;
	dssdesc.frontFaceStencil = stencildesc;
	dssdesc.backFaceStencil = stencildesc;
	[dssdesc setLabel:@"Stencil Write"];
	m_dss_stencil_write = MRCTransfer([m_dev.dev newDepthStencilStateWithDescriptor:dssdesc]);
	dssdesc.frontFaceStencil.depthStencilPassOperation = MTLStencilOperationZero;
	dssdesc.backFaceStencil.depthStencilPassOperation = MTLStencilOperationZero;
	[dssdesc setLabel:@"Stencil Zero"];
	m_dss_stencil_zero = MRCTransfer([m_dev.dev newDepthStencilStateWithDescriptor:dssdesc]);
	stencildesc.stencilCompareFunction = MTLCompareFunctionEqual;
	stencildesc.readMask = 1;
	stencildesc.writeMask = 1;
	for (size_t i = 0; i < std::size(m_dss_hw); i++)
	{
		GSHWDrawConfig::DepthStencilSelector sel;
		sel.key = i;
		if (sel.date)
		{
			if (sel.date_one)
				stencildesc.depthStencilPassOperation = MTLStencilOperationZero;
			else
				stencildesc.depthStencilPassOperation = MTLStencilOperationKeep;
			dssdesc.frontFaceStencil = stencildesc;
			dssdesc.backFaceStencil = stencildesc;
		}
		else
		{
			dssdesc.frontFaceStencil = nil;
			dssdesc.backFaceStencil = nil;
		}
		dssdesc.depthWriteEnabled = sel.zwe ? YES : NO;
		static constexpr MTLCompareFunction ztst[] =
		{
			MTLCompareFunctionNever,
			MTLCompareFunctionAlways,
			MTLCompareFunctionGreaterEqual,
			MTLCompareFunctionGreater,
		};
		static constexpr const char* ztstname[] =
		{
			"DepthNever",
			"DepthAlways",
			"DepthGEq",
			"DepthEq",
		};
		const char* datedesc = sel.date ? (sel.date_one ? " DATE_ONE" : " DATE") : "";
		const char* zwedesc = sel.zwe ? " ZWE" : "";
		dssdesc.depthCompareFunction = ztst[sel.ztst];
		[dssdesc setLabel:[NSString stringWithFormat:@"%s%s%s", ztstname[sel.ztst], zwedesc, datedesc]];
		m_dss_hw[i] = MRCTransfer([m_dev.dev newDepthStencilStateWithDescriptor:dssdesc]);
	}

	// Init HW Vertex Shaders
	for (size_t i = 0; i < std::size(m_hw_vs); i++)
	{
		VSSelector sel;
		sel.key = i;
		if (sel.point_size && sel.expand != GSShader::VSExpand::None)
			continue;
		setFnConstantB(m_fn_constants, sel.fst,        GSMTLConstantIndex_FST);
		setFnConstantB(m_fn_constants, sel.iip,        GSMTLConstantIndex_IIP);
		setFnConstantB(m_fn_constants, sel.point_size, GSMTLConstantIndex_VS_POINT_SIZE);
		setFnConstantI(m_fn_constants, sel.expand,     GSMTLConstantIndex_VS_EXPAND_TYPE);
		m_hw_vs[i] = LoadShader(sel.expand == GSShader::VSExpand::None ? @"vs_main" : @"vs_main_expand");
	}

	// Init pipelines
	auto vs_convert = LoadShader(@"vs_convert");
	auto fs_triangle = LoadShader(@"fs_triangle");
	auto ps_copy = LoadShader(@"ps_copy");
	auto ps_copy_rta_correct = LoadShader(@"ps_rta_correction");
	auto pdesc = [[MTLRenderPipelineDescriptor new] autorelease];
	// FS Triangle Pipelines
	pdesc.colorAttachments[0].pixelFormat = ConvertPixelFormat(GSTexture::Format::Color);
	m_colclip_resolve_pipeline = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_colclip_resolve"), @"ColorClip Resolve");
	m_fxaa_pipeline = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_fxaa"), @"fxaa");
	m_shadeboost_pipeline = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_shadeboost"), @"shadeboost");
	m_clut_pipeline[0] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_convert_clut_4"), @"4-bit CLUT Update");
	m_clut_pipeline[1] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_convert_clut_8"), @"8-bit CLUT Update");
	pdesc.colorAttachments[0].pixelFormat = ConvertPixelFormat(GSTexture::Format::ColorClip);
	m_colclip_init_pipeline = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_colclip_init"), @"ColorClip Init");
	m_colclip_clear_pipeline = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_clear"), @"ColorClip Clear");
	pdesc.colorAttachments[0].pixelFormat = MTLPixelFormatInvalid;
	pdesc.stencilAttachmentPixelFormat = MTLPixelFormatDepth32Float_Stencil8;
	m_datm_pipeline[0] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_datm0"), @"datm0");
	m_datm_pipeline[1] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_datm1"), @"datm1");
	m_datm_pipeline[2] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_datm0_rta_correction"), @"datm0 rta");
	m_datm_pipeline[3] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_datm1_rta_correction"), @"datm1 rta");
	m_stencil_clear_pipeline = MakePipeline(pdesc, fs_triangle, nil, @"Stencil Clear");
	pdesc.colorAttachments[0].pixelFormat = ConvertPixelFormat(GSTexture::Format::PrimID);
	pdesc.stencilAttachmentPixelFormat = MTLPixelFormatInvalid;
	pdesc.depthAttachmentPixelFormat = MTLPixelFormatDepth32Float_Stencil8;
	m_primid_init_pipeline[1][0] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_primid_init_datm0"), @"PrimID DATM0 Clear");
	m_primid_init_pipeline[1][1] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_primid_init_datm1"), @"PrimID DATM1 Clear");
	m_primid_init_pipeline[1][2] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_primid_rta_init_datm0"), @"PrimID DATM0 RTA Clear");
	m_primid_init_pipeline[1][3] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_primid_rta_init_datm1"), @"PrimID DATM1 RTA Clear");
	pdesc.depthAttachmentPixelFormat = MTLPixelFormatInvalid;
	m_primid_init_pipeline[0][0] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_primid_init_datm0"), @"PrimID DATM0 Clear");
	m_primid_init_pipeline[0][1] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_primid_init_datm1"), @"PrimID DATM1 Clear");
	m_primid_init_pipeline[0][2] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_primid_rta_init_datm0"), @"PrimID DATM0 RTA Clear");
	m_primid_init_pipeline[0][3] = MakePipeline(pdesc, fs_triangle, LoadShader(@"ps_primid_rta_init_datm1"), @"PrimID DATM1 RTA Clear");

	pdesc.colorAttachments[0].pixelFormat = ConvertPixelFormat(GSTexture::Format::Color);
	applyAttribute(pdesc.vertexDescriptor, 0, MTLVertexFormatFloat2, offsetof(ConvertShaderVertex, pos),    0);
	applyAttribute(pdesc.vertexDescriptor, 1, MTLVertexFormatFloat2, offsetof(ConvertShaderVertex, texpos), 0);
	pdesc.vertexDescriptor.layouts[0].stride = sizeof(ConvertShaderVertex);

	for (size_t i = 0; i < std::size(m_interlace_pipeline); i++)
	{
		NSString* name = [NSString stringWithFormat:@"ps_interlace%zu", i];
		m_interlace_pipeline[i] = MakePipeline(pdesc, vs_convert, LoadShader(name), name);
	}

	m_convert_pipeline.resize(ShaderConvertSelector::NUM_TOTAL_SHADERS);
	for (u32 i = 0; i < ShaderConvertSelector::NUM_TOTAL_SHADERS; i++)
	{
		const ShaderConvertSelector shader = ShaderConvertSelector::Get(i);
		if (ConvertShaderNotNeeded(shader.Shader()))
			continue;
		NSString* shader_name = [NSString stringWithCString:shader.EntryPoint() encoding:NSUTF8StringEncoding];
		if (shader.DepthOutput())
		{
			pdesc.colorAttachments[0].pixelFormat = MTLPixelFormatInvalid;
			pdesc.depthAttachmentPixelFormat = ConvertPixelFormat(GSTexture::Format::DepthStencil);
		}
		else
		{
			pdesc.colorAttachments[0].pixelFormat = ConvertPixelFormat(shader.OutputFormat());
			pdesc.depthAttachmentPixelFormat = MTLPixelFormatInvalid;
		}
		NSString* name = shader_name;
		if (shader.VariableWriteMask())
			name = [name stringByAppendingString:[NSString stringWithFormat:@" Mask=%x", shader.Mask()]];
		if (shader.Biln())
			name = [name stringByAppendingString:@" Biln"];
		if (shader.Float32Output())
			name = [name stringByAppendingString:shader.DepthOutput() ? @" → Depth" : @" → Float"];

		const u32 scmask = shader.Mask();
		MTLColorWriteMask mask = MTLColorWriteMaskNone;
		if (scmask & 1) mask |= MTLColorWriteMaskRed;
		if (scmask & 2) mask |= MTLColorWriteMaskGreen;
		if (scmask & 4) mask |= MTLColorWriteMaskBlue;
		if (scmask & 8) mask |= MTLColorWriteMaskAlpha;
		pdesc.colorAttachments[0].writeMask = mask;
		setFnConstantB(m_fn_constants, shader.Biln(),        GSMTLConstantIndex_BILN);
		setFnConstantB(m_fn_constants, shader.DepthOutput(), GSMTLConstantIndex_DEPTH_OUT);
		m_convert_pipeline[shader.Index()] = MakePipeline(pdesc, vs_convert, LoadShader(shader_name), name);
	}
	pdesc.colorAttachments[0].writeMask = MTLColorWriteMaskAll;
	pdesc.depthAttachmentPixelFormat = MTLPixelFormatInvalid;
	for (size_t i = 0; i < std::size(m_present_pipeline); i++)
	{
		PresentShader conv = static_cast<PresentShader>(i);
		NSString* name = [NSString stringWithCString:ShaderEntryPoint(conv) encoding:NSUTF8StringEncoding];
		pdesc.colorAttachments[0].pixelFormat = layer_px_fmt;
		m_present_pipeline[i] = MakePipeline(pdesc, vs_convert, LoadShader(name), [NSString stringWithFormat:@"present_%s", ShaderEntryPoint(conv) + 3]);
	}

	pdesc.colorAttachments[0].pixelFormat = ConvertPixelFormat(GSTexture::Format::Color);
	pdesc.colorAttachments[0].blendingEnabled = YES;
	pdesc.colorAttachments[0].rgbBlendOperation = MTLBlendOperationAdd;
	pdesc.colorAttachments[0].sourceRGBBlendFactor = MTLBlendFactorSourceAlpha;
	pdesc.colorAttachments[0].destinationRGBBlendFactor = MTLBlendFactorOneMinusSourceAlpha;
	for (size_t i = 0; i < std::size(m_merge_pipeline); i++)
	{
		bool mmod = i & 1;
		bool amod = i & 2;
		NSString* name = [NSString stringWithFormat:@"ps_merge%d", mmod];
		NSString* pipename = [NSString stringWithFormat:@"Merge%s%s", mmod ? " MMOD" : "", amod ? " AMOD" : ""];
		pdesc.colorAttachments[0].writeMask = amod ? MTLColorWriteMaskRed | MTLColorWriteMaskGreen | MTLColorWriteMaskBlue : MTLColorWriteMaskAll;
		m_merge_pipeline[i] = MakePipeline(pdesc, vs_convert, LoadShader(name), pipename);
	}
	pdesc.colorAttachments[0].writeMask = MTLColorWriteMaskAll;

	applyAttribute(pdesc.vertexDescriptor, 0, MTLVertexFormatFloat2,           offsetof(ImDrawVert, pos), 0);
	applyAttribute(pdesc.vertexDescriptor, 1, MTLVertexFormatFloat2,           offsetof(ImDrawVert, uv),  0);
	applyAttribute(pdesc.vertexDescriptor, 2, MTLVertexFormatUChar4Normalized, offsetof(ImDrawVert, col), 0);
	pdesc.vertexDescriptor.layouts[0].stride = sizeof(ImDrawVert);
	pdesc.colorAttachments[0].pixelFormat = layer_px_fmt;
	m_imgui_pipeline = MakePipeline(pdesc, LoadShader(@"vs_imgui"), LoadShader(@"ps_imgui"), @"imgui");

	[initCommands commit];
	return true;
}}

void GSDeviceMTL::Destroy()
{ @autoreleasepool {
	FlushEncoders();
	std::lock_guard<std::mutex> guard(m_backref->first);
	m_backref->second = nullptr;

	GSDevice::Destroy();
	GSDeviceMTL::DestroySurface();
	m_queue = nullptr;
	m_dev.Reset();
}}

void GSDeviceMTL::DestroySurface()
{
	if (!m_layer)
		return;
	OnMainThread([this]{ DetachSurfaceOnMainThread(); });
	m_layer = nullptr;
}

bool GSDeviceMTL::UpdateWindow()
{
	DestroySurface();

	if (!AcquireWindow(false))
		return false;

	if (m_window_info.type == WindowInfo::Type::Surfaceless)
		return true;

	OnMainThread([this] { AttachSurfaceOnMainThread(); });
	return true;
}

bool GSDeviceMTL::SupportsExclusiveFullscreen() const { return false; }

static const char* GetROVString(const GSMTLDevice::Features& features)
{
	if (!features.rov)
		return "Unsupported";
	if (features.rov_requires_r32 && features.rov_requires_rt)
		return "With R32 Reinterpret and RT Texture";
	if (features.rov_requires_r32)
		return "With R32 Reinterpret";
	if (features.rov_requires_rt)
		return "With RT Texture";
	return "Supported";
}

std::string GSDeviceMTL::GetDriverInfo() const
{ @autoreleasepool {
	std::string desc([[m_dev.dev description] UTF8String]);
	desc += "\n    Texture Swizzle:     " + std::string(m_dev.features.texture_swizzle     ? "Supported" : "Unsupported");
	desc += "\n    Unified Memory:      " + std::string(m_dev.features.unified_memory      ? "Supported" : "Unsupported");
	desc += "\n    Framebuffer Fetch:   " + std::string(m_dev.features.framebuffer_fetch   ? "Supported" : "Unsupported");
	desc += "\n    Memoryless Textures: " + std::string(m_dev.features.memoryless_textures ? "Supported" : "Unsupported");
	desc += "\n    Primitive ID:        " + std::string(m_dev.features.primid              ? "Supported" : "Unsupported");
	desc += "\n    Depth Feedback:      " + std::string(m_dev.features.depth_feedback      ? "Supported" : "Unsupported");
	desc += "\n    ROV:                 " + std::string(GetROVString(m_dev.features));
	desc += "\n    Shader Version:      " + std::string(to_string(m_dev.features.shader_version));
	desc += "\n    Max Texture Size:    " + std::to_string(m_dev.features.max_texsize);
	return desc;
}}

void GSDeviceMTL::ResizeWindow(u32 new_window_width, u32 new_window_height, float new_window_scale)
{
	m_window_info.surface_scale = new_window_scale;
	if (!m_layer ||
		(m_window_info.surface_width == new_window_width && m_window_info.surface_height == new_window_height))
	{
		return;
	}

	m_window_info.surface_width = new_window_width;
	m_window_info.surface_height = new_window_height;
	@autoreleasepool
	{
		[m_layer setDrawableSize:CGSizeMake(new_window_width, new_window_height)];
	}
}

void GSDeviceMTL::UpdateTexture(id<MTLTexture> texture, u32 x, u32 y, u32 width, u32 height, const void* data, u32 data_stride)
{
	id<MTLCommandBuffer> cmdbuf = [m_queue commandBuffer];
	id<MTLBlitCommandEncoder> enc = [cmdbuf blitCommandEncoder];
	size_t bytes = data_stride * height;
	MRCOwned<id<MTLBuffer>> buf = MRCTransfer([m_dev.dev newBufferWithLength:bytes options:m_resource_options_shared_wc]);
	memcpy([buf contents], data, bytes);
	[enc copyFromBuffer:buf
	       sourceOffset:0
	  sourceBytesPerRow:data_stride
	sourceBytesPerImage:bytes
	         sourceSize:MTLSizeMake(width, height, 1)
	          toTexture:texture
	   destinationSlice:0
	   destinationLevel:0
	  destinationOrigin:MTLOriginMake(0, 0, 0)];
	[enc endEncoding];
	[cmdbuf commit];
}

static bool s_capture_next = false;

GSDevice::PresentResult GSDeviceMTL::DoBeginPresent(bool frame_skip)
{ @autoreleasepool {
	if (m_capture_start_frame && FrameNo() == m_capture_start_frame)
		s_capture_next = true;
	if (frame_skip || m_window_info.type == WindowInfo::Type::Surfaceless || !g_gs_device)
	{
		ImGui::EndFrame();
		return PresentResult::FrameSkipped;
	}
	id<MTLCommandBuffer> buf = GetRenderCmdBuf();
	m_current_drawable = MRCRetain([m_layer nextDrawable]);
	EndRenderPass();
	if (!m_current_drawable)
	{
		[buf pushDebugGroup:@"Present Skipped"];
		[buf popDebugGroup];
		FlushEncoders();
		ImGui::EndFrame();
		return PresentResult::FrameSkipped;
	}
	[m_pass_desc colorAttachments][0].texture = [m_current_drawable texture];
	id<MTLRenderCommandEncoder> enc = [buf renderCommandEncoderWithDescriptor:m_pass_desc];
	[enc setLabel:@"Present"];
	m_current_render.encoder = MRCRetain(enc);
	return PresentResult::OK;
}}

void GSDeviceMTL::EndPresent()
{ @autoreleasepool {
	pxAssertMsg(m_current_render.encoder && m_current_render_cmdbuf, "DoBeginPresent cmdbuf was destroyed");
	ImGui::Render();
	RenderImGui(ImGui::GetDrawData());
	EndRenderPass();
	if (m_current_drawable)
	{
		const bool use_present_drawable = m_use_present_drawable == UsePresentDrawable::Always ||
			(m_use_present_drawable == UsePresentDrawable::IfVsync && m_vsync_mode == GSVSyncMode::FIFO);

		if (use_present_drawable)
			[m_current_render_cmdbuf presentDrawable:m_current_drawable];
		else
			[m_current_render_cmdbuf addScheduledHandler:[drawable = std::move(m_current_drawable)](id<MTLCommandBuffer>){
				[drawable present];
			}];
	}
	FlushEncoders();
	FrameCompleted();
	m_current_drawable = nullptr;
	if (m_capture_start_frame)
	{
		if (@available(macOS 10.15, iOS 13, *))
		{
			static NSString* const path = @"/tmp/PCSX2MTLCapture.gputrace";
			static u32 frames;
			if (frames)
			{
				--frames;
				if (!frames)
				{
					[[MTLCaptureManager sharedCaptureManager] stopCapture];
				Console.WriteLn("Metal Trace Capture to /tmp/PCSX2MTLCapture.gputrace finished");
#if !TARGET_OS_IPHONE
				[[NSWorkspace sharedWorkspace] selectFile:path
				                 inFileViewerRootedAtPath:@"/tmp/"];
#endif
				}
			}
			else if (s_capture_next)
			{
				s_capture_next = false;
				MTLCaptureManager* mgr = [MTLCaptureManager sharedCaptureManager];
				if ([mgr supportsDestination:MTLCaptureDestinationGPUTraceDocument])
				{
					MTLCaptureDescriptor* desc = [[MTLCaptureDescriptor new] autorelease];
					[desc setCaptureObject:m_dev.dev];
					if ([[NSFileManager defaultManager] fileExistsAtPath:path])
						[[NSFileManager defaultManager] removeItemAtPath:path error:nil];
					[desc setOutputURL:[NSURL fileURLWithPath:path]];
					[desc setDestination:MTLCaptureDestinationGPUTraceDocument];
					NSError* err = nullptr;
					[mgr startCaptureWithDescriptor:desc error:&err];
					if (err)
					{
						Console.Error("Metal Trace Capture failed: %s", [[err localizedDescription] UTF8String]);
					}
					else
					{
						Console.WriteLn("Metal Trace Capture to /tmp/PCSX2MTLCapture.gputrace started");
						frames = 2;
					}
				}
				else
				{
					Console.Error("Metal Trace Capture Failed: MTLCaptureManager doesn't support GPU trace documents! (Did you forget to run with METAL_CAPTURE_ENABLED=1?)");
				}
			}
		}
	}
}}

void GSDeviceMTL::SetVSyncMode(GSVSyncMode mode, bool allow_present_throttle)
{
	m_allow_present_throttle = allow_present_throttle;

	if (m_vsync_mode == mode)
		return;

	m_vsync_mode = (mode == GSVSyncMode::Mailbox) ? GSVSyncMode::FIFO : mode;
#if !TARGET_OS_IPHONE
	[m_layer setDisplaySyncEnabled:m_vsync_mode == GSVSyncMode::FIFO];
#endif
}

bool GSDeviceMTL::SetGPUTimingEnabled(bool enabled)
{
	if (enabled == m_gpu_timing_enabled)
		return true;
	if (@available(macOS 10.15, iOS 10.3, *))
	{
		std::lock_guard<std::mutex> l(m_mtx);
		m_gpu_timing_enabled = enabled;
		m_accumulated_gpu_time = 0;
		m_last_gpu_time_end = 0;
		return true;
	}
	return false;
}

float GSDeviceMTL::GetAndResetAccumulatedGPUTime()
{
	std::lock_guard<std::mutex> l(m_mtx);
	float time = m_accumulated_gpu_time * 1000;
	m_accumulated_gpu_time = 0;
	return time;
}

void GSDeviceMTL::AccumulateCommandBufferTime(id<MTLCommandBuffer> buffer)
{
	std::lock_guard<std::mutex> l(m_mtx);
	if (!m_gpu_timing_enabled)
		return;
	// We do the check before enabling m_gpu_timing_enabled
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wunguarded-availability"
	// It's unlikely, but command buffers can overlap or run out of order
	// This doesn't handle every case (fully out of order), but it should at least handle overlapping
	double begin = std::max(m_last_gpu_time_end, [buffer GPUStartTime]);
	double end = [buffer GPUEndTime];
	if (end > begin)
	{
		m_accumulated_gpu_time += end - begin;
		m_last_gpu_time_end = end;
	}
#pragma clang diagnostic pop
}

std::unique_ptr<GSDownloadTexture> GSDeviceMTL::CreateDownloadTexture(u32 width, u32 height, GSTexture::Format format)
{
	return GSDownloadTextureMTL::Create(this, width, height, format);
}

void GSDeviceMTL::ClearSamplerCache()
{ @autoreleasepool {
	std::fill(std::begin(m_sampler_hw), std::end(m_sampler_hw), nullptr);
	m_sampler_hw[SamplerSelector::Linear().key] = CreateSampler(m_dev.dev, SamplerSelector::Linear());
	m_sampler_hw[SamplerSelector::Point().key] = CreateSampler(m_dev.dev, SamplerSelector::Point());
}}

void GSDeviceMTL::DoCopyRect(GSTexture* sTex, GSTexture* dTex, const GSVector4i& r, u32 destX, u32 destY)
{ @autoreleasepool {
	// Empty rect, abort copy.
	if (r.rempty())
	{
		GL_INS("Metal: DoCopyRect rect empty.");
		return;
	}
	
	GSTextureMTL* sT = static_cast<GSTextureMTL*>(sTex);
	GSTextureMTL* dT = static_cast<GSTextureMTL*>(dTex);
	const GSVector4i dst_rect(0, 0, dT->GetWidth(), dT->GetHeight());
	const bool full_draw_copy = dst_rect.eq(r);

	// Source is cleared, if destination is a render target, we can carry the clear forward.
	if (sT->GetState() == GSTexture::State::Cleared)
	{
		if (dT->IsRenderTargetOrDepthStencil() && ProcessClearsBeforeCopy(sTex, dTex, full_draw_copy))
			return;

		// Commit clear for the source texture.
		sT->FlushClears();
	}

	g_perfmon.Put(GSPerfMon::TextureCopies, 1);

	// Commit clear for the destination texture.
	GSVector2i dsize = dTex->GetSize();
	if (r.width() < dsize.x || r.height() < dsize.y)
		dT->FlushClears();
	else
		dT->SetState(GSTexture::State::Dirty);

	EndRenderPass();

	sT->m_last_read  = m_current_draw;
	dT->m_last_write = m_current_draw;

	id<MTLCommandBuffer> cmdbuf = GetRenderCmdBuf();
	id<MTLBlitCommandEncoder> encoder = [cmdbuf blitCommandEncoder];
	[encoder setLabel:@"DoCopyRect"];
	[encoder copyFromTexture:sT->GetTexture()
	             sourceSlice:0
	             sourceLevel:0
	            sourceOrigin:MTLOriginMake(r.x, r.y, 0)
	              sourceSize:MTLSizeMake(r.width(), r.height(), 1)
	               toTexture:dT->GetTexture()
	        destinationSlice:0
	        destinationLevel:0
	       destinationOrigin:MTLOriginMake((int)destX, (int)destY, 0)];
	[encoder endEncoding];
}}

void GSDeviceMTL::BeginStretchRect(NSString* name, GSTexture* dTex, MTLLoadAction action)
{
	if (dTex->GetFormat() == GSTexture::Format::DepthStencil)
		BeginRenderPass(name, nullptr, MTLLoadActionDontCare, dTex, action);
	else
		BeginRenderPass(name, dTex, action, nullptr, MTLLoadActionDontCare);

	FlushDebugEntries(m_current_render.encoder);
	MREClearScissor();
	DepthStencilSelector dsel = DepthStencilSelector::NoDepth();
	dsel.zwe = dTex->GetFormat() == GSTexture::Format::DepthStencil;
	MRESetDSS(dsel);
}

void GSDeviceMTL::DoStretchRect(GSTexture* sTex, const GSVector4& sRect, GSTexture* dTex, const GSVector4& dRect, id<MTLRenderPipelineState> pipeline, std::optional<Filter> filter, LoadAction load_action, const void* frag_uniform, size_t frag_uniform_len)
{
	FlushClears(sTex);

	GSVector2i ds = dTex->GetSize();

	bool covers_target = static_cast<int>(dRect.x) <= 0
	                  && static_cast<int>(dRect.y) <= 0
	                  && static_cast<int>(dRect.z) >= ds.x
	                  && static_cast<int>(dRect.w) >= ds.y;
	bool dontcare = load_action == LoadAction::DontCare || (load_action == LoadAction::DontCareIfFull && covers_target);
	MTLLoadAction action = dontcare ? MTLLoadActionDontCare : MTLLoadActionLoad;

	BeginStretchRect(@"StretchRect", dTex, action);

	MRESetPipeline(pipeline);
	MRESetTexture(sTex, GSMTLTextureIndexNonHW);

	if (frag_uniform && frag_uniform_len)
		[m_current_render.encoder setFragmentBytes:frag_uniform length:frag_uniform_len atIndex:GSMTLBufferIndexUniforms];

	if (filter)
		MRESetSampler(*filter == Biln ? SamplerSelector::Linear() : SamplerSelector::Point());

	DrawStretchRect(sRect, dRect, GSVector2(static_cast<float>(ds.x), static_cast<float>(ds.y)));
}

static std::array<GSVector4, 4> CalcStrechRectPoints(const GSVector4& sRect, const GSVector4& dRect, const GSVector2& ds)
{
	static_assert(sizeof(GSDeviceMTL::ConvertShaderVertex) == sizeof(GSVector4), "Using GSVector4 as a ConvertShaderVertex");
	GSVector4 dst = dRect;
	dst /= GSVector4(ds.x, ds.y, ds.x, ds.y);
	dst *= GSVector4(2, -2, 2, -2);
	dst += GSVector4(-1, 1, -1, 1);
	return {
		dst.xyxy(sRect),
		dst.zyzy(sRect),
		dst.xwxw(sRect),
		dst.zwzw(sRect)
	};
}

void GSDeviceMTL::DrawStretchRect(const GSVector4& sRect, const GSVector4& dRect, const GSVector2& ds)
{
	std::array<GSVector4, 4> vertices = CalcStrechRectPoints(sRect, dRect, ds);

	[m_current_render.encoder setVertexBytes:&vertices length:sizeof(vertices) atIndex:GSMTLBufferIndexVertices];

	[m_current_render.encoder drawPrimitives:MTLPrimitiveTypeTriangleStrip
	                             vertexStart:0
	                             vertexCount:4];
	g_perfmon.Put(GSPerfMon::TextureCopies, 1);
	g_perfmon.Put(GSPerfMon::DrawCalls, 1);
}

void GSDeviceMTL::RenderCopy(GSTexture* sTex, id<MTLRenderPipelineState> pipeline, const GSVector4i& rect)
{
	// FS Triangle encoder uses vertex ID alone to make a FS triangle, which we then scissor to the desired rectangle
	MRESetScissor(rect);
	MRESetPipeline(pipeline);
	MRESetTexture(sTex, GSMTLTextureIndexNonHW);
	[m_current_render.encoder drawPrimitives:MTLPrimitiveTypeTriangle vertexStart:0 vertexCount:3];
	g_perfmon.Put(GSPerfMon::TextureCopies, 1);
	g_perfmon.Put(GSPerfMon::DrawCalls, 1);
}

void GSDeviceMTL::DoStretchRect(GSTexture* sTex, const GSVector4& sRect, GSTexture* dTex, const GSVector4& dRect,
	ShaderConvertSelector shader, Filter filter)
{ @autoreleasepool {

	const LoadAction load_action = (shader.Mask() == 0xf) ? LoadAction::DontCareIfFull : LoadAction::Load;
	id<MTLRenderPipelineState> pipeline = GetConvertPipeline(shader);
	pxAssertRel(pipeline, fmt::format("No pipeline for {}", ShaderEntryPoint(shader.Shader())).c_str());
	std::optional<Filter> filter_if_needed = shader.SupportsBilinear() ? std::nullopt : std::make_optional(filter);

	DoStretchRect(sTex, sRect, dTex, dRect, pipeline, filter_if_needed, load_action, nullptr, 0);
}}

static_assert(sizeof(DisplayConstantBuffer) == sizeof(GSMTLPresentPSUniform));
static_assert(offsetof(DisplayConstantBuffer, SourceRect)          == offsetof(GSMTLPresentPSUniform, source_rect));
static_assert(offsetof(DisplayConstantBuffer, TargetRect)          == offsetof(GSMTLPresentPSUniform, target_rect));
static_assert(offsetof(DisplayConstantBuffer, TargetSize)          == offsetof(GSMTLPresentPSUniform, target_size));
static_assert(offsetof(DisplayConstantBuffer, TargetResolution)    == offsetof(GSMTLPresentPSUniform, target_resolution));
static_assert(offsetof(DisplayConstantBuffer, RcpTargetResolution) == offsetof(GSMTLPresentPSUniform, rcp_target_resolution));
static_assert(offsetof(DisplayConstantBuffer, SourceResolution)    == offsetof(GSMTLPresentPSUniform, source_resolution));
static_assert(offsetof(DisplayConstantBuffer, RcpSourceResolution) == offsetof(GSMTLPresentPSUniform, rcp_source_resolution));
static_assert(offsetof(DisplayConstantBuffer, TimeAndPad.x)        == offsetof(GSMTLPresentPSUniform, time));

void GSDeviceMTL::PresentRect(GSTexture* sTex, const GSVector4& sRect, GSTexture* dTex, const GSVector4& dRect, PresentShader shader, float shaderTime, Filter filter)
{ @autoreleasepool {
	GSVector2i ds = dTex ? dTex->GetSize() : GetWindowSize();
	DisplayConstantBuffer cb;
	cb.SetSource(sRect, sTex->GetSize());
	cb.SetTarget(dRect, ds);
	cb.SetTime(shaderTime);
	id<MTLRenderPipelineState> pipe = m_present_pipeline[static_cast<int>(shader)];

	if (dTex)
	{
		DoStretchRect(sTex, sRect, dTex, dRect, pipe, filter, LoadAction::DontCareIfFull, &cb, sizeof(cb));
	}
	else
	{
		// !dTex → Use current draw encoder
		[m_current_render.encoder setRenderPipelineState:pipe];
		[m_current_render.encoder setFragmentSamplerState:m_sampler_hw[filter == Biln ? SamplerSelector::Linear().key : SamplerSelector::Point().key] atIndex:0];
		[m_current_render.encoder setFragmentTexture:static_cast<GSTextureMTL*>(sTex)->GetTexture() atIndex:0];
		[m_current_render.encoder setFragmentBytes:&cb length:sizeof(cb) atIndex:GSMTLBufferIndexUniforms];
		DrawStretchRect(sRect, dRect, GSVector2(static_cast<float>(ds.x), static_cast<float>(ds.y)));
	}
}}

void GSDeviceMTL::DoDrawMultiStretchRects(const MultiStretchRect* rects, u32 num_rects, GSTexture* dTex, ShaderConvertSelector shader)
{ @autoreleasepool {
	BeginStretchRect(@"MultiStretchRect", dTex, MTLLoadActionLoad);

	id<MTLRenderPipelineState> pipeline = nullptr;
	GSTexture* sTex = rects[0].src;
	Filter filter = rects[0].filter;
	u8 wmask = rects[0].wmask.wrgba;

	const GSVector2 ds(static_cast<float>(dTex->GetWidth()), static_cast<float>(dTex->GetHeight()));
	const Map allocation = Allocate(m_vertex_upload_buf, sizeof(ConvertShaderVertex) * 4 * num_rects);
	std::array<GSVector4, 4>* write = static_cast<std::array<GSVector4, 4>*>(allocation.cpu_buffer);
	const id<MTLRenderCommandEncoder> enc = m_current_render.encoder;
	[enc setVertexBuffer:allocation.gpu_buffer
	              offset:allocation.gpu_offset
	             atIndex:GSMTLBufferIndexVertices];
	u32 start = 0;

	auto flush = [&](u32 i) {
		const u32 end = i * 4;
		const u32 vertex_count = end - start;
		const u32 index_count = vertex_count + (vertex_count >> 1); // 6 indices per 4 vertices
		id<MTLRenderPipelineState> new_pipeline = GetConvertPipeline(shader.SetMask(wmask));
		if (new_pipeline != pipeline)
		{
			pipeline = new_pipeline;
			pxAssertRel(pipeline, fmt::format("No pipeline for {}", ShaderEntryPoint(shader.Shader())).c_str());
			MRESetPipeline(pipeline);
		}
		MRESetSampler(filter == Biln ? SamplerSelector::Linear() : SamplerSelector::Point());
		MRESetTexture(sTex, GSMTLTextureIndexNonHW);
		[enc drawIndexedPrimitives:MTLPrimitiveTypeTriangle
		                indexCount:index_count
		                 indexType:MTLIndexTypeUInt16
		               indexBuffer:m_expand_index_buffer
		         indexBufferOffset:0
		             instanceCount:1
		                baseVertex:start
		              baseInstance:0];
		start = end;
	};

	for (u32 i = 0; i < num_rects; i++)
	{
		const MultiStretchRect& rect = rects[i];
		if (rect.src != sTex || rect.filter != filter || rect.wmask.wrgba != wmask)
		{
			flush(i);
			sTex = rect.src;
			filter = rect.filter;
			wmask = rect.wmask.wrgba;
		}
		*write++ = CalcStrechRectPoints(rect.src_rect, rect.dst_rect, ds);
	}

	flush(num_rects);
}}

void GSDeviceMTL::DoUpdateCLUTTexture(GSTexture* sTex, float sScale, u32 offsetX, u32 offsetY, GSTexture* dTex, u32 dOffset, u32 dSize)
{
	GSMTLCLUTConvertPSUniform uniform = { sScale, {offsetX, offsetY}, dOffset };

	const bool is_clut4 = dSize == 16;
	const GSVector4i dRect(0, 0, dSize, 1);

	BeginRenderPass(@"CLUT Update", dTex, MTLLoadActionDontCare, nullptr, MTLLoadActionDontCare);
	[m_current_render.encoder setFragmentBytes:&uniform length:sizeof(uniform) atIndex:GSMTLBufferIndexUniforms];
	RenderCopy(sTex, m_clut_pipeline[!is_clut4], dRect);
}

void GSDeviceMTL::DoConvertToIndexedTexture(GSTexture* sTex, float sScale, u32 offsetX, u32 offsetY, u32 SBW, u32 SPSM, GSTexture* dTex, u32 DBW, u32 DPSM)
{ @autoreleasepool {
	const ShaderConvert shader = ((SPSM & 0xE) == 0) ? ShaderConvert::RGBA_TO_8I : ShaderConvert::RGB5A1_TO_8I;
	id<MTLRenderPipelineState> pipeline = GetConvertPipeline(shader);
	if (!pipeline)
		[NSException raise:@"StretchRect Missing Pipeline" format:@"No pipeline for %d", static_cast<int>(shader)];

	GSMTLIndexedConvertPSUniform uniform = { sScale, SBW, DBW, SPSM };

	const GSVector4 dRect(0, 0, dTex->GetWidth(), dTex->GetHeight());
	DoStretchRect(sTex, GSVector4::zero(), dTex, dRect, pipeline, Nearest, LoadAction::DontCareIfFull, &uniform, sizeof(uniform));
}}

void GSDeviceMTL::DoFilteredDownsampleTexture(GSTexture* sTex, GSTexture* dTex, u32 downsample_factor, const GSVector2i& clamp_min, const GSVector4& dRect)
{ @autoreleasepool {
	const ShaderConvert shader = ShaderConvert::DOWNSAMPLE_COPY;
	id<MTLRenderPipelineState> pipeline = GetConvertPipeline(shader);
	if (!pipeline)
		[NSException raise:@"StretchRect Missing Pipeline" format:@"No pipeline for %d", static_cast<int>(shader)];

	GSMTLDownsamplePSUniform uniform = { {static_cast<uint>(clamp_min.x), static_cast<uint>(clamp_min.x)}, downsample_factor,
	  static_cast<float>(downsample_factor * downsample_factor), (GSConfig.UserHacks_NativeScaling > GSNativeScaling::Aggressive) ? 2.0f : 1.0f };

	DoStretchRect(sTex, GSVector4::zero(), dTex, dRect, pipeline, Nearest, LoadAction::DontCareIfFull, &uniform, sizeof(uniform));
}}

static id<MTLTexture> CreateDSAsRTTexture(id<MTLDevice> dev, NSUInteger width, NSUInteger height, MTLStorageMode storage, NSString* name)
{
	MTLTextureDescriptor *desc = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatR32Float width:width height:height mipmapped:false];
	[desc setUsage:MTLTextureUsageRenderTarget];
	[desc setStorageMode:storage];
	id<MTLTexture> result = [dev newTextureWithDescriptor:desc];
	[result setLabel:name];
	return result;
}

void GSDeviceMTL::DoBeginDSAsRT(GSTexture* ds, const GSVector4i& drawarea)
{
	if (!m_features.framebuffer_fetch)
		return GSDevice::DoBeginDSAsRT(ds, drawarea);
	u32 needed_width = ds->GetWidth();
	u32 needed_height = ds->GetHeight();
	u32 current_width = static_cast<u32>([m_ds_as_rt_texture width]);
	u32 current_height = static_cast<u32>([m_ds_as_rt_texture height]);
	if (m_dev.features.memoryless_textures)
	{
		if (needed_width > current_width || needed_height > current_height) [[unlikely]] @autoreleasepool
		{
			u32 width = std::max(needed_width, current_width);
			u32 height = std::max(needed_height, current_height);
			[m_ds_as_rt_texture release];
			m_ds_as_rt_texture = CreateDSAsRTTexture(m_dev.dev, width, height, MTLStorageModeMemoryless, @"DS as RT");
		}
	}
	else
	{
		if (needed_width == current_width && needed_height == current_height)
			return;
		if (m_ds_as_rt_gstexture)
			Recycle(m_ds_as_rt_gstexture);
		m_ds_as_rt_gstexture = CreateRenderTarget(needed_width, needed_height, GSTexture::Format::DepthColor, false, true);
		m_ds_as_rt_texture = static_cast<GSTextureMTL*>(m_ds_as_rt_gstexture)->GetTexture();
		@autoreleasepool
		{
			NSString* name = [NSString stringWithFormat:@"DS as RT %dx%d", needed_width, needed_height];
			[m_ds_as_rt_texture setLabel:name];
		}
	}
}

void GSDeviceMTL::FlushClears(GSTexture* tex)
{
	if (tex)
		static_cast<GSTextureMTL*>(tex)->FlushClears();
}

// MARK: - MainRenderEncoder Operations

static MTLBlendFactor ConvertBlendFactor(GSDevice::BlendFactor generic)
{
	switch (generic)
	{
		case GSDevice::SRC_COLOR:       return MTLBlendFactorSourceColor;
		case GSDevice::INV_SRC_COLOR:   return MTLBlendFactorOneMinusSourceColor;
		case GSDevice::DST_COLOR:       return MTLBlendFactorDestinationColor;
		case GSDevice::INV_DST_COLOR:   return MTLBlendFactorOneMinusBlendColor;
		case GSDevice::SRC1_COLOR:      return MTLBlendFactorSource1Color;
		case GSDevice::INV_SRC1_COLOR:  return MTLBlendFactorOneMinusSource1Color;
		case GSDevice::SRC_ALPHA:       return MTLBlendFactorSourceAlpha;
		case GSDevice::INV_SRC_ALPHA:   return MTLBlendFactorOneMinusSourceAlpha;
		case GSDevice::DST_ALPHA:       return MTLBlendFactorDestinationAlpha;
		case GSDevice::INV_DST_ALPHA:   return MTLBlendFactorOneMinusDestinationAlpha;
		case GSDevice::SRC1_ALPHA:      return MTLBlendFactorSource1Alpha;
		case GSDevice::INV_SRC1_ALPHA:  return MTLBlendFactorOneMinusSource1Alpha;
		case GSDevice::CONST_COLOR:     return MTLBlendFactorBlendColor;
		case GSDevice::INV_CONST_COLOR: return MTLBlendFactorOneMinusBlendColor;
		case GSDevice::CONST_ONE:       return MTLBlendFactorOne;
		case GSDevice::CONST_ZERO:      return MTLBlendFactorZero;
	}
}

static MTLBlendOperation ConvertBlendOp(GSDevice::BlendOp generic)
{
	switch (generic)
	{
		case GSDevice::OP_ADD:          return MTLBlendOperationAdd;
		case GSDevice::OP_SUBTRACT:     return MTLBlendOperationSubtract;
		case GSDevice::OP_REV_SUBTRACT: return MTLBlendOperationReverseSubtract;
	}
}

void GSDeviceMTL::MRESetHWPipelineState(GSHWDrawConfig::VSSelector vssel, GSHWDrawConfig::PSSelector pssel, GSHWDrawConfig::BlendState blend, GSHWDrawConfig::ColorMaskSelector cms)
{
	PipelineSelectorExtrasMTL extras(blend, m_current_render.color_target, cms, m_current_render.depth_target, m_current_render.stencil_target, m_current_render.has.rt1_depth);
	PipelineSelectorMTL fullsel(vssel, pssel, extras);
	if (m_current_render.has.pipeline_sel && fullsel == m_current_render.pipeline_sel)
		return;
	m_current_render.pipeline_sel = fullsel;
	m_current_render.has.pipeline_sel = true;
	auto idx = m_hw_pipeline.find(fullsel);
	if (idx != m_hw_pipeline.end())
	{
		[m_current_render.encoder setRenderPipelineState:idx->second];
		return;
	}

	bool primid_tracking_init = pssel.date == 1 || pssel.date == 2;

	VSSelector vssel_mtl;
	vssel_mtl.fst = vssel.fst;
	vssel_mtl.iip = vssel.iip;
	vssel_mtl.point_size = vssel.point_size;
	vssel_mtl.expand = vssel.expand;
	id<MTLFunction> vs = m_hw_vs[vssel_mtl.key];

	id<MTLFunction> ps;
	auto idx2 = m_hw_ps.find(pssel);
	if (idx2 != m_hw_ps.end())
	{
		ps = idx2->second;
	}
	else
	{
		setFnConstantB(m_fn_constants, pssel.fst,                   GSMTLConstantIndex_FST);
		setFnConstantB(m_fn_constants, pssel.iip,                   GSMTLConstantIndex_IIP);
		setFnConstantI(m_fn_constants, pssel.aem_fmt,               GSMTLConstantIndex_PS_AEM_FMT);
		setFnConstantI(m_fn_constants, pssel.pal_fmt,               GSMTLConstantIndex_PS_PAL_FMT);
		setFnConstantI(m_fn_constants, pssel.dst_fmt,               GSMTLConstantIndex_PS_DST_FMT);
		setFnConstantI(m_fn_constants, pssel.depth_fmt,             GSMTLConstantIndex_PS_DEPTH_FMT);
		setFnConstantB(m_fn_constants, pssel.aem,                   GSMTLConstantIndex_PS_AEM);
		setFnConstantB(m_fn_constants, pssel.fba,                   GSMTLConstantIndex_PS_FBA);
		setFnConstantB(m_fn_constants, pssel.fog,                   GSMTLConstantIndex_PS_FOG);
		setFnConstantI(m_fn_constants, pssel.date,                  GSMTLConstantIndex_PS_DATE);
		setFnConstantI(m_fn_constants, pssel.atst,                  GSMTLConstantIndex_PS_ATST);
		setFnConstantI(m_fn_constants, pssel.afail,                 GSMTLConstantIndex_PS_AFAIL);
		setFnConstantI(m_fn_constants, pssel.ztst,                  GSMTLConstantIndex_PS_ZTST);
		setFnConstantI(m_fn_constants, pssel.tfx,                   GSMTLConstantIndex_PS_TFX);
		setFnConstantB(m_fn_constants, pssel.tcc,                   GSMTLConstantIndex_PS_TCC);
		setFnConstantI(m_fn_constants, pssel.wms,                   GSMTLConstantIndex_PS_WMS);
		setFnConstantI(m_fn_constants, pssel.wmt,                   GSMTLConstantIndex_PS_WMT);
		setFnConstantB(m_fn_constants, pssel.adjs,                  GSMTLConstantIndex_PS_ADJS);
		setFnConstantB(m_fn_constants, pssel.adjt,                  GSMTLConstantIndex_PS_ADJT);
		setFnConstantB(m_fn_constants, pssel.ltf,                   GSMTLConstantIndex_PS_LTF);
		setFnConstantB(m_fn_constants, pssel.shuffle,               GSMTLConstantIndex_PS_SHUFFLE);
		setFnConstantB(m_fn_constants, pssel.shuffle_same,          GSMTLConstantIndex_PS_SHUFFLE_SAME);
		setFnConstantI(m_fn_constants, pssel.process_ba,            GSMTLConstantIndex_PS_PROCESS_BA);
		setFnConstantI(m_fn_constants, pssel.process_rg,            GSMTLConstantIndex_PS_PROCESS_RG);
		setFnConstantB(m_fn_constants, pssel.shuffle_across,        GSMTLConstantIndex_PS_SHUFFLE_ACROSS);
		setFnConstantB(m_fn_constants, pssel.real16src,             GSMTLConstantIndex_PS_READ16_SRC);
		setFnConstantB(m_fn_constants, pssel.write_rg,              GSMTLConstantIndex_PS_WRITE_RG);
		setFnConstantB(m_fn_constants, pssel.fbmask,                GSMTLConstantIndex_PS_FBMASK);
		setFnConstantI(m_fn_constants, pssel.blend_a,               GSMTLConstantIndex_PS_BLEND_A);
		setFnConstantI(m_fn_constants, pssel.blend_b,               GSMTLConstantIndex_PS_BLEND_B);
		setFnConstantI(m_fn_constants, pssel.blend_c,               GSMTLConstantIndex_PS_BLEND_C);
		setFnConstantI(m_fn_constants, pssel.blend_d,               GSMTLConstantIndex_PS_BLEND_D);
		setFnConstantI(m_fn_constants, pssel.blend_hw,              GSMTLConstantIndex_PS_BLEND_HW);
		setFnConstantB(m_fn_constants, pssel.a_masked,              GSMTLConstantIndex_PS_A_MASKED);
		setFnConstantB(m_fn_constants, pssel.colclip_hw,            GSMTLConstantIndex_PS_COLCLIP_HW);
		setFnConstantB(m_fn_constants, pssel.rta_correction,        GSMTLConstantIndex_PS_RTA_CORRECTION);
		setFnConstantB(m_fn_constants, pssel.rta_source_correction, GSMTLConstantIndex_PS_RTA_SRC_CORRECTION);
		setFnConstantB(m_fn_constants, pssel.colclip,               GSMTLConstantIndex_PS_COLCLIP);
		setFnConstantI(m_fn_constants, pssel.blend_mix,             GSMTLConstantIndex_PS_BLEND_MIX);
		setFnConstantB(m_fn_constants, pssel.round_inv,             GSMTLConstantIndex_PS_ROUND_INV);
		setFnConstantB(m_fn_constants, pssel.fixed_one_a,           GSMTLConstantIndex_PS_FIXED_ONE_A);
		setFnConstantB(m_fn_constants, pssel.pabe,                  GSMTLConstantIndex_PS_PABE);
		setFnConstantB(m_fn_constants, pssel.no_color,              GSMTLConstantIndex_PS_NO_COLOR);
		setFnConstantB(m_fn_constants, pssel.no_color1,             GSMTLConstantIndex_PS_NO_COLOR1);
		setFnConstantI(m_fn_constants, pssel.channel,               GSMTLConstantIndex_PS_CHANNEL);
		setFnConstantI(m_fn_constants, pssel.dither,                GSMTLConstantIndex_PS_DITHER);
		setFnConstantI(m_fn_constants, pssel.dither_adjust,         GSMTLConstantIndex_PS_DITHER_ADJUST);
		setFnConstantB(m_fn_constants, pssel.zclamp,                GSMTLConstantIndex_PS_ZCLAMP);
		setFnConstantB(m_fn_constants, pssel.zfloor,                GSMTLConstantIndex_PS_ZFLOOR);
		setFnConstantB(m_fn_constants, pssel.tcoffsethack,          GSMTLConstantIndex_PS_TCOFFSETHACK);
		setFnConstantB(m_fn_constants, pssel.urban_chaos_hle,       GSMTLConstantIndex_PS_URBAN_CHAOS_HLE);
		setFnConstantB(m_fn_constants, pssel.tales_of_abyss_hle,    GSMTLConstantIndex_PS_TALES_OF_ABYSS_HLE);
		setFnConstantB(m_fn_constants, pssel.tex_is_fb,             GSMTLConstantIndex_PS_TEX_IS_FB);
		setFnConstantB(m_fn_constants, pssel.automatic_lod,         GSMTLConstantIndex_PS_AUTOMATIC_LOD);
		setFnConstantB(m_fn_constants, pssel.manual_lod,            GSMTLConstantIndex_PS_MANUAL_LOD);
		setFnConstantB(m_fn_constants, pssel.region_rect,           GSMTLConstantIndex_PS_REGION_RECT);
		setFnConstantI(m_fn_constants, pssel.scanmsk,               GSMTLConstantIndex_PS_SCANMSK);
		setFnConstantI(m_fn_constants, pssel.aa1,                   GSMTLConstantIndex_PS_AA1);
		setFnConstantB(m_fn_constants, pssel.abe,                   GSMTLConstantIndex_PS_ABE);
		setFnConstantI(m_fn_constants, pssel.sw_aniso,              GSMTLConstantIndex_PS_SW_ANISO);
		setFnConstantB(m_fn_constants, pssel.rov_color,             GSMTLConstantIndex_PS_ROV_COLOR);
		setFnConstantI(m_fn_constants, pssel.rov_depth,             GSMTLConstantIndex_PS_ROV_DEPTH);
		bool eft = pssel.HasColorROV() && !pssel.HasDepthROV() && !pssel.HasDepthOutput();
		auto newps = LoadShader(eft ? @"ps_main_rov_eft" : @"ps_main");
		ps = newps;
		m_hw_ps.insert(std::make_pair(pssel, std::move(newps)));
	}

	MRCOwned<MTLRenderPipelineDescriptor*> pdesc = MRCTransfer([MTLRenderPipelineDescriptor new]);
	if (vssel_mtl.point_size)
		[pdesc setInputPrimitiveTopology:MTLPrimitiveTopologyClassPoint];
	if (vssel_mtl.expand == GSShader::VSExpand::None)
		[pdesc setVertexDescriptor:m_hw_vertex];
	else
		[pdesc setInputPrimitiveTopology:MTLPrimitiveTopologyClassTriangle];
	MTLRenderPipelineColorAttachmentDescriptor* color = [[pdesc colorAttachments] objectAtIndexedSubscript:0];
	color.pixelFormat = ConvertPixelFormat(extras.rt);
	[pdesc setDepthAttachmentPixelFormat:extras.has_depth ? MTLPixelFormatDepth32Float_Stencil8 : MTLPixelFormatInvalid];
	[pdesc setStencilAttachmentPixelFormat:extras.has_stencil ? MTLPixelFormatDepth32Float_Stencil8 : MTLPixelFormatInvalid];
	color.writeMask = extras.writemask;
	if (primid_tracking_init)
	{
		color.blendingEnabled = YES;
		color.rgbBlendOperation = MTLBlendOperationMin;
		color.sourceRGBBlendFactor = MTLBlendFactorOne;
		color.destinationRGBBlendFactor = MTLBlendFactorOne;
		color.writeMask = MTLColorWriteMaskRed;
	}
	else if (blend.IsEffective(cms))
	{
		color.blendingEnabled = YES;
		color.rgbBlendOperation = ConvertBlendOp(extras.blend_op);
		color.sourceRGBBlendFactor = ConvertBlendFactor(extras.src_factor);
		color.destinationRGBBlendFactor = ConvertBlendFactor(extras.dst_factor);
		color.sourceAlphaBlendFactor = ConvertBlendFactor(extras.src_factor_alpha);
		color.destinationAlphaBlendFactor = ConvertBlendFactor(extras.dst_factor_alpha);
	}
	if (extras.has_rt1)
	{
		MTLRenderPipelineColorAttachmentDescriptor* color1 = [[pdesc colorAttachments] objectAtIndexedSubscript:1];
		[color1 setPixelFormat:MTLPixelFormatR32Float];
	}
	if (m_dev.features.rov_requires_rt && extras.rt == GSTexture::Format::Invalid && !extras.has_depth && !extras.has_stencil)
	{
		// Dummy texture
		[color setPixelFormat:MTLPixelFormatR8Unorm];
		[color setWriteMask:MTLColorWriteMaskNone];
	}
	NSString* pname = [NSString stringWithFormat:@"HW Render %x.%llx.%llx.%x", vssel_mtl.key, pssel.key_hi, pssel.key_lo, extras.fullkey];
	auto pipeline = MakePipeline(pdesc, vs, ps, pname);

	[m_current_render.encoder setRenderPipelineState:pipeline];
	m_hw_pipeline.insert(std::make_pair(fullsel, std::move(pipeline)));
}

void GSDeviceMTL::MRESetDSS(DepthStencilSelector sel)
{
	if (!m_current_render.depth_target || m_current_render.depth_sel.key == sel.key)
		return;
	[m_current_render.encoder setDepthStencilState:m_dss_hw[sel.key]];
	m_current_render.depth_sel = sel;
}

void GSDeviceMTL::MRESetDSS(id<MTLDepthStencilState> dss)
{
	[m_current_render.encoder setDepthStencilState:dss];
	m_current_render.depth_sel.key = -1;
}

void GSDeviceMTL::MRESetSampler(SamplerSelector sel)
{
	if (m_current_render.has.sampler && m_current_render.sampler_sel.key == sel.key)
		return;
	if (!m_sampler_hw[sel.key]) [[unlikely]]
		m_sampler_hw[sel.key] = CreateSampler(m_dev.dev, sel);
	[m_current_render.encoder setFragmentSamplerState:m_sampler_hw[sel.key] atIndex:0];
	m_current_render.sampler_sel = sel;
	m_current_render.has.sampler = true;
}

static void textureBarrier(id<MTLRenderCommandEncoder> enc)
{
#if TARGET_OS_IPHONE
#if !TARGET_OS_SIMULATOR
	[enc memoryBarrierWithScope:MTLBarrierScopeTextures
	                afterStages:MTLRenderStageFragment
	               beforeStages:MTLRenderStageFragment];
#endif
#else
	[enc memoryBarrierWithScope:MTLBarrierScopeRenderTargets
	                afterStages:MTLRenderStageFragment
	               beforeStages:MTLRenderStageFragment];
#endif
}

void GSDeviceMTL::MRESetTexture(GSTexture* tex, int pos)
{
	if (!tex || tex == m_current_render.tex[pos])
		return;
	m_current_render.tex[pos] = tex;
	GSTextureMTL* mtex = static_cast<GSTextureMTL*>(tex);
	[m_current_render.encoder setFragmentTexture:mtex->GetTexture() atIndex:pos];
	mtex->m_last_read = m_current_draw;
}

void GSDeviceMTL::MRESetTexture(id<MTLTexture> tex, int pos)
{
	GSTexture* gstex = reinterpret_cast<GSTexture*>(tex); // No one actually dereferences this
	if (!tex || gstex == m_current_render.tex[pos])
		return;
	m_current_render.tex[pos] = gstex;
	[m_current_render.encoder setFragmentTexture:tex atIndex:pos];
}

void GSDeviceMTL::MRESetVertices(id<MTLBuffer> buffer, size_t offset)
{
	if (m_current_render.vertex_buffer != buffer)
	{
		m_current_render.vertex_buffer = buffer;
		[m_current_render.encoder setVertexBuffer:buffer offset:offset atIndex:GSMTLBufferIndexHWVertices];
	}
	else
	{
		[m_current_render.encoder setVertexBufferOffset:offset atIndex:GSMTLBufferIndexHWVertices];
	}
}

void GSDeviceMTL::MRESetVSIndices(id<MTLBuffer> buffer, size_t offset)
{
	if (m_current_render.vs_index_buffer != buffer)
	{
		m_current_render.vs_index_buffer = buffer;
		[m_current_render.encoder setVertexBuffer:buffer offset:offset atIndex:GSMTLBufferIndexHWIndices];
	}
	else
	{
		[m_current_render.encoder setVertexBufferOffset:offset atIndex:GSMTLBufferIndexHWIndices];
	}
}

void GSDeviceMTL::MRESetScissor(const GSVector4i& scissor)
{
	if (m_current_render.has.scissor && (m_current_render.scissor == scissor).alltrue())
		return;
	MTLScissorRect r;
	r.x = scissor.x;
	r.y = scissor.y;
	r.width = scissor.width();
	r.height = scissor.height();
	[m_current_render.encoder setScissorRect:r];
	m_current_render.scissor = scissor;
	m_current_render.has.scissor = true;
}

void GSDeviceMTL::MREClearScissor()
{
	assert(!m_current_render.is_full_rov());
	if (!m_current_render.has.scissor)
		return;
	m_current_render.has.scissor = false;
	GSVector4i size = GSVector4i(0);
	if (m_current_render.color_target)   size = size.max_u32(GSVector4i(m_current_render.color_target  ->GetSize()));
	if (m_current_render.depth_target)   size = size.max_u32(GSVector4i(m_current_render.depth_target  ->GetSize()));
	if (m_current_render.stencil_target) size = size.max_u32(GSVector4i(m_current_render.stencil_target->GetSize()));
	MTLScissorRect r;
	r.x = 0;
	r.y = 0;
	r.width = size.x;
	r.height = size.y;
	[m_current_render.encoder setScissorRect:r];
}

void GSDeviceMTL::MRESetCB(const GSHWDrawConfig::VSConstantBuffer& cb)
{
	if (m_current_render.has.cb_vs && m_current_render.cb_vs == cb)
		return;
	[m_current_render.encoder setVertexBytes:&cb length:sizeof(cb) atIndex:GSMTLBufferIndexHWUniforms];
	m_current_render.has.cb_vs = true;
	m_current_render.cb_vs = cb;
}

void GSDeviceMTL::MRESetCB(const GSHWDrawConfig::PSConstantBuffer& cb)
{
	if (m_current_render.has.cb_ps && m_current_render.cb_ps == cb)
		return;
	[m_current_render.encoder setFragmentBytes:&cb length:sizeof(cb) atIndex:GSMTLBufferIndexHWUniforms];
	m_current_render.has.cb_ps = true;
	m_current_render.cb_ps = cb;
}

void GSDeviceMTL::MRESetBlendColor(u8 color)
{
	if (m_current_render.has.blend_color && m_current_render.blend_color == color)
		return;
	float fc = static_cast<float>(color) / 128.f;
	[m_current_render.encoder setBlendColorRed:fc green:fc blue:fc alpha:fc];
	m_current_render.has.blend_color = true;
	m_current_render.blend_color = color;
}

void GSDeviceMTL::MRESetPipeline(id<MTLRenderPipelineState> pipe)
{
	[m_current_render.encoder setRenderPipelineState:pipe];
	m_current_render.has.pipeline_sel = false;
}

// MARK: - HW Render

// Metal can't import GSDevice.h, but we should make sure the structs are at least compatible
static_assert(sizeof(GSVertex) == sizeof(GSMTLMainVertex));
static_assert(offsetof(GSVertex, ST)      == offsetof(GSMTLMainVertex, st));
static_assert(offsetof(GSVertex, RGBAQ.R) == offsetof(GSMTLMainVertex, rgba));
static_assert(offsetof(GSVertex, RGBAQ.Q) == offsetof(GSMTLMainVertex, q));
static_assert(offsetof(GSVertex, XYZ.X)   == offsetof(GSMTLMainVertex, xy));
static_assert(offsetof(GSVertex, XYZ.Z)   == offsetof(GSMTLMainVertex, z));
static_assert(offsetof(GSVertex, UV)      == offsetof(GSMTLMainVertex, uv));
static_assert(offsetof(GSVertex, FOG)     == offsetof(GSMTLMainVertex, fog));

static_assert(sizeof(GSHWDrawConfig::VSConstantBuffer) == sizeof(GSMTLMainVSUniform));
static_assert(sizeof(GSHWDrawConfig::PSConstantBuffer) == sizeof(GSMTLMainPSUniform));
static_assert(offsetof(GSHWDrawConfig::VSConstantBuffer, vertex_scale)     == offsetof(GSMTLMainVSUniform, vertex_scale));
static_assert(offsetof(GSHWDrawConfig::VSConstantBuffer, vertex_offset)    == offsetof(GSMTLMainVSUniform, vertex_offset));
static_assert(offsetof(GSHWDrawConfig::VSConstantBuffer, texture_scale)    == offsetof(GSMTLMainVSUniform, texture_scale));
static_assert(offsetof(GSHWDrawConfig::VSConstantBuffer, texture_offset)   == offsetof(GSMTLMainVSUniform, texture_offset));
static_assert(offsetof(GSHWDrawConfig::VSConstantBuffer, point_size)       == offsetof(GSMTLMainVSUniform, point_size));
static_assert(offsetof(GSHWDrawConfig::VSConstantBuffer, max_depth)        == offsetof(GSMTLMainVSUniform, max_depth));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, FogColor_AREF.x)  == offsetof(GSMTLMainPSUniform, fog_color));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, FogColor_AREF.a)  == offsetof(GSMTLMainPSUniform, aref));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, WH)               == offsetof(GSMTLMainPSUniform, wh));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, TA_MaxDepth_Af.x) == offsetof(GSMTLMainPSUniform, ta));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, TA_MaxDepth_Af.z) == offsetof(GSMTLMainPSUniform, max_depth));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, TA_MaxDepth_Af.w) == offsetof(GSMTLMainPSUniform, alpha_fix));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, FbMask)           == offsetof(GSMTLMainPSUniform, fbmask));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, HalfTexel)        == offsetof(GSMTLMainPSUniform, half_texel));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, MinMax)           == offsetof(GSMTLMainPSUniform, uv_min_max));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, STRange)          == offsetof(GSMTLMainPSUniform, st_range));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, ChannelShuffle)   == offsetof(GSMTLMainPSUniform, channel_shuffle));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, ChannelShuffleOffset) == offsetof(GSMTLMainPSUniform, channel_shuffle_offset));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, TCOffsetHack)     == offsetof(GSMTLMainPSUniform, tc_offset));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, STScale)          == offsetof(GSMTLMainPSUniform, st_scale));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, DitherMatrix)     == offsetof(GSMTLMainPSUniform, dither_matrix));
static_assert(offsetof(GSHWDrawConfig::PSConstantBuffer, ScaleFactor)      == offsetof(GSMTLMainPSUniform, scale_factor));

void GSDeviceMTL::SetupDestinationAlpha(GSTexture* rt, GSTexture* ds, const GSVector4i& r, SetDATM datm)
{
	FlushClears(rt);
	BeginRenderPass(@"Destination Alpha Setup", nullptr, MTLLoadActionDontCare, nullptr, MTLLoadActionDontCare, ds, MTLLoadActionDontCare);
	[m_current_render.encoder setStencilReferenceValue:1];
	MRESetDSS(m_dss_stencil_zero);
	RenderCopy(nullptr, m_stencil_clear_pipeline, r);
	MRESetDSS(m_dss_stencil_write);
	RenderCopy(rt, m_datm_pipeline[static_cast<u8>(datm)], r);
}

static id<MTLTexture> getTexture(GSTexture* tex)
{
	return tex ? static_cast<GSTextureMTL*>(tex)->GetTexture() : nil;
}

static bool usesStencil(GSHWDrawConfig::DestinationAlphaMode dstalpha)
{
	switch (dstalpha)
	{
		case GSHWDrawConfig::DestinationAlphaMode::Stencil:
		case GSHWDrawConfig::DestinationAlphaMode::StencilOne:
			return true;
		default:
			return false;
	}
}

void GSDeviceMTL::MREInitHWDraw(GSHWDrawConfig& config, const Map& verts)
{
	MRESetScissor(config.scissor);
	MRESetTexture(config.tex, GSMTLTextureIndexTex);
	MRESetTexture(config.pal, GSMTLTextureIndexPalette);
	MRESetSampler(config.sampler);
	MRESetCB(config.cb_vs);
	MRESetCB(config.cb_ps);
	MRESetVertices(verts.gpu_buffer, verts.gpu_offset);
	if (config.vs.UseVSExpandIndexBuffer())
		MRESetVSIndices(verts.gpu_buffer, verts.gpu_offset + config.nverts * sizeof(*config.verts));
}

__fi void GSDeviceMTL::PrepareROVTexture(GSTexture** ptex)
{
	GSTextureMTL* tex = static_cast<GSTextureMTL*>(*ptex);
	FlushClears(tex);
	tex->m_last_read  = m_current_draw;
	tex->m_last_write = m_current_draw;
	*ptex = nullptr;
}

void GSDeviceMTL::DoRenderHW(GSHWDrawConfig& config)
{ @autoreleasepool {
	if (config.tex && (config.ds == config.tex || config.rt == config.tex))
		EndRenderPass(); // Barrier

	size_t vertsize = config.nverts * sizeof(*config.verts);
	size_t idxsize = config.vs.UseFixedExpandIndexBuffer() ? 0 : (config.nindices * sizeof(*config.indices));
	Map allocation = Allocate(m_vertex_upload_buf, vertsize + idxsize);
	memcpy(allocation.cpu_buffer, config.verts, vertsize);

	id<MTLBuffer> index_buffer = nil;
	size_t index_buffer_offset = 0;
	if (!config.vs.UseFixedExpandIndexBuffer())
	{
		memcpy(static_cast<u8*>(allocation.cpu_buffer) + vertsize, config.indices, idxsize);
		if (config.vs.UseVSExpandIndexBuffer())
		{
			// VS expand index buffer is bound to the VS instead of the input assembler
			u32 expand = GetExpansionFactor(config.vs.expand);
			config.nindices *= expand;
			config.indices_per_prim *= expand;
		}
		else
		{
			index_buffer = allocation.gpu_buffer;
			index_buffer_offset = allocation.gpu_offset + vertsize;
		}
	}
	else
	{
		index_buffer = m_expand_index_buffer;
	}

	FlushClears(config.tex);
	FlushClears(config.pal);

	GSTexture* stencil = nullptr;
	GSTexture* primid_tex = nullptr;
	GSTexture* rt = config.rt;
	GSTexture* colclip_rt = g_gs_device->GetColorClipTexture();
	
	if (colclip_rt)
	{
		if (config.colclip_mode == GSHWDrawConfig::ColClipMode::EarlyResolve)
		{
			BeginRenderPass(@"ColorClip Resolve", config.rt, MTLLoadActionLoad, nullptr, MTLLoadActionDontCare);
			RenderCopy(colclip_rt, m_colclip_resolve_pipeline, config.colclip_update_area);

			Recycle(colclip_rt);
			
			g_gs_device->SetColorClipTexture(nullptr);
			
			colclip_rt = nullptr;
		}
		else
			config.ps.colclip_hw = 1;
	}
	
	if (config.ps.colclip_hw)
	{
		if (!colclip_rt)
		{
			config.colclip_update_area = config.drawarea;
			
			GSVector2i size = config.rt->GetSize();
			rt = colclip_rt = CreateFeedbackTarget(size.x, size.y, GSTexture::Format::ColorClip, false);
			
			g_gs_device->SetColorClipTexture(colclip_rt);
			
			const GSVector4i copy_rect = (config.colclip_mode == GSHWDrawConfig::ColClipMode::ConvertOnly) ? GSVector4i::loadh(size) : config.drawarea;
			
			switch (config.rt->GetState())
			{
				case GSTexture::State::Dirty:
					BeginRenderPass(@"ColorClip Init", colclip_rt, MTLLoadActionDontCare, nullptr, MTLLoadActionDontCare);
					RenderCopy(config.rt, m_colclip_init_pipeline, copy_rect);
					break;

				case GSTexture::State::Cleared:
				{
					BeginRenderPass(@"ColorClip Clear", colclip_rt, MTLLoadActionDontCare, nullptr, MTLLoadActionDontCare);
					GSVector4 color = GSVector4::rgba32(config.rt->GetClearColor()) / GSVector4::cxpr(65535, 65535, 65535, 255);
					[m_current_render.encoder setFragmentBytes:&color length:sizeof(color) atIndex:GSMTLBufferIndexUniforms];
					RenderCopy(nullptr, m_colclip_clear_pipeline, copy_rect);
					break;
				}

				case GSTexture::State::Invalidated:
					break;
			}
		}

		rt = colclip_rt;
	}
	
	switch (config.destination_alpha)
	{
		case GSHWDrawConfig::DestinationAlphaMode::Off:
		case GSHWDrawConfig::DestinationAlphaMode::Full:
			break; // No setup
		case GSHWDrawConfig::DestinationAlphaMode::PrimIDTracking:
		{
			FlushClears(rt);
			GSVector2i size = rt->GetSize();
			primid_tex = CreateRenderTarget(size.x, size.y, GSTexture::Format::PrimID);
			DepthStencilSelector dsel = config.depth;
			dsel.zwe = 0;
			GSTexture* depth = dsel.key == DepthStencilSelector::NoDepth().key ? nullptr : config.ds;
			BeginRenderPass(@"PrimID Destination Alpha Init", primid_tex, MTLLoadActionDontCare, depth, MTLLoadActionLoad);
			RenderCopy(rt, m_primid_init_pipeline[static_cast<bool>(depth)][static_cast<u8>(config.datm)], config.drawarea);
			MRESetDSS(dsel);
			pxAssert(config.ps.date == 1 || config.ps.date == 2);
			if (config.ps.tex_is_fb)
				MRESetTexture(rt, GSMTLTextureIndexRenderTarget);
			config.require_one_barrier = false; // Ending render pass is our barrier
			pxAssert(config.require_full_barrier == false && config.drawlist == nullptr);
			MRESetHWPipelineState(config.vs, config.ps, {}, {});
			MREInitHWDraw(config, allocation);
			SendHWDraw(config, m_current_render.encoder, index_buffer, index_buffer_offset, false, false);
			config.ps.date = 3;
			break;
		}
		case GSHWDrawConfig::DestinationAlphaMode::StencilOne:
			BeginRenderPass(@"Destination Alpha Stencil Clear", nullptr, MTLLoadActionDontCare, nullptr, MTLLoadActionDontCare, config.ds, MTLLoadActionDontCare);
			[m_current_render.encoder setStencilReferenceValue:1];
			MRESetDSS(m_dss_stencil_write);
			RenderCopy(nullptr, m_stencil_clear_pipeline, config.drawarea);
			stencil = config.ds;
			break;
		case GSHWDrawConfig::DestinationAlphaMode::Stencil:
			SetupDestinationAlpha(rt, config.ds, config.drawarea, config.datm);
			stencil = config.ds;
			break;
	}

	// Try to reduce render pass restarts
	if (!config.ps.HasColorROV() && !config.ds && m_current_render.color_target == rt && stencil == m_current_render.stencil_target && m_current_render.depth_target != config.tex)
		config.ds = m_current_render.depth_target;
	if (!rt && config.ds == m_current_render.depth_target && m_current_render.color_target != config.tex)
		rt = m_current_render.color_target;
	if (!rt && !config.ds)
	{
		// If we were rendering depth-only and depth gets cleared by the above check, that turns into rendering nothing, which should be a no-op
		pxAssertMsg(0, "DoRenderHW was given a completely useless draw call!");
		[m_current_render.encoder insertDebugSignpost:@"Skipped no-color no-depth draw"];
		if (primid_tex)
			Recycle(primid_tex);
		return;
	}

	const bool feedback_depth = config.ps.IsFeedbackLoopDepth();
	const bool rt1 = feedback_depth && !m_features.depth_feedback && !config.ps.HasDepthROV();
	GSTexture* rt_bind = rt;
	GSTexture* ds_bind = config.ds;
	GSTexture* rt_size = rt_bind ? rt_bind : ds_bind;
	if (config.ps.HasColorROV() && rt_bind) PrepareROVTexture(&rt_bind);
	if (config.ps.HasDepthROV() && ds_bind) PrepareROVTexture(&ds_bind);
	if (!rt_bind && !ds_bind && !stencil)
		BeginFullROV(@"RenderHWROV", rt_size->GetWidth(), rt_size->GetHeight());
	else
		BeginRenderPass(@"DoRenderHW", rt_bind, MTLLoadActionLoad, ds_bind, MTLLoadActionLoad, stencil, MTLLoadActionLoad, rt1);
	id<MTLRenderCommandEncoder> mtlenc = m_current_render.encoder;
	FlushDebugEntries(mtlenc);
	if (usesStencil(config.destination_alpha))
		[mtlenc setStencilReferenceValue:1];
	MREInitHWDraw(config, allocation);
	if (config.ps.HasColorROV() && m_dev.features.rov_requires_r32)
		MRESetTexture(reinterpret_cast<GSTextureMTL*>(rt)->GetROVTexture(), GSMTLTextureIndexRenderTarget);
	else if (config.require_one_barrier || config.require_full_barrier || config.ps.HasColorROV())
		MRESetTexture(rt, GSMTLTextureIndexRenderTarget);
	if (config.ps.HasDepthROV())
	{
		MRESetTexture(config.ds, GSMTLTextureIndexDepthTarget);
	}
	else if (feedback_depth)
	{
		GSTexture* tex = !m_features.depth_feedback && !m_features.framebuffer_fetch ? m_ds_as_rt : config.ds;
		MRESetTexture(tex, GSMTLTextureIndexDepthTarget);
	}
	if (primid_tex)
		MRESetTexture(primid_tex, GSMTLTextureIndexPrimIDs);
	if (config.blend.constant_enable)
		MRESetBlendColor(config.blend.constant);
	MRESetHWPipelineState(config.vs, config.ps, config.blend, config.colormask);
	MRESetDSS(config.depth);

	SendHWDraw(config, mtlenc, index_buffer, index_buffer_offset, config.require_one_barrier, config.require_full_barrier);

	if (config.alpha_second_pass.enable)
	{
		if (config.alpha_second_pass.ps_aref != config.cb_ps.FogColor_AREF.a)
		{
			config.cb_ps.FogColor_AREF.a = config.alpha_second_pass.ps_aref;
			MRESetCB(config.cb_ps);
		}
		MRESetHWPipelineState(config.vs, config.alpha_second_pass.ps, config.blend, config.alpha_second_pass.colormask);
		MRESetDSS(config.alpha_second_pass.depth);
		SendHWDraw(config, mtlenc, index_buffer, index_buffer_offset, config.alpha_second_pass.require_one_barrier, config.alpha_second_pass.require_full_barrier);
	}

	if (colclip_rt)
	{
		config.colclip_update_area = config.colclip_update_area.runion(config.drawarea);

		if ((config.colclip_mode == GSHWDrawConfig::ColClipMode::ResolveOnly || config.colclip_mode == GSHWDrawConfig::ColClipMode::ConvertAndResolve))
		{
			BeginRenderPass(@"ColorClip Resolve", config.rt, MTLLoadActionLoad, nullptr, MTLLoadActionDontCare);
			RenderCopy(colclip_rt, m_colclip_resolve_pipeline, config.colclip_update_area);

			Recycle(colclip_rt);
			
			SetColorClipTexture(nullptr);
		}
	}

	if (primid_tex)
		Recycle(primid_tex);
}}

static void EncodeDraw(id<MTLRenderCommandEncoder> enc, MTLPrimitiveType topology, size_t count, id<MTLBuffer> indices, size_t off, size_t base_vertex)
{
	if (indices)
	{
		[enc drawIndexedPrimitives:topology
		                indexCount:count
		                 indexType:MTLIndexTypeUInt16
		               indexBuffer:indices
		         indexBufferOffset:off + base_vertex * sizeof(uint16_t)];
	}
	else
	{
		[enc drawPrimitives:topology
		        vertexStart:base_vertex
		        vertexCount:count];
	}
}

void GSDeviceMTL::SendHWDraw(GSHWDrawConfig& config, id<MTLRenderCommandEncoder> enc, id<MTLBuffer> buffer, size_t off, bool one_barrier, bool full_barrier)
{
	MTLPrimitiveType topology;
	switch (config.topology)
	{
		case GSHWDrawConfig::Topology::Point:    topology = MTLPrimitiveTypePoint;    break;
		case GSHWDrawConfig::Topology::Line:     topology = MTLPrimitiveTypeLine;     break;
		case GSHWDrawConfig::Topology::Triangle: topology = MTLPrimitiveTypeTriangle; break;
	}

	if (!m_features.texture_barrier) [[unlikely]]
	{
		EncodeDraw(enc, topology, config.nindices, buffer, off, 0);
		g_perfmon.Put(GSPerfMon::DrawCalls, 1);
		return;
	}


	if (full_barrier)
	{
		pxAssert(config.drawlist && !config.drawlist->empty());

		[enc pushDebugGroup:[NSString stringWithFormat:@"Full barrier split draw (%d primitives in %zu groups)", config.nindices / config.indices_per_prim, config.drawlist->size()]];
#if defined(_DEBUG)
		// Check how draw call is split.
		std::map<size_t, size_t> frequency;
		for (const auto& it : *config.drawlist)
			++frequency[it];

		std::string message;
		for (const auto& it : frequency)
			message += " " + std::to_string(it.first) + "(" + std::to_string(it.second) + ")";

		[enc insertDebugSignpost:[NSString stringWithFormat:@"Split single draw (%d primitives) into %zu draws: consecutive draws(frequency):%s",
			config.nindices / config.indices_per_prim, config.drawlist->size(), message.c_str()]];
#endif


		g_perfmon.Put(GSPerfMon::DrawCalls, config.drawlist->size());
		g_perfmon.Put(GSPerfMon::Barriers, config.drawlist->size());

		const u32 indices_per_prim = config.indices_per_prim;
		const u32 draw_list_size = static_cast<u32>(config.drawlist->size());

		for (u32 n = 0, p = 0; n < draw_list_size; n++)
		{
			const size_t count = config.drawlist->at(n) * indices_per_prim;
			textureBarrier(enc);
			EncodeDraw(enc, topology, count, buffer, off, p);
			p += count;
		}

		[enc popDebugGroup];
		return;
	}
	else if (one_barrier)
	{
		// One barrier needed
		textureBarrier(enc);
		g_perfmon.Put(GSPerfMon::Barriers, 1);
	}

	EncodeDraw(enc, topology, config.nindices, buffer, off, 0);
	g_perfmon.Put(GSPerfMon::DrawCalls, 1);
}

// tbh I'm not a fan of the current debug groups
// not much useful information and makes things harder to find
// good to turn on if you're debugging tc stuff though
#ifndef MTL_ENABLE_DEBUG
	#define MTL_ENABLE_DEBUG 0
#endif

void GSDeviceMTL::PushDebugGroup(const char* fmt, ...)
{
#if MTL_ENABLE_DEBUG
	va_list va;
	va_start(va, fmt);
	MRCOwned<NSString*> nsfmt = MRCTransfer([[NSString alloc] initWithUTF8String:fmt]);
	m_debug_entries.emplace_back(DebugEntry::Push, MRCTransfer([[NSString alloc] initWithFormat:nsfmt arguments:va]));
	va_end(va);
#endif
}

void GSDeviceMTL::PopDebugGroup()
{
#if MTL_ENABLE_DEBUG
	m_debug_entries.emplace_back(DebugEntry::Pop, nullptr);
#endif
}

void GSDeviceMTL::InsertDebugMessage(DebugMessageCategory category, const char* fmt, ...)
{
#if MTL_ENABLE_DEBUG
	va_list va;
	va_start(va, fmt);
	MRCOwned<NSString*> nsfmt = MRCTransfer([[NSString alloc] initWithUTF8String:fmt]);
	m_debug_entries.emplace_back(DebugEntry::Insert, MRCTransfer([[NSString alloc] initWithFormat:nsfmt arguments:va]));
	va_end(va);
#endif
}

void GSDeviceMTL::ProcessDebugEntry(id<MTLCommandEncoder> enc, const DebugEntry& entry)
{
	switch (entry.op)
	{
		case DebugEntry::Push:
			[enc pushDebugGroup:entry.str];
			m_debug_group_level++;
			break;
		case DebugEntry::Pop:
			[enc popDebugGroup];
			if (m_debug_group_level > 0)
				m_debug_group_level--;
			break;
		case DebugEntry::Insert:
			[enc insertDebugSignpost:entry.str];
			break;
	}
}

void GSDeviceMTL::FlushDebugEntries(id<MTLCommandEncoder> enc)
{
#if MTL_ENABLE_DEBUG
	if (!m_debug_entries.empty())
	{
		for (const DebugEntry& entry : m_debug_entries)
		{
			ProcessDebugEntry(enc, entry);
		}
		m_debug_entries.clear();
	}
#endif
}

void GSDeviceMTL::EndDebugGroup(id<MTLCommandEncoder> enc)
{
#if MTL_ENABLE_DEBUG
	if (!m_debug_entries.empty() && m_debug_group_level)
	{
		auto begin = m_debug_entries.begin();
		auto cur = begin;
		auto end = m_debug_entries.end();
		while (cur != end && m_debug_group_level)
		{
			ProcessDebugEntry(enc, *cur);
			cur++;
		}
		m_debug_entries.erase(begin, cur);
	}
#endif
}

static simd::float2 ToSimd(const ImVec2& vec)
{
	return simd::make_float2(vec.x, vec.y);
}

static simd::float4 ToSimd(const ImVec4& vec)
{
	return simd::make_float4(vec.x, vec.y, vec.z, vec.w);
}

void GSDeviceMTL::RenderImGui(ImDrawData* data)
{
	if (data->CmdListsCount == 0)
		return;
	UpdateImGuiTextures();
	simd::float4 transform;
	transform.xy = 2.f / simd::make_float2(data->DisplaySize.x, -data->DisplaySize.y);
	transform.zw = ToSimd(data->DisplayPos) * -transform.xy + simd::make_float2(-1, 1);
	id<MTLRenderCommandEncoder> enc = m_current_render.encoder;
	[enc pushDebugGroup:@"ImGui"];

	Map map = Allocate(m_vertex_upload_buf, data->TotalVtxCount * sizeof(ImDrawVert) + data->TotalIdxCount * sizeof(ImDrawIdx));
	size_t vtx_off = 0;
	size_t idx_off = data->TotalVtxCount * sizeof(ImDrawVert);

	[enc setRenderPipelineState:m_imgui_pipeline];
	[enc setVertexBuffer:map.gpu_buffer offset:map.gpu_offset atIndex:GSMTLBufferIndexVertices];
	[enc setVertexBytes:&transform length:sizeof(transform) atIndex:GSMTLBufferIndexUniforms];

	simd::uint4 last_scissor = simd::make_uint4(0, 0, GetWindowWidth(), GetWindowHeight());
	simd::float2 fb_size = simd_float(last_scissor.zw);
	simd::float2 clip_off   = ToSimd(data->DisplayPos);       // (0,0) unless using multi-viewports
	simd::float2 clip_scale = ToSimd(data->FramebufferScale); // (1,1) unless using retina display which are often (2,2)
	ImTextureID last_tex = reinterpret_cast<ImTextureID>(nullptr);

	for (int i = 0; i < data->CmdListsCount; i++)
	{
		const ImDrawList* cmd_list = data->CmdLists[i];
		size_t vtx_size = cmd_list->VtxBuffer.Size * sizeof(ImDrawVert);
		size_t idx_size = cmd_list->IdxBuffer.Size * sizeof(ImDrawIdx);
		memcpy(static_cast<char*>(map.cpu_buffer) + vtx_off, cmd_list->VtxBuffer.Data, vtx_size);
		memcpy(static_cast<char*>(map.cpu_buffer) + idx_off, cmd_list->IdxBuffer.Data, idx_size);

		for (const ImDrawCmd& cmd : cmd_list->CmdBuffer)
		{
			if (cmd.UserCallback)
				[NSException raise:@"Unimplemented" format:@"UserCallback not implemented"];
			if (!cmd.ElemCount)
				continue;

			simd::float4 clip_rect = (ToSimd(cmd.ClipRect) - clip_off.xyxy) * clip_scale.xyxy;
			simd::float2 clip_min = clip_rect.xy;
			simd::float2 clip_max = clip_rect.zw;
			clip_min = simd::max(clip_min, simd::float2(0));
			clip_max = simd::min(clip_max, fb_size);
			if (simd::any(clip_min >= clip_max))
				continue;
			simd::uint4 scissor = simd::make_uint4(simd_uint(clip_min), simd_uint(clip_max - clip_min));
			ImTextureID tex = cmd.GetTexID();
			if (simd::any(scissor != last_scissor))
			{
				last_scissor = scissor;
				[enc setScissorRect:(MTLScissorRect){ .x = scissor.x, .y = scissor.y, .width = scissor.z, .height = scissor.w }];
			}
			if (tex != last_tex)
			{
				last_tex = tex;
				[enc setFragmentTexture:(__bridge id<MTLTexture>)tex atIndex:0];
			}

			[enc setVertexBufferOffset:map.gpu_offset + vtx_off + cmd.VtxOffset * sizeof(ImDrawVert) atIndex:0];
			[enc drawIndexedPrimitives:MTLPrimitiveTypeTriangle
			                indexCount:cmd.ElemCount
			                 indexType:sizeof(ImDrawIdx) == 2 ? MTLIndexTypeUInt16 : MTLIndexTypeUInt32
			               indexBuffer:map.gpu_buffer
			         indexBufferOffset:map.gpu_offset + idx_off + cmd.IdxOffset * sizeof(ImDrawIdx)];
		}

		vtx_off += vtx_size;
		idx_off += idx_size;
	}

	[enc popDebugGroup];
}

#endif // __APPLE__