diff --git a/Common/GPU/Metal/MetalRenderManager.h b/Common/GPU/Metal/MetalRenderManager.h
index 4f1bcb0b6f..f30d2b8676 100644
--- a/Common/GPU/Metal/MetalRenderManager.h
+++ b/Common/GPU/Metal/MetalRenderManager.h
@@ -48,6 +48,8 @@
 // bottom. Vertex data goes above all of them.
 enum {
 	METAL_VERTEX_BUFFER_INDEX = 16,
+	// Push constants (only compute shaders have them) go in a buffer of their own, set with setBytes.
+	METAL_PUSH_CONSTANT_INDEX = 15,
 	METAL_MAX_TEXTURES = 8,
 	METAL_MAX_BUFFERS = 8,
 };
@@ -183,6 +185,7 @@ enum class MetalCmd : uint8_t {
 	DRAW,
 	DRAW_INDEXED,
 	COPY,
+	DISPATCH,
 };
 
 // Objects are unretained. They're kept alive by their owners or the deferred release list until the GPU is
@@ -211,6 +214,7 @@ struct MetalCommand {
 		} drawIndexed;
 		uint32_t pass;  // Index into the chunk's passes.
 		uint32_t copy;  // Index into the chunk's copies.
+		uint32_t dispatch;  // Index into the chunk's dispatches.
 	};
 };
 
@@ -255,6 +259,19 @@ struct MetalRecordedCopy {
 	MTLBlitOption options;
 };
 
