From 8f5c63dd3c0f6aa64613bcb2d6047932897321ec Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 12:53:51 +0100 Subject: [PATCH 01/25] Preserve parent iOS deployment target --- Engine/CMakeLists.txt | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/Engine/CMakeLists.txt b/Engine/CMakeLists.txt index 5f2c9d65..6d462978 100644 --- a/Engine/CMakeLists.txt +++ b/Engine/CMakeLists.txt @@ -8,7 +8,9 @@ option(TEMPEST_BUILD_SHARED "Build shared Tempest." ON) set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS ON) set(CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/cmake") -set(CMAKE_OSX_DEPLOYMENT_TARGET 12.0) +if(NOT IOS) + set(CMAKE_OSX_DEPLOYMENT_TARGET 12.0) +endif() if(MSVC) option(TEMPEST_BUILD_DIRECTX12 "Build directx12 support" ON ) @@ -302,7 +304,9 @@ file(GLOB_RECURSE SOURCES ) if(APPLE OR IOS) - set(CMAKE_OSX_DEPLOYMENT_TARGET 12.0) + if(NOT IOS) + set(CMAKE_OSX_DEPLOYMENT_TARGET 12.0) + endif() enable_language(OBJCXX) file(GLOB_RECURSE ObjCSOURCES "*.mm" From c6a02d976ed72d46a031dfe68c9bbbf7bbb4af92 Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 13:04:41 +0100 Subject: [PATCH 02/25] Harden iOS window lifetime --- Engine/system/api/iosapi.mm | 18 ++++++++++++++++-- 1 file changed, 16 insertions(+), 2 deletions(-) diff --git a/Engine/system/api/iosapi.mm b/Engine/system/api/iosapi.mm index 356496ac..c3ec42f0 100644 --- a/Engine/system/api/iosapi.mm +++ b/Engine/system/api/iosapi.mm @@ -197,6 +197,10 @@ - (void)touchesEnded:(NSSet *)touches withEvent:(UIEvent *)ex { swapContext(); } } + +- (void)touchesCancelled:(NSSet *)touches withEvent:(UIEvent *)ex { + [self touchesEnded:touches withEvent:ex]; + } @end static TempestWindow* mainWindow = nullptr; @@ -211,7 +215,9 @@ @implementation ViewController { } -(id)init { - fullScreen = true; + self = [super init]; + if(self!=nil) + fullScreen = true; return self; } @@ -318,7 +324,8 @@ - (void)applicationWillTerminate:(UIApplication *)application { static Fiber mainContext; static Fiber appleContext; static Fiber* currentContext = nullptr; -alignas(16) static char appleStack[1*1024*1024]={}; +// The engine and its script VM share this manually-swapped stack on iOS. +alignas(16) static char appleStack[8*1024*1024]={}; static void appleMain(void*); inline static void createAppleSubContext() { @@ -398,6 +405,13 @@ static void appleMain(void*) { } void iOSApi::implDestroyWindow(SystemApi::Window *w) { + auto wx = reinterpret_cast(w); + if(wx==nullptr) + return; + wx->owner = nullptr; + wx->hasPendingFrame.store(false); + [wx->displayLink invalidate]; + wx->displayLink = nil; } void iOSApi::implExit() { From d0c8dc260fe77236a259f59918b40080fc160052 Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 13:22:29 +0100 Subject: [PATCH 03/25] Add safe iOS scene lifecycle --- Engine/include/Tempest/IOSRuntime | 1 + Engine/system/api/iosapi.mm | 286 ++++++++++++++++++++++++------ Engine/system/iosruntime.h | 22 +++ 3 files changed, 259 insertions(+), 50 deletions(-) create mode 100644 Engine/include/Tempest/IOSRuntime create mode 100644 Engine/system/iosruntime.h diff --git a/Engine/include/Tempest/IOSRuntime b/Engine/include/Tempest/IOSRuntime new file mode 100644 index 00000000..ce7bea6e --- /dev/null +++ b/Engine/include/Tempest/IOSRuntime @@ -0,0 +1 @@ +#include "../system/iosruntime.h" diff --git a/Engine/system/api/iosapi.mm b/Engine/system/api/iosapi.mm index c3ec42f0..333f77d4 100644 --- a/Engine/system/api/iosapi.mm +++ b/Engine/system/api/iosapi.mm @@ -1,6 +1,7 @@ #include "iosapi.h" #include +#include #include #ifdef __IOS__ @@ -36,6 +37,17 @@ static uintptr_t alignDown(uintptr_t val, uintptr_t align) { static void swapContext(); static void drawFrame(); +static void resumeEngineFromUIKit(); + +@class TempestWindow; + +static TempestWindow* mainWindow = nil; +static std::atomic_bool isRunning{true}; +static std::atomic_bool isEngineReady{false}; +static std::atomic_bool isApplicationActive{false}; +static uint64_t lifecycleGeneration = 0; +static bool activationResumePending = false; +static uint32_t preferredFrameRate = 0; @interface TempestWindow : UIWindow { @public Tempest::Window* owner; @@ -123,17 +135,21 @@ - (void)layoutSubviews { frame.origin.y = 0; [self.rootViewController.view setFrame: frame]; - if(owner==nullptr) + if(owner==nullptr || !isEngineReady.load() || !isApplicationActive.load()) return; new (&event.size) SizeEvent(int32_t(frame.size.width*scale), int32_t(frame.size.height*scale)); curentEvent = Event::Resize; - swapContext(); + activationResumePending = false; + resumeEngineFromUIKit(); } - (void)drawFrame { hasPendingFrame.store(true); - swapContext(); + if(!isEngineReady.load() || !isApplicationActive.load()) + return; + activationResumePending = false; + resumeEngineFromUIKit(); // drawFrame(); } @@ -153,7 +169,7 @@ - (void)touchesBegan:(NSSet *)touches withEvent:(UIEvent *)ex { Event::MouseDown ); curentEvent = Event::MouseDown; - swapContext(); + resumeEngineFromUIKit(); } } @@ -173,7 +189,7 @@ - (void)touchesMoved:(NSSet *)touches withEvent:(UIEvent *)ex { Event::MouseMove ); curentEvent = Event::MouseMove; - swapContext(); + resumeEngineFromUIKit(); } } @@ -194,7 +210,7 @@ - (void)touchesEnded:(NSSet *)touches withEvent:(UIEvent *)ex { Event::MouseUp ); curentEvent = Event::MouseUp; - swapContext(); + resumeEngineFromUIKit(); } } @@ -203,8 +219,6 @@ - (void)touchesCancelled:(NSSet *)touches withEvent:(UIEvent *)ex { } @end -static TempestWindow* mainWindow = nullptr; - @interface ViewController:UIViewController{} -(id)init; @@ -257,61 +271,210 @@ -(bool)isFullscreen { } @end -@interface AppDelegate : NSObject { +static void applyPreferredFrameRate(CADisplayLink* displayLink) { + if(displayLink==nil) + return; + if(@available(iOS 15.0, *)) { + if(preferredFrameRate==0) + displayLink.preferredFrameRateRange = CAFrameRateRangeDefault; + else + displayLink.preferredFrameRateRange = + CAFrameRateRangeMake(preferredFrameRate,preferredFrameRate,preferredFrameRate); + } + else { + displayLink.preferredFramesPerSecond = NSInteger(preferredFrameRate); + } + } + +static void invalidateDisplayLink(TempestWindow* window) { + if(window==nil) + return; + window->hasPendingFrame.store(false); + [window->displayLink invalidate]; + window->displayLink = nil; + isEngineReady.store(false); } + +static void createDisplayLink(TempestWindow* window) { + if(window==nil || window.windowScene==nil || window->owner==nullptr) + return; + if(window->displayLink==nil) { + window->displayLink = [CADisplayLink displayLinkWithTarget:window + selector:@selector(drawFrame)]; + applyPreferredFrameRate(window->displayLink); + [window->displayLink addToRunLoop:[NSRunLoop currentRunLoop] + forMode:NSRunLoopCommonModes]; + } + window->displayLink.paused = !isApplicationActive.load(); + window->hasPendingFrame.store(true); + isEngineReady.store(true); + } + +static TempestWindow* attachWindowToScene(UIWindowScene* windowScene) { + if(mainWindow==nil) { + mainWindow = [[TempestWindow alloc] initWithWindowScene:windowScene]; + ViewController* controller = [[ViewController alloc] init]; + mainWindow.rootViewController = controller; + [controller release]; + mainWindow.autoresizingMask = UIViewAutoresizingFlexibleWidth | + UIViewAutoresizingFlexibleHeight; + mainWindow.backgroundColor = [UIColor blackColor]; + mainWindow->owner = nullptr; + mainWindow->displayLink = nil; + mainWindow->hasPendingFrame.store(false); + mainWindow->curentEvent = Event::Type::NoEvent; + } + else { + mainWindow.windowScene = windowScene; + } + + mainWindow.frame = windowScene.coordinateSpace.bounds; + mainWindow.contentScaleFactor = windowScene.screen.scale; + createDisplayLink(mainWindow); + return mainWindow; + } + +static void detachWindowFromScene(TempestWindow* window) { + if(window==nil) + return; + invalidateDisplayLink(window); + window.hidden = YES; + window.windowScene = nil; + } + +@interface SceneDelegate : UIResponder { + TempestWindow* window; + uint64_t activationGeneration; + bool connected; + } +@property(nonatomic, retain) TempestWindow* window; @end -static bool isApplicationActive = false; +@implementation SceneDelegate +@synthesize window; -@implementation AppDelegate -- (BOOL)application:(UIApplication *)application didFinishLaunchingWithOptions:(NSDictionary *)launchOptions { - (void)application; - (void)launchOptions; +- (void)scene:(UIScene *)scene + willConnectToSession:(UISceneSession *)session + options:(UISceneConnectionOptions *)connectionOptions { + (void)session; + (void)connectionOptions; + if(![scene isKindOfClass:[UIWindowScene class]]) + return; - CGRect frame = [ [ UIScreen mainScreen ] bounds ]; - TempestWindow * window = [ [ TempestWindow alloc ] initWithFrame: frame]; - window.contentScaleFactor = [UIScreen mainScreen].scale; - window.rootViewController = [ViewController new]; - window.autoresizingMask = UIViewAutoresizingFlexibleWidth | UIViewAutoresizingFlexibleHeight; - window.backgroundColor = [ UIColor blackColor ]; + [NSObject cancelPreviousPerformRequestsWithTarget:self]; + activationResumePending = false; + isApplicationActive.store(false); + connected = true; + activationGeneration = ++lifecycleGeneration; + self.window = attachWindowToScene((UIWindowScene*)scene); + [self.window makeKeyAndVisible]; + } - window->owner = nullptr; - window->displayLink = nullptr; - window->hasPendingFrame.store(false); - window->curentEvent = Event::Type::NoEvent; - - mainWindow = window; - [ window makeKeyAndVisible ]; // possible switch here - return YES; +- (void)sceneDidBecomeActive:(UIScene *)scene { + if(!connected || scene!=self.window.windowScene) + return; + [NSObject cancelPreviousPerformRequestsWithTarget:self]; + isApplicationActive.store(true); + if(self.window->displayLink!=nil) + self.window->displayLink.paused = NO; + activationResumePending = true; + activationGeneration = ++lifecycleGeneration; + [self performSelector:@selector(resumeEngineIfCurrent:) + withObject:[NSNumber numberWithUnsignedLongLong:activationGeneration] + afterDelay:0.1]; + } + +- (void)sceneWillResignActive:(UIScene *)scene { + if(!connected || scene!=self.window.windowScene) + return; + [NSObject cancelPreviousPerformRequestsWithTarget:self]; + activationGeneration = ++lifecycleGeneration; + activationResumePending = false; + isApplicationActive.store(false); + self.window->hasPendingFrame.store(false); + if(self.window->displayLink!=nil) + self.window->displayLink.paused = YES; } -- (UIInterfaceOrientationMask)application:(UIApplication *)application - supportedInterfaceOrientationsForWindow:(UIWindow *)window { - return UIInterfaceOrientationMaskAll; +- (void)sceneDidDisconnect:(UIScene *)scene { + if(!connected || scene!=self.window.windowScene) + return; + [NSObject cancelPreviousPerformRequestsWithTarget:self]; + activationGeneration = ++lifecycleGeneration; + activationResumePending = false; + isApplicationActive.store(false); + detachWindowFromScene(self.window); + connected = false; + self.window = nil; + } + +- (void)resumeEngineIfCurrent:(NSNumber*)generation { + if(!connected || !activationResumePending || !isApplicationActive.load() || + generation.unsignedLongLongValue!=activationGeneration || + activationGeneration!=lifecycleGeneration || self.window!=mainWindow) + return; + activationResumePending = false; + resumeEngineFromUIKit(); + } + +- (void)dealloc { + [NSObject cancelPreviousPerformRequestsWithTarget:self]; + if(connected && self.window==mainWindow) { + activationGeneration = ++lifecycleGeneration; + activationResumePending = false; + isApplicationActive.store(false); + detachWindowFromScene(self.window); + connected = false; + } + self.window = nil; + [super dealloc]; } +@end -- (void)applicationWillResignActive:(UIApplication *)application { - (void)application; - isApplicationActive = false; - swapContext(); +@interface AppDelegate : NSObject { } +@end -- (void)applicationDidEnterBackground:(UIApplication *)application { +@implementation AppDelegate +- (BOOL)application:(UIApplication *)application didFinishLaunchingWithOptions:(NSDictionary *)launchOptions { (void)application; + (void)launchOptions; + return YES; } -- (void)applicationWillEnterForeground:(UIApplication *)application { +- (UISceneConfiguration *)application:(UIApplication *)application + configurationForConnectingSceneSession:(UISceneSession *)connectingSceneSession + options:(UISceneConnectionOptions *)options { (void)application; + (void)options; + UISceneConfiguration* configuration = + [UISceneConfiguration configurationWithName:@"Tempest Scene" + sessionRole:connectingSceneSession.role]; + configuration.delegateClass = [SceneDelegate class]; + return configuration; } -- (void)applicationDidBecomeActive:(UIApplication *)application { +- (UIInterfaceOrientationMask)application:(UIApplication *)application + supportedInterfaceOrientationsForWindow:(UIWindow *)window { (void)application; - isApplicationActive = true; - swapContext(); + (void)window; + return UIInterfaceOrientationMaskAll; } - (void)applicationWillTerminate:(UIApplication *)application { (void)application; + ++lifecycleGeneration; + activationResumePending = false; + isApplicationActive.store(false); + isRunning.store(false); + preferredFrameRate = 0; + if(mainWindow!=nil) { + invalidateDisplayLink(mainWindow); + mainWindow->owner = nullptr; + mainWindow.windowScene = nil; + [mainWindow release]; + mainWindow = nil; + } } @end @@ -320,7 +483,6 @@ - (void)applicationWillTerminate:(UIApplication *)application { jmp_buf jmp = {}; }; -static std::atomic_bool isRunning{true}; static Fiber mainContext; static Fiber appleContext; static Fiber* currentContext = nullptr; @@ -362,6 +524,27 @@ inline static void swapContext() { std::atomic_thread_fence(std::memory_order_seq_cst); } +static void resumeEngineFromUIKit() { + if(currentContext==&appleContext) + swapContext(); + } + +void Tempest::iOS::yieldToUIKit() { + if(![NSThread isMainThread] || currentContext!=&mainContext || + !isApplicationActive.load() || mainWindow==nil || + mainWindow->displayLink==nil) + return; + swapContext(); + } + +void Tempest::iOS::setPreferredFrameRate(uint32_t framesPerSecond) { + if(![NSThread isMainThread]) + return; + preferredFrameRate = framesPerSecond; + if(mainWindow!=nil) + applyPreferredFrameRate(mainWindow->displayLink); + } + static void drawFrame() { auto cb = (mainWindow->owner); @autoreleasepool { @@ -381,13 +564,15 @@ static void appleMain(void*) { } static SystemApi::Window* createWindow(Tempest::Window *owner, uint32_t w, uint32_t h, SystemApi::ShowMode mode) { + (void)w; + (void)h; + (void)mode; auto window = mainWindow; - + if(window==nil) + return nullptr; + window->owner = owner; - window->displayLink = [CADisplayLink displayLinkWithTarget:window selector:@selector(drawFrame)]; - //by adding the display link to the run loop our draw method will be called 60 times per second - [window->displayLink addToRunLoop:[NSRunLoop currentRunLoop] forMode:NSDefaultRunLoopMode]; - window->hasPendingFrame.store(true); + createDisplayLink(window); return reinterpret_cast(window); } @@ -409,13 +594,14 @@ static void appleMain(void*) { if(wx==nullptr) return; wx->owner = nullptr; - wx->hasPendingFrame.store(false); - [wx->displayLink invalidate]; - wx->displayLink = nil; + invalidateDisplayLink(wx); } void iOSApi::implExit() { ::isRunning.store(false); + activationResumePending = false; + isEngineReady.store(false); + invalidateDisplayLink(mainWindow); } Tempest::Rect iOSApi::implWindowClientRect(Window* w) { @@ -492,7 +678,7 @@ static void appleMain(void*) { break; } default: - if(isApplicationActive && mainWindow->hasPendingFrame.load()) { + if(isApplicationActive.load() && mainWindow->hasPendingFrame.load()) { mainWindow->hasPendingFrame.store(false); iOSApi::dispatchRender(wnd); } diff --git a/Engine/system/iosruntime.h b/Engine/system/iosruntime.h new file mode 100644 index 00000000..a3d65179 --- /dev/null +++ b/Engine/system/iosruntime.h @@ -0,0 +1,22 @@ +#pragma once + +#include "platform.h" + +#include + +#if defined(__IOS__) + +namespace Tempest::iOS { + +// Temporarily return control to UIKit while running on Tempest's iOS engine +// fiber. Calls from worker threads or before the native window exists are +// ignored. +void yieldToUIKit(); + +// Request a fixed display-link cadence. Zero restores the native/default +// cadence selected by the system. +void setPreferredFrameRate(uint32_t framesPerSecond); + +} + +#endif From 74b4593b40101ff90be9f07b15a17f9fdbcbda46 Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 13:16:43 +0100 Subject: [PATCH 04/25] Prefer working directory files on iOS --- Engine/io/rfile.mm | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/Engine/io/rfile.mm b/Engine/io/rfile.mm index 07a8cc08..694e2ee7 100644 --- a/Engine/io/rfile.mm +++ b/Engine/io/rfile.mm @@ -5,7 +5,9 @@ #include #include -#import +#include + +#import using namespace Tempest; @@ -16,7 +18,12 @@ throw std::system_error(Tempest::SystemErrc::UnableToOpenFile); return ret; } - + + // Relative paths may refer to user-provided files in the process working + // directory. Packaged application resources remain the fallback. + if(void* ret = fopen(cstr,"rb")) + return ret; + @autoreleasepool { NSString *dir = [[NSBundle mainBundle] resourcePath]; std::string full = [dir UTF8String]; From 7c1de076542c0b769c5007a7af8c39857f31188a Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 13:26:19 +0100 Subject: [PATCH 05/25] Add RG16F texture format --- Engine/formats/image/pixmapcodeccommon.cpp | 1 + Engine/formats/pixmap.cpp | 3 ++ Engine/gapi/abstractgraphicsapi.h | 2 ++ Engine/gapi/directx12/dxcommandbuffer.cpp | 2 +- Engine/gapi/directx12/dxdevice.h | 2 ++ Engine/gapi/metal/mtdevice.h | 2 ++ Engine/gapi/metal/mtdevice.mm | 7 +++-- Engine/gapi/vulkan/vdevice.h | 2 ++ Tests/tests/gapi/directx_test.cpp | 6 ++++ Tests/tests/gapi/gapi_test_common.h | 35 ++++++++++++++++++++++ Tests/tests/gapi/metal_test.cpp | 6 ++++ Tests/tests/gapi/vulkan_test.cpp | 6 ++++ Tests/tests/pixmap_test.cpp | 15 ++++++++++ Tests/tests/utils/imagevalidator.cpp | 1 + 14 files changed, 86 insertions(+), 4 deletions(-) diff --git a/Engine/formats/image/pixmapcodeccommon.cpp b/Engine/formats/image/pixmapcodeccommon.cpp index dff1c736..5fb5a925 100644 --- a/Engine/formats/image/pixmapcodeccommon.cpp +++ b/Engine/formats/image/pixmapcodeccommon.cpp @@ -259,6 +259,7 @@ bool PixmapCodecCommon::save(ODevice &f, const char *ext, const uint8_t* cdata, break; case TextureFormat::R11G11B10UF: case TextureFormat::RGBA16F: + case TextureFormat::RG16F: // hdr or exr? break; } diff --git a/Engine/formats/pixmap.cpp b/Engine/formats/pixmap.cpp index 1cb96208..d7436fdf 100644 --- a/Engine/formats/pixmap.cpp +++ b/Engine/formats/pixmap.cpp @@ -455,6 +455,7 @@ size_t Pixmap::blockSizeForFormat(TextureFormat frm) { //--- case TextureFormat::R11G11B10UF: return 4; case TextureFormat::RGBA16F: return 8; + case TextureFormat::RG16F: return 4; } return 0; } @@ -495,6 +496,7 @@ uint8_t Pixmap::componentCount(TextureFormat frm) { //--- case TextureFormat::R11G11B10UF: return 3; case TextureFormat::RGBA16F: return 4; + case TextureFormat::RG16F: return 2; } return 0; } @@ -526,6 +528,7 @@ Size Pixmap::blockCount(TextureFormat frm, uint32_t w, uint32_t h) { case TextureFormat::Depth32F: case TextureFormat::R11G11B10UF: case TextureFormat::RGBA16F: + case TextureFormat::RG16F: return Size(w,h); case TextureFormat::DXT1: case TextureFormat::DXT3: diff --git a/Engine/gapi/abstractgraphicsapi.h b/Engine/gapi/abstractgraphicsapi.h index 51fe4ef8..e7483f1d 100644 --- a/Engine/gapi/abstractgraphicsapi.h +++ b/Engine/gapi/abstractgraphicsapi.h @@ -128,6 +128,7 @@ namespace Tempest { DXT5, R11G11B10UF, RGBA16F, + RG16F, Last }; @@ -159,6 +160,7 @@ namespace Tempest { case DXT5: return "DXT5"; case R11G11B10UF: return "R11G11B10UF"; case RGBA16F: return "RGBA16F"; + case RG16F: return "RG16F"; case Last: break; } diff --git a/Engine/gapi/directx12/dxcommandbuffer.cpp b/Engine/gapi/directx12/dxcommandbuffer.cpp index 0a0580cd..3a9886dd 100644 --- a/Engine/gapi/directx12/dxcommandbuffer.cpp +++ b/Engine/gapi/directx12/dxcommandbuffer.cpp @@ -974,6 +974,7 @@ DxCompPipeline& DxCommandBuffer::copyShader(DXGI_FORMAT format, int32_t& bitCnt, case DXGI_FORMAT_R16G16_UINT: case DXGI_FORMAT_R16G16_SNORM: case DXGI_FORMAT_R16G16_SINT: + case DXGI_FORMAT_R16G16_FLOAT: bitCnt = 16; compCnt = 2; return *dev.copyS.handler; @@ -1341,4 +1342,3 @@ void DxCommandBuffer::generateMipmap(AbstractGraphicsApi::Texture& dstTex, #endif - diff --git a/Engine/gapi/directx12/dxdevice.h b/Engine/gapi/directx12/dxdevice.h index 1c98db76..8e2bf624 100644 --- a/Engine/gapi/directx12/dxdevice.h +++ b/Engine/gapi/directx12/dxdevice.h @@ -121,6 +121,8 @@ inline DXGI_FORMAT nativeFormat(TextureFormat f) { return DXGI_FORMAT_R11G11B10_FLOAT; case TextureFormat::RGBA16F: return DXGI_FORMAT_R16G16B16A16_FLOAT; + case TextureFormat::RG16F: + return DXGI_FORMAT_R16G16_FLOAT; } return DXGI_FORMAT_UNKNOWN; } diff --git a/Engine/gapi/metal/mtdevice.h b/Engine/gapi/metal/mtdevice.h index 8e1d4464..1d877773 100644 --- a/Engine/gapi/metal/mtdevice.h +++ b/Engine/gapi/metal/mtdevice.h @@ -74,6 +74,8 @@ inline MTL::PixelFormat nativeFormat(TextureFormat frm) { return MTL::PixelFormatRG11B10Float; case RGBA16F: return MTL::PixelFormatRGBA16Float; + case RG16F: + return MTL::PixelFormatRG16Float; } return MTL::PixelFormatInvalid; } diff --git a/Engine/gapi/metal/mtdevice.mm b/Engine/gapi/metal/mtdevice.mm index 70f7a546..19c2b3bd 100644 --- a/Engine/gapi/metal/mtdevice.mm +++ b/Engine/gapi/metal/mtdevice.mm @@ -221,20 +221,20 @@ TextureFormat::R16, TextureFormat::RG16, TextureFormat::RGBA16, TextureFormat::R32F, TextureFormat::RG32F, TextureFormat::RGBA32F, TextureFormat::R32U, TextureFormat::RG32U, TextureFormat::RGBA32U, - TextureFormat::R11G11B10UF, TextureFormat::RGBA16F, + TextureFormat::R11G11B10UF, TextureFormat::RGBA16F, TextureFormat::RG16F, }; static const TextureFormat att[] = {TextureFormat::R8, TextureFormat::RG8, TextureFormat::RGBA8, TextureFormat::R16, TextureFormat::RG16, TextureFormat::RGBA16, TextureFormat::R32F, TextureFormat::RG32F, TextureFormat::RGBA32F, - TextureFormat::R11G11B10UF, TextureFormat::RGBA16F, + TextureFormat::R11G11B10UF, TextureFormat::RGBA16F, TextureFormat::RG16F, }; static const TextureFormat sso[] = {TextureFormat::R8, TextureFormat::RG8, TextureFormat::RGBA8, TextureFormat::R16, TextureFormat::RG16, TextureFormat::RGBA16, TextureFormat::R32U, TextureFormat::RG32U, TextureFormat::RGBA32U, TextureFormat::R32F, TextureFormat::RGBA32F, - TextureFormat::R11G11B10UF, TextureFormat::RGBA16F, + TextureFormat::R11G11B10UF, TextureFormat::RGBA16F, TextureFormat::RG16F, }; static const TextureFormat ds[] = {TextureFormat::Depth16, TextureFormat::Depth32F}; @@ -258,6 +258,7 @@ storBit |= uint64_t(1) << TextureFormat::RGBA32U; // 16 bit storBit |= uint64_t(1) << TextureFormat::RGBA16F; + storBit |= uint64_t(1) << TextureFormat::RG16F; // 8 bit storBit |= uint64_t(1) << TextureFormat::R8; [[fallthrough]]; diff --git a/Engine/gapi/vulkan/vdevice.h b/Engine/gapi/vulkan/vdevice.h index 613295b4..110b84d6 100644 --- a/Engine/gapi/vulkan/vdevice.h +++ b/Engine/gapi/vulkan/vdevice.h @@ -102,6 +102,8 @@ inline VkFormat nativeFormat(TextureFormat f) { return VK_FORMAT_B10G11R11_UFLOAT_PACK32; case TextureFormat::RGBA16F: return VK_FORMAT_R16G16B16A16_SFLOAT; + case TextureFormat::RG16F: + return VK_FORMAT_R16G16_SFLOAT; } return VK_FORMAT_UNDEFINED; } diff --git a/Tests/tests/gapi/directx_test.cpp b/Tests/tests/gapi/directx_test.cpp index 4e8eb75a..f67f9d74 100644 --- a/Tests/tests/gapi/directx_test.cpp +++ b/Tests/tests/gapi/directx_test.cpp @@ -122,6 +122,12 @@ TEST(DirectX12Api,Draw) { #endif } +TEST(DirectX12Api,RG16F) { +#if defined(_MSC_VER) + GapiTestCommon::FormatRG16F(); +#endif + } + TEST(DirectX12Api,DepthWrite) { #if defined(_MSC_VER) GapiTestCommon::DepthWrite("DirectX12Api_DepthWrite.png"); diff --git a/Tests/tests/gapi/gapi_test_common.h b/Tests/tests/gapi/gapi_test_common.h index b387d171..eadc082e 100644 --- a/Tests/tests/gapi/gapi_test_common.h +++ b/Tests/tests/gapi/gapi_test_common.h @@ -561,6 +561,41 @@ void Draw(const char* outImage) { } } +template +void FormatRG16F() { + using namespace Tempest; + + try { + GraphicsApi api{ApiFlags::Validation}; + Device device(api); + + if(!device.properties().hasAttachFormat(TextureFormat::RG16F)) { + Log::d("Skipping RG16F testcase: no format support"); + return; + } + + auto tex = device.attachment(TextureFormat::RG16F,4,4); + auto cmd = device.commandBuffer(); + { + auto enc = cmd.startEncoding(device); + enc.setFramebuffer({{tex,Vec4(0.25f,0.5f,0.f,1.f),Tempest::Preserve}}); + } + + auto sync = device.submit(cmd); + sync.wait(); + + auto pm = device.readPixels(tex); + EXPECT_EQ(pm.format(),TextureFormat::RG16F); + EXPECT_EQ(pm.bpp(),4); + EXPECT_EQ(pm.dataSize(),64); + } + catch(std::system_error& e) { + if(e.code()==Tempest::GraphicsErrc::NoDevice) + Log::d("Skipping graphics testcase: ",e.what()); else + throw; + } + } + template void InstanceIndex(const char* outImage) { using namespace Tempest; diff --git a/Tests/tests/gapi/metal_test.cpp b/Tests/tests/gapi/metal_test.cpp index de095aff..d6d7f1c9 100644 --- a/Tests/tests/gapi/metal_test.cpp +++ b/Tests/tests/gapi/metal_test.cpp @@ -102,6 +102,12 @@ TEST(MetalApi,Draw) { #endif } +TEST(MetalApi,RG16F) { +#if defined(__OSX__) + GapiTestCommon::FormatRG16F(); +#endif + } + TEST(MetalApi,DepthWrite) { #if defined(__OSX__) GapiTestCommon::DepthWrite("MetalApi_DepthWrite.png"); diff --git a/Tests/tests/gapi/vulkan_test.cpp b/Tests/tests/gapi/vulkan_test.cpp index c3b5851d..e0477c06 100644 --- a/Tests/tests/gapi/vulkan_test.cpp +++ b/Tests/tests/gapi/vulkan_test.cpp @@ -103,6 +103,12 @@ TEST(VulkanApi,Draw) { #endif } +TEST(VulkanApi,RG16F) { +#if !defined(__OSX__) + GapiTestCommon::FormatRG16F(); +#endif + } + TEST(VulkanApi,DepthWrite) { #if !defined(__OSX__) GapiTestCommon::DepthWrite("VulkanApi_DepthWrite.png"); diff --git a/Tests/tests/pixmap_test.cpp b/Tests/tests/pixmap_test.cpp index b46c266c..dbe945be 100644 --- a/Tests/tests/pixmap_test.cpp +++ b/Tests/tests/pixmap_test.cpp @@ -8,6 +8,21 @@ using namespace testing; using namespace Tempest; +TEST(main,PixmapFormatRG16F) { + EXPECT_EQ(uint8_t(TextureFormat::RGBA16F),25); + EXPECT_STREQ(formatName(TextureFormat::RG16F),"RG16F"); + + Pixmap pm(3,2,TextureFormat::RG16F); + EXPECT_EQ(pm.format(),TextureFormat::RG16F); + EXPECT_EQ(pm.bpp(),4); + EXPECT_EQ(pm.dataSize(),24); + EXPECT_EQ(Pixmap::componentCount(TextureFormat::RG16F),2); + + const auto blocks = Pixmap::blockCount(TextureFormat::RG16F,3,2); + EXPECT_EQ(blocks.w,3); + EXPECT_EQ(blocks.h,2); + } + TEST(main,PixmapIO_0) { Pixmap pm("assets/pixmap_io/rgba.png"); EXPECT_EQ(pm.w(), 256); diff --git a/Tests/tests/utils/imagevalidator.cpp b/Tests/tests/utils/imagevalidator.cpp index 4e7ad967..48537048 100644 --- a/Tests/tests/utils/imagevalidator.cpp +++ b/Tests/tests/utils/imagevalidator.cpp @@ -64,6 +64,7 @@ ImageValidator::Pixel ImageValidator::at(uint32_t x, uint32_t y) const { case TextureFormat::R11G11B10UF: case TextureFormat::RGBA16F: + case TextureFormat::RG16F: assert(false); break; } From 1a26c082b9779d167bb3dbc7d5ace0e748871091 Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 13:42:09 +0100 Subject: [PATCH 06/25] Fix iOS autorelease-pool lifetime across fiber yields --- Engine/system/api/iosapi.mm | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/Engine/system/api/iosapi.mm b/Engine/system/api/iosapi.mm index 333f77d4..e2b5ef90 100644 --- a/Engine/system/api/iosapi.mm +++ b/Engine/system/api/iosapi.mm @@ -652,7 +652,10 @@ static void appleMain(void*) { return; } - @autoreleasepool { + // The engine and UIKit fibers share one OS thread. An Objective-C pool + // pushed on the engine fiber can be invalidated while UIKit runs and then + // trigger AutoreleasePoolPage::badPop when this stack resumes. + { auto& wnd = *mainWindow->owner; auto eType = mainWindow->curentEvent; mainWindow->curentEvent = Event::Type::NoEvent; From 7c02b187ee26330cd8b10bf09ddf3c469c2eb7e9 Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 13:46:27 +0100 Subject: [PATCH 07/25] Add opt-in iOS display policy controls --- Engine/system/api/iosapi.mm | 92 +++++++++++++++++++++++++++++++++---- Engine/system/iosruntime.h | 11 +++++ 2 files changed, 94 insertions(+), 9 deletions(-) diff --git a/Engine/system/api/iosapi.mm b/Engine/system/api/iosapi.mm index e2b5ef90..40d3ea5a 100644 --- a/Engine/system/api/iosapi.mm +++ b/Engine/system/api/iosapi.mm @@ -47,7 +47,18 @@ static uintptr_t alignDown(uintptr_t val, uintptr_t align) { static std::atomic_bool isApplicationActive{false}; static uint64_t lifecycleGeneration = 0; static bool activationResumePending = false; -static uint32_t preferredFrameRate = 0; +static bool idleTimerDisabled = false; + +enum class FrameRateMode : uint8_t { + SystemDefault, + Fixed, + Range, + }; + +static FrameRateMode frameRateMode = FrameRateMode::SystemDefault; +static uint32_t frameRateMinimum = 0; +static uint32_t frameRateMaximum = 0; +static uint32_t frameRatePreferred = 0; @interface TempestWindow : UIWindow { @public Tempest::Window* owner; @@ -275,17 +286,34 @@ static void applyPreferredFrameRate(CADisplayLink* displayLink) { if(displayLink==nil) return; if(@available(iOS 15.0, *)) { - if(preferredFrameRate==0) - displayLink.preferredFrameRateRange = CAFrameRateRangeDefault; - else - displayLink.preferredFrameRateRange = - CAFrameRateRangeMake(preferredFrameRate,preferredFrameRate,preferredFrameRate); + switch(frameRateMode) { + case FrameRateMode::SystemDefault: + displayLink.preferredFrameRateRange = CAFrameRateRangeDefault; + break; + case FrameRateMode::Fixed: + displayLink.preferredFrameRateRange = + CAFrameRateRangeMake(frameRatePreferred,frameRatePreferred, + frameRatePreferred); + break; + case FrameRateMode::Range: + displayLink.preferredFrameRateRange = + CAFrameRateRangeMake(frameRateMinimum,frameRateMaximum, + frameRatePreferred); + break; + } } else { - displayLink.preferredFramesPerSecond = NSInteger(preferredFrameRate); + displayLink.preferredFramesPerSecond = + frameRateMode==FrameRateMode::SystemDefault ? 0 : + NSInteger(frameRatePreferred); } } +static void applyIdleTimerPreference() { + [UIApplication sharedApplication].idleTimerDisabled = + isApplicationActive.load() && idleTimerDisabled ? YES : NO; + } + static void invalidateDisplayLink(TempestWindow* window) { if(window==nil) return; @@ -375,6 +403,7 @@ - (void)sceneDidBecomeActive:(UIScene *)scene { return; [NSObject cancelPreviousPerformRequestsWithTarget:self]; isApplicationActive.store(true); + applyIdleTimerPreference(); if(self.window->displayLink!=nil) self.window->displayLink.paused = NO; activationResumePending = true; @@ -391,6 +420,7 @@ - (void)sceneWillResignActive:(UIScene *)scene { activationGeneration = ++lifecycleGeneration; activationResumePending = false; isApplicationActive.store(false); + applyIdleTimerPreference(); self.window->hasPendingFrame.store(false); if(self.window->displayLink!=nil) self.window->displayLink.paused = YES; @@ -403,6 +433,7 @@ - (void)sceneDidDisconnect:(UIScene *)scene { activationGeneration = ++lifecycleGeneration; activationResumePending = false; isApplicationActive.store(false); + applyIdleTimerPreference(); detachWindowFromScene(self.window); connected = false; self.window = nil; @@ -423,6 +454,7 @@ - (void)dealloc { activationGeneration = ++lifecycleGeneration; activationResumePending = false; isApplicationActive.store(false); + applyIdleTimerPreference(); detachWindowFromScene(self.window); connected = false; } @@ -467,7 +499,12 @@ - (void)applicationWillTerminate:(UIApplication *)application { activationResumePending = false; isApplicationActive.store(false); isRunning.store(false); - preferredFrameRate = 0; + idleTimerDisabled = false; + applyIdleTimerPreference(); + frameRateMode = FrameRateMode::SystemDefault; + frameRateMinimum = 0; + frameRateMaximum = 0; + frameRatePreferred = 0; if(mainWindow!=nil) { invalidateDisplayLink(mainWindow); mainWindow->owner = nullptr; @@ -540,11 +577,48 @@ static void resumeEngineFromUIKit() { void Tempest::iOS::setPreferredFrameRate(uint32_t framesPerSecond) { if(![NSThread isMainThread]) return; - preferredFrameRate = framesPerSecond; + frameRateMode = framesPerSecond==0 ? FrameRateMode::SystemDefault : + FrameRateMode::Fixed; + frameRateMinimum = framesPerSecond; + frameRateMaximum = framesPerSecond; + frameRatePreferred = framesPerSecond; if(mainWindow!=nil) applyPreferredFrameRate(mainWindow->displayLink); } +void Tempest::iOS::setPreferredFrameRateRange(uint32_t minimumFramesPerSecond, + uint32_t maximumFramesPerSecond, + uint32_t preferredFramesPerSecond) { + if(![NSThread isMainThread]) + return; + if(maximumFramesPerSecond==0) { + setPreferredFrameRate(0); + return; + } + if(minimumFramesPerSecond==0) + minimumFramesPerSecond = 1; + if(maximumFramesPerSecondmaximumFramesPerSecond) + preferredFramesPerSecond = maximumFramesPerSecond; + + frameRateMode = FrameRateMode::Range; + frameRateMinimum = minimumFramesPerSecond; + frameRateMaximum = maximumFramesPerSecond; + frameRatePreferred = preferredFramesPerSecond; + if(mainWindow!=nil) + applyPreferredFrameRate(mainWindow->displayLink); + } + +void Tempest::iOS::setIdleTimerDisabled(bool disabled) { + if(![NSThread isMainThread]) + return; + idleTimerDisabled = disabled; + applyIdleTimerPreference(); + } + static void drawFrame() { auto cb = (mainWindow->owner); @autoreleasepool { diff --git a/Engine/system/iosruntime.h b/Engine/system/iosruntime.h index a3d65179..d9a17dde 100644 --- a/Engine/system/iosruntime.h +++ b/Engine/system/iosruntime.h @@ -17,6 +17,17 @@ void yieldToUIKit(); // cadence selected by the system. void setPreferredFrameRate(uint32_t framesPerSecond); +// Request an adaptive display-link range. Values are normalized to +// minimum <= preferred <= maximum. A zero maximum restores the system +// default cadence. +void setPreferredFrameRateRange(uint32_t minimumFramesPerSecond, + uint32_t maximumFramesPerSecond, + uint32_t preferredFramesPerSecond); + +// Control the iOS idle timer. The preference survives scene deactivation and +// is applied again when the scene becomes active. +void setIdleTimerDisabled(bool disabled); + } #endif From dc4cb2108d539b8c944d924665fc0f59da0d5b65 Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 13:38:19 +0100 Subject: [PATCH 08/25] Add portable spatial scaler API --- Engine/gapi/abstractgraphicsapi.cpp | 9 + Engine/gapi/abstractgraphicsapi.h | 20 ++ Engine/graphics/device.cpp | 5 +- Engine/graphics/device.h | 4 +- Engine/graphics/encoder.cpp | 12 + Engine/graphics/encoder.h | 4 +- Engine/graphics/spatialscaler.cpp | 21 ++ Engine/graphics/spatialscaler.h | 36 +++ Engine/include/Tempest/SpatialScaler | 1 + Tests/tests/spatialscaler_test.cpp | 328 +++++++++++++++++++++++++++ 10 files changed, 437 insertions(+), 3 deletions(-) create mode 100644 Engine/graphics/spatialscaler.cpp create mode 100644 Engine/graphics/spatialscaler.h create mode 100644 Engine/include/Tempest/SpatialScaler create mode 100644 Tests/tests/spatialscaler_test.cpp diff --git a/Engine/gapi/abstractgraphicsapi.cpp b/Engine/gapi/abstractgraphicsapi.cpp index a0826cd4..dd13ee94 100644 --- a/Engine/gapi/abstractgraphicsapi.cpp +++ b/Engine/gapi/abstractgraphicsapi.cpp @@ -93,6 +93,10 @@ void AbstractGraphicsApi::CommandBuffer::dispatchMeshIndirect(const Buffer& indi throw std::system_error(Tempest::GraphicsErrc::UnsupportedExtension); } +bool AbstractGraphicsApi::CommandBuffer::spatialUpscale(SpatialScaler&, Texture&, Texture&) { + return false; + } + AbstractGraphicsApi::AccelerationStructure* AbstractGraphicsApi::createBottomAccelerationStruct(Device* d, const RtGeometry* geom, size_t geomSize) { throw std::system_error(Tempest::GraphicsErrc::UnsupportedExtension); } @@ -102,6 +106,11 @@ AbstractGraphicsApi::AccelerationStructure* throw std::system_error(Tempest::GraphicsErrc::UnsupportedExtension); } +AbstractGraphicsApi::SpatialScaler* + AbstractGraphicsApi::createSpatialScaler(Device*, const SpatialScalerDesc&) { + return nullptr; + } + bool Detail::Bindings::operator ==(const Bindings &other) const { for(size_t i=0; i; @@ -655,6 +674,7 @@ namespace Tempest { virtual AccelerationStructure* createBottomAccelerationStruct(Device* d, const RtGeometry* geom, size_t geomSize); virtual AccelerationStructure* createTopAccelerationStruct(Device* d, const RtInstance* geom, AccelerationStructure*const* as, size_t geomSize); + virtual SpatialScaler* createSpatialScaler(Device* d, const SpatialScalerDesc& desc); virtual void readPixels (Device* d, Pixmap& out, const PTexture t, TextureFormat frm, const uint32_t w, const uint32_t h, uint32_t mip, bool storageImg) = 0; diff --git a/Engine/graphics/device.cpp b/Engine/graphics/device.cpp index 2d1fb36a..e956909f 100644 --- a/Engine/graphics/device.cpp +++ b/Engine/graphics/device.cpp @@ -259,6 +259,10 @@ StorageImage Tempest::Device::image2d(TextureFormat frm, const Size sz, const bo return image2d(frm,sz.w,sz.h,mips); } +SpatialScaler Device::spatialScaler(const SpatialScalerDesc& desc) { + return SpatialScaler(api.createSpatialScaler(dev,desc)); + } + ZBuffer Device::zbuffer(TextureFormat frm, const Size sz) { if(sz.w<0 || sz.h<0) throw std::system_error(Tempest::GraphicsErrc::TooLargeTexture, std::to_string(std::min(sz.w,sz.h))); @@ -410,4 +414,3 @@ Detail::VideoBuffer Device::createVideoBuffer(const void *data, size_t size, Mem Detail::VideoBuffer buf(api.createBuffer(dev,data,size,usage,flg), size); return buf; } - diff --git a/Engine/graphics/device.h b/Engine/graphics/device.h index bc06d6d4..0c053f97 100644 --- a/Engine/graphics/device.h +++ b/Engine/graphics/device.h @@ -15,6 +15,7 @@ #include #include #include +#include #include #include #include @@ -131,6 +132,8 @@ class Device { ZBuffer zbuffer (TextureFormat frm, const Size sz); StorageImage image2d (TextureFormat frm, const Size sz, const bool mips = false); + SpatialScaler spatialScaler(const SpatialScalerDesc& desc); + AccelerationStructure blas(const std::vector& geom); AccelerationStructure blas(std::initializer_list geom); AccelerationStructure blas(const RtGeometry* geom, size_t geomSize); @@ -264,4 +267,3 @@ inline AccelerationStructure Device::blas(const VertexBuffer& vbo, const Inde } } - diff --git a/Engine/graphics/encoder.cpp b/Engine/graphics/encoder.cpp index 959dac5d..8a5c1210 100644 --- a/Engine/graphics/encoder.cpp +++ b/Engine/graphics/encoder.cpp @@ -323,3 +323,15 @@ void Encoder::generateMipmaps(Attachment& tex) { uint32_t w = tex.w(), h = tex.h(); impl->generateMipmap(*textureCast(tex).impl.handler,w,h,mipCount(w,h)); } + +bool Encoder::spatialUpscale(const SpatialScaler& scaler, const Attachment& input, StorageImage& output) { + if(scaler.isEmpty() || input.isEmpty() || output.isEmpty()) + return false; + if(state.stage==Rendering) + impl->endRendering(); + state = State(); + + auto& src = *textureCast(input).impl.handler; + auto& dst = *output.tImpl.impl.handler; + return impl->spatialUpscale(*scaler.impl.handler,src,dst); + } diff --git a/Engine/graphics/encoder.h b/Engine/graphics/encoder.h index 457666da..9f567317 100644 --- a/Engine/graphics/encoder.h +++ b/Engine/graphics/encoder.h @@ -7,6 +7,7 @@ #include #include #include +#include namespace Tempest { @@ -119,6 +120,8 @@ class Encoder { void generateMipmaps(Attachment& tex); + bool spatialUpscale(const SpatialScaler& scaler, const Attachment& input, StorageImage& output); + private: explicit Encoder(CommandBuffer* ow); @@ -147,4 +150,3 @@ class Encoder { }; } - diff --git a/Engine/graphics/spatialscaler.cpp b/Engine/graphics/spatialscaler.cpp new file mode 100644 index 00000000..b83a623c --- /dev/null +++ b/Engine/graphics/spatialscaler.cpp @@ -0,0 +1,21 @@ +#include "spatialscaler.h" + +using namespace Tempest; + +SpatialScaler::SpatialScaler(SpatialScaler&& other) noexcept + :impl(other.impl.handler) { + other.impl.handler = nullptr; + } + +SpatialScaler::~SpatialScaler() { + delete impl.handler; + } + +SpatialScaler& SpatialScaler::operator=(SpatialScaler&& other) noexcept { + if(this==&other) + return *this; + delete impl.handler; + impl.handler = other.impl.handler; + other.impl.handler = nullptr; + return *this; + } diff --git a/Engine/graphics/spatialscaler.h b/Engine/graphics/spatialscaler.h new file mode 100644 index 00000000..08c91e5b --- /dev/null +++ b/Engine/graphics/spatialscaler.h @@ -0,0 +1,36 @@ +#pragma once + +#include + +#include "../utility/dptr.h" + +namespace Tempest { + +class Device; +class CommandBuffer; +template +class Encoder; + +class SpatialScaler final { + public: + SpatialScaler() = default; + SpatialScaler(SpatialScaler&& other) noexcept; + SpatialScaler(const SpatialScaler&) = delete; + ~SpatialScaler(); + + SpatialScaler& operator=(SpatialScaler&& other) noexcept; + SpatialScaler& operator=(const SpatialScaler&) = delete; + + bool isEmpty() const { return impl.handler==nullptr; } + explicit operator bool() const { return !isEmpty(); } + + private: + explicit SpatialScaler(AbstractGraphicsApi::SpatialScaler* scaler):impl(scaler) {} + + Detail::DPtr impl; + + friend class Tempest::Device; + friend class Encoder; + }; + +} diff --git a/Engine/include/Tempest/SpatialScaler b/Engine/include/Tempest/SpatialScaler new file mode 100644 index 00000000..b3b25cc3 --- /dev/null +++ b/Engine/include/Tempest/SpatialScaler @@ -0,0 +1 @@ +#include "../graphics/spatialscaler.h" diff --git a/Tests/tests/spatialscaler_test.cpp b/Tests/tests/spatialscaler_test.cpp new file mode 100644 index 00000000..8e8049de --- /dev/null +++ b/Tests/tests/spatialscaler_test.cpp @@ -0,0 +1,328 @@ +#include +#include +#include + +#include + +#include +#include + +using namespace Tempest; + +namespace { + +struct MockStats { + int scalerCreateAttempts = 0; + int scalerCreated = 0; + int scalerDestroyed = 0; + int scalerEncoded = 0; + int renderingBegun = 0; + int renderingEnded = 0; + bool commandSupportsScaler = true; + + AbstractGraphicsApi::Texture* scalerInput = nullptr; + AbstractGraphicsApi::Texture* scalerOutput = nullptr; + }; + +struct MockDevice final : AbstractGraphicsApi::Device { + void waitIdle() override {} + }; + +struct MockShader final : AbstractGraphicsApi::Shader {}; + +struct MockPipeline final : AbstractGraphicsApi::Pipeline { + IVec3 workGroupSize() const override { return {1,1,1}; } + size_t sizeofBuffer(size_t, size_t) const override { return 0; } + }; + +struct MockCompPipeline final : AbstractGraphicsApi::CompPipeline { + IVec3 workGroupSize() const override { return {1,1,1}; } + size_t sizeofBuffer(size_t, size_t) const override { return 0; } + }; + +struct MockBuffer final : AbstractGraphicsApi::Buffer { + void update(const void*, size_t, size_t) override {} + void read(void*, size_t, size_t) override {} + }; + +struct MockTexture final : AbstractGraphicsApi::Texture { + explicit MockTexture(NonUniqResId id):id(id) {} + + uint32_t mipCount() const override { return 1; } + NonUniqResId syncId() const override { return id; } + + NonUniqResId id; + }; + +struct MockSpatialScaler final : AbstractGraphicsApi::SpatialScaler { + explicit MockSpatialScaler(MockStats& stats):stats(stats) { + ++stats.scalerCreated; + } + + ~MockSpatialScaler() override { + ++stats.scalerDestroyed; + } + + MockStats& stats; + }; + +class MockCommandBuffer final : public AbstractGraphicsApi::CommandBuffer { + public: + explicit MockCommandBuffer(MockStats& stats):stats(stats) {} + + void beginRendering(const Detail::FrameBufferDesc&, size_t, uint32_t, uint32_t) override { + ++stats.renderingBegun; + } + void endRendering() override { + ++stats.renderingEnded; + } + + void barrier(const AbstractGraphicsApi::SyncDesc&, + const AbstractGraphicsApi::BarrierDesc*, size_t) override {} + + void generateMipmap(AbstractGraphicsApi::Texture&, uint32_t, uint32_t, uint32_t) override {} + void copy(AbstractGraphicsApi::Buffer&, size_t, AbstractGraphicsApi::Texture&, + uint32_t, uint32_t, uint32_t) override {} + + bool isRecording() const override { return recording; } + void begin() override { recording = true; } + void end() override { recording = false; } + void reset() override { recording = false; } + + void setPipeline(AbstractGraphicsApi::Pipeline&) override {} + void setComputePipeline(AbstractGraphicsApi::CompPipeline&) override {} + void setBinding(size_t, AbstractGraphicsApi::Texture*, uint32_t, + const ComponentMapping&, const Sampler&) override {} + void setBinding(size_t, AbstractGraphicsApi::Buffer*, size_t) override {} + void setBinding(size_t, AbstractGraphicsApi::DescArray*) override {} + void setBinding(size_t, AbstractGraphicsApi::AccelerationStructure*) override {} + void setBinding(size_t, const Sampler&) override {} + + void setViewport(const Rect&) override {} + void setScissor(const Rect&) override {} + + void draw(const AbstractGraphicsApi::Buffer*, size_t, size_t, size_t, + size_t, size_t) override {} + void drawIndexed(const AbstractGraphicsApi::Buffer*, size_t, size_t, + const AbstractGraphicsApi::Buffer&, Detail::IndexClass, + size_t, size_t, size_t, size_t) override {} + void drawIndirect(const AbstractGraphicsApi::Buffer&, size_t) override {} + void dispatch(size_t, size_t, size_t) override {} + void dispatchIndirect(const AbstractGraphicsApi::Buffer&, size_t) override {} + + bool spatialUpscale(AbstractGraphicsApi::SpatialScaler& scaler, + AbstractGraphicsApi::Texture& input, + AbstractGraphicsApi::Texture& output) override { + if(!stats.commandSupportsScaler) + return AbstractGraphicsApi::CommandBuffer::spatialUpscale(scaler,input,output); + stats.scalerInput = &input; + stats.scalerOutput = &output; + ++stats.scalerEncoded; + return true; + } + + private: + MockStats& stats; + bool recording = false; + }; + +class MockApi final : public AbstractGraphicsApi { + public: + MockApi(MockStats& stats, bool supportsScaler, bool supportsCommandScaler = true) + :stats(stats),supportsScaler(supportsScaler) { + stats.commandSupportsScaler = supportsCommandScaler; + } + + std::vector devices() const override { return {Props()}; } + + protected: + Device* createDevice(std::string_view) override { return new MockDevice(); } + Swapchain* createSwapchain(SystemApi::Window*, AbstractGraphicsApi::Device*) override { return nullptr; } + + PPipeline createPipeline(Device*, const RenderState&, Topology, + const Shader* const*, size_t) override { + return PPipeline(new MockPipeline()); + } + + PCompPipeline createComputePipeline(Device*, Shader*) override { + return PCompPipeline(new MockCompPipeline()); + } + + PShader createShader(Device*, const void*, size_t) override { + return PShader(new MockShader()); + } + + CommandBuffer* createCommandBuffer(Device*) override { + return new MockCommandBuffer(stats); + } + + DescArray* createDescriptors(Device*, AbstractGraphicsApi::Texture**, size_t, uint32_t) override { + return new DescArray(); + } + DescArray* createDescriptors(Device*, AbstractGraphicsApi::Texture**, size_t, uint32_t, + const Sampler&) override { + return new DescArray(); + } + DescArray* createDescriptors(Device*, AbstractGraphicsApi::Buffer**, size_t) override { + return new DescArray(); + } + + PBuffer createBuffer(Device*, const void*, size_t, MemUsage, BufferHeap) override { + return PBuffer(new MockBuffer()); + } + + PTexture createTexture(Device*, const Pixmap&, TextureFormat, uint32_t) override { + return newTexture(); + } + PTexture createTexture(Device*, uint32_t, uint32_t, uint32_t, TextureFormat) override { + return newTexture(); + } + PTexture createStorage(Device*, uint32_t, uint32_t, uint32_t, TextureFormat) override { + return newTexture(); + } + PTexture createStorage(Device*, uint32_t, uint32_t, uint32_t, uint32_t, + TextureFormat) override { + return newTexture(); + } + + SpatialScaler* createSpatialScaler(Device* device, const SpatialScalerDesc& desc) override { + ++stats.scalerCreateAttempts; + if(!supportsScaler) + return AbstractGraphicsApi::createSpatialScaler(device,desc); + return new MockSpatialScaler(stats); + } + + void readPixels(Device*, Pixmap&, const PTexture, TextureFormat, + uint32_t, uint32_t, uint32_t, bool) override {} + void readBytes(Device*, Buffer*, void*, size_t) override {} + void present(Device*, Swapchain*) override {} + std::shared_ptr submit(Device*, CommandBuffer*) override { return {}; } + + void getCaps(Device*, Props& caps) override { + const uint64_t rgba8 = uint64_t(1) << uint64_t(TextureFormat::RGBA8); + caps.setSamplerFormats(rgba8); + caps.setAttachFormats(rgba8); + caps.setStorageFormats(rgba8); + } + + private: + PTexture newTexture() { + const auto id = NonUniqResId(uint32_t(1) << nextTextureId++); + return PTexture(new MockTexture(id)); + } + + MockStats& stats; + bool supportsScaler; + uint32_t nextTextureId = 0; + }; + +SpatialScalerDesc scalerDesc() { + SpatialScalerDesc desc; + desc.inputFormat = TextureFormat::RGBA8; + desc.outputFormat = TextureFormat::RGBA8; + desc.inputWidth = 2; + desc.inputHeight = 2; + desc.outputWidth = 4; + desc.outputHeight = 4; + return desc; + } + +} + +TEST(SpatialScaler, UnsupportedReturnsEmpty) { + MockStats stats; + MockApi api(stats,false); + Device device(api); + + auto scaler = device.spatialScaler(scalerDesc()); + EXPECT_TRUE(scaler.isEmpty()); + EXPECT_FALSE(bool(scaler)); + EXPECT_EQ(stats.scalerCreateAttempts,1); + EXPECT_EQ(stats.scalerCreated,0); + EXPECT_EQ(stats.scalerDestroyed,0); + + auto input = device.attachment(TextureFormat::RGBA8,2,2); + auto output = device.image2d(TextureFormat::RGBA8,4,4); + auto cmd = device.commandBuffer(); + auto encoder = cmd.startEncoding(device); + EXPECT_FALSE(encoder.spatialUpscale(scaler,input,output)); + EXPECT_EQ(stats.scalerEncoded,0); + } + +TEST(SpatialScaler, UnsupportedCommandReturnsFalse) { + MockStats stats; + MockApi api(stats,true,false); + Device device(api); + + auto scaler = device.spatialScaler(scalerDesc()); + auto input = device.attachment(TextureFormat::RGBA8,2,2); + auto output = device.image2d(TextureFormat::RGBA8,4,4); + auto cmd = device.commandBuffer(); + auto encoder = cmd.startEncoding(device); + + EXPECT_FALSE(scaler.isEmpty()); + EXPECT_FALSE(encoder.spatialUpscale(scaler,input,output)); + EXPECT_EQ(stats.scalerEncoded,0); + } + +TEST(SpatialScaler, OwnsAndDestroysBackendObject) { + MockStats stats; + MockApi api(stats,true); + Device device(api); + + { + auto scaler = device.spatialScaler(scalerDesc()); + EXPECT_FALSE(scaler.isEmpty()); + EXPECT_EQ(stats.scalerCreated,1); + EXPECT_EQ(stats.scalerDestroyed,0); + } + EXPECT_EQ(stats.scalerDestroyed,1); + } + +TEST(SpatialScaler, MoveLeavesSourceEmptyAndReleasesDestination) { + MockStats stats; + MockApi api(stats,true); + Device device(api); + + { + auto first = device.spatialScaler(scalerDesc()); + SpatialScaler second(std::move(first)); + EXPECT_TRUE(first.isEmpty()); + EXPECT_FALSE(second.isEmpty()); + + auto third = device.spatialScaler(scalerDesc()); + EXPECT_EQ(stats.scalerCreated,2); + third = std::move(second); + EXPECT_TRUE(second.isEmpty()); + EXPECT_FALSE(third.isEmpty()); + EXPECT_EQ(stats.scalerDestroyed,1); + + third = std::move(third); + EXPECT_FALSE(third.isEmpty()); + } + EXPECT_EQ(stats.scalerDestroyed,2); + } + +TEST(SpatialScaler, EncoderUsesPublicResources) { + MockStats stats; + MockApi api(stats,true); + Device device(api); + + auto scaler = device.spatialScaler(scalerDesc()); + auto input = device.attachment(TextureFormat::RGBA8,2,2); + auto output = device.image2d(TextureFormat::RGBA8,4,4); + auto cmd = device.commandBuffer(); + + { + auto encoder = cmd.startEncoding(device); + encoder.setFramebuffer({{input,Vec4(),Tempest::Preserve}}); + EXPECT_TRUE(encoder.spatialUpscale(scaler,input,output)); + } + + EXPECT_EQ(stats.scalerEncoded,1); + EXPECT_EQ(stats.renderingBegun,1); + EXPECT_EQ(stats.renderingEnded,1); + EXPECT_NE(stats.scalerInput,nullptr); + EXPECT_NE(stats.scalerOutput,nullptr); + EXPECT_NE(stats.scalerInput,stats.scalerOutput); + } From fe2e12a066dcec52809d14bf37ecb620d2edb13e Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 13:49:48 +0100 Subject: [PATCH 09/25] Clamp iOS frame rates to the active screen --- Engine/system/api/iosapi.mm | 23 +++++++++++++++++++---- 1 file changed, 19 insertions(+), 4 deletions(-) diff --git a/Engine/system/api/iosapi.mm b/Engine/system/api/iosapi.mm index 40d3ea5a..8f06f7b8 100644 --- a/Engine/system/api/iosapi.mm +++ b/Engine/system/api/iosapi.mm @@ -286,20 +286,35 @@ static void applyPreferredFrameRate(CADisplayLink* displayLink) { if(displayLink==nil) return; if(@available(iOS 15.0, *)) { + uint32_t screenMaximum = 60; + if(mainWindow!=nil && mainWindow.screen!=nil && + mainWindow.screen.maximumFramesPerSecond>0) + screenMaximum = uint32_t(mainWindow.screen.maximumFramesPerSecond); switch(frameRateMode) { case FrameRateMode::SystemDefault: displayLink.preferredFrameRateRange = CAFrameRateRangeDefault; break; case FrameRateMode::Fixed: + { + const uint32_t rate = frameRatePreferredminimum ? + frameRatePreferred : minimum) : maximum; displayLink.preferredFrameRateRange = - CAFrameRateRangeMake(frameRateMinimum,frameRateMaximum, - frameRatePreferred); + CAFrameRateRangeMake(minimum,maximum,preferred); break; + } } } else { From 691525beeeedd7bf7b2271da98df5e0e5adbd8e8 Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 14:17:11 +0100 Subject: [PATCH 10/25] Add configurable Metal swapchain buffer count --- Engine/gapi/metal/mtswapchain.h | 2 +- Engine/gapi/metal/mtswapchain.mm | 7 +++++-- Engine/gapi/metalapi.cpp | 14 ++++++++++++-- Engine/gapi/metalapi.h | 8 +++++++- Tests/tests/gapi/metal_test.cpp | 16 ++++++++++++++++ 5 files changed, 41 insertions(+), 6 deletions(-) diff --git a/Engine/gapi/metal/mtswapchain.h b/Engine/gapi/metal/mtswapchain.h index 0f250bd1..eac96f9c 100644 --- a/Engine/gapi/metal/mtswapchain.h +++ b/Engine/gapi/metal/mtswapchain.h @@ -18,7 +18,7 @@ class MtDevice; class MtSwapchain : public AbstractGraphicsApi::Swapchain { public: - MtSwapchain(MtDevice& dev, SystemApi::Window* w); + MtSwapchain(MtDevice& dev, SystemApi::Window* w, uint32_t bufferCount); ~MtSwapchain(); void reset() override; diff --git a/Engine/gapi/metal/mtswapchain.mm b/Engine/gapi/metal/mtswapchain.mm index e0eaf26c..c9f3e50a 100644 --- a/Engine/gapi/metal/mtswapchain.mm +++ b/Engine/gapi/metal/mtswapchain.mm @@ -92,7 +92,7 @@ static CGRect windowRect(UIWindow* wnd) { #endif // note : MoltenVK supports NSView, UIView, CAMetalLayer, so we should align to it -MtSwapchain::MtSwapchain(MtDevice& dev, SystemApi::Window *w) +MtSwapchain::MtSwapchain(MtDevice& dev, SystemApi::Window *w, uint32_t bufferCount) :dev(dev), pimpl(new Impl()) { NSObject* obj = reinterpret_cast(w); if([obj isKindOfClass : [SysWindow class]]) @@ -117,7 +117,10 @@ static CGRect windowRect(UIWindow* wnd) { [lay setContentsScale:dpi]; #if defined(__IOS__) // Swapchain takes too much memory on 2GB iPhone - lay.maximumDrawableCount = 2; + lay.maximumDrawableCount = bufferCount==0 ? 2 : bufferCount; +#elif defined(__OSX__) + if(bufferCount!=0) + lay.maximumDrawableCount = bufferCount; #endif lay.pixelFormat = MTLPixelFormatBGRA8Unorm; lay.allowsNextDrawableTimeout = NO; diff --git a/Engine/gapi/metalapi.cpp b/Engine/gapi/metalapi.cpp index e498d252..82a68fae 100644 --- a/Engine/gapi/metalapi.cpp +++ b/Engine/gapi/metalapi.cpp @@ -23,10 +23,20 @@ #include +#include + using namespace Tempest; using namespace Tempest::Detail; -MetalApi::MetalApi(ApiFlags f) { +MetalApi::MetalApi(ApiFlags f) + :MetalApi(f,Options{}) { + } + +MetalApi::MetalApi(ApiFlags f, const Options& options) + :swapchainBufferCount(options.swapchainBufferCount) { + if(swapchainBufferCount!=0 && swapchainBufferCount!=2 && swapchainBufferCount!=3) + throw std::invalid_argument("Metal swapchain buffer count must be 0, 2, or 3"); + if((f & ApiFlags::Validation)==ApiFlags::Validation) { setenv("METAL_DEVICE_WRAPPER_TYPE","1",1); setenv("METAL_DEBUG_ERROR_MODE", "5",0); @@ -69,7 +79,7 @@ AbstractGraphicsApi::Device* MetalApi::createDevice(std::string_view gpuName) { AbstractGraphicsApi::Swapchain *MetalApi::createSwapchain(SystemApi::Window *w, AbstractGraphicsApi::Device* d) { auto& dev = *reinterpret_cast(d); - return new MtSwapchain(dev,w); + return new MtSwapchain(dev,w,swapchainBufferCount); } AbstractGraphicsApi::PPipeline MetalApi::createPipeline(AbstractGraphicsApi::Device *d, diff --git a/Engine/gapi/metalapi.h b/Engine/gapi/metalapi.h index 98cde1a0..88277c71 100644 --- a/Engine/gapi/metalapi.h +++ b/Engine/gapi/metalapi.h @@ -6,7 +6,12 @@ namespace Tempest { class MetalApi : public AbstractGraphicsApi { public: + struct Options { + uint32_t swapchainBufferCount = 0; + }; + explicit MetalApi(ApiFlags f=ApiFlags::NoFlags); + MetalApi(ApiFlags f, const Options& options); ~MetalApi(); std::vector devices() const override; @@ -45,7 +50,8 @@ class MetalApi : public AbstractGraphicsApi { void getCaps(Device *d, Props& caps) override; private: - bool validation = false; + bool validation = false; + uint32_t swapchainBufferCount = 0; }; } diff --git a/Tests/tests/gapi/metal_test.cpp b/Tests/tests/gapi/metal_test.cpp index d6d7f1c9..15b4174d 100644 --- a/Tests/tests/gapi/metal_test.cpp +++ b/Tests/tests/gapi/metal_test.cpp @@ -19,6 +19,22 @@ TEST(MetalApi,MetalApi) { #endif } +TEST(MetalApi,SwapchainBufferCountOptions) { +#if defined(__OSX__) + for(uint32_t count : {0u,2u,3u}) { + MetalApi::Options options; + options.swapchainBufferCount = count; + EXPECT_NO_THROW(MetalApi(ApiFlags::NoFlags,options)); + } + + for(uint32_t count : {1u,4u}) { + MetalApi::Options options; + options.swapchainBufferCount = count; + EXPECT_THROW(MetalApi(ApiFlags::NoFlags,options),std::invalid_argument); + } +#endif + } + TEST(MetalApi,Vbo) { #if defined(__OSX__) GapiTestCommon::Vbo(); From 4e77354181cc62f3078028485936fe78ae4031b4 Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 14:32:14 +0100 Subject: [PATCH 11/25] Add optional Metal shader module cache --- Engine/gapi/metal/mtdevice.h | 5 +- Engine/gapi/metal/mtdevice.mm | 4 +- Engine/gapi/metal/mtshadermodulecache.h | 136 ++++++++++++++++++ Engine/gapi/metalapi.cpp | 9 +- Engine/gapi/metalapi.h | 4 + Tests/tests/CMakeLists.txt | 1 + Tests/tests/gapi/metal_shader_cache_test.cpp | 137 +++++++++++++++++++ Tests/tests/gapi/metal_test.cpp | 29 ++++ 8 files changed, 319 insertions(+), 6 deletions(-) create mode 100644 Engine/gapi/metal/mtshadermodulecache.h create mode 100644 Tests/tests/gapi/metal_shader_cache_test.cpp diff --git a/Engine/gapi/metal/mtdevice.h b/Engine/gapi/metal/mtdevice.h index 1d877773..422fb1eb 100644 --- a/Engine/gapi/metal/mtdevice.h +++ b/Engine/gapi/metal/mtdevice.h @@ -12,6 +12,7 @@ #include "gapi/shaderreflection.h" #include "gapi/metal/mtsync.h" #include "gapi/metal/mtsamplercache.h" +#include "gapi/metal/mtshadermodulecache.h" #include "nsptr.h" class MTLDevice; @@ -245,7 +246,7 @@ inline MTL::RenderStages nativeFormat(ShaderReflection::Stage st) { class MtDevice : public AbstractGraphicsApi::Device { public: - MtDevice(std::string_view name, bool validation); + MtDevice(std::string_view name, bool validation, size_t shaderModuleCacheSize); ~MtDevice(); static const uint32_t MaxFences = 32; @@ -287,6 +288,8 @@ class MtDevice : public AbstractGraphicsApi::Device { Props prop; MtSamplerCache samplers; + ShaderModuleCache + shaderModules; bool validation = false; static void deductProps(AbstractGraphicsApi::Props& prop, MTL::Device& dev); diff --git a/Engine/gapi/metal/mtdevice.mm b/Engine/gapi/metal/mtdevice.mm index 19c2b3bd..464b3e09 100644 --- a/Engine/gapi/metal/mtdevice.mm +++ b/Engine/gapi/metal/mtdevice.mm @@ -35,8 +35,8 @@ return std::min(MTL::LanguageVersion3_1, opt->languageVersion()); } -MtDevice::MtDevice(std::string_view name, bool validation) - : impl(mkDevice(name)), samplers(*impl), validation(validation) { +MtDevice::MtDevice(std::string_view name, bool validation, size_t shaderModuleCacheSize) + : impl(mkDevice(name)), samplers(*impl), shaderModules(shaderModuleCacheSize), validation(validation) { if(impl.get()==nullptr) throw std::system_error(Tempest::GraphicsErrc::NoDevice); diff --git a/Engine/gapi/metal/mtshadermodulecache.h b/Engine/gapi/metal/mtshadermodulecache.h new file mode 100644 index 00000000..23a868ea --- /dev/null +++ b/Engine/gapi/metal/mtshadermodulecache.h @@ -0,0 +1,136 @@ +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace Tempest { +namespace Detail { + +struct ShaderModuleHash { + size_t operator()(const void* source, size_t size) const noexcept { + // FNV-1a is used only to select a bucket. Cache hits always compare the + // complete shader module, so hash collisions cannot alias two modules. + constexpr size_t offset = sizeof(size_t)==8 ? size_t(14695981039346656037ull) : size_t(2166136261u); + constexpr size_t prime = sizeof(size_t)==8 ? size_t(1099511628211ull) : size_t(16777619u); + + auto* bytes = static_cast(source); + size_t hash = offset; + for(size_t i=0; i +class ShaderModuleCache final { + public: + explicit ShaderModuleCache(size_t capacity, Hash hash = Hash{}) + : capacity(capacity), hash(std::move(hash)) { + } + + ShaderModuleCache(const ShaderModuleCache&) = delete; + ShaderModuleCache& operator=(const ShaderModuleCache&) = delete; + + template + Value getOrCreate(const void* source, size_t size, Factory&& factory) { + if(capacity==0) + return std::forward(factory)(); + + const size_t hashValue = hash(source,size); + { + std::lock_guard guard(sync); + auto at = find(hashValue,source,size); + if(at!=entries.end()) { + entries.splice(entries.begin(),entries,at); + return entries.front().value; + } + } + + std::vector key(size); + if(size>0) + std::memcpy(key.data(),source,size); + + // Shader translation and Metal compilation are intentionally outside + // the cache lock. Concurrent misses may compile the same module; the + // second lookup below selects a single shared cache entry. + Value candidate = std::forward(factory)(); + Value evicted; + { + std::lock_guard guard(sync); + auto at = find(hashValue,key.data(),key.size()); + if(at!=entries.end()) { + entries.splice(entries.begin(),entries,at); + return entries.front().value; + } + + entries.push_front(Entry{hashValue,std::move(key),candidate}); + try { + buckets.emplace(hashValue,entries.begin()); + } + catch(...) { + entries.pop_front(); + throw; + } + + if(entries.size()>capacity) { + auto last = std::prev(entries.end()); + eraseBucket(last); + evicted = std::move(last->value); + entries.erase(last); + } + } + // Keep destruction of an evicted module outside the cache lock. + return candidate; + } + + private: + struct Entry { + size_t hash = 0; + std::vector source; + Value value; + }; + + using EntryList = std::list; + using Iterator = typename EntryList::iterator; + + Iterator find(size_t hashValue, const void* source, size_t size) { + auto range = buckets.equal_range(hashValue); + for(auto i=range.first; i!=range.second; ++i) { + auto at = i->second; + if(at->source.size()!=size) + continue; + if(size==0 || std::memcmp(at->source.data(),source,size)==0) + return at; + } + return entries.end(); + } + + void eraseBucket(Iterator entry) { + auto range = buckets.equal_range(entry->hash); + for(auto i=range.first; i!=range.second; ++i) { + if(i->second==entry) { + buckets.erase(i); + return; + } + } + } + + const size_t capacity; + Hash hash; + + std::mutex sync; + EntryList entries; + std::unordered_multimap buckets; + }; + +} +} diff --git a/Engine/gapi/metalapi.cpp b/Engine/gapi/metalapi.cpp index 82a68fae..aace080e 100644 --- a/Engine/gapi/metalapi.cpp +++ b/Engine/gapi/metalapi.cpp @@ -33,7 +33,8 @@ MetalApi::MetalApi(ApiFlags f) } MetalApi::MetalApi(ApiFlags f, const Options& options) - :swapchainBufferCount(options.swapchainBufferCount) { + :swapchainBufferCount(options.swapchainBufferCount), + shaderModuleCacheSize(options.shaderModuleCacheSize) { if(swapchainBufferCount!=0 && swapchainBufferCount!=2 && swapchainBufferCount!=3) throw std::invalid_argument("Metal swapchain buffer count must be 0, 2, or 3"); @@ -73,7 +74,7 @@ std::vector MetalApi::devices() const { } AbstractGraphicsApi::Device* MetalApi::createDevice(std::string_view gpuName) { - return new MtDevice(gpuName,validation); + return new MtDevice(gpuName,validation,shaderModuleCacheSize); } AbstractGraphicsApi::Swapchain *MetalApi::createSwapchain(SystemApi::Window *w, @@ -103,7 +104,9 @@ AbstractGraphicsApi::PCompPipeline MetalApi::createComputePipeline(AbstractGraph AbstractGraphicsApi::PShader MetalApi::createShader(AbstractGraphicsApi::Device *d, const void *source, size_t src_size) { auto& dx = *reinterpret_cast(d); - return PShader(new MtShader(dx,source,src_size)); + return dx.shaderModules.getOrCreate(source,src_size,[&dx,source,src_size]() { + return PShader(new MtShader(dx,source,src_size)); + }); } AbstractGraphicsApi::PBuffer MetalApi::createBuffer(AbstractGraphicsApi::Device *d, const void *mem, size_t size, diff --git a/Engine/gapi/metalapi.h b/Engine/gapi/metalapi.h index 88277c71..fa5de35a 100644 --- a/Engine/gapi/metalapi.h +++ b/Engine/gapi/metalapi.h @@ -8,6 +8,9 @@ class MetalApi : public AbstractGraphicsApi { public: struct Options { uint32_t swapchainBufferCount = 0; + // Maximum number of compiled Metal shader modules kept per device. + // Zero disables caching. + size_t shaderModuleCacheSize = 0; }; explicit MetalApi(ApiFlags f=ApiFlags::NoFlags); @@ -52,6 +55,7 @@ class MetalApi : public AbstractGraphicsApi { private: bool validation = false; uint32_t swapchainBufferCount = 0; + size_t shaderModuleCacheSize = 0; }; } diff --git a/Tests/tests/CMakeLists.txt b/Tests/tests/CMakeLists.txt index b2ac8bad..6d892e9b 100644 --- a/Tests/tests/CMakeLists.txt +++ b/Tests/tests/CMakeLists.txt @@ -13,6 +13,7 @@ add_test(${PROJECT_NAME} COMMAND ${PROJECT_NAME}) target_include_directories(${PROJECT_NAME} PRIVATE .) target_include_directories(${PROJECT_NAME} PRIVATE "${CMAKE_SOURCE_DIR}/../../Engine/include") +target_include_directories(${PROJECT_NAME} PRIVATE "${CMAKE_SOURCE_DIR}/../../Engine") add_definitions(-DGTEST_LANG_CXX11=1) if(MSVC) diff --git a/Tests/tests/gapi/metal_shader_cache_test.cpp b/Tests/tests/gapi/metal_shader_cache_test.cpp new file mode 100644 index 00000000..2f3db065 --- /dev/null +++ b/Tests/tests/gapi/metal_shader_cache_test.cpp @@ -0,0 +1,137 @@ +#include "gapi/metal/mtshadermodulecache.h" + +#include + +#include +#include +#include +#include +#include +#include +#include + +using namespace Tempest::Detail; + +namespace { + +using Module = std::shared_ptr; + +struct ConstantHash { + size_t operator()(const void*,size_t) const noexcept { + return 1; + } + }; + +} + +TEST(MetalShaderModuleCache,Disabled) { + ShaderModuleCache cache(0); + const uint32_t source = 1; + int compiled = 0; + auto factory = [&]() { + ++compiled; + return std::make_shared(compiled); + }; + + auto first = cache.getOrCreate(&source,sizeof(source),factory); + auto second = cache.getOrCreate(&source,sizeof(source),factory); + EXPECT_NE(first,second); + EXPECT_EQ(compiled,2); + } + +TEST(MetalShaderModuleCache,ReusesModule) { + ShaderModuleCache cache(2); + const uint32_t source = 1; + int compiled = 0; + auto factory = [&]() { + ++compiled; + return std::make_shared(compiled); + }; + + auto first = cache.getOrCreate(&source,sizeof(source),factory); + auto second = cache.getOrCreate(&source,sizeof(source),factory); + EXPECT_EQ(first,second); + EXPECT_EQ(compiled,1); + } + +TEST(MetalShaderModuleCache,EvictsLeastRecentlyUsed) { + ShaderModuleCache cache(2); + const uint32_t a = 1; + const uint32_t b = 2; + const uint32_t c = 3; + int compiled = 0; + auto load = [&](const uint32_t& source) { + return cache.getOrCreate(&source,sizeof(source),[&]() { + ++compiled; + return std::make_shared(compiled); + }); + }; + + auto firstA = load(a); + auto firstB = load(b); + EXPECT_EQ(firstA,load(a)); + auto firstC = load(c); + EXPECT_EQ(compiled,3); + + auto secondB = load(b); + EXPECT_NE(firstB,secondB); + EXPECT_EQ(compiled,4); + + // Eviction drops only the cache's reference. Client-held modules stay alive. + EXPECT_EQ(*firstB,2); + EXPECT_EQ(*firstC,3); + } + +TEST(MetalShaderModuleCache,ComparesBytesOnHashCollision) { + ShaderModuleCache cache(2); + const uint32_t a = 1; + const uint32_t b = 2; + int compiled = 0; + auto load = [&](const uint32_t& source) { + return cache.getOrCreate(&source,sizeof(source),[&]() { + ++compiled; + return std::make_shared(compiled); + }); + }; + + auto firstA = load(a); + auto firstB = load(b); + EXPECT_NE(firstA,firstB); + EXPECT_EQ(firstA,load(a)); + EXPECT_EQ(firstB,load(b)); + EXPECT_EQ(compiled,2); + } + +TEST(MetalShaderModuleCache,ConcurrentMissesCompileOutsideLock) { + ShaderModuleCache cache(4); + const uint32_t sources[] = {1,1,2,2}; + constexpr int threadCount = 4; + std::mutex gateSync; + std::condition_variable gate; + int entered = 0; + std::atomic_bool timedOut{false}; + std::vector results(threadCount); + std::vector threads; + threads.reserve(threadCount); + + for(int i=0; i lock(gateSync); + ++entered; + gate.notify_all(); + if(!gate.wait_for(lock,std::chrono::seconds(2),[&]() { return entered==threadCount; })) + timedOut.store(true); + return std::make_shared(i); + }); + }); + } + for(auto& thread:threads) + thread.join(); + + EXPECT_FALSE(timedOut.load()); + EXPECT_EQ(entered,threadCount); + EXPECT_EQ(results[0],results[1]); + EXPECT_EQ(results[2],results[3]); + EXPECT_NE(results[0],results[2]); + } diff --git a/Tests/tests/gapi/metal_test.cpp b/Tests/tests/gapi/metal_test.cpp index 15b4174d..347f4956 100644 --- a/Tests/tests/gapi/metal_test.cpp +++ b/Tests/tests/gapi/metal_test.cpp @@ -35,6 +35,35 @@ TEST(MetalApi,SwapchainBufferCountOptions) { #endif } +TEST(MetalApi,ShaderModuleCacheOptions) { +#if defined(__OSX__) + try { + MetalApi::Options options; + options.shaderModuleCacheSize = 1; + MetalApi api(ApiFlags::Validation,options); + Device device(api); + + auto firstVert = device.shader("shader/simple_test.vert.sprv"); + auto secondVert = device.shader("shader/simple_test.vert.sprv"); + auto frag = device.shader("shader/simple_test.frag.sprv"); + auto firstPso = device.pipeline(Topology::Triangles,RenderState(),firstVert,frag); + + // Loading the vertex shader after eviction also verifies that client-held + // shader modules remain valid when the cache drops its own reference. + auto thirdVert = device.shader("shader/simple_test.vert.sprv"); + auto secondPso = device.pipeline(Topology::Triangles,RenderState(),thirdVert,frag); + (void)secondVert; + (void)firstPso; + (void)secondPso; + } + catch(std::system_error& e) { + if(e.code()==Tempest::GraphicsErrc::NoDevice) + Log::d("Skipping graphics testcase: ", e.what()); else + throw; + } +#endif + } + TEST(MetalApi,Vbo) { #if defined(__OSX__) GapiTestCommon::Vbo(); From 422b6d38f494b64b6de78919c89558ec731b9962 Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 14:25:47 +0100 Subject: [PATCH 12/25] Preserve AbstractGraphicsApi vtable order --- Engine/gapi/abstractgraphicsapi.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Engine/gapi/abstractgraphicsapi.h b/Engine/gapi/abstractgraphicsapi.h index 4f64469e..abfd9fc8 100644 --- a/Engine/gapi/abstractgraphicsapi.h +++ b/Engine/gapi/abstractgraphicsapi.h @@ -674,7 +674,6 @@ namespace Tempest { virtual AccelerationStructure* createBottomAccelerationStruct(Device* d, const RtGeometry* geom, size_t geomSize); virtual AccelerationStructure* createTopAccelerationStruct(Device* d, const RtInstance* geom, AccelerationStructure*const* as, size_t geomSize); - virtual SpatialScaler* createSpatialScaler(Device* d, const SpatialScalerDesc& desc); virtual void readPixels (Device* d, Pixmap& out, const PTexture t, TextureFormat frm, const uint32_t w, const uint32_t h, uint32_t mip, bool storageImg) = 0; @@ -684,6 +683,7 @@ namespace Tempest { virtual auto submit (Device *d, CommandBuffer* cmd) -> std::shared_ptr = 0; virtual void getCaps(Device *d, Props& caps)=0; + virtual SpatialScaler* createSpatialScaler(Device* d, const SpatialScalerDesc& desc); friend class Tempest::Device; }; From 6c5f63e97d1164b78ad3015a7a20baf5ef62b43c Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 14:29:37 +0100 Subject: [PATCH 13/25] Add portable temporal scaler API --- Engine/gapi/abstractgraphicsapi.cpp | 10 + Engine/gapi/abstractgraphicsapi.h | 25 ++ Engine/graphics/device.cpp | 4 + Engine/graphics/device.h | 2 + Engine/graphics/encoder.cpp | 16 + Engine/graphics/encoder.h | 4 +- Engine/graphics/temporalscaler.cpp | 21 ++ Engine/graphics/temporalscaler.h | 36 ++ Engine/include/Tempest/TemporalScaler | 1 + Tests/tests/temporalscaler_test.cpp | 471 ++++++++++++++++++++++++++ 10 files changed, 589 insertions(+), 1 deletion(-) create mode 100644 Engine/graphics/temporalscaler.cpp create mode 100644 Engine/graphics/temporalscaler.h create mode 100644 Engine/include/Tempest/TemporalScaler create mode 100644 Tests/tests/temporalscaler_test.cpp diff --git a/Engine/gapi/abstractgraphicsapi.cpp b/Engine/gapi/abstractgraphicsapi.cpp index dd13ee94..bfc50dc7 100644 --- a/Engine/gapi/abstractgraphicsapi.cpp +++ b/Engine/gapi/abstractgraphicsapi.cpp @@ -97,6 +97,11 @@ bool AbstractGraphicsApi::CommandBuffer::spatialUpscale(SpatialScaler&, Texture& return false; } +bool AbstractGraphicsApi::CommandBuffer::temporalUpscale(TemporalScaler&, Texture&, Texture&, + Texture&, Texture&, const TemporalScalerArgs&) { + return false; + } + AbstractGraphicsApi::AccelerationStructure* AbstractGraphicsApi::createBottomAccelerationStruct(Device* d, const RtGeometry* geom, size_t geomSize) { throw std::system_error(Tempest::GraphicsErrc::UnsupportedExtension); } @@ -111,6 +116,11 @@ AbstractGraphicsApi::SpatialScaler* return nullptr; } +AbstractGraphicsApi::TemporalScaler* + AbstractGraphicsApi::createTemporalScaler(Device*, const TemporalScalerDesc&) { + return nullptr; + } + bool Detail::Bindings::operator ==(const Bindings &other) const { for(size_t i=0; i; @@ -684,6 +708,7 @@ namespace Tempest { virtual void getCaps(Device *d, Props& caps)=0; virtual SpatialScaler* createSpatialScaler(Device* d, const SpatialScalerDesc& desc); + virtual TemporalScaler* createTemporalScaler(Device* d, const TemporalScalerDesc& desc); friend class Tempest::Device; }; diff --git a/Engine/graphics/device.cpp b/Engine/graphics/device.cpp index e956909f..a5dcf6eb 100644 --- a/Engine/graphics/device.cpp +++ b/Engine/graphics/device.cpp @@ -263,6 +263,10 @@ SpatialScaler Device::spatialScaler(const SpatialScalerDesc& desc) { return SpatialScaler(api.createSpatialScaler(dev,desc)); } +TemporalScaler Device::temporalScaler(const TemporalScalerDesc& desc) { + return TemporalScaler(api.createTemporalScaler(dev,desc)); + } + ZBuffer Device::zbuffer(TextureFormat frm, const Size sz) { if(sz.w<0 || sz.h<0) throw std::system_error(Tempest::GraphicsErrc::TooLargeTexture, std::to_string(std::min(sz.w,sz.h))); diff --git a/Engine/graphics/device.h b/Engine/graphics/device.h index 0c053f97..f140c611 100644 --- a/Engine/graphics/device.h +++ b/Engine/graphics/device.h @@ -16,6 +16,7 @@ #include #include #include +#include #include #include #include @@ -133,6 +134,7 @@ class Device { StorageImage image2d (TextureFormat frm, const Size sz, const bool mips = false); SpatialScaler spatialScaler(const SpatialScalerDesc& desc); + TemporalScaler temporalScaler(const TemporalScalerDesc& desc); AccelerationStructure blas(const std::vector& geom); AccelerationStructure blas(std::initializer_list geom); diff --git a/Engine/graphics/encoder.cpp b/Engine/graphics/encoder.cpp index 8a5c1210..8e31c696 100644 --- a/Engine/graphics/encoder.cpp +++ b/Engine/graphics/encoder.cpp @@ -335,3 +335,19 @@ bool Encoder::spatialUpscale(const SpatialScaler& scaler, const A auto& dst = *output.tImpl.impl.handler; return impl->spatialUpscale(*scaler.impl.handler,src,dst); } + +bool Encoder::temporalUpscale(TemporalScaler& scaler, const Attachment& input, + const ZBuffer& depth, const Attachment& motion, + StorageImage& output, const TemporalScalerArgs& args) { + if(scaler.isEmpty() || input.isEmpty() || depth.isEmpty() || motion.isEmpty() || output.isEmpty()) + return false; + if(state.stage==Rendering) + impl->endRendering(); + state = State(); + + auto& src = *textureCast(input).impl.handler; + auto& dep = *textureCast(depth).impl.handler; + auto& mov = *textureCast(motion).impl.handler; + auto& dst = *textureCast(output).impl.handler; + return impl->temporalUpscale(*scaler.impl.handler,src,dep,mov,dst,args); + } diff --git a/Engine/graphics/encoder.h b/Engine/graphics/encoder.h index 9f567317..dc4ea0fc 100644 --- a/Engine/graphics/encoder.h +++ b/Engine/graphics/encoder.h @@ -8,6 +8,7 @@ #include #include #include +#include namespace Tempest { @@ -121,6 +122,8 @@ class Encoder { void generateMipmaps(Attachment& tex); bool spatialUpscale(const SpatialScaler& scaler, const Attachment& input, StorageImage& output); + bool temporalUpscale(TemporalScaler& scaler, const Attachment& input, const ZBuffer& depth, + const Attachment& motion, StorageImage& output, const TemporalScalerArgs& args); private: explicit Encoder(CommandBuffer* ow); @@ -149,4 +152,3 @@ class Encoder { friend class CommandBuffer; }; } - diff --git a/Engine/graphics/temporalscaler.cpp b/Engine/graphics/temporalscaler.cpp new file mode 100644 index 00000000..29853b0a --- /dev/null +++ b/Engine/graphics/temporalscaler.cpp @@ -0,0 +1,21 @@ +#include "temporalscaler.h" + +using namespace Tempest; + +TemporalScaler::TemporalScaler(TemporalScaler&& other) noexcept + :impl(other.impl.handler) { + other.impl.handler = nullptr; + } + +TemporalScaler::~TemporalScaler() { + delete impl.handler; + } + +TemporalScaler& TemporalScaler::operator=(TemporalScaler&& other) noexcept { + if(this==&other) + return *this; + delete impl.handler; + impl.handler = other.impl.handler; + other.impl.handler = nullptr; + return *this; + } diff --git a/Engine/graphics/temporalscaler.h b/Engine/graphics/temporalscaler.h new file mode 100644 index 00000000..ed60fa1c --- /dev/null +++ b/Engine/graphics/temporalscaler.h @@ -0,0 +1,36 @@ +#pragma once + +#include + +#include "../utility/dptr.h" + +namespace Tempest { + +class Device; +class CommandBuffer; +template +class Encoder; + +class TemporalScaler final { + public: + TemporalScaler() = default; + TemporalScaler(TemporalScaler&& other) noexcept; + TemporalScaler(const TemporalScaler&) = delete; + ~TemporalScaler(); + + TemporalScaler& operator=(TemporalScaler&& other) noexcept; + TemporalScaler& operator=(const TemporalScaler&) = delete; + + bool isEmpty() const { return impl.handler==nullptr; } + explicit operator bool() const { return !isEmpty(); } + + private: + explicit TemporalScaler(AbstractGraphicsApi::TemporalScaler* scaler):impl(scaler) {} + + Detail::DPtr impl; + + friend class Tempest::Device; + friend class Encoder; + }; + +} diff --git a/Engine/include/Tempest/TemporalScaler b/Engine/include/Tempest/TemporalScaler new file mode 100644 index 00000000..2aa7e1fa --- /dev/null +++ b/Engine/include/Tempest/TemporalScaler @@ -0,0 +1 @@ +#include "../graphics/temporalscaler.h" diff --git a/Tests/tests/temporalscaler_test.cpp b/Tests/tests/temporalscaler_test.cpp new file mode 100644 index 00000000..85759c3c --- /dev/null +++ b/Tests/tests/temporalscaler_test.cpp @@ -0,0 +1,471 @@ +#include +#include +#include + +#include + +#include +#include + +using namespace Tempest; + +namespace { + +struct MockStats { + int scalerCreateAttempts = 0; + int scalerCreated = 0; + int scalerDestroyed = 0; + int scalerEncoded = 0; + int renderingBegun = 0; + int renderingEnded = 0; + int computePipelinesSet = 0; + bool commandSupportsScaler = true; + + TemporalScalerDesc scalerDesc; + TemporalScalerArgs scalerArgs; + + AbstractGraphicsApi::Texture* scalerInput = nullptr; + AbstractGraphicsApi::Texture* scalerDepth = nullptr; + AbstractGraphicsApi::Texture* scalerMotion = nullptr; + AbstractGraphicsApi::Texture* scalerOutput = nullptr; + }; + +struct MockDevice final : AbstractGraphicsApi::Device { + void waitIdle() override {} + }; + +struct MockShader final : AbstractGraphicsApi::Shader {}; + +struct MockPipeline final : AbstractGraphicsApi::Pipeline { + IVec3 workGroupSize() const override { return {1,1,1}; } + size_t sizeofBuffer(size_t, size_t) const override { return 0; } + }; + +struct MockCompPipeline final : AbstractGraphicsApi::CompPipeline { + IVec3 workGroupSize() const override { return {1,1,1}; } + size_t sizeofBuffer(size_t, size_t) const override { return 0; } + }; + +struct MockBuffer final : AbstractGraphicsApi::Buffer { + void update(const void*, size_t, size_t) override {} + void read(void*, size_t, size_t) override {} + }; + +struct MockTexture final : AbstractGraphicsApi::Texture { + explicit MockTexture(NonUniqResId id):id(id) {} + + uint32_t mipCount() const override { return 1; } + NonUniqResId syncId() const override { return id; } + + NonUniqResId id; + }; + +struct MockTemporalScaler final : AbstractGraphicsApi::TemporalScaler { + explicit MockTemporalScaler(MockStats& stats):stats(stats) { + ++stats.scalerCreated; + } + + ~MockTemporalScaler() override { + ++stats.scalerDestroyed; + } + + MockStats& stats; + }; + +class MockCommandBuffer final : public AbstractGraphicsApi::CommandBuffer { + public: + explicit MockCommandBuffer(MockStats& stats):stats(stats) {} + + void beginRendering(const Detail::FrameBufferDesc&, size_t, uint32_t, uint32_t) override { + ++stats.renderingBegun; + } + void endRendering() override { + ++stats.renderingEnded; + } + + void barrier(const AbstractGraphicsApi::SyncDesc&, + const AbstractGraphicsApi::BarrierDesc*, size_t) override {} + + void generateMipmap(AbstractGraphicsApi::Texture&, uint32_t, uint32_t, uint32_t) override {} + void copy(AbstractGraphicsApi::Buffer&, size_t, AbstractGraphicsApi::Texture&, + uint32_t, uint32_t, uint32_t) override {} + + bool isRecording() const override { return recording; } + void begin() override { recording = true; } + void end() override { recording = false; } + void reset() override { recording = false; } + + void setPipeline(AbstractGraphicsApi::Pipeline&) override {} + void setComputePipeline(AbstractGraphicsApi::CompPipeline&) override { + ++stats.computePipelinesSet; + } + void setBinding(size_t, AbstractGraphicsApi::Texture*, uint32_t, + const ComponentMapping&, const Sampler&) override {} + void setBinding(size_t, AbstractGraphicsApi::Buffer*, size_t) override {} + void setBinding(size_t, AbstractGraphicsApi::DescArray*) override {} + void setBinding(size_t, AbstractGraphicsApi::AccelerationStructure*) override {} + void setBinding(size_t, const Sampler&) override {} + + void setViewport(const Rect&) override {} + void setScissor(const Rect&) override {} + + void draw(const AbstractGraphicsApi::Buffer*, size_t, size_t, size_t, + size_t, size_t) override {} + void drawIndexed(const AbstractGraphicsApi::Buffer*, size_t, size_t, + const AbstractGraphicsApi::Buffer&, Detail::IndexClass, + size_t, size_t, size_t, size_t) override {} + void drawIndirect(const AbstractGraphicsApi::Buffer&, size_t) override {} + void dispatch(size_t, size_t, size_t) override {} + void dispatchIndirect(const AbstractGraphicsApi::Buffer&, size_t) override {} + + bool temporalUpscale(AbstractGraphicsApi::TemporalScaler& scaler, + AbstractGraphicsApi::Texture& input, + AbstractGraphicsApi::Texture& depth, + AbstractGraphicsApi::Texture& motion, + AbstractGraphicsApi::Texture& output, + const TemporalScalerArgs& args) override { + if(!stats.commandSupportsScaler) + return AbstractGraphicsApi::CommandBuffer::temporalUpscale(scaler,input,depth,motion,output,args); + stats.scalerInput = &input; + stats.scalerDepth = &depth; + stats.scalerMotion = &motion; + stats.scalerOutput = &output; + stats.scalerArgs = args; + ++stats.scalerEncoded; + return true; + } + + private: + MockStats& stats; + bool recording = false; + }; + +class MockApi final : public AbstractGraphicsApi { + public: + MockApi(MockStats& stats, bool supportsScaler, bool supportsCommandScaler = true) + :stats(stats),supportsScaler(supportsScaler) { + stats.commandSupportsScaler = supportsCommandScaler; + } + + std::vector devices() const override { return {Props()}; } + + protected: + Device* createDevice(std::string_view) override { return new MockDevice(); } + Swapchain* createSwapchain(SystemApi::Window*, AbstractGraphicsApi::Device*) override { return nullptr; } + + PPipeline createPipeline(Device*, const RenderState&, Topology, + const Shader* const*, size_t) override { + return PPipeline(new MockPipeline()); + } + + PCompPipeline createComputePipeline(Device*, Shader*) override { + return PCompPipeline(new MockCompPipeline()); + } + + PShader createShader(Device*, const void*, size_t) override { + return PShader(new MockShader()); + } + + CommandBuffer* createCommandBuffer(Device*) override { + return new MockCommandBuffer(stats); + } + + DescArray* createDescriptors(Device*, AbstractGraphicsApi::Texture**, size_t, uint32_t) override { + return new DescArray(); + } + DescArray* createDescriptors(Device*, AbstractGraphicsApi::Texture**, size_t, uint32_t, + const Sampler&) override { + return new DescArray(); + } + DescArray* createDescriptors(Device*, AbstractGraphicsApi::Buffer**, size_t) override { + return new DescArray(); + } + + PBuffer createBuffer(Device*, const void*, size_t, MemUsage, BufferHeap) override { + return PBuffer(new MockBuffer()); + } + + PTexture createTexture(Device*, const Pixmap&, TextureFormat, uint32_t) override { + return newTexture(); + } + PTexture createTexture(Device*, uint32_t, uint32_t, uint32_t, TextureFormat) override { + return newTexture(); + } + PTexture createStorage(Device*, uint32_t, uint32_t, uint32_t, TextureFormat) override { + return newTexture(); + } + PTexture createStorage(Device*, uint32_t, uint32_t, uint32_t, uint32_t, + TextureFormat) override { + return newTexture(); + } + + TemporalScaler* createTemporalScaler(Device* device, const TemporalScalerDesc& desc) override { + ++stats.scalerCreateAttempts; + stats.scalerDesc = desc; + if(!supportsScaler) + return AbstractGraphicsApi::createTemporalScaler(device,desc); + return new MockTemporalScaler(stats); + } + + void readPixels(Device*, Pixmap&, const PTexture, TextureFormat, + uint32_t, uint32_t, uint32_t, bool) override {} + void readBytes(Device*, Buffer*, void*, size_t) override {} + void present(Device*, Swapchain*) override {} + std::shared_ptr submit(Device*, CommandBuffer*) override { return {}; } + + void getCaps(Device*, Props& caps) override { + const uint64_t rgba8 = uint64_t(1) << uint64_t(TextureFormat::RGBA8); + const uint64_t rgba16f = uint64_t(1) << uint64_t(TextureFormat::RGBA16F); + const uint64_t rg32f = uint64_t(1) << uint64_t(TextureFormat::RG32F); + const uint64_t depth32 = uint64_t(1) << uint64_t(TextureFormat::Depth32F); + caps.setSamplerFormats(rgba8|rgba16f|rg32f|depth32); + caps.setAttachFormats(rgba8|rgba16f|rg32f); + caps.setDepthFormats(depth32); + caps.setStorageFormats(rgba8|rgba16f); + } + + private: + PTexture newTexture() { + const auto id = NonUniqResId(uint32_t(1) << nextTextureId++); + return PTexture(new MockTexture(id)); + } + + MockStats& stats; + bool supportsScaler; + uint32_t nextTextureId = 0; + }; + +TemporalScalerDesc scalerDesc() { + TemporalScalerDesc desc; + desc.inputFormat = TextureFormat::RGBA16F; + desc.depthFormat = TextureFormat::Depth32F; + desc.motionFormat = TextureFormat::RG32F; + desc.outputFormat = TextureFormat::RGBA8; + desc.inputWidth = 960; + desc.inputHeight = 540; + desc.outputWidth = 1920; + desc.outputHeight = 1080; + desc.autoExposure = false; + return desc; + } + +} + +TEST(TemporalScaler, UnsupportedReturnsEmpty) { + MockStats stats; + MockApi api(stats,false); + Device device(api); + + auto scaler = device.temporalScaler(scalerDesc()); + EXPECT_TRUE(scaler.isEmpty()); + EXPECT_FALSE(bool(scaler)); + EXPECT_EQ(stats.scalerCreateAttempts,1); + EXPECT_EQ(stats.scalerCreated,0); + EXPECT_EQ(stats.scalerDestroyed,0); + + auto input = device.attachment(TextureFormat::RGBA16F,2,2); + auto depth = device.zbuffer(TextureFormat::Depth32F,2,2); + auto motion = device.attachment(TextureFormat::RG32F,2,2); + auto output = device.image2d(TextureFormat::RGBA16F,4,4); + auto cmd = device.commandBuffer(); + auto encoder = cmd.startEncoding(device); + EXPECT_FALSE(encoder.temporalUpscale(scaler,input,depth,motion,output,{})); + EXPECT_EQ(stats.scalerEncoded,0); + } + +TEST(TemporalScaler, UnsupportedCommandReturnsFalse) { + MockStats stats; + MockApi api(stats,true,false); + Device device(api); + + auto scaler = device.temporalScaler(scalerDesc()); + auto input = device.attachment(TextureFormat::RGBA16F,2,2); + auto depth = device.zbuffer(TextureFormat::Depth32F,2,2); + auto motion = device.attachment(TextureFormat::RG32F,2,2); + auto output = device.image2d(TextureFormat::RGBA16F,4,4); + auto cmd = device.commandBuffer(); + auto encoder = cmd.startEncoding(device); + + EXPECT_FALSE(scaler.isEmpty()); + EXPECT_FALSE(encoder.temporalUpscale(scaler,input,depth,motion,output,{})); + EXPECT_EQ(stats.scalerEncoded,0); + } + +TEST(TemporalScaler, EveryEmptyResourceReturnsFalse) { + MockStats stats; + MockApi api(stats,true); + Device device(api); + + auto scaler = device.temporalScaler(scalerDesc()); + auto input = device.attachment(TextureFormat::RGBA16F,2,2); + auto depth = device.zbuffer(TextureFormat::Depth32F,2,2); + auto motion = device.attachment(TextureFormat::RG32F,2,2); + auto output = device.image2d(TextureFormat::RGBA16F,4,4); + auto cmd = device.commandBuffer(); + auto encoder = cmd.startEncoding(device); + + TemporalScaler emptyScaler; + Attachment emptyInput; + ZBuffer emptyDepth; + Attachment emptyMotion; + StorageImage emptyOutput; + + EXPECT_FALSE(encoder.temporalUpscale(emptyScaler,input,depth,motion,output,{})); + EXPECT_FALSE(encoder.temporalUpscale(scaler,emptyInput,depth,motion,output,{})); + EXPECT_FALSE(encoder.temporalUpscale(scaler,input,emptyDepth,motion,output,{})); + EXPECT_FALSE(encoder.temporalUpscale(scaler,input,depth,emptyMotion,output,{})); + EXPECT_FALSE(encoder.temporalUpscale(scaler,input,depth,motion,emptyOutput,{})); + EXPECT_EQ(stats.scalerEncoded,0); + } + +TEST(TemporalScaler, DeviceForwardsDescriptor) { + MockStats stats; + MockApi api(stats,true); + Device device(api); + + const auto expected = scalerDesc(); + auto scaler = device.temporalScaler(expected); + + EXPECT_TRUE(bool(scaler)); + EXPECT_EQ(stats.scalerDesc.inputFormat,expected.inputFormat); + EXPECT_EQ(stats.scalerDesc.depthFormat,expected.depthFormat); + EXPECT_EQ(stats.scalerDesc.motionFormat,expected.motionFormat); + EXPECT_EQ(stats.scalerDesc.outputFormat,expected.outputFormat); + EXPECT_EQ(stats.scalerDesc.inputWidth,expected.inputWidth); + EXPECT_EQ(stats.scalerDesc.inputHeight,expected.inputHeight); + EXPECT_EQ(stats.scalerDesc.outputWidth,expected.outputWidth); + EXPECT_EQ(stats.scalerDesc.outputHeight,expected.outputHeight); + EXPECT_EQ(stats.scalerDesc.autoExposure,expected.autoExposure); + } + +TEST(TemporalScaler, OwnsAndDestroysBackendObject) { + MockStats stats; + MockApi api(stats,true); + Device device(api); + + { + auto scaler = device.temporalScaler(scalerDesc()); + EXPECT_FALSE(scaler.isEmpty()); + EXPECT_EQ(stats.scalerCreated,1); + EXPECT_EQ(stats.scalerDestroyed,0); + } + EXPECT_EQ(stats.scalerDestroyed,1); + } + +TEST(TemporalScaler, MoveLeavesSourceEmptyAndReleasesDestination) { + MockStats stats; + MockApi api(stats,true); + Device device(api); + + { + auto first = device.temporalScaler(scalerDesc()); + TemporalScaler second(std::move(first)); + EXPECT_TRUE(first.isEmpty()); + EXPECT_FALSE(second.isEmpty()); + + auto third = device.temporalScaler(scalerDesc()); + EXPECT_EQ(stats.scalerCreated,2); + third = std::move(second); + EXPECT_TRUE(second.isEmpty()); + EXPECT_FALSE(third.isEmpty()); + EXPECT_EQ(stats.scalerDestroyed,1); + + third = std::move(third); + EXPECT_FALSE(third.isEmpty()); + } + EXPECT_EQ(stats.scalerDestroyed,2); + } + +TEST(TemporalScaler, EncoderForwardsResourcesAndArgsAndEndsRendering) { + MockStats stats; + MockApi api(stats,true); + Device device(api); + + auto scaler = device.temporalScaler(scalerDesc()); + auto input = device.attachment(TextureFormat::RGBA16F,2,2); + auto depth = device.zbuffer(TextureFormat::Depth32F,2,2); + auto motion = device.attachment(TextureFormat::RG32F,2,2); + auto output = device.image2d(TextureFormat::RGBA16F,4,4); + auto cmd = device.commandBuffer(); + + TemporalScalerArgs args; + args.jitterOffsetX = 0.25f; + args.jitterOffsetY = -0.5f; + args.motionVectorScaleX = 2.f; + args.motionVectorScaleY = -3.f; + args.resetHistory = true; + args.depthReversed = true; + + { + auto encoder = cmd.startEncoding(device); + encoder.setFramebuffer({{input,Vec4(),Tempest::Preserve}}); + EXPECT_TRUE(encoder.temporalUpscale(scaler,input,depth,motion,output,args)); + } + + EXPECT_EQ(stats.scalerEncoded,1); + EXPECT_EQ(stats.renderingBegun,1); + EXPECT_EQ(stats.renderingEnded,1); + auto* recordedInput = dynamic_cast(stats.scalerInput); + auto* recordedDepth = dynamic_cast(stats.scalerDepth); + auto* recordedMotion = dynamic_cast(stats.scalerMotion); + auto* recordedOutput = dynamic_cast(stats.scalerOutput); + ASSERT_NE(recordedInput,nullptr); + ASSERT_NE(recordedDepth,nullptr); + ASSERT_NE(recordedMotion,nullptr); + ASSERT_NE(recordedOutput,nullptr); + EXPECT_EQ(recordedInput->syncId(),NonUniqResId(1)); + EXPECT_EQ(recordedDepth->syncId(),NonUniqResId(2)); + EXPECT_EQ(recordedMotion->syncId(),NonUniqResId(4)); + EXPECT_EQ(recordedOutput->syncId(),NonUniqResId(8)); + EXPECT_FLOAT_EQ(stats.scalerArgs.jitterOffsetX,args.jitterOffsetX); + EXPECT_FLOAT_EQ(stats.scalerArgs.jitterOffsetY,args.jitterOffsetY); + EXPECT_FLOAT_EQ(stats.scalerArgs.motionVectorScaleX,args.motionVectorScaleX); + EXPECT_FLOAT_EQ(stats.scalerArgs.motionVectorScaleY,args.motionVectorScaleY); + EXPECT_EQ(stats.scalerArgs.resetHistory,args.resetHistory); + EXPECT_EQ(stats.scalerArgs.depthReversed,args.depthReversed); + } + +TEST(TemporalScaler, RejectsThreeDimensionalOutput) { + MockStats stats; + MockApi api(stats,true); + Device device(api); + + auto scaler = device.temporalScaler(scalerDesc()); + auto input = device.attachment(TextureFormat::RGBA16F,2,2); + auto depth = device.zbuffer(TextureFormat::Depth32F,2,2); + auto motion = device.attachment(TextureFormat::RG32F,2,2); + auto output = device.image3d(TextureFormat::RGBA16F,4,4,2); + auto cmd = device.commandBuffer(); + auto encoder = cmd.startEncoding(device); + + EXPECT_THROW(encoder.temporalUpscale(scaler,input,depth,motion,output,{}),BadTextureCastException); + EXPECT_EQ(stats.scalerEncoded,0); + } + +TEST(TemporalScaler, ClearsComputePipelineCache) { + MockStats stats; + MockApi api(stats,true); + Device device(api); + + const uint32_t shaderCode = 0; + auto shader = device.shader(&shaderCode,sizeof(shaderCode)); + auto pipeline = device.pipeline(shader); + auto scaler = device.temporalScaler(scalerDesc()); + auto input = device.attachment(TextureFormat::RGBA16F,2,2); + auto depth = device.zbuffer(TextureFormat::Depth32F,2,2); + auto motion = device.attachment(TextureFormat::RG32F,2,2); + auto output = device.image2d(TextureFormat::RGBA16F,4,4); + auto cmd = device.commandBuffer(); + + { + auto encoder = cmd.startEncoding(device); + encoder.setPipeline(pipeline); + EXPECT_EQ(stats.computePipelinesSet,1); + EXPECT_TRUE(encoder.temporalUpscale(scaler,input,depth,motion,output,{})); + encoder.setPipeline(pipeline); + EXPECT_EQ(stats.computePipelinesSet,2); + } + + EXPECT_EQ(stats.renderingEnded,0); + EXPECT_EQ(stats.scalerEncoded,1); + } From 066ef1ed89ca66b10617ba2ad94a1b1db8b5282f Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 14:25:22 +0100 Subject: [PATCH 14/25] Add optional MetalFX spatial scaler backend --- Engine/CMakeLists.txt | 12 +++ Engine/gapi/metal/metal_cpp.mm | 6 ++ Engine/gapi/metal/mtcommandbuffer.cpp | 15 ++++ Engine/gapi/metal/mtcommandbuffer.h | 8 ++ Engine/gapi/metal/mtspatialscaler.h | 30 +++++++ Engine/gapi/metal/mtspatialscaler.mm | 108 ++++++++++++++++++++++++++ Engine/gapi/metalapi.cpp | 7 ++ Engine/gapi/metalapi.h | 1 + Tests/tests/gapi/metal_test.cpp | 73 +++++++++++++++++ 9 files changed, 260 insertions(+) create mode 100644 Engine/gapi/metal/mtspatialscaler.h create mode 100644 Engine/gapi/metal/mtspatialscaler.mm diff --git a/Engine/CMakeLists.txt b/Engine/CMakeLists.txt index 6d462978..835e9446 100644 --- a/Engine/CMakeLists.txt +++ b/Engine/CMakeLists.txt @@ -190,6 +190,12 @@ if(TEMPEST_BUILD_METAL) add_definitions(-DTEMPEST_BUILD_METAL) add_subdirectory("thirdparty/metal-cpp" EXCLUDE_FROM_ALL) target_include_directories(${PROJECT_NAME} PRIVATE "thirdparty/metal-cpp") + if(APPLE) + find_library(TEMPEST_METALFX_FRAMEWORK NAMES MetalFX) + if(TEMPEST_METALFX_FRAMEWORK) + target_compile_definitions(${PROJECT_NAME} PRIVATE TEMPEST_BUILD_METALFX) + endif() + endif() endif() ### Spirv-cross @@ -331,8 +337,14 @@ if(WIN32) elseif(IOS) #NOTE: use -DCMAKE_SYSTEM_NAME=iOS on first configure target_link_libraries(${PROJECT_NAME} PRIVATE "-framework UiKit" "-framework Foundation" "-framework QuartzCore" "-framework Metal") + if(TEMPEST_METALFX_FRAMEWORK) + target_link_libraries(${PROJECT_NAME} PRIVATE "-weak_framework MetalFX") + endif() elseif(APPLE) target_link_libraries(${PROJECT_NAME} PRIVATE "-framework AppKit" "-framework QuartzCore" "-framework Metal") + if(TEMPEST_METALFX_FRAMEWORK) + target_link_libraries(${PROJECT_NAME} PRIVATE "-weak_framework MetalFX") + endif() elseif(UNIX) target_link_libraries(${PROJECT_NAME} PRIVATE X11 Xcursor) endif() diff --git a/Engine/gapi/metal/metal_cpp.mm b/Engine/gapi/metal/metal_cpp.mm index 7fa7137b..6516b982 100644 --- a/Engine/gapi/metal/metal_cpp.mm +++ b/Engine/gapi/metal/metal_cpp.mm @@ -1,6 +1,12 @@ #define NS_PRIVATE_IMPLEMENTATION #define CA_PRIVATE_IMPLEMENTATION #define MTL_PRIVATE_IMPLEMENTATION +#if defined(TEMPEST_BUILD_METALFX) +#define MTLFX_PRIVATE_IMPLEMENTATION +#endif #include #include #include +#if defined(TEMPEST_BUILD_METALFX) +#include +#endif diff --git a/Engine/gapi/metal/mtcommandbuffer.cpp b/Engine/gapi/metal/mtcommandbuffer.cpp index 2997f449..75c67069 100644 --- a/Engine/gapi/metal/mtcommandbuffer.cpp +++ b/Engine/gapi/metal/mtcommandbuffer.cpp @@ -10,6 +10,9 @@ #include "mttexture.h" #include "mtswapchain.h" #include "mtaccelerationstructure.h" +#if defined(TEMPEST_BUILD_METALFX) +#include "mtspatialscaler.h" +#endif using namespace Tempest; using namespace Tempest::Detail; @@ -796,4 +799,16 @@ void MtCommandBuffer::copy(AbstractGraphicsApi::Buffer& dest, size_t offset, offset, bpp*width,bpp*width*height); } +#if defined(TEMPEST_BUILD_METALFX) +bool MtCommandBuffer::spatialUpscale(AbstractGraphicsApi::SpatialScaler& scaler, + AbstractGraphicsApi::Texture& input, + AbstractGraphicsApi::Texture& output) { + setEncoder(E_None,nullptr); + auto& sx = reinterpret_cast(scaler); + auto& src = reinterpret_cast(input); + auto& dst = reinterpret_cast(output); + return sx.encode(*impl,src,dst); + } +#endif + #endif diff --git a/Engine/gapi/metal/mtcommandbuffer.h b/Engine/gapi/metal/mtcommandbuffer.h index 0d49c76c..8f49f9c4 100644 --- a/Engine/gapi/metal/mtcommandbuffer.h +++ b/Engine/gapi/metal/mtcommandbuffer.h @@ -20,6 +20,9 @@ class MtPipeline; class MtCompPipeline; class MtDescriptorArray; class MtTopAccelerationStructure; +#if defined(TEMPEST_BUILD_METALFX) +class MtSpatialScaler; +#endif class MtCommandBuffer : public AbstractGraphicsApi::CommandBuffer { public: @@ -66,6 +69,11 @@ class MtCommandBuffer : public AbstractGraphicsApi::CommandBuffer { void generateMipmap(AbstractGraphicsApi::Texture& image, uint32_t texWidth, uint32_t texHeight, uint32_t mipLevels) override; void copy (AbstractGraphicsApi::Buffer& dst, size_t offset, AbstractGraphicsApi::Texture& src, uint32_t width, uint32_t height, uint32_t mip) override; +#if defined(TEMPEST_BUILD_METALFX) + bool spatialUpscale(AbstractGraphicsApi::SpatialScaler& scaler, + AbstractGraphicsApi::Texture& input, + AbstractGraphicsApi::Texture& output) override; +#endif private: enum EncType:uint8_t { diff --git a/Engine/gapi/metal/mtspatialscaler.h b/Engine/gapi/metal/mtspatialscaler.h new file mode 100644 index 00000000..48727550 --- /dev/null +++ b/Engine/gapi/metal/mtspatialscaler.h @@ -0,0 +1,30 @@ +#pragma once + +#if defined(TEMPEST_BUILD_METALFX) + +#include +#include + +#include "nsptr.h" + +namespace Tempest { +namespace Detail { + +class MtDevice; +class MtTexture; + +class MtSpatialScaler final : public AbstractGraphicsApi::SpatialScaler { + public: + MtSpatialScaler(MtDevice& device, const SpatialScalerDesc& desc); + + bool isValid() const { return impl!=nullptr; } + bool encode(MTL::CommandBuffer& cmd, MtTexture& input, MtTexture& output); + + private: + NsPtr impl; + }; + +} +} + +#endif diff --git a/Engine/gapi/metal/mtspatialscaler.mm b/Engine/gapi/metal/mtspatialscaler.mm new file mode 100644 index 00000000..b9a12024 --- /dev/null +++ b/Engine/gapi/metal/mtspatialscaler.mm @@ -0,0 +1,108 @@ +#if defined(TEMPEST_BUILD_METALFX) + +#include "mtspatialscaler.h" + +#include +#include + +#include "mtdevice.h" +#include "mttexture.h" + +using namespace Tempest; +using namespace Tempest::Detail; + +namespace { + +MTLFX::SpatialScalerColorProcessingMode nativeColorMode(SpatialScalerColorMode mode) { + switch(mode) { + case SpatialScalerColorMode::Perceptual: + return MTLFX::SpatialScalerColorProcessingModePerceptual; + case SpatialScalerColorMode::Linear: + return MTLFX::SpatialScalerColorProcessingModeLinear; + case SpatialScalerColorMode::HDR: + return MTLFX::SpatialScalerColorProcessingModeHDR; + } + return MTLFX::SpatialScalerColorProcessingModePerceptual; + } + +bool isMetalFxAvailable() { +#if TARGET_OS_IPHONE + if(@available(iOS 16.0, *)) + return true; +#else + if(@available(macOS 13.0, *)) + return true; +#endif + return false; + } + +} + +MtSpatialScaler::MtSpatialScaler(MtDevice& device, const SpatialScalerDesc& cfg) { + const auto inputFormat = nativeFormat(cfg.inputFormat); + const auto outputFormat = nativeFormat(cfg.outputFormat); + if(!isMetalFxAvailable() || inputFormat==MTL::PixelFormatInvalid || outputFormat==MTL::PixelFormatInvalid) + return; + if(cfg.inputWidth==0 || cfg.inputHeight==0 || cfg.outputWidth==0 || cfg.outputHeight==0) + return; + + auto pool = NsPtr::init(); + if(!MTLFX::SpatialScalerDescriptor::supportsDevice(device.impl.get())) + return; + + auto raw = MTLFX::SpatialScalerDescriptor::alloc(); + if(raw==nullptr) + return; + auto desc = NsPtr(raw->init()); + if(desc==nullptr) + return; + + desc->setColorTextureFormat(inputFormat); + desc->setOutputTextureFormat(outputFormat); + desc->setInputWidth(cfg.inputWidth); + desc->setInputHeight(cfg.inputHeight); + desc->setOutputWidth(cfg.outputWidth); + desc->setOutputHeight(cfg.outputHeight); + desc->setColorProcessingMode(nativeColorMode(cfg.colorMode)); + + impl = NsPtr(desc->newSpatialScaler(device.impl.get())); + } + +bool MtSpatialScaler::encode(MTL::CommandBuffer& cmd, MtTexture& input, MtTexture& output) { + if(impl==nullptr || input.impl==nullptr || output.impl==nullptr) + return false; + if(output.impl->storageMode()!=MTL::StorageModePrivate) + return false; + if(input.impl->pixelFormat()!=impl->colorTextureFormat() || + output.impl->pixelFormat()!=impl->outputTextureFormat()) + return false; + if(input.impl->width()!=impl->inputWidth() || input.impl->height()!=impl->inputHeight() || + output.impl->width()!=impl->outputWidth() || output.impl->height()!=impl->outputHeight()) + return false; + + const auto inputUsage = impl->colorTextureUsage(); + const auto outputUsage = impl->outputTextureUsage(); + if((input.impl->usage()&inputUsage)!=inputUsage || + (output.impl->usage()&outputUsage)!=outputUsage) + return false; + + impl->setInputContentWidth(input.impl->width()); + impl->setInputContentHeight(input.impl->height()); + impl->setColorTexture(input.impl.get()); + impl->setOutputTexture(output.impl.get()); + impl->encodeToCommandBuffer(&cmd); + return true; + } + +AbstractGraphicsApi::SpatialScaler* + MetalApi::createSpatialScaler(AbstractGraphicsApi::Device* device, const SpatialScalerDesc& desc) { + auto& dev = *reinterpret_cast(device); + auto* scaler = new MtSpatialScaler(dev,desc); + if(!scaler->isValid()) { + delete scaler; + return nullptr; + } + return scaler; + } + +#endif diff --git a/Engine/gapi/metalapi.cpp b/Engine/gapi/metalapi.cpp index aace080e..978603b1 100644 --- a/Engine/gapi/metalapi.cpp +++ b/Engine/gapi/metalapi.cpp @@ -49,6 +49,13 @@ MetalApi::MetalApi(ApiFlags f, const Options& options) MetalApi::~MetalApi() { } +#if !defined(TEMPEST_BUILD_METALFX) +AbstractGraphicsApi::SpatialScaler* + MetalApi::createSpatialScaler(AbstractGraphicsApi::Device*, const SpatialScalerDesc&) { + return nullptr; + } +#endif + std::vector MetalApi::devices() const { #if defined(__OSX__) auto dev = MTL::CopyAllDevices(); diff --git a/Engine/gapi/metalapi.h b/Engine/gapi/metalapi.h index fa5de35a..8393991c 100644 --- a/Engine/gapi/metalapi.h +++ b/Engine/gapi/metalapi.h @@ -33,6 +33,7 @@ class MetalApi : public AbstractGraphicsApi { PTexture createTexture(Device* d, const uint32_t w, const uint32_t h, uint32_t mips, TextureFormat frm) override; PTexture createStorage(Device* d, const uint32_t w, const uint32_t h, uint32_t mips, TextureFormat frm) override; PTexture createStorage(Device* d, const uint32_t w, const uint32_t h, const uint32_t depth, uint32_t mips, TextureFormat frm) override; + SpatialScaler* createSpatialScaler(Device* d, const SpatialScalerDesc& desc) override; AccelerationStructure* createBottomAccelerationStruct(Device* d, const RtGeometry* geom, size_t size) override; AccelerationStructure* createTopAccelerationStruct(Device* d, const RtInstance* inst, AccelerationStructure*const* as, size_t size) override; diff --git a/Tests/tests/gapi/metal_test.cpp b/Tests/tests/gapi/metal_test.cpp index 347f4956..fcae11e4 100644 --- a/Tests/tests/gapi/metal_test.cpp +++ b/Tests/tests/gapi/metal_test.cpp @@ -8,11 +8,26 @@ #include #include +#include + #include "gapi_test_common.h" using namespace testing; using namespace Tempest; +namespace { + +float unpackUnsignedFloat(uint32_t value, uint32_t mantissaBits) { + const uint32_t mantissaMask = (1u << mantissaBits)-1u; + const uint32_t mantissa = value & mantissaMask; + const uint32_t exponent = (value >> mantissaBits) & 0x1Fu; + if(exponent==0) + return std::ldexp(float(mantissa),-14-int(mantissaBits)); + return std::ldexp(1.f+float(mantissa)/float(1u << mantissaBits),int(exponent)-15); + } + +} + TEST(MetalApi,MetalApi) { #if defined(__OSX__) GapiTestCommon::init(); @@ -64,6 +79,64 @@ TEST(MetalApi,ShaderModuleCacheOptions) { #endif } +TEST(MetalApi,SpatialScaler) { +#if defined(__OSX__) + try { + MetalApi api{ApiFlags::Validation}; + Device device(api); + + SpatialScalerDesc desc; + desc.inputFormat = TextureFormat::R11G11B10UF; + desc.outputFormat = TextureFormat::R11G11B10UF; + desc.inputWidth = 32; + desc.inputHeight = 32; + desc.outputWidth = 64; + desc.outputHeight = 64; + desc.colorMode = SpatialScalerColorMode::HDR; + + auto scaler = device.spatialScaler(desc); + if(scaler.isEmpty()) { + Log::d("Skipping MetalFX spatial scaler testcase: unsupported device or system"); + return; + } + + auto input = device.attachment(desc.inputFormat,desc.inputWidth,desc.inputHeight); + auto output = device.image2d(desc.outputFormat,desc.outputWidth,desc.outputHeight); + auto cmd = device.commandBuffer(); + { + auto enc = cmd.startEncoding(device); + enc.setFramebuffer({{input,Vec4(0.25f,0.5f,0.75f,1.f),Tempest::Preserve}}); + EXPECT_TRUE(enc.spatialUpscale(scaler,input,output)); + } + + auto sync = device.submit(cmd); + sync.wait(); + auto result = device.readPixels(output); + EXPECT_EQ(result.w(),desc.outputWidth); + EXPECT_EQ(result.h(),desc.outputHeight); + ASSERT_EQ(result.format(),TextureFormat::R11G11B10UF); + ASSERT_EQ(result.dataSize(),size_t(desc.outputWidth)*desc.outputHeight*sizeof(uint32_t)); + + const auto* pixels = reinterpret_cast(result.data()); + double average[3] = {}; + for(size_t i=0; i> 11 & 0x7FFu,6); + average[2] += unpackUnsignedFloat(pixels[i] >> 22 & 0x3FFu,5); + } + const double pixelCount = double(result.w())*result.h(); + EXPECT_NEAR(average[0]/pixelCount,0.25,0.03); + EXPECT_NEAR(average[1]/pixelCount,0.50,0.03); + EXPECT_NEAR(average[2]/pixelCount,0.75,0.03); + } + catch(std::system_error& e) { + if(e.code()==Tempest::GraphicsErrc::NoDevice) + Log::d("Skipping MetalFX spatial scaler testcase: ", e.what()); else + throw; + } +#endif + } + TEST(MetalApi,Vbo) { #if defined(__OSX__) GapiTestCommon::Vbo(); From 03e756824370771d98836f27ef96ba00678abc41 Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 15:25:08 +0100 Subject: [PATCH 15/25] Preserve parent Apple deployment target --- Engine/CMakeLists.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/Engine/CMakeLists.txt b/Engine/CMakeLists.txt index 835e9446..eb28408c 100644 --- a/Engine/CMakeLists.txt +++ b/Engine/CMakeLists.txt @@ -8,7 +8,7 @@ option(TEMPEST_BUILD_SHARED "Build shared Tempest." ON) set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS ON) set(CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/cmake") -if(NOT IOS) +if(NOT CMAKE_OSX_DEPLOYMENT_TARGET) set(CMAKE_OSX_DEPLOYMENT_TARGET 12.0) endif() @@ -310,7 +310,7 @@ file(GLOB_RECURSE SOURCES ) if(APPLE OR IOS) - if(NOT IOS) + if(NOT CMAKE_OSX_DEPLOYMENT_TARGET) set(CMAKE_OSX_DEPLOYMENT_TARGET 12.0) endif() enable_language(OBJCXX) From 03fe9c1490e4dd9b6237742b88bbd297679c44cd Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 15:25:08 +0100 Subject: [PATCH 16/25] Add optional MetalFX temporal scaler backend --- Engine/gapi/metal/mtcommandbuffer.cpp | 33 ++- Engine/gapi/metal/mtcommandbuffer.h | 12 + Engine/gapi/metal/mtmetalfx.h | 21 ++ Engine/gapi/metal/mtspatialscaler.h | 5 + Engine/gapi/metal/mtspatialscaler.mm | 14 +- Engine/gapi/metal/mttemporalscaler.h | 42 ++++ Engine/gapi/metal/mttemporalscaler.mm | 143 ++++++++++++ Engine/gapi/metalapi.cpp | 8 + Engine/gapi/metalapi.h | 3 +- Tests/tests/gapi/metal_test.cpp | 317 ++++++++++++++++++++++++++ 10 files changed, 590 insertions(+), 8 deletions(-) create mode 100644 Engine/gapi/metal/mtmetalfx.h create mode 100644 Engine/gapi/metal/mttemporalscaler.h create mode 100644 Engine/gapi/metal/mttemporalscaler.mm diff --git a/Engine/gapi/metal/mtcommandbuffer.cpp b/Engine/gapi/metal/mtcommandbuffer.cpp index 75c67069..f91e634c 100644 --- a/Engine/gapi/metal/mtcommandbuffer.cpp +++ b/Engine/gapi/metal/mtcommandbuffer.cpp @@ -13,6 +13,9 @@ #if defined(TEMPEST_BUILD_METALFX) #include "mtspatialscaler.h" #endif +#if defined(TEMPEST_BUILD_METALFX_TEMPORAL) +#include "mttemporalscaler.h" +#endif using namespace Tempest; using namespace Tempest::Detail; @@ -804,10 +807,32 @@ bool MtCommandBuffer::spatialUpscale(AbstractGraphicsApi::SpatialScaler& scaler, AbstractGraphicsApi::Texture& input, AbstractGraphicsApi::Texture& output) { setEncoder(E_None,nullptr); - auto& sx = reinterpret_cast(scaler); - auto& src = reinterpret_cast(input); - auto& dst = reinterpret_cast(output); - return sx.encode(*impl,src,dst); + auto* sx = dynamic_cast(&scaler); + auto* src = dynamic_cast(&input); + auto* dst = dynamic_cast(&output); + if(sx==nullptr || src==nullptr || dst==nullptr || !sx->belongsTo(device)) + return false; + return sx->encode(*impl,*src,*dst); + } +#endif + +#if defined(TEMPEST_BUILD_METALFX_TEMPORAL) +bool MtCommandBuffer::temporalUpscale(AbstractGraphicsApi::TemporalScaler& scaler, + AbstractGraphicsApi::Texture& input, + AbstractGraphicsApi::Texture& depth, + AbstractGraphicsApi::Texture& motion, + AbstractGraphicsApi::Texture& output, + const TemporalScalerArgs& args) { + setEncoder(E_None,nullptr); + auto* sx = dynamic_cast(&scaler); + auto* src = dynamic_cast(&input); + auto* dep = dynamic_cast(&depth); + auto* mov = dynamic_cast(&motion); + auto* dst = dynamic_cast(&output); + if(sx==nullptr || src==nullptr || dep==nullptr || mov==nullptr || dst==nullptr || + !sx->belongsTo(device)) + return false; + return sx->encode(*impl,*src,*dep,*mov,*dst,args); } #endif diff --git a/Engine/gapi/metal/mtcommandbuffer.h b/Engine/gapi/metal/mtcommandbuffer.h index 8f49f9c4..32f35689 100644 --- a/Engine/gapi/metal/mtcommandbuffer.h +++ b/Engine/gapi/metal/mtcommandbuffer.h @@ -7,6 +7,7 @@ #include "mtfbolayout.h" #include "mtpipelinelay.h" #include "nsptr.h" +#include "mtmetalfx.h" namespace Tempest { @@ -23,6 +24,9 @@ class MtTopAccelerationStructure; #if defined(TEMPEST_BUILD_METALFX) class MtSpatialScaler; #endif +#if defined(TEMPEST_BUILD_METALFX_TEMPORAL) +class MtTemporalScaler; +#endif class MtCommandBuffer : public AbstractGraphicsApi::CommandBuffer { public: @@ -74,6 +78,14 @@ class MtCommandBuffer : public AbstractGraphicsApi::CommandBuffer { AbstractGraphicsApi::Texture& input, AbstractGraphicsApi::Texture& output) override; #endif +#if defined(TEMPEST_BUILD_METALFX_TEMPORAL) + bool temporalUpscale(AbstractGraphicsApi::TemporalScaler& scaler, + AbstractGraphicsApi::Texture& input, + AbstractGraphicsApi::Texture& depth, + AbstractGraphicsApi::Texture& motion, + AbstractGraphicsApi::Texture& output, + const TemporalScalerArgs& args) override; +#endif private: enum EncType:uint8_t { diff --git a/Engine/gapi/metal/mtmetalfx.h b/Engine/gapi/metal/mtmetalfx.h new file mode 100644 index 00000000..61cf82b0 --- /dev/null +++ b/Engine/gapi/metal/mtmetalfx.h @@ -0,0 +1,21 @@ +#pragma once + +#if defined(__APPLE__) +#include +#include + +#if __has_include() +#if TARGET_OS_IPHONE +#if defined(__IPHONE_OS_VERSION_MAX_ALLOWED) && __IPHONE_OS_VERSION_MAX_ALLOWED >= 160000 +#define TEMPEST_METALFX_TEMPORAL_SDK_AVAILABLE +#endif +#elif defined(__MAC_OS_X_VERSION_MAX_ALLOWED) && __MAC_OS_X_VERSION_MAX_ALLOWED >= 130000 +#define TEMPEST_METALFX_TEMPORAL_SDK_AVAILABLE +#endif +#endif + +#endif + +#if defined(TEMPEST_BUILD_METALFX) && defined(TEMPEST_METALFX_TEMPORAL_SDK_AVAILABLE) +#define TEMPEST_BUILD_METALFX_TEMPORAL +#endif diff --git a/Engine/gapi/metal/mtspatialscaler.h b/Engine/gapi/metal/mtspatialscaler.h index 48727550..638422df 100644 --- a/Engine/gapi/metal/mtspatialscaler.h +++ b/Engine/gapi/metal/mtspatialscaler.h @@ -5,6 +5,8 @@ #include #include +#include + #include "nsptr.h" namespace Tempest { @@ -18,9 +20,12 @@ class MtSpatialScaler final : public AbstractGraphicsApi::SpatialScaler { MtSpatialScaler(MtDevice& device, const SpatialScalerDesc& desc); bool isValid() const { return impl!=nullptr; } + bool belongsTo(const MtDevice& device) const { return owner==&device; } bool encode(MTL::CommandBuffer& cmd, MtTexture& input, MtTexture& output); private: + MtDevice* owner = nullptr; + std::mutex sync; NsPtr impl; }; diff --git a/Engine/gapi/metal/mtspatialscaler.mm b/Engine/gapi/metal/mtspatialscaler.mm index b9a12024..80bd72d0 100644 --- a/Engine/gapi/metal/mtspatialscaler.mm +++ b/Engine/gapi/metal/mtspatialscaler.mm @@ -5,6 +5,8 @@ #include #include +#include + #include "mtdevice.h" #include "mttexture.h" @@ -38,7 +40,8 @@ bool isMetalFxAvailable() { } -MtSpatialScaler::MtSpatialScaler(MtDevice& device, const SpatialScalerDesc& cfg) { +MtSpatialScaler::MtSpatialScaler(MtDevice& device, const SpatialScalerDesc& cfg) + :owner(&device) { const auto inputFormat = nativeFormat(cfg.inputFormat); const auto outputFormat = nativeFormat(cfg.outputFormat); if(!isMetalFxAvailable() || inputFormat==MTL::PixelFormatInvalid || outputFormat==MTL::PixelFormatInvalid) @@ -69,8 +72,11 @@ bool isMetalFxAvailable() { } bool MtSpatialScaler::encode(MTL::CommandBuffer& cmd, MtTexture& input, MtTexture& output) { + std::lock_guard guard(sync); if(impl==nullptr || input.impl==nullptr || output.impl==nullptr) return false; + if(cmd.device()!=owner->impl.get() || &input.dev!=owner || &output.dev!=owner) + return false; if(output.impl->storageMode()!=MTL::StorageModePrivate) return false; if(input.impl->pixelFormat()!=impl->colorTextureFormat() || @@ -96,8 +102,10 @@ bool isMetalFxAvailable() { AbstractGraphicsApi::SpatialScaler* MetalApi::createSpatialScaler(AbstractGraphicsApi::Device* device, const SpatialScalerDesc& desc) { - auto& dev = *reinterpret_cast(device); - auto* scaler = new MtSpatialScaler(dev,desc); + auto* dev = dynamic_cast(device); + if(dev==nullptr) + return nullptr; + auto* scaler = new MtSpatialScaler(*dev,desc); if(!scaler->isValid()) { delete scaler; return nullptr; diff --git a/Engine/gapi/metal/mttemporalscaler.h b/Engine/gapi/metal/mttemporalscaler.h new file mode 100644 index 00000000..64d42467 --- /dev/null +++ b/Engine/gapi/metal/mttemporalscaler.h @@ -0,0 +1,42 @@ +#pragma once + +#include "mtmetalfx.h" + +#if defined(TEMPEST_BUILD_METALFX_TEMPORAL) + +#include +#include + +#include + +#include "nsptr.h" + +namespace Tempest { +namespace Detail { + +class MtDevice; +class MtTexture; + +class MtTemporalScaler final : public AbstractGraphicsApi::TemporalScaler { + public: + MtTemporalScaler(MtDevice& device, const TemporalScalerDesc& desc); + + bool isValid() const { return impl!=nullptr; } + bool belongsTo(const MtDevice& device) const { return owner==&device; } + bool encode(MTL::CommandBuffer& cmd, MtTexture& input, MtTexture& depth, + MtTexture& motion, MtTexture& output, const TemporalScalerArgs& args); + + private: + // A temporal scaler owns one history. The mutex makes validate/set/encode + // atomic with respect to other CPU threads. Callers must still record and + // submit history-dependent frames in one explicit order; resetHistory starts + // a new history and must be submitted before the frames that consume it. + MtDevice* owner = nullptr; + std::mutex sync; + NsPtr impl; + }; + +} +} + +#endif diff --git a/Engine/gapi/metal/mttemporalscaler.mm b/Engine/gapi/metal/mttemporalscaler.mm new file mode 100644 index 00000000..216171f5 --- /dev/null +++ b/Engine/gapi/metal/mttemporalscaler.mm @@ -0,0 +1,143 @@ +#include "mttemporalscaler.h" + +#if defined(TEMPEST_BUILD_METALFX_TEMPORAL) + +#include +#include + +#include +#include + +#include "mtdevice.h" +#include "mttexture.h" + +using namespace Tempest; +using namespace Tempest::Detail; + +namespace { + +bool isMetalFxTemporalAvailable() { +#if TARGET_OS_IPHONE + if(@available(iOS 16.0, *)) + return true; +#else + if(@available(macOS 13.0, *)) + return true; +#endif + return false; + } + +bool hasUsage(const MTL::Texture& texture, MTL::TextureUsage required) { + return (texture.usage()&required)==required; + } + +bool hasFormatAndSize(const MTL::Texture& texture, MTL::PixelFormat format, + NS::UInteger width, NS::UInteger height) { + return texture.textureType()==MTL::TextureType2D && texture.sampleCount()==1 && + texture.pixelFormat()==format && texture.width()==width && texture.height()==height; + } + +} + +MtTemporalScaler::MtTemporalScaler(MtDevice& device, const TemporalScalerDesc& cfg) + :owner(&device) { + const auto inputFormat = nativeFormat(cfg.inputFormat); + const auto depthFormat = nativeFormat(cfg.depthFormat); + const auto motionFormat = nativeFormat(cfg.motionFormat); + const auto outputFormat = nativeFormat(cfg.outputFormat); + if(!isMetalFxTemporalAvailable() || inputFormat==MTL::PixelFormatInvalid || + depthFormat==MTL::PixelFormatInvalid || motionFormat==MTL::PixelFormatInvalid || + outputFormat==MTL::PixelFormatInvalid) + return; + if(cfg.inputWidth==0 || cfg.inputHeight==0 || cfg.outputWidth==0 || cfg.outputHeight==0) + return; + if(cfg.inputWidth>device.prop.tex2d.maxSize || cfg.inputHeight>device.prop.tex2d.maxSize || + cfg.outputWidth>device.prop.tex2d.maxSize || cfg.outputHeight>device.prop.tex2d.maxSize) + return; + + auto pool = NsPtr::init(); + if(!MTLFX::TemporalScalerDescriptor::supportsDevice(device.impl.get())) + return; + + auto raw = MTLFX::TemporalScalerDescriptor::alloc(); + if(raw==nullptr) + return; + auto desc = NsPtr(raw->init()); + if(desc==nullptr) + return; + + desc->setColorTextureFormat(inputFormat); + desc->setDepthTextureFormat(depthFormat); + desc->setMotionTextureFormat(motionFormat); + desc->setOutputTextureFormat(outputFormat); + desc->setInputWidth(cfg.inputWidth); + desc->setInputHeight(cfg.inputHeight); + desc->setOutputWidth(cfg.outputWidth); + desc->setOutputHeight(cfg.outputHeight); + desc->setAutoExposureEnabled(cfg.autoExposure); + + auto scaler = NsPtr(desc->newTemporalScaler(device.impl.get())); + if(scaler==nullptr) + return; + if(scaler->colorTextureFormat()!=inputFormat || scaler->depthTextureFormat()!=depthFormat || + scaler->motionTextureFormat()!=motionFormat || scaler->outputTextureFormat()!=outputFormat || + scaler->inputWidth()!=cfg.inputWidth || scaler->inputHeight()!=cfg.inputHeight || + scaler->outputWidth()!=cfg.outputWidth || scaler->outputHeight()!=cfg.outputHeight) + return; + impl = std::move(scaler); + } + +bool MtTemporalScaler::encode(MTL::CommandBuffer& cmd, MtTexture& input, MtTexture& depth, + MtTexture& motion, MtTexture& output, const TemporalScalerArgs& args) { + std::lock_guard guard(sync); + if(impl==nullptr || input.impl==nullptr || depth.impl==nullptr || + motion.impl==nullptr || output.impl==nullptr) + return false; + if(cmd.device()!=owner->impl.get() || &input.dev!=owner || &depth.dev!=owner || + &motion.dev!=owner || &output.dev!=owner) + return false; + + if(!hasFormatAndSize(*input.impl,impl->colorTextureFormat(),impl->inputWidth(),impl->inputHeight()) || + !hasFormatAndSize(*depth.impl,impl->depthTextureFormat(),impl->inputWidth(),impl->inputHeight()) || + !hasFormatAndSize(*motion.impl,impl->motionTextureFormat(),impl->inputWidth(),impl->inputHeight()) || + !hasFormatAndSize(*output.impl,impl->outputTextureFormat(),impl->outputWidth(),impl->outputHeight())) + return false; + if(output.impl->storageMode()!=MTL::StorageModePrivate) + return false; + + if(!hasUsage(*input.impl,impl->colorTextureUsage()) || + !hasUsage(*depth.impl,impl->depthTextureUsage()) || + !hasUsage(*motion.impl,impl->motionTextureUsage()) || + !hasUsage(*output.impl,impl->outputTextureUsage())) + return false; + + impl->setInputContentWidth(input.impl->width()); + impl->setInputContentHeight(input.impl->height()); + impl->setColorTexture(input.impl.get()); + impl->setDepthTexture(depth.impl.get()); + impl->setMotionTexture(motion.impl.get()); + impl->setOutputTexture(output.impl.get()); + impl->setJitterOffsetX(args.jitterOffsetX); + impl->setJitterOffsetY(args.jitterOffsetY); + impl->setMotionVectorScaleX(args.motionVectorScaleX); + impl->setMotionVectorScaleY(args.motionVectorScaleY); + impl->setReset(args.resetHistory); + impl->setDepthReversed(args.depthReversed); + impl->encodeToCommandBuffer(&cmd); + return true; + } + +AbstractGraphicsApi::TemporalScaler* + MetalApi::createTemporalScaler(AbstractGraphicsApi::Device* device, const TemporalScalerDesc& desc) { + auto* dev = dynamic_cast(device); + if(dev==nullptr) + return nullptr; + auto* scaler = new MtTemporalScaler(*dev,desc); + if(!scaler->isValid()) { + delete scaler; + return nullptr; + } + return scaler; + } + +#endif diff --git a/Engine/gapi/metalapi.cpp b/Engine/gapi/metalapi.cpp index 978603b1..2307776a 100644 --- a/Engine/gapi/metalapi.cpp +++ b/Engine/gapi/metalapi.cpp @@ -20,6 +20,7 @@ #include "gapi/metal/mtsync.h" #include "gapi/metal/mtswapchain.h" #include "gapi/metal/mtaccelerationstructure.h" +#include "gapi/metal/mtmetalfx.h" #include @@ -56,6 +57,13 @@ AbstractGraphicsApi::SpatialScaler* } #endif +#if !defined(TEMPEST_BUILD_METALFX_TEMPORAL) +AbstractGraphicsApi::TemporalScaler* + MetalApi::createTemporalScaler(AbstractGraphicsApi::Device*, const TemporalScalerDesc&) { + return nullptr; + } +#endif + std::vector MetalApi::devices() const { #if defined(__OSX__) auto dev = MTL::CopyAllDevices(); diff --git a/Engine/gapi/metalapi.h b/Engine/gapi/metalapi.h index 8393991c..f8ec892a 100644 --- a/Engine/gapi/metalapi.h +++ b/Engine/gapi/metalapi.h @@ -33,7 +33,8 @@ class MetalApi : public AbstractGraphicsApi { PTexture createTexture(Device* d, const uint32_t w, const uint32_t h, uint32_t mips, TextureFormat frm) override; PTexture createStorage(Device* d, const uint32_t w, const uint32_t h, uint32_t mips, TextureFormat frm) override; PTexture createStorage(Device* d, const uint32_t w, const uint32_t h, const uint32_t depth, uint32_t mips, TextureFormat frm) override; - SpatialScaler* createSpatialScaler(Device* d, const SpatialScalerDesc& desc) override; + SpatialScaler* createSpatialScaler(Device* d, const SpatialScalerDesc& desc) override; + TemporalScaler* createTemporalScaler(Device* d, const TemporalScalerDesc& desc) override; AccelerationStructure* createBottomAccelerationStruct(Device* d, const RtGeometry* geom, size_t size) override; AccelerationStructure* createTopAccelerationStruct(Device* d, const RtInstance* inst, AccelerationStructure*const* as, size_t size) override; diff --git a/Tests/tests/gapi/metal_test.cpp b/Tests/tests/gapi/metal_test.cpp index fcae11e4..5fdc85ed 100644 --- a/Tests/tests/gapi/metal_test.cpp +++ b/Tests/tests/gapi/metal_test.cpp @@ -8,7 +8,18 @@ #include #include +#include #include +#include +#include +#include + +#if defined(__OSX__) +#include "../../../Engine/gapi/metal/mtmetalfx.h" +#if defined(TEMPEST_METALFX_TEMPORAL_SDK_AVAILABLE) +#include +#endif +#endif #include "gapi_test_common.h" @@ -26,6 +37,65 @@ float unpackUnsignedFloat(uint32_t value, uint32_t mantissaBits) { return std::ldexp(1.f+float(mantissa)/float(1u << mantissaBits),int(exponent)-15); } +float unpackHalf(uint16_t value) { + const uint32_t sign = value >> 15; + const uint32_t exponent = (value >> 10) & 0x1Fu; + const uint32_t mantissa = value & 0x3FFu; + float ret = 0.f; + if(exponent==0) + ret = std::ldexp(float(mantissa),-24); else + if(exponent==0x1Fu) + ret = mantissa==0 ? std::numeric_limits::infinity() : + std::numeric_limits::quiet_NaN(); else + ret = std::ldexp(float(0x400u+mantissa),int(exponent)-25); + return sign!=0 ? -ret : ret; + } + +std::array averageRgba16F(const Pixmap& image) { + std::array average = {}; + const auto* pixels = reinterpret_cast(image.data()); + const size_t count = size_t(image.w())*image.h(); + for(size_t i=0; i(dlsym(RTLD_DEFAULT,"MTLCreateSystemDefaultDevice")); + auto lookupClass = reinterpret_cast (dlsym(RTLD_DEFAULT,"objc_lookUpClass")); + auto registerSel = reinterpret_cast (dlsym(RTLD_DEFAULT,"sel_registerName")); + auto sendBool = reinterpret_cast (dlsym(RTLD_DEFAULT,"objc_msgSend")); + auto sendVoid = reinterpret_cast (dlsym(RTLD_DEFAULT,"objc_msgSend")); + if(createDevice==nullptr || lookupClass==nullptr || registerSel==nullptr || + sendBool==nullptr || sendVoid==nullptr) + return false; + auto scalerClass = lookupClass("MTLFXTemporalScalerDescriptor"); + if(scalerClass==nullptr) + return false; + + auto device = createDevice(); + if(device==nullptr) + return false; + + const bool supported = sendBool(scalerClass,registerSel("supportsDevice:"),device); + sendVoid(device,registerSel("release")); + return supported; +#else + return false; +#endif + } +#endif + } TEST(MetalApi,MetalApi) { @@ -102,6 +172,24 @@ TEST(MetalApi,SpatialScaler) { auto input = device.attachment(desc.inputFormat,desc.inputWidth,desc.inputHeight); auto output = device.image2d(desc.outputFormat,desc.outputWidth,desc.outputHeight); + + Device otherDevice(api); + auto otherScaler = otherDevice.spatialScaler(desc); + auto otherInput = otherDevice.attachment(desc.inputFormat,desc.inputWidth,desc.inputHeight); + auto otherOutput = otherDevice.image2d(desc.outputFormat,desc.outputWidth,desc.outputHeight); + ASSERT_FALSE(otherScaler.isEmpty()); + { + auto otherCmd = otherDevice.commandBuffer(); + auto enc = otherCmd.startEncoding(otherDevice); + EXPECT_FALSE(enc.spatialUpscale(scaler,otherInput,otherOutput)); + } + { + auto localCmd = device.commandBuffer(); + auto enc = localCmd.startEncoding(device); + EXPECT_FALSE(enc.spatialUpscale(scaler,otherInput,otherOutput)); + EXPECT_FALSE(enc.spatialUpscale(otherScaler,input,output)); + } + auto cmd = device.commandBuffer(); { auto enc = cmd.startEncoding(device); @@ -137,6 +225,235 @@ TEST(MetalApi,SpatialScaler) { #endif } +TEST(MetalApi,TemporalScalerDescriptorValidation) { +#if defined(__OSX__) + try { + MetalApi api; + Device device(api); + + TemporalScalerDesc desc; + desc.inputFormat = TextureFormat::RGBA16F; + desc.depthFormat = TextureFormat::Depth32F; + desc.motionFormat = TextureFormat::RG32F; + desc.outputFormat = TextureFormat::RGBA16F; + desc.inputWidth = 64; + desc.inputHeight = 64; + desc.outputWidth = 128; + desc.outputHeight = 128; + + auto invalid = desc; + invalid.inputFormat = TextureFormat::Undefined; + EXPECT_TRUE(device.temporalScaler(invalid).isEmpty()); + invalid = desc; + invalid.depthFormat = TextureFormat::Undefined; + EXPECT_TRUE(device.temporalScaler(invalid).isEmpty()); + invalid = desc; + invalid.motionFormat = TextureFormat::Undefined; + EXPECT_TRUE(device.temporalScaler(invalid).isEmpty()); + invalid = desc; + invalid.outputFormat = TextureFormat::Undefined; + EXPECT_TRUE(device.temporalScaler(invalid).isEmpty()); + + invalid = desc; + invalid.inputWidth = 0; + EXPECT_TRUE(device.temporalScaler(invalid).isEmpty()); + invalid = desc; + invalid.inputHeight = 0; + EXPECT_TRUE(device.temporalScaler(invalid).isEmpty()); + invalid = desc; + invalid.outputWidth = 0; + EXPECT_TRUE(device.temporalScaler(invalid).isEmpty()); + invalid = desc; + invalid.outputHeight = 0; + EXPECT_TRUE(device.temporalScaler(invalid).isEmpty()); + + invalid = desc; + invalid.inputWidth = device.properties().tex2d.maxSize+1; + EXPECT_TRUE(device.temporalScaler(invalid).isEmpty()); + invalid = desc; + invalid.outputHeight = device.properties().tex2d.maxSize+1; + EXPECT_TRUE(device.temporalScaler(invalid).isEmpty()); + + if(metalFxTemporalSupported()) { + auto manualExposure = desc; + manualExposure.autoExposure = false; + EXPECT_FALSE(device.temporalScaler(manualExposure).isEmpty()); + } + } + catch(std::system_error& e) { + if(e.code()==Tempest::GraphicsErrc::NoDevice) + Log::d("Skipping MetalFX temporal descriptor testcase: ", e.what()); else + throw; + } +#endif + } + +TEST(MetalApi,TemporalScaler) { +#if defined(__OSX__) + try { + MetalApi api{ApiFlags::Validation}; + Device device(api); + + TemporalScalerDesc desc; + desc.inputFormat = TextureFormat::RGBA16F; + desc.depthFormat = TextureFormat::Depth32F; + desc.motionFormat = TextureFormat::RG32F; + desc.outputFormat = TextureFormat::RGBA16F; + desc.inputWidth = 64; + desc.inputHeight = 64; + desc.outputWidth = 128; + desc.outputHeight = 128; + desc.autoExposure = true; + + if(!metalFxTemporalSupported()) { + Log::d("Skipping MetalFX temporal scaler testcase: unsupported device or system"); + return; + } + auto scaler = device.temporalScaler(desc); + ASSERT_FALSE(scaler.isEmpty()); + + auto input = device.attachment(desc.inputFormat,desc.inputWidth,desc.inputHeight); + auto depth = device.zbuffer(desc.depthFormat,desc.inputWidth,desc.inputHeight); + auto motion = device.attachment(desc.motionFormat,desc.inputWidth,desc.inputHeight); + auto output = device.image2d(desc.outputFormat,desc.outputWidth,desc.outputHeight); + auto wrongInput = device.attachment(desc.inputFormat,desc.inputWidth/2,desc.inputHeight); + auto wrongDepth = device.zbuffer(desc.depthFormat,desc.inputWidth,desc.inputHeight/2); + auto wrongMotion = device.attachment(desc.motionFormat,desc.inputWidth/2,desc.inputHeight); + auto wrongOutput = device.image2d(desc.outputFormat,desc.outputWidth/2,desc.outputHeight); + auto wrongInputFormat = device.attachment(TextureFormat::R11G11B10UF,desc.inputWidth,desc.inputHeight); + auto wrongDepthFormat = device.zbuffer(TextureFormat::Depth16,desc.inputWidth,desc.inputHeight); + auto wrongMotionFormat = device.attachment(TextureFormat::RGBA16F,desc.inputWidth,desc.inputHeight); + auto wrongOutputFormat = device.image2d(TextureFormat::R11G11B10UF,desc.outputWidth,desc.outputHeight); + + Device otherDevice(api); + auto otherScaler = otherDevice.temporalScaler(desc); + auto otherInput = otherDevice.attachment(desc.inputFormat,desc.inputWidth,desc.inputHeight); + auto otherDepth = otherDevice.zbuffer(desc.depthFormat,desc.inputWidth,desc.inputHeight); + auto otherMotion = otherDevice.attachment(desc.motionFormat,desc.inputWidth,desc.inputHeight); + auto otherOutput = otherDevice.image2d(desc.outputFormat,desc.outputWidth,desc.outputHeight); + ASSERT_FALSE(otherScaler.isEmpty()); + { + auto otherCmd = otherDevice.commandBuffer(); + auto enc = otherCmd.startEncoding(otherDevice); + EXPECT_FALSE(enc.temporalUpscale(scaler,otherInput,otherDepth,otherMotion,otherOutput,{})); + } + + auto cmd = device.commandBuffer(); + { + auto enc = cmd.startEncoding(device); + EXPECT_FALSE(enc.temporalUpscale(scaler,wrongInput,depth,motion,output,{})); + EXPECT_FALSE(enc.temporalUpscale(scaler,input,wrongDepth,motion,output,{})); + EXPECT_FALSE(enc.temporalUpscale(scaler,input,depth,wrongMotion,output,{})); + EXPECT_FALSE(enc.temporalUpscale(scaler,input,depth,motion,wrongOutput,{})); + EXPECT_FALSE(enc.temporalUpscale(scaler,wrongInputFormat,depth,motion,output,{})); + EXPECT_FALSE(enc.temporalUpscale(scaler,input,wrongDepthFormat,motion,output,{})); + EXPECT_FALSE(enc.temporalUpscale(scaler,input,depth,wrongMotionFormat,output,{})); + EXPECT_FALSE(enc.temporalUpscale(scaler,input,depth,motion,wrongOutputFormat,{})); + EXPECT_FALSE(enc.temporalUpscale(scaler,otherInput,otherDepth,otherMotion,otherOutput,{})); + EXPECT_FALSE(enc.temporalUpscale(otherScaler,input,depth,motion,output,{})); + + enc.setFramebuffer({{input,Vec4(0.25f,0.5f,0.75f,1.f),Tempest::Preserve}, + {motion,Vec4(1.f/float(desc.inputWidth), + -1.f/float(desc.inputHeight),0.f,0.f),Tempest::Preserve}}, + {depth,0.75f,Tempest::Preserve}); + + TemporalScalerArgs args; + args.jitterOffsetX = 0.25f; + args.jitterOffsetY = -0.25f; + args.motionVectorScaleX = float(desc.inputWidth); + args.motionVectorScaleY = float(desc.inputHeight); + args.resetHistory = true; + args.depthReversed = false; + EXPECT_TRUE(enc.temporalUpscale(scaler,input,depth,motion,output,args)); + } + + auto sync = device.submit(cmd); + sync.wait(); + auto result = device.readPixels(output); + EXPECT_EQ(result.w(),desc.outputWidth); + EXPECT_EQ(result.h(),desc.outputHeight); + ASSERT_EQ(result.format(),desc.outputFormat); + ASSERT_EQ(result.dataSize(),size_t(desc.outputWidth)*desc.outputHeight*4*sizeof(uint16_t)); + const auto firstAverage = averageRgba16F(result); + for(auto component:firstAverage) + EXPECT_TRUE(std::isfinite(component)); + EXPECT_NEAR(firstAverage[0],0.25,0.12); + EXPECT_NEAR(firstAverage[1],0.50,0.12); + EXPECT_NEAR(firstAverage[2],0.75,0.12); + EXPECT_NEAR(firstAverage[3],1.00,0.08); + + TemporalScalerArgs secondArgs; + secondArgs.jitterOffsetX = -0.125f; + secondArgs.jitterOffsetY = 0.375f; + secondArgs.motionVectorScaleX = float(desc.inputWidth); + secondArgs.motionVectorScaleY = float(desc.inputHeight); + secondArgs.resetHistory = false; + secondArgs.depthReversed = true; + auto secondCmd = device.commandBuffer(); + { + auto enc = secondCmd.startEncoding(device); + enc.setFramebuffer({{input,Vec4(0.75f,0.25f,0.5f,1.f),Tempest::Preserve}, + {motion,Vec4(-1.f/float(desc.inputWidth), + 1.f/float(desc.inputHeight),0.f,0.f),Tempest::Preserve}}, + {depth,0.6f,Tempest::Preserve}); + EXPECT_TRUE(enc.temporalUpscale(scaler,input,depth,motion,output,secondArgs)); + } + auto secondSync = device.submit(secondCmd); + secondSync.wait(); + auto secondResult = device.readPixels(output); + ASSERT_EQ(secondResult.format(),desc.outputFormat); + const auto secondAverage = averageRgba16F(secondResult); + for(auto component:secondAverage) + EXPECT_TRUE(std::isfinite(component)); + EXPECT_NEAR(secondAverage[0],0.75,0.15); + EXPECT_NEAR(secondAverage[1],0.25,0.15); + EXPECT_NEAR(secondAverage[2],0.50,0.15); + EXPECT_NEAR(secondAverage[3],1.00,0.08); + + std::array parallelCmd = { + device.commandBuffer(),device.commandBuffer() + }; + std::array parallelOutput = { + device.image2d(desc.outputFormat,desc.outputWidth,desc.outputHeight), + device.image2d(desc.outputFormat,desc.outputWidth,desc.outputHeight) + }; + std::array encoded = {}; + std::array workers; + for(size_t i=0; i(); From 83675dc55123ed7d982573d43d8108ce4405d600 Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 15:39:47 +0100 Subject: [PATCH 17/25] Fix CTest registration for TempestTests --- Tests/tests/CMakeLists.txt | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/Tests/tests/CMakeLists.txt b/Tests/tests/CMakeLists.txt index 6d892e9b..c1a97690 100644 --- a/Tests/tests/CMakeLists.txt +++ b/Tests/tests/CMakeLists.txt @@ -9,7 +9,9 @@ set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/testsuite) # Setup testing enable_testing() add_executable(${PROJECT_NAME}) -add_test(${PROJECT_NAME} COMMAND ${PROJECT_NAME}) +add_test(NAME ${PROJECT_NAME} COMMAND ${PROJECT_NAME}) +set_tests_properties(${PROJECT_NAME} PROPERTIES + WORKING_DIRECTORY "${CMAKE_RUNTIME_OUTPUT_DIRECTORY}") target_include_directories(${PROJECT_NAME} PRIVATE .) target_include_directories(${PROJECT_NAME} PRIVATE "${CMAKE_SOURCE_DIR}/../../Engine/include") From 0f958d5ee2fbe8985d4c53b93a6777a452c47cfb Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 15:46:18 +0100 Subject: [PATCH 18/25] Add optional direct Metal drawable rendering --- Engine/gapi/metal/mtcommandbuffer.cpp | 9 +- Engine/gapi/metal/mtcommandbuffer.h | 4 + Engine/gapi/metal/mtswapchain.h | 43 +- Engine/gapi/metal/mtswapchain.mm | 410 +++++++++++++++--- Engine/gapi/metal/mtswapchainstate.h | 267 ++++++++++++ Engine/gapi/metalapi.cpp | 12 +- Engine/gapi/metalapi.h | 17 +- .../tests/gapi/metal_swapchain_state_test.cpp | 195 +++++++++ Tests/tests/gapi/metal_test.cpp | 78 ++++ 9 files changed, 953 insertions(+), 82 deletions(-) create mode 100644 Engine/gapi/metal/mtswapchainstate.h create mode 100644 Tests/tests/gapi/metal_swapchain_state_test.cpp diff --git a/Engine/gapi/metal/mtcommandbuffer.cpp b/Engine/gapi/metal/mtcommandbuffer.cpp index f91e634c..b606ca20 100644 --- a/Engine/gapi/metal/mtcommandbuffer.cpp +++ b/Engine/gapi/metal/mtcommandbuffer.cpp @@ -71,6 +71,7 @@ void MtCommandBuffer::end() { } void MtCommandBuffer::reset() { + swapchainFrames.clear(); auto pool = NsPtr::init(); auto desc = NsPtr::init(); desc->setRetainedReferences(false); @@ -100,7 +101,13 @@ void MtCommandBuffer::beginRendering(const FrameBufferDesc& fbo, size_t fboSize, auto clr = desc->colorAttachments()->object(i); if(fbo.sw[i]!=nullptr) { auto& s = *reinterpret_cast(fbo.sw[i]); - clr->setTexture(s.img[fbo.imgId[i]].tex.get()); + auto frame = s.acquireRenderTarget(fbo.imgId[i]); + clr->setTexture(frame->texture.get()); + bool known = false; + for(auto& existing:swapchainFrames) + known = known || existing.get()==frame.get(); + if(!known) + swapchainFrames.push_back(std::move(frame)); curFbo.colorFormat[curFbo.numColors] = s.format(); } else { auto& t = *reinterpret_cast(fbo.att[i]); diff --git a/Engine/gapi/metal/mtcommandbuffer.h b/Engine/gapi/metal/mtcommandbuffer.h index 32f35689..ee5b781b 100644 --- a/Engine/gapi/metal/mtcommandbuffer.h +++ b/Engine/gapi/metal/mtcommandbuffer.h @@ -3,6 +3,8 @@ #include #include +#include + #include "gapi/shaderreflection.h" #include "mtfbolayout.h" #include "mtpipelinelay.h" @@ -27,6 +29,7 @@ class MtSpatialScaler; #if defined(TEMPEST_BUILD_METALFX_TEMPORAL) class MtTemporalScaler; #endif +class MtSwapchainFrame; class MtCommandBuffer : public AbstractGraphicsApi::CommandBuffer { public: @@ -133,6 +136,7 @@ class MtCommandBuffer : public AbstractGraphicsApi::CommandBuffer { NsPtr encBlit; std::vector usedResources; + std::vector> swapchainFrames; MtFboLayout curFbo; Push pushData; diff --git a/Engine/gapi/metal/mtswapchain.h b/Engine/gapi/metal/mtswapchain.h index eac96f9c..e998b3d7 100644 --- a/Engine/gapi/metal/mtswapchain.h +++ b/Engine/gapi/metal/mtswapchain.h @@ -2,10 +2,14 @@ #include #include "utility/spinlock.h" +#include "mtswapchainstate.h" #include "nsptr.h" #include +#include +#include + namespace CA { class MetalDrawable; @@ -16,9 +20,27 @@ namespace Detail { class MtDevice; +struct MtSwapchainFrame final { + MtSwapchainFrame(uint64_t generation, uint32_t image, bool direct, + MTL::Texture* texture, CA::MetalDrawable* drawable); + ~MtSwapchainFrame(); + + MtSwapchainFrame(const MtSwapchainFrame&) = delete; + MtSwapchainFrame& operator=(const MtSwapchainFrame&) = delete; + + uint64_t generation = 0; + uint32_t image = 0; + bool direct = false; + NsPtr texture; + NsPtr drawable; + }; + class MtSwapchain : public AbstractGraphicsApi::Swapchain { public: - MtSwapchain(MtDevice& dev, SystemApi::Window* w, uint32_t bufferCount); + using Frame = std::shared_ptr; + + MtSwapchain(MtDevice& dev, SystemApi::Window* w, uint32_t bufferCount, + bool directPreferred); ~MtSwapchain(); void reset() override; @@ -29,26 +51,29 @@ class MtSwapchain : public AbstractGraphicsApi::Swapchain { void present(); NonUniqResId syncId() const override { return NonUniqResId::I_None; } + Frame acquireRenderTarget(uint32_t image); MTL::PixelFormat format() const; + private: struct Image { NsPtr tex; }; - std::vector img; - - private: struct Impl; std::unique_ptr pimpl; - SpinLock sync; + mutable SpinLock sync; MtDevice& dev; Tempest::Size sz; - uint32_t imgCount = 0; - uint32_t currentImg = 0; + std::vector img; + MtSwapchainState state; + MtSwapchainOperationGate operationGate; + Frame activeFrame; + bool directPreferred = false; - NsPtr mkTexture(); - void nextDrawable(); + NsPtr mkTexture(const Tempest::Size& size); + Frame acquireRenderTargetImpl(uint32_t image); + Frame acquireCopyFrame(const MtSwapchainState::Ticket& ticket); }; } diff --git a/Engine/gapi/metal/mtswapchain.mm b/Engine/gapi/metal/mtswapchain.mm index c9f3e50a..bb23dba7 100644 --- a/Engine/gapi/metal/mtswapchain.mm +++ b/Engine/gapi/metal/mtswapchain.mm @@ -22,6 +22,9 @@ #import #import +#include +#include + using namespace Tempest; using namespace Tempest::Detail; @@ -63,6 +66,56 @@ - (CALayer *)makeBackingLayer { } }; +template +static NsPtr strongRef(T* ptr) { + if(ptr!=nullptr) + ptr->retain(); + return NsPtr(ptr); + } + +MtSwapchainFrame::MtSwapchainFrame(uint64_t generation, uint32_t image, + bool direct, MTL::Texture* texture, + CA::MetalDrawable* drawable) + :generation(generation), image(image), direct(direct), + texture(strongRef(texture)), drawable(strongRef(drawable)) { + } + +MtSwapchainFrame::~MtSwapchainFrame() { + } + +namespace { + +struct PresentFrame final { + MtSwapchain::Frame rendered; + NsPtr drawable; + + PresentFrame(MtSwapchain::Frame rendered, CA::MetalDrawable* drawable) + :rendered(std::move(rendered)), drawable(strongRef(drawable)) { + } + }; + +class DeviceSubmission final { + public: + explicit DeviceSubmission(MtDevice* device) :device(device) { + device->onSubmit(); + } + + ~DeviceSubmission() { + finish(); + } + + void finish() noexcept { + if(!finished.exchange(true,std::memory_order_acq_rel)) + device->onFinish(); + } + + private: + MtDevice* device = nullptr; + std::atomic_bool finished{false}; + }; + +} + static float backingScaleFactor(SysWindow* w) { #if defined(__OSX__) return [w screen].backingScaleFactor; @@ -92,8 +145,9 @@ static CGRect windowRect(UIWindow* wnd) { #endif // note : MoltenVK supports NSView, UIView, CAMetalLayer, so we should align to it -MtSwapchain::MtSwapchain(MtDevice& dev, SystemApi::Window *w, uint32_t bufferCount) - :dev(dev), pimpl(new Impl()) { +MtSwapchain::MtSwapchain(MtDevice& dev, SystemApi::Window *w, + uint32_t bufferCount, bool directPreferred) + :pimpl(new Impl()), dev(dev), directPreferred(directPreferred) { NSObject* obj = reinterpret_cast(w); if([obj isKindOfClass : [SysWindow class]]) pimpl->wnd = reinterpret_cast(w); @@ -123,94 +177,321 @@ static CGRect windowRect(UIWindow* wnd) { lay.maximumDrawableCount = bufferCount; #endif lay.pixelFormat = MTLPixelFormatBGRA8Unorm; - lay.allowsNextDrawableTimeout = NO; + lay.allowsNextDrawableTimeout = directPreferred ? YES : NO; lay.framebufferOnly = NO; reset(); } MtSwapchain::~MtSwapchain() { + Frame retiredFrame; + std::vector retiredImages; + auto exclusive = operationGate.blockNewOperations(); + { + std::lock_guard guard(sync); + state.reset(0); + activeFrame.swap(retiredFrame); + img.swap(retiredImages); + sz = {0,0}; + } + exclusive.wait(); + dev.waitIdle(); + + // Releasing Metal objects can enter the Objective-C runtime. Keep it out of + // the swapchain spinlock, including during destruction. + retiredFrame.reset(); + retiredImages.clear(); + if(pimpl->view!=nil) [pimpl->view release]; } void MtSwapchain::reset() { + Frame retiredFrame; + std::vector retiredImages; + // First stop new CPU operations, then invalidate all tickets. Active + // acquire/present operations can unwind without contending on the gate. + auto exclusive = operationGate.blockNewOperations(); + { + std::lock_guard guard(sync); + state.reset(0); + activeFrame.swap(retiredFrame); + img.swap(retiredImages); + sz = {0,0}; + } + + // No layer/device call is made under the swapchain spinlock. + exclusive.wait(); dev.waitIdle(); // pending commands - std::lock_guard guard(sync); + retiredFrame.reset(); + retiredImages.clear(); // https://developer.apple.com/documentation/quartzcore/cametallayer?language=objc CAMetalLayer* lay = pimpl->metalLayer(); auto wrect = windowRect(pimpl->wnd); // auto lrect = lay.frame; lay.drawableSize = wrect.size; - sz = {int(wrect.size.width), int(wrect.size.height)}; - imgCount = uint32_t(lay.maximumDrawableCount); - - img.resize(imgCount); - for(size_t i=0; i newImages(imageCount); + if(!directPreferred) { + // Preserve the established Copy path: all private back buffers are + // allocated eagerly during reset. + for(auto& image:newImages) + image.tex = mkTexture(newSize); + } - currentImg = 0; + { + std::lock_guard guard(sync); + sz = newSize; + img.swap(newImages); + state.reset(imageCount); + } } uint32_t MtSwapchain::currentBackBufferIndex() { - return currentImg; + std::lock_guard guard(sync); + return state.currentImage(); } -void MtSwapchain::present() { - auto pool = NsPtr::init(); - - CA::MetalLayer* lay = reinterpret_cast(pimpl->metalLayer()); - uint32_t i = currentImg; - auto drawable = lay->nextDrawable(); - if(drawable==nullptr) +MtSwapchain::Frame MtSwapchain::acquireCopyFrame(const MtSwapchainState::Ticket& ticket) { + Tempest::Size expected; + MTL::Texture* existing = nullptr; + bool valid = false; + { + std::lock_guard guard(sync); + valid = state.isAcquiring(ticket) && ticket.image guard(sync); - auto dr = drawable->texture(); - if(dr->width()!=img[i].tex->width() || dr->height()!=img[i].tex->height()) { + if(existing!=nullptr) + return std::make_shared(ticket.generation,ticket.image, + false,existing,nullptr); + + // Direct mode allocates its private fallback only when drawable acquisition + // fails. Texture allocation is intentionally outside the state lock. + auto created = mkTexture(expected); + MTL::Texture* selected = nullptr; + { + std::lock_guard guard(sync); + valid = state.isAcquiring(ticket) && ticket.image(ticket.generation,ticket.image, + false,selected,nullptr); + } + +MtSwapchain::Frame MtSwapchain::acquireRenderTarget(uint32_t image) { + auto operation = operationGate.startOperation(); + return acquireRenderTargetImpl(image); + } + +MtSwapchain::Frame MtSwapchain::acquireRenderTargetImpl(uint32_t image) { + MtSwapchainState::Acquire acquire; + Tempest::Size expected; + Frame reuse; + bool invalid = false; + { + std::lock_guard guard(sync); + acquire = state.beginAcquire(); + if(acquire.result==MtSwapchainState::Acquire::Result::Reuse) { + invalid = activeFrame==nullptr || image!=acquire.ticket.image; + if(!invalid) + reuse = activeFrame; + } + else if(acquire.result!=MtSwapchainState::Acquire::Result::Start || + image!=acquire.ticket.image) { + if(acquire.result==MtSwapchainState::Acquire::Result::Start) + state.cancelAcquire(acquire.ticket); + invalid = true; + } + else { + expected = sz; + } } - - auto desc = NsPtr::init(); - //desc->setRetainedReferences(true); - desc->setErrorOptions(MTL::CommandBufferErrorOptionEncoderExecutionStatus); - - auto cmd = dev.queue->commandBuffer(desc.get()); - auto enc = cmd->blitCommandEncoder(); - - enc->copyFromTexture(img[i].tex.get(), 0, 0, - dr, 0, 0, - 1, 1); - enc->endEncoding(); - cmd->presentDrawable(drawable); - - dev.onSubmit(); - cmd->addCompletedHandler(^(MTL::CommandBuffer* c){ - MTL::CommandBufferStatus s = c->status(); - if(s==MTL::CommandBufferStatusNotEnqueued || - s==MTL::CommandBufferStatusEnqueued || - s==MTL::CommandBufferStatusCommitted || - s==MTL::CommandBufferStatusScheduled) - return; - - if(s!=MTL::CommandBufferStatusCompleted) { - Log::e("swapchain fatal error"); - dev.onFinish(); - dev.waitIdle(); - return; + if(invalid) + throw SwapchainSuboptimal(); + if(reuse!=nullptr) + return reuse; + + Frame frame; + try { + if(directPreferred) { + auto pool = NsPtr::init(); + auto* lay = reinterpret_cast(pimpl->metalLayer()); + auto drawable = strongRef(lay->nextDrawable()); + auto* texture = drawable==nullptr ? nullptr : drawable->texture(); + const bool sizeMatches = texture!=nullptr && + texture->width()==uint32_t(expected.w) && + texture->height()==uint32_t(expected.h); + if(MtSwapchainState::chooseTarget(true,texture!=nullptr,sizeMatches)== + MtSwapchainState::Target::Direct) { + frame = std::make_shared(acquire.ticket.generation, + acquire.ticket.image,true, + texture,drawable.get()); + } } - dev.onFinish(); - }); - cmd->commit(); + if(frame==nullptr) + frame = acquireCopyFrame(acquire.ticket); + + // Prepare the shared ownership before entering the spinlock. Swapping it + // into activeFrame cannot allocate or release the previous frame there. + Frame publishedFrame = frame; + bool published = false; + { + std::lock_guard guard(sync); + published = state.publish(acquire.ticket, + frame->direct ? MtSwapchainState::Target::Direct + : MtSwapchainState::Target::Copy); + if(published) + activeFrame.swap(publishedFrame); + } + if(!published) + throw SwapchainSuboptimal(); + return frame; + } + catch(...) { + { + std::lock_guard guard(sync); + state.cancelAcquire(acquire.ticket); + } + throw; + } + } - nextDrawable(); +void MtSwapchain::present() { + auto operation = operationGate.startOperation(); + MtSwapchainState::Ticket ticket; + Frame frame; + bool acquireCopy = false; + uint32_t copyImage = 0; + bool invalid = false; + { + std::lock_guard guard(sync); + if(state.beginPresent(ticket)) { + if(activeFrame==nullptr) { + state.presentFailed(ticket); + invalid = true; + } + else { + frame = activeFrame; + } + } + else if(!directPreferred && state.currentPhase()==MtSwapchainState::Phase::Idle) { + // The old Copy implementation allowed presenting an untouched private + // back buffer. Preserve that edge case for source/behaviour compatibility. + acquireCopy = true; + copyImage = state.currentImage(); + } + else { + invalid = true; + } + } + if(invalid) + throw SwapchainSuboptimal(); + + if(acquireCopy) { + acquireRenderTargetImpl(copyImage); + { + std::lock_guard guard(sync); + const bool began = state.beginPresent(ticket); + invalid = !began || activeFrame==nullptr; + if(invalid) { + if(began) + state.presentFailed(ticket); + } + else { + frame = activeFrame; + } + } + if(invalid) + throw SwapchainSuboptimal(); + } + + bool committed = false; + std::shared_ptr submission; + try { + auto pool = NsPtr::init(); + CA::MetalDrawable* drawable = frame->drawable.get(); + NsPtr acquired; + if(drawable==nullptr) { + auto* lay = reinterpret_cast(pimpl->metalLayer()); + acquired = strongRef(lay->nextDrawable()); + drawable = acquired.get(); + } + + auto* target = drawable==nullptr ? nullptr : drawable->texture(); + if(target==nullptr || target->width()!=frame->texture->width() || + target->height()!=frame->texture->height()) + throw SwapchainSuboptimal(); + + auto keepAlive = std::make_shared(frame,drawable); + auto desc = NsPtr::init(); + if(desc==nullptr) + throw std::system_error(GraphicsErrc::OutOfVideoMemory); + desc->setRetainedReferences(true); + desc->setErrorOptions(MTL::CommandBufferErrorOptionEncoderExecutionStatus); + + auto* cmd = dev.queue->commandBuffer(desc.get()); + if(cmd==nullptr) + throw std::system_error(GraphicsErrc::OutOfVideoMemory); + + if(!frame->direct) { + auto* enc = cmd->blitCommandEncoder(); + if(enc==nullptr) + throw std::system_error(GraphicsErrc::OutOfVideoMemory); + enc->copyFromTexture(frame->texture.get(), 0, 0, + target, 0, 0, + 1, 1); + enc->endEncoding(); + } + cmd->presentDrawable(drawable); + auto* stableDevice = &dev; + submission = std::make_shared(stableDevice); + cmd->addCompletedHandler(^(MTL::CommandBuffer* c){ + (void)keepAlive; + const MTL::CommandBufferStatus status = c->status(); + if(status!=MTL::CommandBufferStatusCompleted) + Log::e("swapchain fatal error"); + submission->finish(); + }); + cmd->commit(); + committed = true; + + Frame retiredFrame; + { + std::lock_guard guard(sync); + if(state.presentCommitted(ticket)) + activeFrame.swap(retiredFrame); + } + } + catch(...) { + if(!committed) { + if(submission!=nullptr) + submission->finish(); + { + std::lock_guard guard(sync); + state.presentFailed(ticket); + } + } + throw; + } } -NsPtr MtSwapchain::mkTexture() { +NsPtr MtSwapchain::mkTexture(const Tempest::Size& size) { auto pool = NsPtr::init(); auto desc = NsPtr::init(); if(desc==nullptr) @@ -218,8 +499,8 @@ static CGRect windowRect(UIWindow* wnd) { desc->setTextureType(MTL::TextureType2D); desc->setPixelFormat(MTL::PixelFormatBGRA8Unorm); - desc->setWidth(sz.w); - desc->setHeight(sz.h); + desc->setWidth(size.w); + desc->setHeight(size.h); desc->setMipmapLevelCount(1); desc->setCpuCacheMode(MTL::CPUCacheModeDefaultCache); desc->setStorageMode(MTL::StorageModePrivate); @@ -232,19 +513,18 @@ static CGRect windowRect(UIWindow* wnd) { return impl; } -void MtSwapchain::nextDrawable() { - currentImg = (currentImg+1) % img.size(); - } - uint32_t MtSwapchain::imageCount() const { - return imgCount; + std::lock_guard guard(sync); + return state.imageCount(); } uint32_t MtSwapchain::w() const { + std::lock_guard guard(sync); return sz.w; } uint32_t MtSwapchain::h() const { + std::lock_guard guard(sync); return sz.h; } diff --git a/Engine/gapi/metal/mtswapchainstate.h b/Engine/gapi/metal/mtswapchainstate.h new file mode 100644 index 00000000..0a8bb286 --- /dev/null +++ b/Engine/gapi/metal/mtswapchainstate.h @@ -0,0 +1,267 @@ +#pragma once + +#include +#include +#include +#include + +namespace Tempest::Detail { + +/** + * Serializes swapchain reset/destruction against CPU-side acquire/present + * operations. The exclusive side can invalidate the render generation before + * waiting, while new operations remain blocked until layer reconfiguration is + * complete. + */ +class MtSwapchainOperationGate final { + public: + class Operation final { + public: + Operation() = default; + Operation(const Operation&) = delete; + Operation& operator=(const Operation&) = delete; + + Operation(Operation&& other) noexcept :gate(other.gate) { + other.gate = nullptr; + } + Operation& operator=(Operation&& other) noexcept { + if(this==&other) + return *this; + finish(); + gate = other.gate; + other.gate = nullptr; + return *this; + } + ~Operation() { finish(); } + + void finish() noexcept { + if(gate==nullptr) + return; + auto* owner = gate; + gate = nullptr; + owner->finishOperation(); + } + + private: + explicit Operation(MtSwapchainOperationGate* gate) :gate(gate) {} + MtSwapchainOperationGate* gate = nullptr; + + friend class MtSwapchainOperationGate; + }; + + class Exclusive final { + public: + Exclusive() = default; + Exclusive(const Exclusive&) = delete; + Exclusive& operator=(const Exclusive&) = delete; + + Exclusive(Exclusive&& other) noexcept :gate(other.gate) { + other.gate = nullptr; + } + Exclusive& operator=(Exclusive&& other) noexcept { + if(this==&other) + return *this; + release(); + gate = other.gate; + other.gate = nullptr; + return *this; + } + ~Exclusive() { release(); } + + void wait() { + if(gate!=nullptr) + gate->waitForOperations(); + } + + void release() noexcept { + if(gate==nullptr) + return; + auto* owner = gate; + gate = nullptr; + owner->releaseExclusive(); + } + + private: + explicit Exclusive(MtSwapchainOperationGate* gate) :gate(gate) {} + MtSwapchainOperationGate* gate = nullptr; + + friend class MtSwapchainOperationGate; + }; + + Operation startOperation() { + std::unique_lock guard(sync); + changed.wait(guard,[this](){ return !exclusive; }); + ++active; + return Operation(this); + } + + Exclusive blockNewOperations() { + std::unique_lock guard(sync); + changed.wait(guard,[this](){ return !exclusive; }); + exclusive = true; + return Exclusive(this); + } + + private: + void finishOperation() noexcept { + std::lock_guard guard(sync); + assert(active>0); + --active; + changed.notify_all(); + } + + void waitForOperations() { + std::unique_lock guard(sync); + changed.wait(guard,[this](){ return active==0; }); + } + + void releaseExclusive() noexcept { + std::lock_guard guard(sync); + exclusive = false; + changed.notify_all(); + } + + std::mutex sync; + std::condition_variable changed; + uint32_t active = 0; + bool exclusive = false; + }; + +/** + * Small, platform-independent state machine used by MtSwapchain. + * + * CAMetalLayer calls deliberately live outside this class (and outside the + * swapchain lock). Tickets make their results invalid as soon as reset() + * starts a new generation. + */ +class MtSwapchainState final { + public: + enum class Target:uint8_t { + Copy, + Direct, + }; + + enum class Phase:uint8_t { + Idle, + Acquiring, + Ready, + Presenting, + }; + + struct Ticket { + uint64_t generation = 0; + uint64_t serial = 0; + uint32_t image = 0; + + bool operator==(const Ticket& other) const noexcept { + return generation==other.generation && serial==other.serial && + image==other.image; + } + bool operator!=(const Ticket& other) const noexcept { + return !(*this==other); + } + }; + + struct Acquire { + enum class Result:uint8_t { + Start, + Reuse, + Busy, + }; + + Result result = Result::Busy; + Ticket ticket; + }; + + void reset(uint32_t count) noexcept { + ++generation; + ++serial; + images = count; + current = 0; + phase = Phase::Idle; + active = {}; + target = Target::Copy; + } + + Acquire beginAcquire() noexcept { + if(images==0) + return {}; + if(phase==Phase::Ready) + return {Acquire::Result::Reuse,active}; + if(phase!=Phase::Idle) + return {}; + + active = {generation,++serial,current}; + phase = Phase::Acquiring; + return {Acquire::Result::Start,active}; + } + + bool publish(const Ticket& ticket, Target renderTarget) noexcept { + if(phase!=Phase::Acquiring || ticket!=active || + ticket.generation!=generation || ticket.image!=current) + return false; + target = renderTarget; + phase = Phase::Ready; + return true; + } + + bool cancelAcquire(const Ticket& ticket) noexcept { + if(phase!=Phase::Acquiring || ticket!=active) + return false; + phase = Phase::Idle; + active = {}; + return true; + } + + bool isAcquiring(const Ticket& ticket) const noexcept { + return phase==Phase::Acquiring && ticket==active && + ticket.generation==generation && ticket.image==current; + } + + bool beginPresent(Ticket& ticket) noexcept { + if(phase!=Phase::Ready) + return false; + phase = Phase::Presenting; + ticket = active; + return true; + } + + bool presentFailed(const Ticket& ticket) noexcept { + if(phase!=Phase::Presenting || ticket!=active) + return false; + phase = Phase::Ready; + return true; + } + + bool presentCommitted(const Ticket& ticket) noexcept { + if(phase!=Phase::Presenting || ticket!=active || images==0) + return false; + current = (current+1)%images; + phase = Phase::Idle; + active = {}; + return true; + } + + uint32_t currentImage() const noexcept { return current; } + uint32_t imageCount() const noexcept { return images; } + uint64_t currentGeneration() const noexcept { return generation; } + Phase currentPhase() const noexcept { return phase; } + Target currentTarget() const noexcept { return target; } + + static Target chooseTarget(bool directPreferred, bool drawableAvailable, + bool drawableSizeMatches) noexcept { + return directPreferred && drawableAvailable && drawableSizeMatches + ? Target::Direct : Target::Copy; + } + + private: + uint64_t generation = 0; + uint64_t serial = 0; + uint32_t images = 0; + uint32_t current = 0; + Phase phase = Phase::Idle; + Target target = Target::Copy; + Ticket active; + }; + +} diff --git a/Engine/gapi/metalapi.cpp b/Engine/gapi/metalapi.cpp index 2307776a..a7200563 100644 --- a/Engine/gapi/metalapi.cpp +++ b/Engine/gapi/metalapi.cpp @@ -35,9 +35,13 @@ MetalApi::MetalApi(ApiFlags f) MetalApi::MetalApi(ApiFlags f, const Options& options) :swapchainBufferCount(options.swapchainBufferCount), - shaderModuleCacheSize(options.shaderModuleCacheSize) { + shaderModuleCacheSize(options.shaderModuleCacheSize), + swapchainRenderMode(options.swapchainRenderMode) { if(swapchainBufferCount!=0 && swapchainBufferCount!=2 && swapchainBufferCount!=3) throw std::invalid_argument("Metal swapchain buffer count must be 0, 2, or 3"); + if(swapchainRenderMode!=SwapchainRenderMode::Copy && + swapchainRenderMode!=SwapchainRenderMode::Direct) + throw std::invalid_argument("Unknown Metal swapchain render mode"); if((f & ApiFlags::Validation)==ApiFlags::Validation) { setenv("METAL_DEVICE_WRAPPER_TYPE","1",1); @@ -95,7 +99,8 @@ AbstractGraphicsApi::Device* MetalApi::createDevice(std::string_view gpuName) { AbstractGraphicsApi::Swapchain *MetalApi::createSwapchain(SystemApi::Window *w, AbstractGraphicsApi::Device* d) { auto& dev = *reinterpret_cast(d); - return new MtSwapchain(dev,w,swapchainBufferCount); + return new MtSwapchain(dev,w,swapchainBufferCount, + swapchainRenderMode==SwapchainRenderMode::Direct); } AbstractGraphicsApi::PPipeline MetalApi::createPipeline(AbstractGraphicsApi::Device *d, @@ -239,8 +244,11 @@ std::shared_ptr MetalApi::submit(Device* d, CommandB throw DeviceLostException(); MTL::CommandBuffer& cmd = *cx.impl; + auto swapchainFrames = std::make_shared>(); + swapchainFrames->swap(cx.swapchainFrames); dx->onSubmit(); cmd.addCompletedHandler(^(MTL::CommandBuffer* c){ + (void)swapchainFrames; const MTL::CommandBufferStatus s = c->status(); dx->signalFence(*pfence, s, MTL::CommandBufferError(c->error()->code()), c->error()); if(s==MTL::CommandBufferStatusCompleted || s==MTL::CommandBufferStatusError) diff --git a/Engine/gapi/metalapi.h b/Engine/gapi/metalapi.h index f8ec892a..bea43c05 100644 --- a/Engine/gapi/metalapi.h +++ b/Engine/gapi/metalapi.h @@ -6,11 +6,17 @@ namespace Tempest { class MetalApi : public AbstractGraphicsApi { public: + enum class SwapchainRenderMode:uint8_t { + Copy, + Direct, + }; + struct Options { - uint32_t swapchainBufferCount = 0; + uint32_t swapchainBufferCount = 0; // Maximum number of compiled Metal shader modules kept per device. // Zero disables caching. - size_t shaderModuleCacheSize = 0; + size_t shaderModuleCacheSize = 0; + SwapchainRenderMode swapchainRenderMode = SwapchainRenderMode::Copy; }; explicit MetalApi(ApiFlags f=ApiFlags::NoFlags); @@ -55,9 +61,10 @@ class MetalApi : public AbstractGraphicsApi { void getCaps(Device *d, Props& caps) override; private: - bool validation = false; - uint32_t swapchainBufferCount = 0; - size_t shaderModuleCacheSize = 0; + bool validation = false; + uint32_t swapchainBufferCount = 0; + size_t shaderModuleCacheSize = 0; + SwapchainRenderMode swapchainRenderMode = SwapchainRenderMode::Copy; }; } diff --git a/Tests/tests/gapi/metal_swapchain_state_test.cpp b/Tests/tests/gapi/metal_swapchain_state_test.cpp new file mode 100644 index 00000000..6d5d33d4 --- /dev/null +++ b/Tests/tests/gapi/metal_swapchain_state_test.cpp @@ -0,0 +1,195 @@ +#include + +#include "../../../Engine/gapi/metal/mtswapchainstate.h" + +#include +#include +#include +#include + +using namespace Tempest::Detail; +using namespace std::chrono_literals; + +TEST(MetalSwapchainGate,AcquireResetOrdering) { + MtSwapchainOperationGate gate; + auto acquire = gate.startOperation(); + + std::promise resetStarted; + std::promise resetPassed; + std::promise releaseReset; + auto releaseResetFuture = releaseReset.get_future(); + auto resetPassedFuture = resetPassed.get_future(); + std::thread reset([&](){ + auto exclusive = gate.blockNewOperations(); + resetStarted.set_value(); + exclusive.wait(); + resetPassed.set_value(); + releaseResetFuture.wait(); + }); + + resetStarted.get_future().wait(); + EXPECT_EQ(resetPassedFuture.wait_for(0ms),std::future_status::timeout); + + std::promise secondAcquireStarted; + auto secondAcquireFuture = secondAcquireStarted.get_future(); + std::thread secondAcquire([&](){ + auto operation = gate.startOperation(); + secondAcquireStarted.set_value(); + }); + EXPECT_EQ(secondAcquireFuture.wait_for(10ms),std::future_status::timeout); + + acquire.finish(); + EXPECT_EQ(resetPassedFuture.wait_for(1s),std::future_status::ready); + EXPECT_EQ(secondAcquireFuture.wait_for(10ms),std::future_status::timeout); + + releaseReset.set_value(); + reset.join(); + EXPECT_EQ(secondAcquireFuture.wait_for(1s),std::future_status::ready); + secondAcquire.join(); + } + +TEST(MetalSwapchainGate,PresentRegistersSubmissionBeforeResetPasses) { + MtSwapchainOperationGate gate; + auto present = gate.startOperation(); + std::atomic_bool submitted{false}; + std::atomic_bool resetObservedSubmit{false}; + std::promise resetStarted; + + std::thread reset([&](){ + auto exclusive = gate.blockNewOperations(); + resetStarted.set_value(); + exclusive.wait(); + resetObservedSubmit.store(submitted.load(std::memory_order_acquire), + std::memory_order_release); + }); + + resetStarted.get_future().wait(); + submitted.store(true,std::memory_order_release); + present.finish(); + reset.join(); + EXPECT_TRUE(resetObservedSubmit.load(std::memory_order_acquire)); + } + +TEST(MetalSwapchainGate,ResetWaitsForAllFramesInFlight) { + MtSwapchainOperationGate gate; + auto first = gate.startOperation(); + auto second = gate.startOperation(); + std::promise resetStarted; + std::promise resetPassed; + auto resetPassedFuture = resetPassed.get_future(); + + std::thread reset([&](){ + auto exclusive = gate.blockNewOperations(); + resetStarted.set_value(); + exclusive.wait(); + resetPassed.set_value(); + }); + + resetStarted.get_future().wait(); + first.finish(); + EXPECT_EQ(resetPassedFuture.wait_for(10ms),std::future_status::timeout); + second.finish(); + EXPECT_EQ(resetPassedFuture.wait_for(1s),std::future_status::ready); + reset.join(); + } + +TEST(MetalSwapchainGate,CompletionOwnsResourcesUntilRelease) { + auto resource = std::make_shared(42); + std::weak_ptr lifetime = resource; + auto completion = [keepAlive=resource]() mutable { + keepAlive.reset(); + }; + resource.reset(); + EXPECT_FALSE(lifetime.expired()); + completion(); + EXPECT_TRUE(lifetime.expired()); + } + +TEST(MetalSwapchainState,SingleAcquireAndReuse) { + MtSwapchainState state; + state.reset(3); + + const auto first = state.beginAcquire(); + EXPECT_EQ(first.result,MtSwapchainState::Acquire::Result::Start); + EXPECT_EQ(first.ticket.image,0u); + EXPECT_EQ(state.beginAcquire().result,MtSwapchainState::Acquire::Result::Busy); + + ASSERT_TRUE(state.publish(first.ticket,MtSwapchainState::Target::Direct)); + const auto reused = state.beginAcquire(); + EXPECT_EQ(reused.result,MtSwapchainState::Acquire::Result::Reuse); + EXPECT_EQ(reused.ticket,first.ticket); + EXPECT_EQ(state.currentTarget(),MtSwapchainState::Target::Direct); + } + +TEST(MetalSwapchainState,CopyFallbackCanBePublished) { + MtSwapchainState state; + state.reset(2); + const auto acquire = state.beginAcquire(); + + ASSERT_TRUE(state.publish(acquire.ticket,MtSwapchainState::Target::Copy)); + EXPECT_EQ(state.currentTarget(),MtSwapchainState::Target::Copy); + EXPECT_EQ(state.currentImage(),0u); + } + +TEST(MetalSwapchainState,MissingOrWrongDrawableSelectsCopyFallback) { + EXPECT_EQ(MtSwapchainState::chooseTarget(true,false,false), + MtSwapchainState::Target::Copy); + EXPECT_EQ(MtSwapchainState::chooseTarget(true,true,false), + MtSwapchainState::Target::Copy); + EXPECT_EQ(MtSwapchainState::chooseTarget(true,true,true), + MtSwapchainState::Target::Direct); + EXPECT_EQ(MtSwapchainState::chooseTarget(false,true,true), + MtSwapchainState::Target::Copy); + } + +TEST(MetalSwapchainState,ResetInvalidatesAcquireAndPresent) { + MtSwapchainState state; + state.reset(2); + const auto acquiring = state.beginAcquire(); + state.reset(3); + EXPECT_FALSE(state.publish(acquiring.ticket,MtSwapchainState::Target::Direct)); + EXPECT_EQ(state.currentImage(),0u); + EXPECT_EQ(state.imageCount(),3u); + + const auto acquired = state.beginAcquire(); + ASSERT_TRUE(state.publish(acquired.ticket,MtSwapchainState::Target::Direct)); + MtSwapchainState::Ticket presenting; + ASSERT_TRUE(state.beginPresent(presenting)); + state.reset(2); + EXPECT_FALSE(state.presentCommitted(presenting)); + EXPECT_EQ(state.currentImage(),0u); + } + +TEST(MetalSwapchainState,IndexAdvancesOnlyAfterCommittedPresent) { + MtSwapchainState state; + state.reset(2); + const auto acquired = state.beginAcquire(); + ASSERT_TRUE(state.publish(acquired.ticket,MtSwapchainState::Target::Direct)); + + MtSwapchainState::Ticket firstPresent; + ASSERT_TRUE(state.beginPresent(firstPresent)); + ASSERT_TRUE(state.presentFailed(firstPresent)); + EXPECT_EQ(state.currentImage(),0u); + + MtSwapchainState::Ticket retry; + ASSERT_TRUE(state.beginPresent(retry)); + ASSERT_TRUE(state.presentCommitted(retry)); + EXPECT_EQ(state.currentImage(),1u); + + MtSwapchainState::Ticket duplicate; + EXPECT_FALSE(state.beginPresent(duplicate)); + EXPECT_FALSE(state.presentCommitted(retry)); + EXPECT_EQ(state.currentImage(),1u); + } + +TEST(MetalSwapchainState,CancelledAcquireIsRetryable) { + MtSwapchainState state; + state.reset(2); + const auto failed = state.beginAcquire(); + ASSERT_TRUE(state.cancelAcquire(failed.ticket)); + + const auto retry = state.beginAcquire(); + EXPECT_EQ(retry.result,MtSwapchainState::Acquire::Result::Start); + EXPECT_NE(retry.ticket.serial,failed.ticket.serial); + ASSERT_TRUE(state.publish(retry.ticket,MtSwapchainState::Target::Copy)); + } diff --git a/Tests/tests/gapi/metal_test.cpp b/Tests/tests/gapi/metal_test.cpp index 5fdc85ed..3aeea3cb 100644 --- a/Tests/tests/gapi/metal_test.cpp +++ b/Tests/tests/gapi/metal_test.cpp @@ -4,6 +4,8 @@ #include #include #include +#include +#include #include #include @@ -94,6 +96,59 @@ bool metalFxTemporalSupported() { return false; #endif } + +class MetalSwapchainWindow final : public Window { + public: + SystemApi::Window* nativeWindow() const { return hwnd(); } + }; + +void swapchainSmoke(MetalApi::SwapchainRenderMode mode) { + MetalApi::Options options; + options.swapchainRenderMode = mode; + MetalApi api(ApiFlags::Validation,options); + Device device(api); + MetalSwapchainWindow window; + window.resize(64,64); + Application::processEvents(); + + Swapchain swapchain(device,window.nativeWindow()); + std::vector fences; + auto renderFrame = [&](const Vec4& clear) { + const uint32_t before = swapchain.currentImage(); + auto command = device.commandBuffer(); + { + auto encoder = command.startEncoding(device); + encoder.setFramebuffer({{swapchain[before],clear,Tempest::Preserve}}); + } + fences.emplace_back(device.submit(command)); + device.present(swapchain); + EXPECT_EQ(swapchain.currentImage(),(before+1)%swapchain.imageCount()); + }; + + renderFrame(Vec4(1.f,0.f,0.f,1.f)); + renderFrame(Vec4(0.f,1.f,0.f,1.f)); + // Reset while both render and presentation command buffers can still be in + // flight. The operation gate and device idle tracking must close this race. + swapchain.reset(); + for(auto& fence:fences) + fence.wait(); + fences.clear(); + + window.resize(96,80); + Application::processEvents(); + swapchain.reset(); + EXPECT_GT(swapchain.w(),0u); + EXPECT_GT(swapchain.h(),0u); + + renderFrame(Vec4(0.f,0.f,1.f,1.f)); + renderFrame(Vec4(0.25f,0.5f,0.75f,1.f)); + if(mode==MetalApi::SwapchainRenderMode::Direct) + EXPECT_THROW(device.present(swapchain),SwapchainSuboptimal); + // Leave the final submissions in flight. MtSwapchain destruction must wait + // for both CPU presentation setup and GPU completion without using `this` + // from a completion handler. + } + #endif } @@ -454,6 +509,29 @@ TEST(MetalApi,TemporalScaler) { #endif } +TEST(MetalApi,SwapchainRenderModeOptions) { + MetalApi::Options defaults; + EXPECT_EQ(defaults.swapchainRenderMode,MetalApi::SwapchainRenderMode::Copy); + + MetalApi::Options direct; + direct.swapchainRenderMode = MetalApi::SwapchainRenderMode::Direct; +#if defined(__OSX__) + EXPECT_NO_THROW(MetalApi(ApiFlags::NoFlags,direct)); + + MetalApi::Options invalid; + invalid.swapchainRenderMode = static_cast(255); + EXPECT_THROW(MetalApi(ApiFlags::NoFlags,invalid),std::invalid_argument); +#endif + } + +TEST(MetalApi,SwapchainCopyAndDirectSmoke) { +#if defined(__OSX__) + Application app; + swapchainSmoke(MetalApi::SwapchainRenderMode::Copy); + swapchainSmoke(MetalApi::SwapchainRenderMode::Direct); +#endif + } + TEST(MetalApi,Vbo) { #if defined(__OSX__) GapiTestCommon::Vbo(); From ab6777c78acc6c4d1b336aee8ed1ff1a5fb050d6 Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 16:22:58 +0100 Subject: [PATCH 19/25] Add optional precompiled Metal shader libraries --- Engine/gapi/metal/mtdevice.h | 8 +- Engine/gapi/metal/mtdevice.mm | 10 +- Engine/gapi/metal/mtprecompiledlibrary.h | 45 +++ Engine/gapi/metal/mtprecompiledlibrary.mm | 157 ++++++++ Engine/gapi/metal/mtsha256.cpp | 133 +++++++ Engine/gapi/metal/mtsha256.h | 35 ++ Engine/gapi/metal/mtshader.cpp | 42 ++- Engine/gapi/metalapi.cpp | 220 ++++++++++- Engine/gapi/metalapi.h | 95 ++++- Tests/tests/CMakeLists.txt | 8 + Tests/tests/gapi/metal_test.cpp | 427 ++++++++++++++++++++++ 11 files changed, 1154 insertions(+), 26 deletions(-) create mode 100644 Engine/gapi/metal/mtprecompiledlibrary.h create mode 100644 Engine/gapi/metal/mtprecompiledlibrary.mm create mode 100644 Engine/gapi/metal/mtsha256.cpp create mode 100644 Engine/gapi/metal/mtsha256.h diff --git a/Engine/gapi/metal/mtdevice.h b/Engine/gapi/metal/mtdevice.h index 422fb1eb..8f662cb4 100644 --- a/Engine/gapi/metal/mtdevice.h +++ b/Engine/gapi/metal/mtdevice.h @@ -4,6 +4,7 @@ #include #include #include +#include #include #include @@ -20,6 +21,8 @@ class MTLDevice; namespace Tempest { namespace Detail { +class MtPrecompiledLibraries; + inline MTL::PixelFormat nativeFormat(TextureFormat frm) { switch(frm) { case Undefined: @@ -246,7 +249,8 @@ inline MTL::RenderStages nativeFormat(ShaderReflection::Stage st) { class MtDevice : public AbstractGraphicsApi::Device { public: - MtDevice(std::string_view name, bool validation, size_t shaderModuleCacheSize); + MtDevice(std::string_view name, bool validation, + std::shared_ptr precompiledOptions = {}); ~MtDevice(); static const uint32_t MaxFences = 32; @@ -288,8 +292,10 @@ class MtDevice : public AbstractGraphicsApi::Device { Props prop; MtSamplerCache samplers; + std::shared_ptr precompiledOptions; ShaderModuleCache shaderModules; + std::unique_ptr precompiledLibraries; bool validation = false; static void deductProps(AbstractGraphicsApi::Props& prop, MTL::Device& dev); diff --git a/Engine/gapi/metal/mtdevice.mm b/Engine/gapi/metal/mtdevice.mm index 464b3e09..6feadf11 100644 --- a/Engine/gapi/metal/mtdevice.mm +++ b/Engine/gapi/metal/mtdevice.mm @@ -1,6 +1,7 @@ #if defined(TEMPEST_BUILD_METAL) #include "mtdevice.h" +#include "mtprecompiledlibrary.h" #include "thirdparty/spirv_cross/spirv_msl.hpp" #include @@ -35,8 +36,11 @@ return std::min(MTL::LanguageVersion3_1, opt->languageVersion()); } -MtDevice::MtDevice(std::string_view name, bool validation, size_t shaderModuleCacheSize) - : impl(mkDevice(name)), samplers(*impl), shaderModules(shaderModuleCacheSize), validation(validation) { +MtDevice::MtDevice(std::string_view name, bool validation, + std::shared_ptr precompiledOptions) + : impl(mkDevice(name)), samplers(*impl), precompiledOptions(std::move(precompiledOptions)), + shaderModules(this->precompiledOptions!=nullptr ? + this->precompiledOptions->shaderModuleCacheSize : 0), validation(validation) { if(impl.get()==nullptr) throw std::system_error(Tempest::GraphicsErrc::NoDevice); @@ -61,6 +65,8 @@ prop.transposedRtMatrix = true; } deductProps(prop,*impl); + if(this->precompiledOptions!=nullptr && !this->precompiledOptions->precompiledLibraries.empty()) + precompiledLibraries = std::make_unique(*impl,*this->precompiledOptions); } MtDevice::~MtDevice() { diff --git a/Engine/gapi/metal/mtprecompiledlibrary.h b/Engine/gapi/metal/mtprecompiledlibrary.h new file mode 100644 index 00000000..91529b1b --- /dev/null +++ b/Engine/gapi/metal/mtprecompiledlibrary.h @@ -0,0 +1,45 @@ +#pragma once + +#include + +#include + +#include "nsptr.h" + +#include +#include + +namespace Tempest { +namespace Detail { + +class MtPrecompiledLibraries final { + public: + MtPrecompiledLibraries(MTL::Device& device, const MetalApi::Options& options); + + NsPtr find(std::string_view canonicalMsl, + const MetalApi::PrecompiledShaderProfile& runtimeProfile) const; + + private: + struct Library { + NsPtr impl; + }; + + struct Entry { + MetalApi::PrecompiledShaderProfile profile; + MetalApi::PrecompiledShaderKey key = {}; + size_t library = 0; + bool duplicate = false; + }; + + static MetalApi::PrecompiledPlatform currentPlatform(); + static bool valid(const MetalApi::PrecompiledShader& shader); + static bool sameGenerationProfile(const MetalApi::PrecompiledShaderProfile& a, + const MetalApi::PrecompiledShaderProfile& b); + static bool sameStage(MTL::FunctionType actual, MetalApi::PrecompiledShaderStage expected); + + std::vector libraries; + std::vector entries; + }; + +} +} diff --git a/Engine/gapi/metal/mtprecompiledlibrary.mm b/Engine/gapi/metal/mtprecompiledlibrary.mm new file mode 100644 index 00000000..bcb8db54 --- /dev/null +++ b/Engine/gapi/metal/mtprecompiledlibrary.mm @@ -0,0 +1,157 @@ +#if defined(TEMPEST_BUILD_METAL) + +#include "mtprecompiledlibrary.h" + +#include +#include + +using namespace Tempest; +using namespace Tempest::Detail; + +MtPrecompiledLibraries::MtPrecompiledLibraries(MTL::Device& device, + const MetalApi::Options& options) { + auto pool = NsPtr::init(); + if(![(id)(void*)(&device) respondsToSelector:@selector(newLibraryWithData:error:)]) + return; + + for(const auto& src:options.precompiledLibraries) { + if(src.data.empty() || src.shaders.empty()) + continue; + if(MetalApi::precompiledLibraryHash(src.data.data(),src.data.size())!=src.dataHash) + continue; + + bool hasEligibleEntry = false; + for(const auto& shader:src.shaders) + hasEligibleEntry |= valid(shader); + if(!hasEligibleEntry) + continue; + + dispatch_data_t data = dispatch_data_create(src.data.data(),src.data.size(),nullptr, + DISPATCH_DATA_DESTRUCTOR_DEFAULT); + if(data==nullptr) + continue; + + NS::Error* err = nullptr; + auto library = NsPtr(device.newLibrary(data,&err)); + dispatch_release(data); + if(library==nullptr || err!=nullptr) + continue; + + const size_t libraryIndex = libraries.size(); + libraries.push_back({std::move(library)}); + for(const auto& shader:src.shaders) { + if(!valid(shader)) + continue; + entries.push_back({shader.profile,shader.key,libraryIndex,false}); + } + } + + for(size_t i=0; i MtPrecompiledLibraries::find( + std::string_view canonicalMsl, + const MetalApi::PrecompiledShaderProfile& runtimeProfile) const { + auto pool = NsPtr::init(); + for(const auto& entry:entries) { + if(entry.duplicate) + continue; + if(!sameGenerationProfile(entry.profile,runtimeProfile)) + continue; + if(MetalApi::precompiledShaderKey(canonicalMsl,entry.profile)!=entry.key) + continue; + if(entry.library>=libraries.size() || libraries[entry.library].impl==nullptr) + continue; + + auto name = NsPtr(NS::String::string(entry.profile.entryPoint.c_str(),NS::UTF8StringEncoding)); + if(name==nullptr) + continue; + name->retain(); + + auto constants = NsPtr::init(); + if(constants==nullptr) + continue; + NS::Error* err = nullptr; + auto fn = NsPtr(libraries[entry.library].impl.get()->newFunction(name.get(),constants.get(),&err)); + if(fn==nullptr || err!=nullptr) + continue; + if(!sameStage(fn->functionType(),runtimeProfile.stage)) + continue; + return fn; + } + return NsPtr(); + } + +MetalApi::PrecompiledPlatform MtPrecompiledLibraries::currentPlatform() { +#if defined(__IOS__) +#if defined(TARGET_OS_SIMULATOR) && TARGET_OS_SIMULATOR + return MetalApi::PrecompiledPlatform::IOSSimulator; +#else + return MetalApi::PrecompiledPlatform::IOSDevice; +#endif +#else + return MetalApi::PrecompiledPlatform::MacOS; +#endif + } + +bool MtPrecompiledLibraries::valid(const MetalApi::PrecompiledShader& shader) { + const auto& profile = shader.profile; + if(profile.schemaVersion!=MetalApi::PrecompiledShaderSchemaVersion || + profile.mslGeneratorVersion!=MetalApi::MslGeneratorVersion || + profile.platform!=currentPlatform() || profile.entryPoint.empty() || + profile.entryPoint.find('\0')!=std::string::npos || profile.mslVersion==0) + return false; + if(uint8_t(profile.stage)>uint8_t(MetalApi::PrecompiledShaderStage::Mesh) || + profile.argumentBuffersTier>2) + return false; + return true; + } + +bool MtPrecompiledLibraries::sameGenerationProfile( + const MetalApi::PrecompiledShaderProfile& a, + const MetalApi::PrecompiledShaderProfile& b) { + return a.schemaVersion==b.schemaVersion && + a.mslGeneratorVersion==b.mslGeneratorVersion && + a.platform==b.platform && a.stage==b.stage && + a.entryPoint==b.entryPoint && + a.mslVersion==b.mslVersion && a.flipVertY==b.flipVertY && + a.bufferSizeBufferIndex==b.bufferSizeBufferIndex && + a.argumentBuffersTier==b.argumentBuffersTier && + a.runtimeArrayRichDescriptor==b.runtimeArrayRichDescriptor && + a.readWriteTextureFences==b.readWriteTextureFences && + a.nativeImageAtomics==b.nativeImageAtomics && + a.r32uiLinearTextureAlignment==b.r32uiLinearTextureAlignment && + a.r32uiAlignmentConstantId==b.r32uiAlignmentConstantId; + } + +bool MtPrecompiledLibraries::sameStage(MTL::FunctionType actual, + MetalApi::PrecompiledShaderStage expected) { + switch(expected) { + case MetalApi::PrecompiledShaderStage::Vertex: + return actual==MTL::FunctionTypeVertex; + case MetalApi::PrecompiledShaderStage::Control: + return actual==MTL::FunctionTypeKernel; + case MetalApi::PrecompiledShaderStage::Evaluate: + return actual==MTL::FunctionTypeVertex; + case MetalApi::PrecompiledShaderStage::Geometry: + return actual==MTL::FunctionTypeVertex; + case MetalApi::PrecompiledShaderStage::Fragment: + return actual==MTL::FunctionTypeFragment; + case MetalApi::PrecompiledShaderStage::Compute: + return actual==MTL::FunctionTypeKernel; + case MetalApi::PrecompiledShaderStage::Task: + return actual==MTL::FunctionTypeObject; + case MetalApi::PrecompiledShaderStage::Mesh: + return actual==MTL::FunctionTypeMesh; + } + return false; + } + +#endif diff --git a/Engine/gapi/metal/mtsha256.cpp b/Engine/gapi/metal/mtsha256.cpp new file mode 100644 index 00000000..f5ce6dfd --- /dev/null +++ b/Engine/gapi/metal/mtsha256.cpp @@ -0,0 +1,133 @@ +#include "mtsha256.h" + +#include +#include + +using namespace Tempest::Detail; + +namespace { + +constexpr uint32_t k[64] = { + 0x428a2f98u,0x71374491u,0xb5c0fbcfu,0xe9b5dba5u,0x3956c25bu,0x59f111f1u,0x923f82a4u,0xab1c5ed5u, + 0xd807aa98u,0x12835b01u,0x243185beu,0x550c7dc3u,0x72be5d74u,0x80deb1feu,0x9bdc06a7u,0xc19bf174u, + 0xe49b69c1u,0xefbe4786u,0x0fc19dc6u,0x240ca1ccu,0x2de92c6fu,0x4a7484aau,0x5cb0a9dcu,0x76f988dau, + 0x983e5152u,0xa831c66du,0xb00327c8u,0xbf597fc7u,0xc6e00bf3u,0xd5a79147u,0x06ca6351u,0x14292967u, + 0x27b70a85u,0x2e1b2138u,0x4d2c6dfcu,0x53380d13u,0x650a7354u,0x766a0abbu,0x81c2c92eu,0x92722c85u, + 0xa2bfe8a1u,0xa81a664bu,0xc24b8b70u,0xc76c51a3u,0xd192e819u,0xd6990624u,0xf40e3585u,0x106aa070u, + 0x19a4c116u,0x1e376c08u,0x2748774cu,0x34b0bcb5u,0x391c0cb3u,0x4ed8aa4au,0x5b9cca4fu,0x682e6ff3u, + 0x748f82eeu,0x78a5636fu,0x84c87814u,0x8cc70208u,0x90befffau,0xa4506cebu,0xbef9a3f7u,0xc67178f2u, + }; + +constexpr uint32_t rotr(uint32_t value, uint32_t count) { + return (value>>count) | (value<<(32u-count)); + } + +uint32_t loadBe(const uint8_t* data) { + return (uint32_t(data[0])<<24u) | (uint32_t(data[1])<<16u) | + (uint32_t(data[2])<<8u) | uint32_t(data[3]); + } + +} + +MtSha256::MtSha256() + :state{0x6a09e667u,0xbb67ae85u,0x3c6ef372u,0xa54ff53au, + 0x510e527fu,0x9b05688cu,0x1f83d9abu,0x5be0cd19u} { + } + +void MtSha256::update(std::string_view data) { + update(data.data(),data.size()); + } + +void MtSha256::update(const void* ptr, size_t size) { + if(finalized || size==0) + return; + + const auto* data = static_cast(ptr); + byteCount += size; + + if(pendingSize>0) { + const size_t take = std::min(size_t(64)-pendingSize,size); + std::memcpy(pending+pendingSize,data,take); + pendingSize += take; + data += take; + size -= take; + if(pendingSize==64) { + transform(pending); + pendingSize = 0; + } + } + + while(size>=64) { + transform(data); + data += 64; + size -= 64; + } + + if(size>0) { + std::memcpy(pending,data,size); + pendingSize = size; + } + } + +MtSha256::Digest MtSha256::finalize() { + if(!finalized) { + const uint64_t bitCount = byteCount*8u; + pending[pendingSize++] = 0x80u; + if(pendingSize>56) { + std::fill(pending+pendingSize,pending+64,0u); + transform(pending); + pendingSize = 0; + } + std::fill(pending+pendingSize,pending+56,0u); + for(size_t i=0; i<8; ++i) + pending[56+i] = uint8_t(bitCount>>(56u-8u*i)); + transform(pending); + pendingSize = 0; + finalized = true; + } + + Digest ret = {}; + for(size_t i=0; i<8; ++i) { + ret[i*4+0] = uint8_t(state[i]>>24u); + ret[i*4+1] = uint8_t(state[i]>>16u); + ret[i*4+2] = uint8_t(state[i]>>8u); + ret[i*4+3] = uint8_t(state[i]); + } + return ret; + } + +MtSha256::Digest MtSha256::hash(const void* data, size_t size) { + MtSha256 hash; + hash.update(data,size); + return hash.finalize(); + } + +MtSha256::Digest MtSha256::hash(std::string_view data) { + return hash(data.data(),data.size()); + } + +void MtSha256::transform(const uint8_t block[64]) { + uint32_t w[64] = {}; + for(size_t i=0; i<16; ++i) + w[i] = loadBe(block+i*4); + for(size_t i=16; i<64; ++i) { + const uint32_t s0 = rotr(w[i-15],7)^rotr(w[i-15],18)^(w[i-15]>>3u); + const uint32_t s1 = rotr(w[i-2],17)^rotr(w[i-2],19)^(w[i-2]>>10u); + w[i] = w[i-16]+s0+w[i-7]+s1; + } + + uint32_t a=state[0], b=state[1], c=state[2], d=state[3]; + uint32_t e=state[4], f=state[5], g=state[6], h=state[7]; + for(size_t i=0; i<64; ++i) { + const uint32_t s1 = rotr(e,6)^rotr(e,11)^rotr(e,25); + const uint32_t choice= (e&f)^((~e)&g); + const uint32_t temp1 = h+s1+choice+k[i]+w[i]; + const uint32_t s0 = rotr(a,2)^rotr(a,13)^rotr(a,22); + const uint32_t major = (a&b)^(a&c)^(b&c); + const uint32_t temp2 = s0+major; + h=g; g=f; f=e; e=d+temp1; + d=c; c=b; b=a; a=temp1+temp2; + } + state[0]+=a; state[1]+=b; state[2]+=c; state[3]+=d; + state[4]+=e; state[5]+=f; state[6]+=g; state[7]+=h; + } diff --git a/Engine/gapi/metal/mtsha256.h b/Engine/gapi/metal/mtsha256.h new file mode 100644 index 00000000..f4fc54a6 --- /dev/null +++ b/Engine/gapi/metal/mtsha256.h @@ -0,0 +1,35 @@ +#pragma once + +#include +#include +#include +#include + +namespace Tempest { +namespace Detail { + +class MtSha256 final { + public: + using Digest = std::array; + + MtSha256(); + + void update(const void* data, size_t size); + void update(std::string_view data); + Digest finalize(); + + static Digest hash(const void* data, size_t size); + static Digest hash(std::string_view data); + + private: + void transform(const uint8_t block[64]); + + uint32_t state[8] = {}; + uint8_t pending[64] = {}; + uint64_t byteCount = 0; + size_t pendingSize = 0; + bool finalized = false; + }; + +} +} diff --git a/Engine/gapi/metal/mtshader.cpp b/Engine/gapi/metal/mtshader.cpp index 87b06f6a..b281f33c 100644 --- a/Engine/gapi/metal/mtshader.cpp +++ b/Engine/gapi/metal/mtshader.cpp @@ -6,6 +6,7 @@ #include #include "mtdevice.h" +#include "mtprecompiledlibrary.h" #include "gapi/shaderreflection.h" #include "thirdparty/spirv_cross/spirv_msl.hpp" @@ -20,16 +21,40 @@ static uint32_t spvVersion(MTL::LanguageVersion v) { return spirv_cross::CompilerMSL::Options::make_msl_version(major,minor,0); } +static MetalApi::PrecompiledShaderStage shaderStage(ShaderReflection::Stage stage) { + switch(stage) { + case ShaderReflection::Stage::Vertex: return MetalApi::PrecompiledShaderStage::Vertex; + case ShaderReflection::Stage::Control: return MetalApi::PrecompiledShaderStage::Control; + case ShaderReflection::Stage::Evaluate: return MetalApi::PrecompiledShaderStage::Evaluate; + case ShaderReflection::Stage::Geometry: return MetalApi::PrecompiledShaderStage::Geometry; + case ShaderReflection::Stage::Fragment: return MetalApi::PrecompiledShaderStage::Fragment; + case ShaderReflection::Stage::Compute: return MetalApi::PrecompiledShaderStage::Compute; + case ShaderReflection::Stage::Task: return MetalApi::PrecompiledShaderStage::Task; + case ShaderReflection::Stage::Mesh: return MetalApi::PrecompiledShaderStage::Mesh; + case ShaderReflection::Stage::None: break; + } + return MetalApi::PrecompiledShaderStage::Compute; + } + MtShader::MtShader(MtDevice& dev, const void* source, size_t srcSize) : Shader(source, srcSize) { auto pool = NsPtr::init(); spirv_cross::CompilerMSL::Options optMSL; + MetalApi::PrecompiledShaderProfile precompiledProfile; #if defined(__OSX__) optMSL.platform = spirv_cross::CompilerMSL::Options::macOS; + precompiledProfile.platform = MetalApi::PrecompiledPlatform::MacOS; #else optMSL.platform = spirv_cross::CompilerMSL::Options::iOS; +#if defined(TARGET_OS_SIMULATOR) && TARGET_OS_SIMULATOR + precompiledProfile.platform = MetalApi::PrecompiledPlatform::IOSSimulator; +#else + precompiledProfile.platform = MetalApi::PrecompiledPlatform::IOSDevice; +#endif #endif optMSL.buffer_size_buffer_index = MSL_BUFFER_LENGTH; + precompiledProfile.stage = shaderStage(stage); + precompiledProfile.bufferSizeBufferIndex = MSL_BUFFER_LENGTH; spirv_cross::CompilerGLSL::Options optGLSL; optGLSL.vertex.flip_vert_y = true; @@ -48,7 +73,8 @@ MtShader::MtShader(MtDevice& dev, const void* source, size_t srcSize) optMSL.readwrite_texture_fences = false; } - if(!dev.useNativeImageAtomic()) { + const bool nativeImageAtomics = dev.useNativeImageAtomic(); + if(!nativeImageAtomics) { const uint32_t align = dev.linearImageAlignment(); optMSL.r32ui_linear_texture_alignment = align; optMSL.r32ui_alignment_constant_id = 0; @@ -70,6 +96,14 @@ MtShader::MtShader(MtDevice& dev, const void* source, size_t srcSize) break; } } + precompiledProfile.mslVersion = optMSL.msl_version; + precompiledProfile.flipVertY = optGLSL.vertex.flip_vert_y; + precompiledProfile.argumentBuffersTier = uint8_t(optMSL.argument_buffers_tier); + precompiledProfile.runtimeArrayRichDescriptor = optMSL.runtime_array_rich_descriptor; + precompiledProfile.readWriteTextureFences = optMSL.readwrite_texture_fences; + precompiledProfile.nativeImageAtomics = nativeImageAtomics; + precompiledProfile.r32uiLinearTextureAlignment= optMSL.r32ui_linear_texture_alignment; + precompiledProfile.r32uiAlignmentConstantId = optMSL.r32ui_alignment_constant_id; comp.set_msl_options (optMSL ); comp.set_common_options(optGLSL); @@ -125,6 +159,12 @@ MtShader::MtShader(MtDevice& dev, const void* source, size_t srcSize) //Log::d(msl); + if(dev.precompiledLibraries!=nullptr) { + impl = dev.precompiledLibraries->find(msl,precompiledProfile); + if(impl!=nullptr) + return; + } + auto opt = NsPtr::init(); NS::Error* err = nullptr; auto str = NsPtr(NS::String::string(msl.c_str(),NS::UTF8StringEncoding)); diff --git a/Engine/gapi/metalapi.cpp b/Engine/gapi/metalapi.cpp index a7200563..66e291da 100644 --- a/Engine/gapi/metalapi.cpp +++ b/Engine/gapi/metalapi.cpp @@ -21,37 +21,220 @@ #include "gapi/metal/mtswapchain.h" #include "gapi/metal/mtaccelerationstructure.h" #include "gapi/metal/mtmetalfx.h" +#include "gapi/metal/mtsha256.h" #include +#include +#include #include +#include using namespace Tempest; using namespace Tempest::Detail; -MetalApi::MetalApi(ApiFlags f) - :MetalApi(f,Options{}) { +namespace { + +void appendU8(MtSha256& hash, uint8_t value) { + hash.update(&value,sizeof(value)); + } + +void appendU32(MtSha256& hash, uint32_t value) { + uint8_t data[4] = {uint8_t(value),uint8_t(value>>8u),uint8_t(value>>16u),uint8_t(value>>24u)}; + hash.update(data,sizeof(data)); + } + +void appendU64(MtSha256& hash, uint64_t value) { + uint8_t data[8] = {}; + for(size_t i=0; i<8; ++i) + data[i] = uint8_t(value>>(i*8u)); + hash.update(data,sizeof(data)); + } + +void appendString(MtSha256& hash, std::string_view value) { + appendU64(hash,value.size()); + hash.update(value); + } + +struct MetalApiOptionsRegistry final { + std::mutex mutex; + std::unordered_map> entries; + }; + +std::atomic publishedOptionsRegistry{nullptr}; + +MetalApiOptionsRegistry& ensureRegistry() { + // Only the opt-in Options constructor reaches this allocation. The process- + // lifetime registry keeps destruction of global MetalApi objects safe. + static auto* instance = []() { + auto* value = new MetalApiOptionsRegistry; + publishedOptionsRegistry.store(value,std::memory_order_release); + return value; + }(); + return *instance; + } + +MetalApiOptionsRegistry* tryGetRegistry() noexcept { + return publishedOptionsRegistry.load(std::memory_order_acquire); + } + +bool enableValidation(ApiFlags f) { + if((f & ApiFlags::Validation)!=ApiFlags::Validation) + return false; + setenv("METAL_DEVICE_WRAPPER_TYPE","1",1); + setenv("METAL_DEBUG_ERROR_MODE", "5",0); + setenv("METAL_ERROR_MODE", "5",0); + return true; } -MetalApi::MetalApi(ApiFlags f, const Options& options) - :swapchainBufferCount(options.swapchainBufferCount), - shaderModuleCacheSize(options.shaderModuleCacheSize), - swapchainRenderMode(options.swapchainRenderMode) { - if(swapchainBufferCount!=0 && swapchainBufferCount!=2 && swapchainBufferCount!=3) +void validateOptions(const MetalApi::Options& options) { + if(options.swapchainBufferCount!=0 && + options.swapchainBufferCount!=2 && options.swapchainBufferCount!=3) throw std::invalid_argument("Metal swapchain buffer count must be 0, 2, or 3"); - if(swapchainRenderMode!=SwapchainRenderMode::Copy && - swapchainRenderMode!=SwapchainRenderMode::Direct) + if(options.swapchainRenderMode!=MetalApi::SwapchainRenderMode::Copy && + options.swapchainRenderMode!=MetalApi::SwapchainRenderMode::Direct) throw std::invalid_argument("Unknown Metal swapchain render mode"); + } + +void registerOptions(const MetalApi* api, const MetalApi::Options& options) { + auto value = std::make_shared(options); + auto& registry = ensureRegistry(); + std::lock_guard guard(registry.mutex); + registry.entries.insert_or_assign(api,std::move(value)); + } + +void eraseRegistrationBestEffort(MetalApiOptionsRegistry& registry, + const MetalApi* api) noexcept { + try { + std::lock_guard guard(registry.mutex); + registry.entries.erase(api); + } + catch(...) { + } + } + +void copyOptionsRegistration(const MetalApi* destination, + const MetalApi* source) noexcept { + if(destination==source) + return; + auto* registry = tryGetRegistry(); + if(registry==nullptr) + return; + try { + std::lock_guard guard(registry->mutex); + const auto found = registry->entries.find(source); + if(found==registry->entries.end()) { + registry->entries.erase(destination); + return; + } + auto value = found->second; + registry->entries.insert_or_assign(destination,std::move(value)); + } + catch(...) { + eraseRegistrationBestEffort(*registry,destination); + } + } - if((f & ApiFlags::Validation)==ApiFlags::Validation) { - setenv("METAL_DEVICE_WRAPPER_TYPE","1",1); - setenv("METAL_DEBUG_ERROR_MODE", "5",0); - setenv("METAL_ERROR_MODE", "5",0); - validation = true; +std::shared_ptr registeredOptions(const MetalApi* api) noexcept { + auto* registry = tryGetRegistry(); + if(registry==nullptr) + return {}; + try { + std::lock_guard guard(registry->mutex); + const auto found = registry->entries.find(api); + if(found==registry->entries.end()) + return {}; + return found->second; } + catch(...) { + return {}; + } + } + +void unregisterOptions(const MetalApi* api) noexcept { + auto* registry = tryGetRegistry(); + if(registry==nullptr) + return; + eraseRegistrationBestEffort(*registry,api); + } + +} + +struct Tempest::Detail::MetalApiAbiProbe final { +#if defined(__clang__) +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Winvalid-offsetof" +#endif + static constexpr size_t validationOffset = __builtin_offsetof(MetalApi,validation); +#if defined(__clang__) +#pragma clang diagnostic pop +#endif + }; + +#if defined(__arm64__) || defined(__aarch64__) +static_assert(sizeof(MetalApi)==16,"MetalApi arm64 ABI size changed"); +static_assert(alignof(MetalApi)==8,"MetalApi arm64 ABI alignment changed"); +static_assert(Detail::MetalApiAbiProbe::validationOffset==8, + "MetalApi validation bool moved from its legacy arm64 offset"); +#endif + +MetalApi::MetalApi(ApiFlags f) + :validation(enableValidation(f)) { + } + +MetalApi::MetalApi(ApiFlags f, const Options& options) { + validateOptions(options); + validation = enableValidation(f); + registerOptions(this,options); + } + +MetalApi::MetalApi(const MetalApi& other) noexcept + :validation(other.validation) { + copyOptionsRegistration(this,&other); + } + +MetalApi& MetalApi::operator=(const MetalApi& other) noexcept { + if(this==&other) + return *this; + copyOptionsRegistration(this,&other); + validation = other.validation; + return *this; } MetalApi::~MetalApi() { + unregisterOptions(this); + } + +MetalApi::PrecompiledShaderKey MetalApi::precompiledShaderKey(std::string_view canonicalMsl, + const PrecompiledShaderProfile& profile) { + static constexpr std::string_view domain = "Tempest.Metal.PrecompiledShader"; + MtSha256 hash; + hash.update(domain); + appendU8(hash,0); + appendU32(hash,profile.schemaVersion); + appendU64(hash,canonicalMsl.size()); + hash.update(canonicalMsl); + appendU32(hash,profile.mslGeneratorVersion); + appendU8(hash,uint8_t(profile.platform)); + appendU8(hash,uint8_t(profile.stage)); + appendString(hash,profile.entryPoint); + appendU32(hash,profile.mslVersion); + appendU8(hash,profile.flipVertY ? 1 : 0); + appendU32(hash,profile.bufferSizeBufferIndex); + appendU8(hash,profile.argumentBuffersTier); + appendU8(hash,profile.runtimeArrayRichDescriptor ? 1 : 0); + appendU8(hash,profile.readWriteTextureFences ? 1 : 0); + appendU8(hash,profile.nativeImageAtomics ? 1 : 0); + appendU32(hash,profile.r32uiLinearTextureAlignment); + appendU32(hash,profile.r32uiAlignmentConstantId); + return hash.finalize(); + } + +MetalApi::PrecompiledLibraryHash MetalApi::precompiledLibraryHash(const void* data, + size_t size) { + if(data==nullptr && size!=0) + return {}; + return MtSha256::hash(data,size); } #if !defined(TEMPEST_BUILD_METALFX) @@ -93,14 +276,17 @@ std::vector MetalApi::devices() const { } AbstractGraphicsApi::Device* MetalApi::createDevice(std::string_view gpuName) { - return new MtDevice(gpuName,validation,shaderModuleCacheSize); + auto options = registeredOptions(this); + return new MtDevice(gpuName,validation,std::move(options)); } AbstractGraphicsApi::Swapchain *MetalApi::createSwapchain(SystemApi::Window *w, AbstractGraphicsApi::Device* d) { auto& dev = *reinterpret_cast(d); - return new MtSwapchain(dev,w,swapchainBufferCount, - swapchainRenderMode==SwapchainRenderMode::Direct); + auto options = registeredOptions(this); + const uint32_t bufferCount = options!=nullptr ? options->swapchainBufferCount : 0; + const auto renderMode = options!=nullptr ? options->swapchainRenderMode : SwapchainRenderMode::Copy; + return new MtSwapchain(dev,w,bufferCount,renderMode==SwapchainRenderMode::Direct); } AbstractGraphicsApi::PPipeline MetalApi::createPipeline(AbstractGraphicsApi::Device *d, diff --git a/Engine/gapi/metalapi.h b/Engine/gapi/metalapi.h index bea43c05..283b257c 100644 --- a/Engine/gapi/metalapi.h +++ b/Engine/gapi/metalapi.h @@ -2,8 +2,20 @@ #include +#include +#include +#include +#include +#include +#include +#include + namespace Tempest { +namespace Detail { +struct MetalApiAbiProbe; +} + class MetalApi : public AbstractGraphicsApi { public: enum class SwapchainRenderMode:uint8_t { @@ -11,18 +23,92 @@ class MetalApi : public AbstractGraphicsApi { Direct, }; + static constexpr uint32_t PrecompiledShaderSchemaVersion = 1; + static constexpr uint32_t MslGeneratorVersion = 1; + + enum class PrecompiledPlatform : uint8_t { + MacOS, + IOSDevice, + IOSSimulator, + }; + + enum class PrecompiledShaderStage : uint8_t { + Vertex, + Control, + Evaluate, + Geometry, + Fragment, + Compute, + Task, + Mesh, + }; + + using PrecompiledShaderKey = std::array; + using PrecompiledLibraryHash = std::array; + + /** + * Exact input profile used when SPIR-V was converted to canonical MSL. + * MslGeneratorVersion pins every SPIRV-Cross default not listed here; it + * must be incremented whenever those defaults or the serialization change. + */ + struct PrecompiledShaderProfile { + uint32_t schemaVersion = PrecompiledShaderSchemaVersion; + uint32_t mslGeneratorVersion = MslGeneratorVersion; + PrecompiledPlatform platform = PrecompiledPlatform::MacOS; + PrecompiledShaderStage stage = PrecompiledShaderStage::Compute; + std::string entryPoint = "main0"; + uint32_t mslVersion = 0; + bool flipVertY = true; + uint32_t bufferSizeBufferIndex = 29; + uint8_t argumentBuffersTier = 0; + bool runtimeArrayRichDescriptor = false; + bool readWriteTextureFences = true; + bool nativeImageAtomics = true; + uint32_t r32uiLinearTextureAlignment = 0; + uint32_t r32uiAlignmentConstantId = 0; + }; + + struct PrecompiledShader { + PrecompiledShaderProfile profile; + PrecompiledShaderKey key = {}; + }; + + struct PrecompiledLibrary { + // One target-specific metallib. MetalApi copies these bytes at construction. + std::vector data; + // Full SHA-256 of data. It is verified before Metal sees the library. + PrecompiledLibraryHash dataHash = {}; + std::vector shaders; + }; + struct Options { uint32_t swapchainBufferCount = 0; // Maximum number of compiled Metal shader modules kept per device. // Zero disables caching. size_t shaderModuleCacheSize = 0; - SwapchainRenderMode swapchainRenderMode = SwapchainRenderMode::Copy; + SwapchainRenderMode swapchainRenderMode = SwapchainRenderMode::Copy; + // Empty by default: runtime MSL compilation remains the only code path. + std::vector precompiledLibraries; }; explicit MetalApi(ApiFlags f=ApiFlags::NoFlags); MetalApi(ApiFlags f, const Options& options); + MetalApi(const MetalApi& other) noexcept; + MetalApi& operator=(const MetalApi& other) noexcept; ~MetalApi(); + /** + * Computes the canonical, full SHA-256 identity of an MSL source/profile + * pair. The serialized profile includes the schema, target, entry point, + * shader stage and every Metal code-generation option above. An artifact + * is used only when this key and the complete runtime profile agree. + */ + static PrecompiledShaderKey precompiledShaderKey(std::string_view canonicalMsl, + const PrecompiledShaderProfile& profile); + + /** Computes the full SHA-256 identity of a target-specific metallib. */ + static PrecompiledLibraryHash precompiledLibraryHash(const void* data, size_t size); + std::vector devices() const override; protected: @@ -61,10 +147,9 @@ class MetalApi : public AbstractGraphicsApi { void getCaps(Device *d, Props& caps) override; private: - bool validation = false; - uint32_t swapchainBufferCount = 0; - size_t shaderModuleCacheSize = 0; - SwapchainRenderMode swapchainRenderMode = SwapchainRenderMode::Copy; + bool validation = false; + + friend struct Detail::MetalApiAbiProbe; }; } diff --git a/Tests/tests/CMakeLists.txt b/Tests/tests/CMakeLists.txt index c1a97690..a347d4fe 100644 --- a/Tests/tests/CMakeLists.txt +++ b/Tests/tests/CMakeLists.txt @@ -151,6 +151,14 @@ target_sources(${PROJECT_NAME} PRIVATE set(BUILD_SHARED_LIBS ${BUILD_SHARED_MOLTEN_TEMPEST}) add_subdirectory("${CMAKE_SOURCE_DIR}/../../Engine" build) +if(APPLE AND TEMPEST_BUILD_METAL AND NOT IOS) + target_compile_definitions(${PROJECT_NAME} PRIVATE TEMPEST_TEST_METAL_PRECOMPILED=1) + target_include_directories(${PROJECT_NAME} PRIVATE + "${CMAKE_SOURCE_DIR}/../../Engine" + "${CMAKE_SOURCE_DIR}/../../Engine/thirdparty/metal-cpp") + target_link_libraries(${PROJECT_NAME} "-framework Metal" "-framework Foundation") +endif() + if(UNIX) target_link_libraries(${PROJECT_NAME} -lpthread) endif() diff --git a/Tests/tests/gapi/metal_test.cpp b/Tests/tests/gapi/metal_test.cpp index 3aeea3cb..70df5706 100644 --- a/Tests/tests/gapi/metal_test.cpp +++ b/Tests/tests/gapi/metal_test.cpp @@ -25,6 +25,24 @@ #include "gapi_test_common.h" +#if defined(__OSX__) && defined(TEMPEST_TEST_METAL_PRECOMPILED) +#include "gapi/metal/mtdevice.h" +#include "gapi/metal/mtprecompiledlibrary.h" +#include "gapi/metal/mtsha256.h" + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#endif + using namespace testing; using namespace Tempest; @@ -152,7 +170,416 @@ void swapchainSmoke(MetalApi::SwapchainRenderMode mode) { #endif } +#if defined(__OSX__) && defined(TEMPEST_TEST_METAL_PRECOMPILED) +namespace { + +std::string hexDigest(const Detail::MtSha256::Digest& digest) { + std::ostringstream out; + out << std::hex << std::setfill('0'); + for(uint8_t value:digest) + out << std::setw(2) << uint32_t(value); + return out.str(); + } + +MetalApi::PrecompiledShaderProfile testProfile() { + MetalApi::PrecompiledShaderProfile ret; + ret.platform = MetalApi::PrecompiledPlatform::MacOS; + ret.stage = MetalApi::PrecompiledShaderStage::Compute; + ret.entryPoint = "main0"; + ret.mslVersion = 20400; + return ret; + } + +struct TemporaryMetallib final { + explicit TemporaryMetallib(uint32_t value=7, bool alternate=false) { + char path[] = "/tmp/tempest-metal-precompiled-XXXXXX"; + const char* created = mkdtemp(path); + if(created==nullptr) + return; + directory = created; + + canonicalMsl = + "#include \n" + "using namespace metal;\n" + "kernel void main0(device uint* output [[buffer(0)]], " + "uint id [[thread_position_in_grid]]) { output[id] = "+ + std::to_string(value)+"; }\n"; + if(alternate) { + canonicalMsl += + "kernel void alternate(device uint* output [[buffer(0)]], " + "uint id [[thread_position_in_grid]]) { output[id] = "+ + std::to_string(value+1)+"; }\n"; + } + + const auto source = directory/"fixture.metal"; + const auto air = directory/"fixture.air"; + const auto library= directory/"fixture.metallib"; + { + std::ofstream out(source); + out << canonicalMsl; + } + + const std::string compile = "xcrun -sdk macosx metal -std=macos-metal2.4 -c \""+ + source.string()+"\" -o \""+air.string()+"\""; + const std::string link = "xcrun -sdk macosx metallib \""+air.string()+ + "\" -o \""+library.string()+"\""; + if(std::system(compile.c_str())!=0 || std::system(link.c_str())!=0) + return; + + std::ifstream in(library,std::ios::binary); + data.assign(std::istreambuf_iterator(in),std::istreambuf_iterator()); + } + + ~TemporaryMetallib() { + if(!directory.empty()) + std::filesystem::remove_all(directory); + } + + bool valid() const { + return !data.empty(); + } + + std::filesystem::path directory; + std::string canonicalMsl; + std::vector data; + }; + +MetalApi::Options optionsFor(const TemporaryMetallib& fixture, + MetalApi::PrecompiledShaderProfile profile = testProfile()) { + MetalApi::PrecompiledShader shader; + shader.profile = std::move(profile); + shader.key = MetalApi::precompiledShaderKey(fixture.canonicalMsl,shader.profile); + + MetalApi::PrecompiledLibrary library; + library.data = fixture.data; + library.dataHash = MetalApi::precompiledLibraryHash(library.data.data(),library.data.size()); + library.shaders.push_back(std::move(shader)); + + MetalApi::Options ret; + ret.precompiledLibraries.push_back(std::move(library)); + return ret; + } + +Detail::NsPtr testMetalDevice() { + setenv("METAL_DEVICE_WRAPPER_TYPE","1",1); + setenv("METAL_DEBUG_ERROR_MODE", "5",0); + setenv("METAL_ERROR_MODE", "5",0); + return Detail::NsPtr(MTL::CreateSystemDefaultDevice()); + } + +std::vector readBinary(const char* filename) { + std::ifstream in(filename,std::ios::binary); + return std::vector(std::istreambuf_iterator(in), + std::istreambuf_iterator()); + } + +class InspectableMetalApi final : public MetalApi { + public: + using MetalApi::MetalApi; + + InspectableMetalApi(const InspectableMetalApi&) noexcept = default; + InspectableMetalApi& operator=(const InspectableMetalApi&) noexcept = default; + + std::unique_ptr createTestDevice() { + return std::unique_ptr( + static_cast(createDevice(std::string_view{}))); + } + }; + +} +#endif + +TEST(MetalApi,PrecompiledSha256Vectors) { +#if defined(__OSX__) && defined(TEMPEST_TEST_METAL_PRECOMPILED) + EXPECT_EQ(hexDigest(Detail::MtSha256::hash("")), + "e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855"); + EXPECT_EQ(hexDigest(Detail::MtSha256::hash("abc")), + "ba7816bf8f01cfea414140de5dae2223b00361a396177a9cb410ff61f20015ad"); + EXPECT_EQ(hexDigest(Detail::MtSha256::hash( + "abcdbcdecdefdefgefghfghighijhijkijkljklmklmnlmnomnopnopq")), + "248d6a61d20638b8e5c026930c3e6039a33ce45964ff2167f6ecedd419db06c1"); + + Detail::MtSha256 incremental; + const std::string block(1000,'a'); + for(size_t i=0; i<1000; ++i) + incremental.update(block); + EXPECT_EQ(hexDigest(incremental.finalize()), + "cdc76e5c9914fb9281a1c7e284d73e67f1809a48a497200e046d39ccc7112cd0"); + EXPECT_EQ(hexDigest(MetalApi::precompiledLibraryHash("abc",3)), + "ba7816bf8f01cfea414140de5dae2223b00361a396177a9cb410ff61f20015ad"); + EXPECT_EQ(sizeof(MetalApi),2*sizeof(void*)); +#endif + } + +TEST(MetalApi,PrecompiledKeyCoversProfile) { +#if defined(__OSX__) && defined(TEMPEST_TEST_METAL_PRECOMPILED) + const std::string msl = "kernel void main0() {}"; + const auto base = testProfile(); + const auto key = MetalApi::precompiledShaderKey(msl,base); + auto differs = [&](auto mutate) { + auto profile = base; + mutate(profile); + EXPECT_NE(MetalApi::precompiledShaderKey(msl,profile),key); + }; + + EXPECT_NE(MetalApi::precompiledShaderKey(msl+"\n",base),key); + differs([](auto& p){ ++p.schemaVersion; }); + differs([](auto& p){ ++p.mslGeneratorVersion; }); + differs([](auto& p){ p.platform=MetalApi::PrecompiledPlatform::IOSDevice; }); + differs([](auto& p){ p.stage=MetalApi::PrecompiledShaderStage::Fragment; }); + differs([](auto& p){ p.entryPoint="alternate"; }); + differs([](auto& p){ ++p.mslVersion; }); + differs([](auto& p){ p.flipVertY=!p.flipVertY; }); + differs([](auto& p){ ++p.bufferSizeBufferIndex; }); + differs([](auto& p){ ++p.argumentBuffersTier; }); + differs([](auto& p){ p.runtimeArrayRichDescriptor=!p.runtimeArrayRichDescriptor; }); + differs([](auto& p){ p.readWriteTextureFences=!p.readWriteTextureFences; }); + differs([](auto& p){ p.nativeImageAtomics=!p.nativeImageAtomics; }); + differs([](auto& p){ ++p.r32uiLinearTextureAlignment; }); + differs([](auto& p){ ++p.r32uiAlignmentConstantId; }); +#endif + } +TEST(MetalApi,PrecompiledLibraryValidHitAndLifetime) { +#if defined(__OSX__) && defined(TEMPEST_TEST_METAL_PRECOMPILED) + TemporaryMetallib fixture; + ASSERT_TRUE(fixture.valid()); + auto device = testMetalDevice(); + ASSERT_NE(device,nullptr); + + std::unique_ptr loader; + { + auto temporaryOptions = optionsFor(fixture); + loader = std::make_unique(*device,temporaryOptions); + temporaryOptions.precompiledLibraries.clear(); + } + + auto function = loader->find(fixture.canonicalMsl,testProfile()); + ASSERT_NE(function,nullptr); + EXPECT_EQ(function->functionType(),MTL::FunctionTypeKernel); + NS::Error* error = nullptr; + auto pipeline = Detail::NsPtr( + device->newComputePipelineState(function.get(),&error)); + EXPECT_EQ(error,nullptr); + EXPECT_NE(pipeline,nullptr); +#endif + } + +TEST(MetalApi,PrecompiledOptionsCopyAndLegacyFallback) { +#if defined(__OSX__) && defined(TEMPEST_TEST_METAL_PRECOMPILED) + static_assert(std::is_copy_constructible_v,"MetalApi lost legacy copying"); + static_assert(std::is_copy_assignable_v,"MetalApi lost legacy assignment"); + static_assert(std::is_nothrow_copy_constructible_v,"MetalApi copy must remain noexcept"); + static_assert(std::is_nothrow_copy_assignable_v,"MetalApi assignment must remain noexcept"); + static_assert(std::is_nothrow_move_constructible_v,"MetalApi move-via-copy must remain noexcept"); +#if defined(__arm64__) || defined(__aarch64__) + static_assert(sizeof(MetalApi)==16,"MetalApi arm64 ABI size changed"); + static_assert(alignof(MetalApi)==8,"MetalApi arm64 ABI alignment changed"); +#endif + + TemporaryMetallib fixture; + ASSERT_TRUE(fixture.valid()); + + auto sourceOptions = optionsFor(fixture); + sourceOptions.swapchainBufferCount = 3; + sourceOptions.shaderModuleCacheSize = 2; + sourceOptions.swapchainRenderMode = MetalApi::SwapchainRenderMode::Direct; + InspectableMetalApi source(ApiFlags::Validation,sourceOptions); + InspectableMetalApi copied(source); + InspectableMetalApi assigned; + assigned = source; + + auto expectRegisteredHit = [&](InspectableMetalApi& api) { + auto device = api.createTestDevice(); + ASSERT_NE(device,nullptr); + ASSERT_NE(device->precompiledOptions,nullptr); + ASSERT_NE(device->precompiledLibraries,nullptr); + EXPECT_TRUE(device->validation); + EXPECT_EQ(device->precompiledOptions->swapchainBufferCount,3u); + EXPECT_EQ(device->precompiledOptions->shaderModuleCacheSize,2u); + EXPECT_EQ(device->precompiledOptions->swapchainRenderMode, + MetalApi::SwapchainRenderMode::Direct); + EXPECT_NE(device->precompiledLibraries->find(fixture.canonicalMsl,testProfile()),nullptr); + }; + expectRegisteredHit(copied); + expectRegisteredHit(assigned); + + std::unique_ptr survivingDevice; + { + InspectableMetalApi scoped(ApiFlags::NoFlags,optionsFor(fixture)); + survivingDevice = scoped.createTestDevice(); + } + ASSERT_NE(survivingDevice,nullptr); + ASSERT_NE(survivingDevice->precompiledOptions,nullptr); + ASSERT_NE(survivingDevice->precompiledLibraries,nullptr); + EXPECT_NE(survivingDevice->precompiledLibraries->find(fixture.canonicalMsl,testProfile()),nullptr); + + InspectableMetalApi legacy; + InspectableMetalApi legacyCopy(legacy); + InspectableMetalApi overwritten(ApiFlags::NoFlags,optionsFor(fixture)); + overwritten = legacy; + auto expectLegacyFallback = [](InspectableMetalApi& api) { + auto device = api.createTestDevice(); + ASSERT_NE(device,nullptr); + EXPECT_EQ(device->precompiledOptions,nullptr); + EXPECT_EQ(device->precompiledLibraries,nullptr); + }; + expectLegacyFallback(legacyCopy); + expectLegacyFallback(overwritten); +#endif + } + +TEST(MetalApi,PrecompiledLibraryMismatchAndDuplicatesFailClosed) { +#if defined(__OSX__) && defined(TEMPEST_TEST_METAL_PRECOMPILED) + TemporaryMetallib fixture; + ASSERT_TRUE(fixture.valid()); + auto device = testMetalDevice(); + ASSERT_NE(device,nullptr); + + auto options = optionsFor(fixture); + Detail::MtPrecompiledLibraries loader(*device,options); + auto runtime = testProfile(); + + auto expectMiss = [&](auto mutate) { + auto mismatch = runtime; + mutate(mismatch); + EXPECT_EQ(loader.find(fixture.canonicalMsl,mismatch),nullptr); + }; + EXPECT_EQ(loader.find(fixture.canonicalMsl+" ",runtime),nullptr); + expectMiss([](auto& p){ ++p.schemaVersion; }); + expectMiss([](auto& p){ ++p.mslGeneratorVersion; }); + expectMiss([](auto& p){ p.platform=MetalApi::PrecompiledPlatform::IOSDevice; }); + expectMiss([](auto& p){ p.stage=MetalApi::PrecompiledShaderStage::Fragment; }); + expectMiss([](auto& p){ p.entryPoint="alternate"; }); + expectMiss([](auto& p){ ++p.mslVersion; }); + expectMiss([](auto& p){ p.flipVertY=!p.flipVertY; }); + expectMiss([](auto& p){ ++p.bufferSizeBufferIndex; }); + expectMiss([](auto& p){ ++p.argumentBuffersTier; }); + expectMiss([](auto& p){ p.runtimeArrayRichDescriptor=!p.runtimeArrayRichDescriptor; }); + expectMiss([](auto& p){ p.readWriteTextureFences=!p.readWriteTextureFences; }); + expectMiss([](auto& p){ p.nativeImageAtomics=!p.nativeImageAtomics; }); + expectMiss([](auto& p){ ++p.r32uiLinearTextureAlignment; }); + expectMiss([](auto& p){ ++p.r32uiAlignmentConstantId; }); + + auto badKey = optionsFor(fixture); + ++badKey.precompiledLibraries[0].shaders[0].key[0]; + Detail::MtPrecompiledLibraries badKeyLoader(*device,badKey); + EXPECT_EQ(badKeyLoader.find(fixture.canonicalMsl,runtime),nullptr); + + auto duplicate = optionsFor(fixture); + duplicate.precompiledLibraries[0].shaders.push_back( + duplicate.precompiledLibraries[0].shaders[0]); + Detail::MtPrecompiledLibraries duplicateLoader(*device,duplicate); + EXPECT_EQ(duplicateLoader.find(fixture.canonicalMsl,runtime),nullptr); + + auto wrongStage = optionsFor(fixture); + wrongStage.precompiledLibraries[0].shaders[0].profile.stage = + MetalApi::PrecompiledShaderStage::Fragment; + auto& stageEntry = wrongStage.precompiledLibraries[0].shaders[0]; + stageEntry.key = MetalApi::precompiledShaderKey(fixture.canonicalMsl,stageEntry.profile); + Detail::MtPrecompiledLibraries wrongStageLoader(*device,wrongStage); + EXPECT_EQ(wrongStageLoader.find(fixture.canonicalMsl,stageEntry.profile),nullptr); + + auto missingEntry = optionsFor(fixture); + missingEntry.precompiledLibraries[0].shaders[0].profile.entryPoint = "missing_entry"; + auto& missing = missingEntry.precompiledLibraries[0].shaders[0]; + missing.key = MetalApi::precompiledShaderKey(fixture.canonicalMsl,missing.profile); + Detail::MtPrecompiledLibraries missingEntryLoader(*device,missingEntry); + EXPECT_EQ(missingEntryLoader.find(fixture.canonicalMsl,missing.profile),nullptr); + + TemporaryMetallib alternateFixture(7,true); + ASSERT_TRUE(alternateFixture.valid()); + auto alternateProfile = testProfile(); + alternateProfile.entryPoint = "alternate"; + auto alternateOptions = optionsFor(alternateFixture,alternateProfile); + Detail::MtPrecompiledLibraries alternateLoader(*device,alternateOptions); + EXPECT_NE(alternateLoader.find(alternateFixture.canonicalMsl,alternateProfile),nullptr); + EXPECT_EQ(alternateLoader.find(alternateFixture.canonicalMsl,testProfile()),nullptr); +#endif + } + +TEST(MetalApi,PrecompiledLibraryBytesHashRejectsValidSubstitution) { +#if defined(__OSX__) && defined(TEMPEST_TEST_METAL_PRECOMPILED) + TemporaryMetallib expected(7); + TemporaryMetallib replacement(9); + ASSERT_TRUE(expected.valid()); + ASSERT_TRUE(replacement.valid()); + auto device = testMetalDevice(); + ASSERT_NE(device,nullptr); + + auto replacementOptions = optionsFor(replacement); + Detail::MtPrecompiledLibraries replacementLoader(*device,replacementOptions); + ASSERT_NE(replacementLoader.find(replacement.canonicalMsl,testProfile()),nullptr); + + auto substituted = optionsFor(expected); + substituted.precompiledLibraries[0].data = replacement.data; + Detail::MtPrecompiledLibraries substitutedLoader(*device,substituted); + EXPECT_EQ(substitutedLoader.find(expected.canonicalMsl,testProfile()),nullptr); +#endif + } + +TEST(MetalApi,PrecompiledLibraryCorruptionFallsBack) { +#if defined(__OSX__) && defined(TEMPEST_TEST_METAL_PRECOMPILED) + TemporaryMetallib fixture; + ASSERT_TRUE(fixture.valid()); + auto options = optionsFor(fixture); + options.precompiledLibraries[0].data.assign({0x01,0x02,0x03,0x04}); + options.precompiledLibraries[0].dataHash = MetalApi::precompiledLibraryHash( + options.precompiledLibraries[0].data.data(), + options.precompiledLibraries[0].data.size()); + + MetalApi api(ApiFlags::Validation,std::move(options)); + Device device(api); + EXPECT_NO_THROW({ + auto shader = device.shader("shader/simple_test.comp.sprv"); + auto pipeline = device.pipeline(shader); + }); +#endif + } + +TEST(MetalApi,PrecompiledLibraryConcurrentLookup) { +#if defined(__OSX__) && defined(TEMPEST_TEST_METAL_PRECOMPILED) + TemporaryMetallib fixture; + ASSERT_TRUE(fixture.valid()); + auto device = testMetalDevice(); + ASSERT_NE(device,nullptr); + auto options = optionsFor(fixture); + Detail::MtPrecompiledLibraries loader(*device,options); + + std::atomic_uint32_t hits = 0; + std::vector threads; + for(size_t i=0; i<8; ++i) { + threads.emplace_back([&]() { + for(size_t r=0; r<16; ++r) { + auto function = loader.find(fixture.canonicalMsl,testProfile()); + if(function!=nullptr) + hits.fetch_add(1,std::memory_order_relaxed); + } + }); + } + for(auto& thread:threads) + thread.join(); + EXPECT_EQ(hits.load(),8u*16u); + + const auto spirv = readBinary("shader/simple_test.comp.sprv"); + ASSERT_FALSE(spirv.empty()); + MetalApi api(ApiFlags::Validation,optionsFor(fixture)); + Device tempestDevice(api); + std::atomic_uint32_t shaders = 0; + threads.clear(); + for(size_t i=0; i<4; ++i) { + threads.emplace_back([&]() { + for(size_t r=0; r<4; ++r) { + auto shader = tempestDevice.shader(spirv.data(),spirv.size()); + shaders.fetch_add(1,std::memory_order_relaxed); + } + }); + } + for(auto& thread:threads) + thread.join(); + EXPECT_EQ(shaders.load(),4u*4u); +#endif + } TEST(MetalApi,MetalApi) { #if defined(__OSX__) GapiTestCommon::init(); From 1160a932366db6ac870d5ba7eaf47e0c19e483af Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 16:47:54 +0100 Subject: [PATCH 20/25] Respect host iOS orientation policy --- Engine/system/api/iosapi.mm | 7 ------- 1 file changed, 7 deletions(-) diff --git a/Engine/system/api/iosapi.mm b/Engine/system/api/iosapi.mm index 8f06f7b8..8d4f17df 100644 --- a/Engine/system/api/iosapi.mm +++ b/Engine/system/api/iosapi.mm @@ -501,13 +501,6 @@ - (UISceneConfiguration *)application:(UIApplication *)application return configuration; } -- (UIInterfaceOrientationMask)application:(UIApplication *)application - supportedInterfaceOrientationsForWindow:(UIWindow *)window { - (void)application; - (void)window; - return UIInterfaceOrientationMaskAll; - } - - (void)applicationWillTerminate:(UIApplication *)application { (void)application; ++lifecycleGeneration; From 179766f721bd7cd01d6b8979e1490476d80648ad Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 16:52:07 +0100 Subject: [PATCH 21/25] Respect host iOS scene configuration --- Engine/system/api/iosapi.mm | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/Engine/system/api/iosapi.mm b/Engine/system/api/iosapi.mm index 8d4f17df..c34bab7d 100644 --- a/Engine/system/api/iosapi.mm +++ b/Engine/system/api/iosapi.mm @@ -494,9 +494,10 @@ - (UISceneConfiguration *)application:(UIApplication *)application options:(UISceneConnectionOptions *)options { (void)application; (void)options; - UISceneConfiguration* configuration = - [UISceneConfiguration configurationWithName:@"Tempest Scene" - sessionRole:connectingSceneSession.role]; + UISceneConfiguration* configuration = connectingSceneSession.configuration; + if(configuration==nil) + configuration = [UISceneConfiguration configurationWithName:nil + sessionRole:connectingSceneSession.role]; configuration.delegateClass = [SceneDelegate class]; return configuration; } From ce76038209032b68e4b4bccaf8494616360d63de Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 17:53:06 +0100 Subject: [PATCH 22/25] Fix scaler tests on MSVC --- Tests/tests/spatialscaler_test.cpp | 4 ++-- Tests/tests/temporalscaler_test.cpp | 8 ++++---- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/Tests/tests/spatialscaler_test.cpp b/Tests/tests/spatialscaler_test.cpp index 8e8049de..80268f9c 100644 --- a/Tests/tests/spatialscaler_test.cpp +++ b/Tests/tests/spatialscaler_test.cpp @@ -322,7 +322,7 @@ TEST(SpatialScaler, EncoderUsesPublicResources) { EXPECT_EQ(stats.scalerEncoded,1); EXPECT_EQ(stats.renderingBegun,1); EXPECT_EQ(stats.renderingEnded,1); - EXPECT_NE(stats.scalerInput,nullptr); - EXPECT_NE(stats.scalerOutput,nullptr); + EXPECT_TRUE(stats.scalerInput!=nullptr); + EXPECT_TRUE(stats.scalerOutput!=nullptr); EXPECT_NE(stats.scalerInput,stats.scalerOutput); } diff --git a/Tests/tests/temporalscaler_test.cpp b/Tests/tests/temporalscaler_test.cpp index 85759c3c..44789485 100644 --- a/Tests/tests/temporalscaler_test.cpp +++ b/Tests/tests/temporalscaler_test.cpp @@ -409,10 +409,10 @@ TEST(TemporalScaler, EncoderForwardsResourcesAndArgsAndEndsRendering) { auto* recordedDepth = dynamic_cast(stats.scalerDepth); auto* recordedMotion = dynamic_cast(stats.scalerMotion); auto* recordedOutput = dynamic_cast(stats.scalerOutput); - ASSERT_NE(recordedInput,nullptr); - ASSERT_NE(recordedDepth,nullptr); - ASSERT_NE(recordedMotion,nullptr); - ASSERT_NE(recordedOutput,nullptr); + ASSERT_TRUE(recordedInput!=nullptr); + ASSERT_TRUE(recordedDepth!=nullptr); + ASSERT_TRUE(recordedMotion!=nullptr); + ASSERT_TRUE(recordedOutput!=nullptr); EXPECT_EQ(recordedInput->syncId(),NonUniqResId(1)); EXPECT_EQ(recordedDepth->syncId(),NonUniqResId(2)); EXPECT_EQ(recordedMotion->syncId(),NonUniqResId(4)); From cd0cf6bdb04e7beb3a6ef1e45a1b39b54589aba4 Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 18:43:11 +0100 Subject: [PATCH 23/25] Use shared host-visible Metal storage on iOS --- Engine/gapi/metal/mtaccelerationstructure.cpp | 6 +++--- Engine/gapi/metal/mtdevice.h | 17 +++++++++++++++++ Engine/gapi/metal/mttexture.cpp | 4 ++-- Engine/gapi/metalapi.cpp | 10 ++-------- 4 files changed, 24 insertions(+), 13 deletions(-) diff --git a/Engine/gapi/metal/mtaccelerationstructure.cpp b/Engine/gapi/metal/mtaccelerationstructure.cpp index 186bcbc1..bc4225ca 100644 --- a/Engine/gapi/metal/mtaccelerationstructure.cpp +++ b/Engine/gapi/metal/mtaccelerationstructure.cpp @@ -100,7 +100,7 @@ MtTopAccelerationStructure::MtTopAccelerationStructure(MtDevice& dx, const RtIns :owner(dx) { auto pool = NsPtr::init(); instances = NsPtr(dx.impl->newBuffer(sizeof(MTL::AccelerationStructureUserIDInstanceDescriptor)*asSize, - MTL::ResourceStorageModeManaged)); + hostVisibleResourceOptions(*dx.impl))); if(instances==nullptr) throw std::system_error(GraphicsErrc::OutOfVideoMemory); @@ -128,7 +128,8 @@ MtTopAccelerationStructure::MtTopAccelerationStructure(MtDevice& dx, const RtIns blas.push_back(ax->impl.get()); } } - instances->didModifyRange(NS::Range(0,instances->length())); + if(instances->storageMode()==MTL::StorageModeManaged) + instances->didModifyRange(NS::Range(0,instances->length())); auto asArray = NsPtr(NS::Array::array(reinterpret_cast(blas.data()), blas.size())); auto desc = NsPtr::init(); @@ -196,4 +197,3 @@ void MtTopAccelerationStructure::implUseResource(MTL::RenderCommandEncoder& cmd, } #endif - diff --git a/Engine/gapi/metal/mtdevice.h b/Engine/gapi/metal/mtdevice.h index 8f662cb4..8404372d 100644 --- a/Engine/gapi/metal/mtdevice.h +++ b/Engine/gapi/metal/mtdevice.h @@ -23,6 +23,23 @@ namespace Detail { class MtPrecompiledLibraries; +inline MTL::StorageMode hostVisibleStorageMode(MTL::Device& dev) { +#ifdef __IOS__ + // Managed storage is unavailable on iOS. The Simulator can still report + // hasUnifiedMemory()==false, so platform support is authoritative here. + (void)dev; + return MTL::StorageModeShared; +#else + return dev.hasUnifiedMemory() ? MTL::StorageModeShared : MTL::StorageModeManaged; +#endif + } + +inline MTL::ResourceOptions hostVisibleResourceOptions(MTL::Device& dev) { + return hostVisibleStorageMode(dev)==MTL::StorageModeShared + ? MTL::ResourceStorageModeShared + : MTL::ResourceStorageModeManaged; + } + inline MTL::PixelFormat nativeFormat(TextureFormat frm) { switch(frm) { case Undefined: diff --git a/Engine/gapi/metal/mttexture.cpp b/Engine/gapi/metal/mttexture.cpp index 3172ea75..ce204219 100644 --- a/Engine/gapi/metal/mttexture.cpp +++ b/Engine/gapi/metal/mttexture.cpp @@ -45,7 +45,7 @@ MtTexture::MtTexture(MtDevice& d, const uint32_t w, const uint32_t h, const uint MtTexture::MtTexture(MtDevice& dev, const Pixmap& pm, uint32_t mipCnt, TextureFormat frm) :dev(dev), mipCnt(mipCnt) { const uint32_t smip = (isCompressedFormat(frm) ? mipCnt : 1); - const MTL::StorageMode smode = dev.impl->hasUnifiedMemory() ? MTL::StorageModeShared : MTL::StorageModeManaged; + const MTL::StorageMode smode = hostVisibleStorageMode(*dev.impl); NsPtr stage = alloc(frm,pm.w(),pm.h(),0,smip,smode,MTL::TextureUsageShaderRead); impl = alloc(frm,pm.w(),pm.h(),0,mipCnt,MTL::StorageModePrivate,MTL::TextureUsageShaderRead); @@ -152,7 +152,7 @@ void MtTexture::readPixels(Pixmap& out, TextureFormat frm, const uint32_t w, con throw std::runtime_error("not implemented"); out = Pixmap(w,h,frm); - const MTL::StorageMode opt = dev.impl->hasUnifiedMemory() ? MTL::StorageModeShared : MTL::StorageModeManaged; + const MTL::StorageMode opt = hostVisibleStorageMode(*dev.impl); NsPtr stage = alloc(frm,w,h,0,1,opt,MTL::TextureUsageShaderRead); auto pool = NsPtr::init(); diff --git a/Engine/gapi/metalapi.cpp b/Engine/gapi/metalapi.cpp index 66e291da..a01d9004 100644 --- a/Engine/gapi/metalapi.cpp +++ b/Engine/gapi/metalapi.cpp @@ -327,18 +327,12 @@ AbstractGraphicsApi::PBuffer MetalApi::createBuffer(AbstractGraphicsApi::Device opt |= MTL::ResourceStorageModePrivate; break; case BufferHeap::Upload: { - if(dx.impl->hasUnifiedMemory()) { - // Shared resources are only available on systems with integrated graphics, - // such as Apple silicon and integrated GPUs on Intel-based Mac computers - opt |= MTL::ResourceStorageModeShared; - } else { - opt |= MTL::ResourceStorageModeManaged; - } + opt |= hostVisibleResourceOptions(*dx.impl); opt |= MTL::ResourceCPUCacheModeWriteCombined; break; } case BufferHeap::Readback: - opt |= MTL::ResourceStorageModeManaged; + opt |= hostVisibleResourceOptions(*dx.impl); opt |= MTL::ResourceCPUCacheModeDefaultCache; break; } From 16fac300106d983630c132c762ecb479b0a599af Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 19:05:35 +0100 Subject: [PATCH 24/25] Initialize brush state for triangle drawing --- Engine/2d/painter.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/Engine/2d/painter.cpp b/Engine/2d/painter.cpp index 66c305a8..70cd779a 100644 --- a/Engine/2d/painter.cpp +++ b/Engine/2d/painter.cpp @@ -85,6 +85,11 @@ void Painter::implSetColor(float r, float g, float b, float a) { void Painter::drawTriangle(int x0, int y0, float u0, float v0, int x1, int y1, float u1, float v1, int x2, int y2, float u2, float v2) { + if(state!=StBrush) { + dev.setTopology(Triangles); + state=StBrush; + implBrush(s.br); + } FPoint trigBuf[4+4+4+4]; implDrawTrig( float(x0), float(y0), s.dU+u0*s.invW,s.dV+v0*s.invH, float(x1), float(y1), s.dU+u1*s.invW,s.dV+v1*s.invH, @@ -95,6 +100,11 @@ void Painter::drawTriangle(int x0, int y0, float u0, float v0, void Painter::drawTriangle(float x0, float y0, float u0, float v0, float x1, float y1, float u1, float v1, float x2, float y2, float u2, float v2) { + if(state!=StBrush) { + dev.setTopology(Triangles); + state=StBrush; + implBrush(s.br); + } FPoint trigBuf[4+4+4+4]; implDrawTrig( x0, y0, s.dU+u0*s.invW,s.dV+v0*s.invH, x1, y1, s.dU+u1*s.invW,s.dV+v1*s.invH, From 047b9498c0bb6690e99cd77abef9e4afa4696dcd Mon Sep 17 00:00:00 2001 From: Patrick Baran Date: Sun, 30 Aug 2026 21:24:45 +0100 Subject: [PATCH 25/25] Guard missing platform font fallback --- Engine/formats/font.cpp | 11 +++++++++-- Tests/tests/font_test.cpp | 17 +++++++++++++++++ 2 files changed, 26 insertions(+), 2 deletions(-) create mode 100644 Tests/tests/font_test.cpp diff --git a/Engine/formats/font.cpp b/Engine/formats/font.cpp index aa7d0d0c..9b2870d3 100644 --- a/Engine/formats/font.cpp +++ b/Engine/formats/font.cpp @@ -261,8 +261,15 @@ struct FontElement::Impl { try { std::lock_guard guard(syncMap); if(fallback==nullptr){ - fallback.reset(new Impl(Detail::getFallbackFont())); - if(stbtt_InitFont(&fallback->info,fallback->data,0)==0) + const std::string path = Detail::getFallbackFont(); + // Only Windows currently provides a platform fallback path. An empty + // Impl has no stb_truetype data, so passing its null buffer to + // stbtt_InitFont is undefined behaviour (and crashes on iOS/macOS). + if(path.empty()) + return nullLater(); + fallback.reset(new Impl(path)); + if(fallback->data==nullptr || fallback->size==0 || + stbtt_InitFont(&fallback->info,fallback->data,0)==0) throw std::system_error(Tempest::SystemErrc::UnableToLoadAsset); } diff --git a/Tests/tests/font_test.cpp b/Tests/tests/font_test.cpp new file mode 100644 index 00000000..4b2dbd12 --- /dev/null +++ b/Tests/tests/font_test.cpp @@ -0,0 +1,17 @@ +#include +#include + +#include + +using namespace Tempest; + +TEST(FontTest, MissingPlatformFallbackIsSafe) { + Font font = Application::defaultFont(); + + // U+10FFFF is intentionally absent from the bundled Roboto font. Platforms + // without a configured system fallback must return empty geometry instead + // of trying to initialize stb_truetype with a null buffer. + const Size size = font.textSize("\xF4\x8F\xBF\xBF"); + EXPECT_GE(size.w,0); + EXPECT_GE(size.h,0); + }