+// A compute dispatch. Only used for uploads, so far (texture scaling). Resources go at their Vulkan binding
+// numbers, like for draws.
+struct MetalRecordedDispatch {
+	void *pipeline;  // id<MTLComputePipelineState>
+	void *textures[METAL_MAX_TEXTURES];
+	void *buffers[METAL_MAX_BUFFERS];
+	uint32_t bufferOffsets[METAL_MAX_BUFFERS];
+	uint8_t pushConstants[16];
+	uint32_t pushConstantSize;
+	MTLSize threadgroups;
+	MTLSize threadsPerThreadgroup;
+};
+
 enum class MetalChunkEnd {
 	FLUSH,
 	PRESENT,
@@ -279,6 +296,7 @@ struct MetalCommandChunk {
 	std::vector<MetalCommand> uploads;  // Run in a command buffer of their own, before the commands.
 	std::vector<MetalRecordedPass> passes;
 	std::vector<MetalRecordedCopy> copies;
+	std::vector<MetalRecordedDispatch> dispatches;
 	MetalChunkEnd end = MetalChunkEnd::FLUSH;
 	bool vsync = true;
 	int newDrawableWidth = 0;  // Non-zero to resize the layer's drawables first.
@@ -291,6 +309,7 @@ struct MetalCommandChunk {
 		uploads.clear();
 		passes.clear();
 		copies.clear();
+		dispatches.clear();
 		end = MetalChunkEnd::FLUSH;
 		vsync = true;
 		newDrawableWidth = 0;
@@ -405,6 +424,8 @@ public:
 	void CopyTextureToBuffer(__unsafe_unretained id<MTLTexture> src, int level, MTLOrigin origin, MTLSize size, __unsafe_unretained id<MTLBuffer> dst, uint32_t offset, uint32_t bytesPerRow, uint32_t bytesPerImage, MTLBlitOption options);
 	void CopyBufferToTexture(__unsafe_unretained id<MTLBuffer> src, uint32_t offset, uint32_t bytesPerRow, uint32_t bytesPerImage, MTLSize size, __unsafe_unretained id<MTLTexture> dst, int slice, int level, MTLOrigin origin, MTLBlitOption options, bool upload);
 	void GenerateMips(__unsafe_unretained id<MTLTexture> texture, bool upload);
+	// Runs in the upload command buffer, in order with the upload copies. Fill in the returned dispatch.
+	MetalRecordedDispatch &RecordUploadDispatch();
 
 	const MetalStats &Stats() const { return lastStats_; }
 
@@ -435,6 +456,7 @@ private:
 	void ReplayPass(const MetalCommandChunk &chunk, MetalStep &step, id<MTLCommandBuffer> cmd);
 	void ReplayCopy(const MetalCommandChunk &chunk, const MetalRecordedCopy &copy, id<MTLCommandBuffer> cmd, bool upload);
 	void EndReplayEncoders();
+	void ReplayDispatch(const MetalRecordedDispatch &dispatch, id<MTLCommandBuffer> cmd);
 	id<MTLTexture> ReplayBackbuffer();
 
 	id<MTLDevice> device_;
@@ -520,6 +542,7 @@ private:
 	id<MTLCommandBuffer> rtUploadCmd_;
 	id<MTLBlitCommandEncoder> rtBlit_;
 	id<MTLBlitCommandEncoder> rtUploadBlit_;
+	id<MTLComputeCommandEncoder> rtUploadCompute_;
 	id<CAMetalDrawable> rtDrawable_;
 	// Replay scratch, reused between chunks.
 	std::vector<MetalStep> rtSteps_;
diff --git a/Common/GPU/Metal/MetalRenderManager.mm b/Common/GPU/Metal/MetalRenderManager.mm
index 256503b91e..a606a20efa 100644
--- a/Common/GPU/Metal/MetalRenderManager.mm
+++ b/Common/GPU/Metal/MetalRenderManager.mm
@@ -187,6 +187,14 @@ bool TranslateVulkanGLSLToMSL(ShaderStage stage, const char *glsl, MetalTranslat
 		for (const auto &resource : resources.storage_images) {
 			addImage(resource);
 		}
+		if (!resources.push_constant_buffers.empty()) {
+			spirv_cross::MSLResourceBinding binding;
+			binding.stage = model;
+			binding.desc_set = spirv_cross::kPushConstDescSet;
+			binding.binding = spirv_cross::kPushConstBinding;
+			binding.msl_buffer = METAL_PUSH_CONSTANT_INDEX;
+			msl.add_msl_resource_binding(binding);
+		}
 		out->readsFramebuffer = !resources.subpass_inputs.empty();
 
 		out->msl = msl.compile();
@@ -688,6 +696,17 @@ void MetalRenderManager::GenerateMips(__unsafe_unretained id<MTLTexture> texture
 	copy.dst = (__bridge void *)texture;
 }
 
+MetalRecordedDispatch &MetalRenderManager::RecordUploadDispatch() {
+	chunk_->dispatches.emplace_back();
+	MetalCommand c{};
+	c.cmd = MetalCmd::DISPATCH;
+	c.dispatch = (uint32_t)chunk_->dispatches.size() - 1;
+	chunk_->uploads.push_back(c);
+	MetalRecordedDispatch &dispatch = chunk_->dispatches.back();
+	dispatch = {};
+	return dispatch;
+}
+
 void MetalRenderManager::SetPipeline(__unsafe_unretained MetalPipelineState *pipeline) {
 	if (state_.pipeline != pipeline) {
 		Record(MetalCmd::PIPELINE).bind.obj = (__bridge void *)pipeline;
@@ -917,6 +936,38 @@ void MetalRenderManager::EndReplayEncoders() {
 		[rtUploadBlit_ endEncoding];
 		rtUploadBlit_ = nil;
 	}
+	if (rtUploadCompute_) {
+		[rtUploadCompute_ endEncoding];
+		rtUploadCompute_ = nil;
+	}
+}
+
+// Metal orders the encoders' accesses to the (tracked) textures, so there are no barriers to place.
+void MetalRenderManager::ReplayDispatch(const MetalRecordedDispatch &dispatch, id<MTLCommandBuffer> cmd) {
+	if (rtUploadBlit_) {
+		[rtUploadBlit_ endEncoding];
+		rtUploadBlit_ = nil;
+	}
+	if (!rtUploadCompute_) {
+		rtUploadCompute_ = [cmd computeCommandEncoder];
+		rtUploadCompute_.label = @"upload compute";
+	}
+	id<MTLComputeCommandEncoder> enc = rtUploadCompute_;
+	[enc setComputePipelineState:(__bridge id<MTLComputePipelineState>)dispatch.pipeline];
+	for (int i = 0; i < METAL_MAX_TEXTURES; i++) {
+		if (dispatch.textures[i]) {
+			[enc setTexture:(__bridge id<MTLTexture>)dispatch.textures[i] atIndex:i];
+		}
+	}
+	for (int i = 0; i < METAL_MAX_BUFFERS; i++) {
+		if (dispatch.buffers[i]) {
+			[enc setBuffer:(__bridge id<MTLBuffer>)dispatch.buffers[i] offset:dispatch.bufferOffsets[i] atIndex:i];
+		}
+	}
+	if (dispatch.pushConstantSize) {
+		[enc setBytes:dispatch.pushConstants length:dispatch.pushConstantSize atIndex:METAL_PUSH_CONSTANT_INDEX];
+	}
+	[enc dispatchThreadgroups:dispatch.threadgroups threadsPerThreadgroup:dispatch.threadsPerThreadgroup];
 }
 
 void MetalRenderManager::Replay(MetalCommandChunk *chunk) {
@@ -933,7 +984,15 @@ void MetalRenderManager::Replay(MetalCommandChunk *chunk) {
 	[cmd enqueue];
 
 	for (const MetalCommand &c : chunk->uploads) {
-		ReplayCopy(*chunk, chunk->copies[c.copy], uploadCmd, true);
+		if (c.cmd == MetalCmd::DISPATCH) {
+			ReplayDispatch(chunk->dispatches[c.dispatch], uploadCmd);
+		} else {
+			if (rtUploadCompute_) {
+				[rtUploadCompute_ endEncoding];
+				rtUploadCompute_ = nil;
+			}
+			ReplayCopy(*chunk, chunk->copies[c.copy], uploadCmd, true);
+		}
 	}
 	BuildSteps(*chunk);
 	MergeSteps(*chunk);
diff --git a/GPU/Metal/TextureCacheMetal.h b/GPU/Metal/TextureCacheMetal.h
index 644023c13e..edff6473eb 100644
--- a/GPU/Metal/TextureCacheMetal.h
+++ b/GPU/Metal/TextureCacheMetal.h
@@ -19,10 +19,14 @@
 
 // Objective-C++ only.
 
+#include <string>
 #include <unordered_map>
+#include <vector>
 
 #include "Common/GPU/Metal/MetalRenderManager.h"
+#include "Common/File/Path.h"
 #include "GPU/Common/TextureCacheCommon.h"
+#include "GPU/Common/TextureShaderCommon.h"
 
 class FramebufferManagerMetal;
 class ShaderManagerMetal;
@@ -41,6 +45,7 @@ public:
 
 	void DeviceLost() override;
 	void DeviceRestore(Draw::DrawContext *draw) override;
+	void NotifyConfigChanged() override;
 
 	void *GetNativeTextureView(const TexCacheEntry *entry, bool flat) const override;
 
@@ -72,6 +77,14 @@ private:
 		return (__bridge id<MTLTexture>)entry->texturePtr;
 	}
 
+	// Hardware texture scaling, with the same compute shaders as the Vulkan backend.
+	void CompileScalingShader();
+	void ClearScalingShaders();
+	bool HasScalingShader() const { return singlePassCS_ != nil || !multipassCS_.empty(); }
+	id<MTLComputePipelineState> CompileComputeShader(const std::string &source, const char *tag);
+	// Records the dispatches that scale a decoded level (RGBA8, packed) into dst, a single level view.
+	bool ScaleBufferToTexture(MetalPushAlloc src, id<MTLTexture> dst, int srcWidth, int srcHeight, int dstWidth, int dstHeight);
+
 	MetalRenderManager *rm_ = nullptr;
 	DrawEngineMetal *drawEngine_ = nullptr;
 
@@ -79,4 +92,10 @@ private:
 
 	id<MTLTexture> boundTexture_;
 	id<MTLSamplerState> boundSampler_;
+
+	std::string textureShader_;
+	id<MTLComputePipelineState> singlePassCS_;
+	std::vector<id<MTLComputePipelineState>> multipassCS_;
+	std::vector<TextureScaleScratchDesc> multipassScratchDescs_;
+	std::vector<TextureScaleStageDesc> multipassStageDescs_;
 };
diff --git a/GPU/Metal/TextureCacheMetal.mm b/GPU/Metal/TextureCacheMetal.mm
index 8a0d9509cf..caacf4e4ab 100644
--- a/GPU/Metal/TextureCacheMetal.mm
+++ b/GPU/Metal/TextureCacheMetal.mm
@@ -21,6 +21,9 @@
 #include <cfloat>
 
 #include "Common/Log.h"
+#include "Common/StringUtils.h"
+#include "Common/File/VFS/VFS.h"
+#include "GPU/Common/PostShader.h"
 #include "Common/GPU/thin3d.h"
 #include "Core/Config.h"
 #include "Core/ConfigValues.h"
@@ -94,6 +97,176 @@ static Draw::DataFormat GetDestFormat(GETextureFormat format, GEPaletteFormat cl
 
 TextureCacheMetal::TextureCacheMetal(Draw::DrawContext *draw, Draw2D *draw2D) : TextureCacheCommon(draw, draw2D) {
 	rm_ = (MetalRenderManager *)draw->GetNativeObject(Draw::NativeObject::RENDER_MANAGER);
+	CompileScalingShader();
+}
+
+void TextureCacheMetal::NotifyConfigChanged() {
+	TextureCacheCommon::NotifyConfigChanged();
+	CompileScalingShader();
+}
+
+static std::string ReadShaderSrc(const Path &filename) {
+	size_t size = 0;
+	uint8_t *data = g_VFS.ReadFile(filename.c_str(), &size);
+	if (!data) {
+		return std::string();
+	}
+	std::string src((const char *)data, size);
+	delete[] data;
+	return src;
+}
+
+void TextureCacheMetal::ClearScalingShaders() {
+	singlePassCS_ = nil;
+	multipassCS_.clear();
+	multipassScratchDescs_.clear();
+	multipassStageDescs_.clear();
+	textureShader_.clear();
+	shaderScaleFactor_ = 0;
+}
+
+id<MTLComputePipelineState> TextureCacheMetal::CompileComputeShader(const std::string &source, const char *tag) {
+	MetalTranslatedShader translated;
+	std::string error;
+	if (!TranslateVulkanGLSLToMSL(ShaderStage::Compute, source.c_str(), &translated, &error)) {
+		ERROR_LOG(Log::G3D, "Failed to translate texture scaling shader %s: %s", tag, error.c_str());
+		return nil;
+	}
+	id<MTLFunction> function = CreateMetalFunction(rm_->Device(), translated, tag, &error);
+	if (!function) {
+		ERROR_LOG(Log::G3D, "Failed to compile texture scaling shader %s: %s", tag, error.c_str());
+		return nil;
+	}
+	NSError *nsError = nil;
+	id<MTLComputePipelineState> pipeline = [rm_->Device() newComputePipelineStateWithFunction:function error:&nsError];
+	if (!pipeline) {
+		ERROR_LOG(Log::G3D, "Failed to create texture scaling pipeline %s: %s", tag, nsError.localizedDescription.UTF8String);
+	}
+	return pipeline;
+}
+
+// Compiles synchronously, which only happens when the setting changes or the game starts.
+void TextureCacheMetal::CompileScalingShader() {
+	if (!rm_) {
+		return;
+	}
+	if (!g_Config.bTexHardwareScaling || g_Config.sTextureShaderName != textureShader_) {
+		ClearScalingShaders();
+	} else if (HasScalingShader()) {
+		return;
+	}
+	if (!g_Config.bTexHardwareScaling) {
+		return;
+	}
+
+	ReloadAllPostShaderInfo(draw_);
+	const TextureShaderInfo *shaderInfo = GetTextureShaderInfo(g_Config.sTextureShaderName);
+	if (!shaderInfo || shaderInfo->computeShaderFiles.empty()) {
+		return;
+	}
+
+	if (shaderInfo->computeShaderFiles.size() > 1) {
+		if (!GetTextureScaleMultipassLayout(shaderInfo->scaleFactor, &multipassScratchDescs_, &multipassStageDescs_)) {
+			ERROR_LOG(Log::G3D, "Unsupported multipass scale factor %d in section '%s'", shaderInfo->scaleFactor, shaderInfo->section.c_str());
+			return;
+		}
+		const size_t expectedShaderCount = shaderInfo->scaleFactor == 2 ? 3 : 4;
+		if (shaderInfo->computeShaderFiles.size() != expectedShaderCount) {
+			ERROR_LOG(Log::G3D, "Expected %d compute shader stages in section '%s', got %d", (int)expectedShaderCount, shaderInfo->section.c_str(), (int)shaderInfo->computeShaderFiles.size());
+			ClearScalingShaders();
+			return;
+		}
+		for (const Path &file : shaderInfo->computeShaderFiles) {
+			id<MTLComputePipelineState> pipeline = CompileComputeShader(ReadShaderSrc(file), shaderInfo->section.c_str());
+			if (!pipeline) {
+				ClearScalingShaders();
+				return;
+			}
+			multipassCS_.push_back(pipeline);
+		}
+	} else {
+		std::string source = StringFromFormat(textureScaleUploadShader, ReadShaderSrc(shaderInfo->computeShaderFiles[0]).c_str());
+		singlePassCS_ = CompileComputeShader(source, shaderInfo->section.c_str());
+		if (!singlePassCS_) {
+			return;
+		}
+	}
+
+	if (!shaderInfo->constantBuffer.empty()) {
+		// Only the smiley test shader has one, and it's going away.
+		ERROR_LOG(Log::G3D, "Texture scaling shaders with constant buffers aren't supported on Metal");
+		ClearScalingShaders();
+		return;
+	}
+
+	textureShader_ = g_Config.sTextureShaderName;
+	shaderScaleFactor_ = shaderInfo->scaleFactor;
+}
+
+bool TextureCacheMetal::ScaleBufferToTexture(MetalPushAlloc src, id<MTLTexture> dst, int srcWidth, int srcHeight, int dstWidth, int dstHeight) {
+	// The bindings of the shaders (Vulkan binding = Metal index): 0 the output image, 1 the input buffer,
+	// 3 the input image. All have 8x8 threadgroups.
+
+	if (singlePassCS_) {
+		MetalRecordedDispatch &dispatch = rm_->RecordUploadDispatch();
+		dispatch.threadsPerThreadgroup = MTLSizeMake(8, 8, 1);
+		dispatch.pipeline = (__bridge void *)singlePassCS_;
+		dispatch.textures[0] = (__bridge void *)dst;
+		dispatch.buffers[1] = (__bridge void *)src.buffer;
+		dispatch.bufferOffsets[1] = src.offset;
+		const int32_t params[2] = { srcWidth, srcHeight };
+		memcpy(dispatch.pushConstants, params, sizeof(params));
+		dispatch.pushConstantSize = sizeof(params);
+		// One thread per input pixel.
+		dispatch.threadgroups = MTLSizeMake((srcWidth + 7) / 8, (srcHeight + 7) / 8, 1);
+		return true;
+	}
+	if (multipassCS_.empty()) {
+		return false;
+	}
+
+	std::vector<id<MTLTexture>> scratch;
+	for (const TextureScaleScratchDesc &desc : multipassScratchDescs_) {
+		MTLTextureDescriptor *td = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatRGBA16Float width:srcWidth * desc.widthScale height:srcHeight * desc.heightScale mipmapped:NO];
+		td.usage = MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite;
+		td.storageMode = MTLStorageModePrivate;
+		id<MTLTexture> tex = [rm_->Device() newTextureWithDescriptor:td];
+		if (!tex) {
+			ERROR_LOG(Log::G3D, "Failed to allocate multipass scratch images");
+			return false;
+		}
+		tex.label = [NSString stringWithUTF8String:desc.tag];
+		// Used by this frame's upload only.
+		rm_->DeferRelease(tex);
+		scratch.push_back(tex);
+	}
+
+	for (const TextureScaleStageDesc &stage : multipassStageDescs_) {
+		if (stage.shaderIndex >= multipassCS_.size()) {
+			return false;
+		}
+		MetalRecordedDispatch &dispatch = rm_->RecordUploadDispatch();
+		dispatch.threadsPerThreadgroup = MTLSizeMake(8, 8, 1);
+		dispatch.pipeline = (__bridge void *)multipassCS_[stage.shaderIndex];
+		if (stage.inputScratch < 0) {
+			dispatch.buffers[1] = (__bridge void *)src.buffer;
+			dispatch.bufferOffsets[1] = src.offset;
+		} else {
+			dispatch.textures[3] = (__bridge void *)scratch[stage.inputScratch];
+		}
+		dispatch.textures[0] = (__bridge void *)(stage.outputScratch < 0 ? dst : scratch[stage.outputScratch]);
+		const int32_t params[4] = {
+			srcWidth * stage.srcWidthScale,
+			srcHeight * stage.srcHeightScale,
+			stage.useFinalOutputSize ? dstWidth : srcWidth * stage.dstWidthScale,
+			stage.useFinalOutputSize ? dstHeight : srcHeight * stage.dstHeightScale,
+		};
+		memcpy(dispatch.pushConstants, params, sizeof(params));
+		dispatch.pushConstantSize = sizeof(params);
+		// One thread per output pixel.
+		dispatch.threadgroups = MTLSizeMake((params[2] + 7) / 8, (params[3] + 7) / 8, 1);
+	}
+	return true;
 }
 
 TextureCacheMetal::~TextureCacheMetal() {
@@ -107,6 +280,7 @@ void TextureCacheMetal::SetFramebufferManager(FramebufferManagerMetal *fbManager
 void TextureCacheMetal::DeviceLost() {
 	TextureCacheCommon::DeviceLost();
 	samplerCache_.clear();
+	ClearScalingShaders();
 	SetBoundTexture(nil);
 	SetBoundSampler(nil);
 	draw_ = nullptr;
@@ -116,6 +290,7 @@ void TextureCacheMetal::DeviceLost() {
 void TextureCacheMetal::DeviceRestore(Draw::DrawContext *draw) {
 	rm_ = (MetalRenderManager *)draw->GetNativeObject(Draw::NativeObject::RENDER_MANAGER);
 	TextureCacheCommon::DeviceRestore(draw);
+	CompileScalingShader();
 }
 
 void TextureCacheMetal::ReleaseTexture(TexCacheEntry *entry, bool delete_them) {
@@ -180,6 +355,8 @@ void TextureCacheMetal::ApplySamplerByKey(const SamplerCacheKey &key) {
 
 void TextureCacheMetal::BuildTexture(TexCacheEntry *const entry) {
 	BuildTexturePlan plan;
+	plan.hardwareScaling = g_Config.bTexHardwareScaling && HasScalingShader();
+	plan.slowScaler = !plan.hardwareScaling;
 	if (!PrepareBuildTexture(plan, entry)) {
 		ReleaseTexture(entry, true);
 		return;
@@ -209,7 +386,6 @@ void TextureCacheMetal::BuildTexture(TexCacheEntry *const entry) {
 	const bool compressed = Draw::DataFormatIsBlockCompressed(dstFmt, &blockSize);
 	const int bpp = (int)Draw::DataFormatSizeInBytes(dstFmt);
 
-	// We don't generate mips (yet), so only create the ones we load.
 	int levels = plan.depth == 1 ? plan.levelsToCreate : 1;
 	const int levelsToLoad = plan.depth == 1 ? std::min(plan.levelsToCreate, plan.levelsToLoad) : 1;
 
@@ -225,7 +401,11 @@ void TextureCacheMetal::BuildTexture(TexCacheEntry *const entry) {
 	td.height = th;
 	td.depth = plan.depth;
 	td.mipmapLevelCount = levels;
-	td.usage = MTLTextureUsageShaderRead;
+	// Scaled with a compute shader, which decodes unscaled and writes each level.
+	const bool computeUpload = plan.hardwareScaling && plan.scaleFactor > 1 && dstFmt == Draw::DataFormat::R8G8B8A8_UNORM &&
+		plan.depth == 1 && !plan.doReplace && !plan.saveTexture;
+
+	td.usage = MTLTextureUsageShaderRead | (computeUpload ? MTLTextureUsageShaderWrite : 0);
 	td.storageMode = MTLStorageModePrivate;
 	td.swizzle = Swizzle16(mtlFmt);
 	id<MTLTexture> texture = [rm_->Device() newTextureWithDescriptor:td];
@@ -256,6 +436,24 @@ void TextureCacheMetal::BuildTexture(TexCacheEntry *const entry) {
 		if (plan.doReplace && compressed) {
 			size = std::max(size, (size_t)plan.replaced->GetLevelDataSizeAfterCopy(plan.baseLevelSrc + i));
 		}
+		if (computeUpload) {
+			const int srcWidth = gstate.getTextureWidth(i);
+			const int srcHeight = gstate.getTextureHeight(i);
+			const int srcStride = srcWidth * 4;
+			const size_t srcSize = (size_t)srcStride * srcHeight;
+			BuildTexturePlan unscaledPlan = plan;
+			unscaledPlan.scaleFactor = 1;
+			MetalPushAlloc alloc = rm_->Push().Allocate(srcSize, 16);
+			LoadTextureLevel(*entry, alloc.ptr, srcSize, srcStride, unscaledPlan, srcLevel, dstFmt, TexDecodeFlags{});
+			id<MTLTexture> view = [texture newTextureViewWithPixelFormat:mtlFmt textureType:MTLTextureType2D levels:NSMakeRange(i, 1) slices:NSMakeRange(0, 1)];
+			if (!ScaleBufferToTexture(alloc, view, srcWidth, srcHeight, mipWidth, mipHeight)) {
+				// Leaves the level empty.
+				WARN_LOG_ONCE(mtlscalefail, Log::G3D, "Hardware texture scaling failed");
+			}
+			rm_->DeferRelease(view);
+			continue;
+		}
+
 		MetalPushAlloc alloc = rm_->Push().Allocate(size, 16);
 		LoadTextureLevel(*entry, alloc.ptr, size, stride, plan, srcLevel, dstFmt, TexDecodeFlags{});
 
diff --git a/UI/GameSettingsScreen.cpp b/UI/GameSettingsScreen.cpp
index de3c6b1fa6..02f61a333d 100644
--- a/UI/GameSettingsScreen.cpp
+++ b/UI/GameSettingsScreen.cpp
@@ -163,8 +163,7 @@ void GameSettingsScreen::PreCreateViews() {
 }
 
 static bool UsingHardwareTextureScaling() {
-	// For now, Vulkan only.
-	return g_Config.bTexHardwareScaling && GetGPUBackend() == GPUBackend::VULKAN && !g_Config.bSoftwareRendering;
+	return g_Config.bTexHardwareScaling && (GetGPUBackend() == GPUBackend::VULKAN || GetGPUBackend() == GPUBackend::METAL) && !g_Config.bSoftwareRendering;
 }
 
 static std::string TextureTranslateName(std::string_view value) {
@@ -519,7 +518,7 @@ void GameSettingsScreen::CreateGraphicsSettings(UI::ViewGroup *graphicsSettings)
 
 	graphicsSettings->Add(new ItemHeader(gr->T("Texture upscaling")));
 
-	if (GetGPUBackend() == GPUBackend::VULKAN) {
+	if (GetGPUBackend() == GPUBackend::VULKAN || GetGPUBackend() == GPUBackend::METAL) {
 		ChoiceWithValueDisplay *textureShaderChoice = graphicsSettings->Add(new ChoiceWithValueDisplay(&g_Config.sTextureShaderName, gr->T("GPU texture upscaler (fast)"), &TextureTranslateName));
 		textureShaderChoice->OnClick.Add([this](UI::EventParams &e) {
 			auto gr = GetI18NCategory(I18NCat::GRAPHICS);