diff --git a/src/eepp/graphics/texture.cpp b/src/eepp/graphics/texture.cpp
index 9c5dc43a0..0d0cfc922 100644
--- a/src/eepp/graphics/texture.cpp
+++ b/src/eepp/graphics/texture.cpp
@@ -530,10 +530,18 @@ void Texture::update( const Uint8* pixels, Uint32 width, Uint32 height, Uint32 x
 
 		{
 			ScopedTexture saver( mTexture );
+			const bool tightlyPackedThreeChannelRows =
+				( pf == Image::PixelFormat::PIXEL_FORMAT_RGB ||
+				  pf == Image::PixelFormat::PIXEL_FORMAT_BGR ) &&
+				( static_cast<size_t>( width ) * 3 ) % 4 != 0;
+			if ( tightlyPackedThreeChannelRows )
+				glPixelStorei( GL_UNPACK_ALIGNMENT, 1 );
 
 			glTexSubImage2D( GL_TEXTURE_2D, 0, x, y, width, height,
 							 (unsigned int)convertPixelFormatToGLFormat( pf ), GL_UNSIGNED_BYTE,
 							 pixels );
+			if ( tightlyPackedThreeChannelRows )
+				glPixelStorei( GL_UNPACK_ALIGNMENT, 4 );
 
 			if ( hasLocalCopy() ) {
 				Image image( pixels, width, height, mChannels );
diff --git a/src/eepp/system/base64.cpp b/src/eepp/system/base64.cpp
index e4cd6a3a6..d41235e0e 100644
--- a/src/eepp/system/base64.cpp
+++ b/src/eepp/system/base64.cpp
@@ -108,40 +108,51 @@ size_t decodeBase64( size_t in_len, const char* in, size_t out_len, unsigned cha
 
 size_t decodeBase64Strict( size_t inLen, const char* input, size_t outLen, unsigned char* output ) {
 	size_t outputOffset = 0;
-	Uint32 value = 0;
-	unsigned bits = 0;
-	size_t padding = 0;
-	for ( size_t offset = 0; offset < inLen; ++offset ) {
-		const Uint8 decoded = base64dec_tab[static_cast<unsigned char>( input[offset] )];
-		if ( decoded == BASE64_PADDING ) {
-			padding = inLen - offset;
-			if ( padding > 2 || inLen % 4 != 0 )
-				return static_cast<size_t>( -1 );
-			for ( size_t remainder = offset; remainder < inLen; ++remainder )
-				if ( input[remainder] != '=' )
-					return static_cast<size_t>( -1 );
+	size_t inputOffset = 0;
+	while ( inputOffset + 4 <= inLen ) {
+		const Uint8 a = base64dec_tab[static_cast<unsigned char>( input[inputOffset] )];
+		const Uint8 b = base64dec_tab[static_cast<unsigned char>( input[inputOffset + 1] )];
+		const Uint8 c = base64dec_tab[static_cast<unsigned char>( input[inputOffset + 2] )];
+		const Uint8 d = base64dec_tab[static_cast<unsigned char>( input[inputOffset + 3] )];
+		if ( ( a | b | c | d ) > 63 )
 			break;
-		}
-		if ( decoded > 63 || padding != 0 )
+		if ( outputOffset + 3 > outLen )
 			return static_cast<size_t>( -1 );
-		value = ( value << 6 ) | decoded;
-		bits += 6;
-		if ( bits >= 8 ) {
-			bits -= 8;
-			if ( outputOffset >= outLen )
-				return static_cast<size_t>( -1 );
-			output[outputOffset++] = static_cast<unsigned char>( ( value >> bits ) & 0xFF );
-			if ( bits == 0 )
-				value = 0;
-		}
+		output[outputOffset] = static_cast<unsigned char>( ( a << 2 ) | ( b >> 4 ) );
+		output[outputOffset + 1] = static_cast<unsigned char>( ( b << 4 ) | ( c >> 2 ) );
+		output[outputOffset + 2] = static_cast<unsigned char>( ( c << 6 ) | d );
+		inputOffset += 4;
+		outputOffset += 3;
 	}
-	const size_t dataLength = inLen - padding;
-	if ( dataLength % 4 == 1 || ( padding == 1 && dataLength % 4 != 3 ) ||
-		 ( padding == 2 && dataLength % 4 != 2 ) )
+
+	const size_t remaining = inLen - inputOffset;
+	if ( remaining == 0 )
+		return outputOffset;
+	if ( remaining == 1 || remaining > 4 )
 		return static_cast<size_t>( -1 );
-	// Reject non-canonical encodings whose unused bits are non-zero. Besides being strict, this
-	// avoids accepting multiple byte strings for the same payload at protocol boundaries.
-	if ( bits != 0 && ( value & ( ( 1u << bits ) - 1u ) ) != 0 )
+
+	const Uint8 a = base64dec_tab[static_cast<unsigned char>( input[inputOffset] )];
+	const Uint8 b = base64dec_tab[static_cast<unsigned char>( input[inputOffset + 1] )];
+	if ( a > 63 || b > 63 || outputOffset >= outLen )
+		return static_cast<size_t>( -1 );
+	output[outputOffset++] = static_cast<unsigned char>( ( a << 2 ) | ( b >> 4 ) );
+
+	if ( remaining == 2 )
+		return ( b & 0x0F ) == 0 ? outputOffset : static_cast<size_t>( -1 );
+
+	const Uint8 c = base64dec_tab[static_cast<unsigned char>( input[inputOffset + 2] )];
+	if ( c == BASE64_PADDING ) {
+		return remaining == 4 && input[inputOffset + 3] == '=' && ( b & 0x0F ) == 0
+				   ? outputOffset
+				   : static_cast<size_t>( -1 );
+	}
+	if ( c > 63 || outputOffset >= outLen )
+		return static_cast<size_t>( -1 );
+	output[outputOffset++] = static_cast<unsigned char>( ( b << 4 ) | ( c >> 2 ) );
+
+	if ( remaining == 3 )
+		return ( c & 0x03 ) == 0 ? outputOffset : static_cast<size_t>( -1 );
+	if ( input[inputOffset + 3] != '=' || ( c & 0x03 ) != 0 )
 		return static_cast<size_t>( -1 );
 	return outputOffset;
 }
diff --git a/src/eepp/system/cpu.cpp b/src/eepp/system/cpu.cpp
index 756aa0a1a..68ad612a6 100644
--- a/src/eepp/system/cpu.cpp
+++ b/src/eepp/system/cpu.cpp
@@ -42,6 +42,12 @@ bool CPU::hasAVX2() {
 		__cpuid( cpuInfo, 0 );
 		if ( cpuInfo[0] < 7 )
 			return false;
+		__cpuid( cpuInfo, 1 );
+		constexpr int OSXSAVE = 1 << 27;
+		constexpr int AVX = 1 << 28;
+		if ( ( cpuInfo[2] & ( OSXSAVE | AVX ) ) != ( OSXSAVE | AVX ) ||
+			 ( _xgetbv( 0 ) & 0x6 ) != 0x6 )
+			return false;
 		__cpuid( cpuInfo, 7 );
 		return ( cpuInfo[1] & ( 1 << 5 ) ) != 0;
 #elif defined( COMPILER_GCC_CLANG )
diff --git a/src/modules/eterm/include/eterm/terminal/kittygraphicsprotocol.hpp b/src/modules/eterm/include/eterm/terminal/kittygraphicsprotocol.hpp
index c137975fb..7972c7547 100644
--- a/src/modules/eterm/include/eterm/terminal/kittygraphicsprotocol.hpp
+++ b/src/modules/eterm/include/eterm/terminal/kittygraphicsprotocol.hpp
@@ -168,19 +168,20 @@ class KittyGraphicsProtocol {
 			Int32 gapMs{ 40 };
 			Uint32 usageHint{ 0 };
 		};
-		std::shared_ptr<std::vector<Uint8>> rgba;
+		std::shared_ptr<std::vector<Uint8>> pixels;
 		std::unordered_map<Uint32, Frame> frames;
 		Sizei size;
+		Uint64 creationSerial{ 0 };
+		EE::System::Clock frameClock;
 		Uint32 imageNumber{ 0 };
 		Uint32 usageHint{ 0 };
-		bool anonymous{ false };
-		Uint64 creationSerial{ 0 };
 		Uint32 currentFrame{ 1 };
 		Uint32 loopCount{ 1 };
 		Uint32 loopsCompleted{ 0 };
-		Uint8 animationState{ 1 };
 		Int32 rootGapMs{ 0 };
-		EE::System::Clock frameClock;
+		Uint8 animationState{ 1 };
+		Uint8 channels{ 4 };
+		bool anonymous{ false };
 	};
 
 	struct PendingTransfer {
@@ -210,6 +211,7 @@ class KittyGraphicsProtocol {
 	KittyImageId allocateImageId();
 	KittyImageId resolveImageId( const KittyGraphicsCommandData& data ) const;
 	bool ensureCapacity( size_t bytes, KittyImageId replacingId, bool addingImage );
+	bool ensureRootRGBA( KittyImageId imageId, Image& image );
 	bool isImagePlaced( KittyImageId imageId ) const;
 	void eraseImage( KittyImageId imageId );
 	std::string response( const KittyGraphicsCommandData& data, KittyGraphicsError error,
diff --git a/src/modules/eterm/include/eterm/terminal/kittygraphicsrenderer.hpp b/src/modules/eterm/include/eterm/terminal/kittygraphicsrenderer.hpp
index e0e9b8173..a165304ee 100644
--- a/src/modules/eterm/include/eterm/terminal/kittygraphicsrenderer.hpp
+++ b/src/modules/eterm/include/eterm/terminal/kittygraphicsrenderer.hpp
@@ -30,6 +30,7 @@ class KittyGraphicsRenderer {
 		TexturePtr texture;
 		std::unordered_map<Uint32, TexturePtr> frames;
 		Sizei size;
+		Uint8 channels{ 4 };
 	};
 
 	std::unordered_map<KittyImageId, GPUImage> mImages;
diff --git a/src/modules/eterm/include/eterm/terminal/terminalgraphics.hpp b/src/modules/eterm/include/eterm/terminal/terminalgraphics.hpp
index 83b4ab6f0..8f059a7b9 100644
--- a/src/modules/eterm/include/eterm/terminal/terminalgraphics.hpp
+++ b/src/modules/eterm/include/eterm/terminal/terminalgraphics.hpp
@@ -64,15 +64,16 @@ enum class TerminalGraphicsUpdateType : Uint8 {
 };
 
 struct TerminalGraphicsUpdate {
-	std::shared_ptr<const std::vector<Uint8>> rgba;
+	std::shared_ptr<const std::vector<Uint8>> pixels;
 	Sizei imageSize;
 	Rect region;
 	Uint64 sequence{ 0 };
 	KittyImageId imageId{ 0 };
 	Uint32 frameNumber{ 1 };
 	TerminalGraphicsUpdateType type{ TerminalGraphicsUpdateType::Resync };
+	Uint8 channels{ 4 };
 
-	size_t payloadBytes() const { return rgba ? rgba->size() : 0; }
+	size_t payloadBytes() const { return pixels ? pixels->size() : 0; }
 };
 
 /** Bounded ordered worker-to-UI mutation queue with explicit overflow recovery. */
diff --git a/src/modules/eterm/src/eterm/terminal/kittygraphicsprotocol.cpp b/src/modules/eterm/src/eterm/terminal/kittygraphicsprotocol.cpp
index 218b36a5e..1d32fdc82 100644
--- a/src/modules/eterm/src/eterm/terminal/kittygraphicsprotocol.cpp
+++ b/src/modules/eterm/src/eterm/terminal/kittygraphicsprotocol.cpp
@@ -9,17 +9,8 @@
 #include <charconv>
 #include <cstring>
 #include <limits>
-#include <thread>
 #include <type_traits>
 
-#if EE_PLATFORM != EE_PLATFORM_WIN && EE_PLATFORM != EE_PLATFORM_EMSCRIPTEN && \
-	EE_PLATFORM != EE_PLATFORM_ANDROID
-#include <fcntl.h>
-#include <sys/mman.h>
-#include <sys/stat.h>
-#include <unistd.h>
-#endif
-
 using namespace EE::System;
 
 namespace eterm { namespace Terminal {
@@ -78,10 +69,12 @@ bool decodeBase64( std::string_view input, bool finalChunk, std::vector<Uint8>&
 	if ( oldSize > MaxTransferBytes || capacity > MaxTransferBytes - oldSize )
 		return false;
 	output.resize( oldSize + capacity );
+	const Base64::DecodeMode mode =
+		finalChunk ? Base64::DecodeMode::NoWhitespaceStrict : Base64::DecodeMode::NoWhitespace;
 	const size_t decodedSize =
-		Base64::decode( input.size(), input.data(), capacity, output.data() + oldSize,
-						Base64::DecodeMode::NoWhitespaceStrict );
-	if ( decodedSize == static_cast<size_t>( -1 ) ) {
+		Base64::decode( input.size(), input.data(), capacity, output.data() + oldSize, mode );
+	if ( decodedSize == static_cast<size_t>( -1 ) ||
+		 ( !finalChunk && decodedSize != input.size() / 4 * 3 ) ) {
 		output.resize( oldSize );
 		return false;
 	}
@@ -89,79 +82,6 @@ bool decodeBase64( std::string_view input, bool finalChunk, std::vector<Uint8>&
 	return true;
 }
 
-bool readSharedMemory( const KittyGraphicsCommandData& data, std::vector<Uint8>& output ) {
-#if EE_PLATFORM == EE_PLATFORM_WIN || EE_PLATFORM == EE_PLATFORM_EMSCRIPTEN || \
-	EE_PLATFORM == EE_PLATFORM_ANDROID
-	(void)data;
-	(void)output;
-	return false;
-#else
-	std::vector<Uint8> decodedName;
-	if ( !decodeBase64( data.payload, true, decodedName ) || decodedName.empty() ||
-		 decodedName.size() > 255 ||
-		 std::find( decodedName.begin() + 1, decodedName.end(), '/' ) != decodedName.end() ||
-		 std::find( decodedName.begin(), decodedName.end(), 0 ) != decodedName.end() )
-		return false;
-	std::string name( decodedName.begin(), decodedName.end() );
-	const int descriptor = shm_open( name.c_str(), O_RDONLY, 0 );
-	if ( descriptor == -1 )
-		return false;
-	// POSIX Kitty transfers are single-use. Unlink immediately after opening so all error paths
-	// still retire the client-owned object while the descriptor keeps its contents alive.
-	shm_unlink( name.c_str() );
-	struct stat status{};
-	const size_t offset = data.dataOffset.value_or( 0 );
-	bool valid = fstat( descriptor, &status ) == 0 && status.st_size >= 0 &&
-				 static_cast<Uint64>( status.st_size ) >= offset;
-	size_t bytes = 0;
-	if ( valid ) {
-		const size_t available = static_cast<size_t>( status.st_size ) - offset;
-		bytes = data.dataSize.value_or( static_cast<Uint32>(
-			std::min<size_t>( available, std::numeric_limits<Uint32>::max() ) ) );
-		valid = bytes <= available && bytes <= 128 * 1024 * 1024;
-	}
-	if ( valid && bytes != 0 ) {
-		const size_t mappingBytes = offset + bytes;
-		valid = mappingBytes >= bytes && mappingBytes <= 128 * 1024 * 1024;
-		void* mapping = valid ? mmap( nullptr, mappingBytes, PROT_READ, MAP_SHARED, descriptor, 0 )
-							  : MAP_FAILED;
-		if ( mapping == MAP_FAILED ) {
-			valid = false;
-		} else {
-			const auto* source = static_cast<const Uint8*>( mapping ) + offset;
-			// mpv reuses the same shm name for every frame. If it reopened the object before
-			// we unlinked it above, its next memcpy can overlap this read. Require consecutive
-			// identical observations after yielding to the writer; otherwise retain the previous
-			// displayed frame instead of publishing visibly torn rows.
-			std::vector<Uint8> snapshot( source, source + bytes );
-			bool stable = false;
-			unsigned stableObservations = 0;
-			constexpr unsigned RequiredStableObservations = 2;
-			constexpr unsigned MaxSnapshotAttempts = 6;
-			for ( unsigned attempt = 0; attempt < MaxSnapshotAttempts; ++attempt ) {
-				std::this_thread::yield();
-				if ( std::memcmp( snapshot.data(), source, bytes ) == 0 ) {
-					if ( ++stableObservations == RequiredStableObservations ) {
-						stable = true;
-						break;
-					}
-				} else {
-					stableObservations = 0;
-					snapshot.assign( source, source + bytes );
-				}
-			}
-			if ( stable )
-				output = std::move( snapshot );
-			else
-				valid = false;
-			munmap( mapping, mappingBytes );
-		}
-	}
-	close( descriptor );
-	return valid && bytes != 0;
-#endif
-}
-
 bool placementContains( const TerminalVisiblePlacement& placement, Vector2i cell ) {
 	const Int64 right = static_cast<Int64>( placement.visibleAnchorCell.x ) + placement.columns;
 	const Int64 bottom = static_cast<Int64>( placement.visibleAnchorCell.y ) + placement.rows;
@@ -402,20 +322,6 @@ KittyGraphicsHandleResult
 KittyGraphicsProtocol::handleTransmit( const KittyGraphicsCommandData& data, bool display,
 									   bool query, bool frame, Vector2i cursor ) {
 	const bool more = data.more.value_or( 0 ) != 0;
-	if ( data.transmission == 's' ) {
-		if ( mPending.active )
-			mPending = {};
-		PendingTransfer transfer;
-		transfer.data = data;
-		transfer.data.payload = {};
-		transfer.display = display;
-		transfer.query = query;
-		transfer.frame = frame;
-		if ( !readSharedMemory( data, transfer.decodedData ) )
-			return { response( data, KittyGraphicsError::DecodeFailed ),
-					 KittyGraphicsError::DecodeFailed, false };
-		return finishTransfer( std::move( transfer ), cursor );
-	}
 	if ( mPending.active ) {
 		if ( frame != mPending.frame || data.format || data.dataSize || data.imageId ||
 			 data.imageNumber || data.usageHint || data.placementId || data.width || data.height ||
@@ -561,7 +467,7 @@ KittyGraphicsHandleResult KittyGraphicsProtocol::finishTransfer( PendingTransfer
 		}
 	}
 
-	if ( format == 24 ) {
+	if ( format == 24 && transfer.frame ) {
 		std::vector<Uint8> rgba = std::move( mPixelScratch );
 		rgba.resize( static_cast<size_t>( width ) * height * 4 );
 		for ( size_t sourceOffset = 0, destinationOffset = 0; sourceOffset < pixels.size();
@@ -580,6 +486,9 @@ KittyGraphicsHandleResult KittyGraphicsProtocol::finishTransfer( PendingTransfer
 		if ( image == mImages.end() )
 			return { response( data, KittyGraphicsError::NotFound ), KittyGraphicsError::NotFound,
 					 false };
+		if ( !ensureRootRGBA( imageId, image->second ) )
+			return { response( data, KittyGraphicsError::NoSpace ), KittyGraphicsError::NoSpace,
+					 false };
 		Uint32 frameNumber = data.rows.value_or( 0 );
 		if ( frameNumber == 0 ) {
 			frameNumber = 2;
@@ -599,9 +508,10 @@ KittyGraphicsHandleResult KittyGraphicsProtocol::finishTransfer( PendingTransfer
 		std::vector<Uint8>* destinationPixels = nullptr;
 		bool createdFrame = false;
 		if ( frameNumber == 1 ) {
-			if ( image->second.rgba.use_count() != 1 )
-				image->second.rgba = std::make_shared<std::vector<Uint8>>( *image->second.rgba );
-			destinationPixels = image->second.rgba.get();
+			if ( image->second.pixels.use_count() != 1 )
+				image->second.pixels =
+					std::make_shared<std::vector<Uint8>>( *image->second.pixels );
+			destinationPixels = image->second.pixels.get();
 		} else {
 			auto frame = image->second.frames.find( frameNumber );
 			if ( frame == image->second.frames.end() ) {
@@ -611,7 +521,7 @@ KittyGraphicsHandleResult KittyGraphicsProtocol::finishTransfer( PendingTransfer
 				if ( data.columns ) {
 					const Uint32 baseFrame = *data.columns;
 					if ( baseFrame == 1 )
-						newFrame.rgba = *image->second.rgba;
+						newFrame.rgba = *image->second.pixels;
 					else {
 						auto base = image->second.frames.find( baseFrame );
 						if ( base == image->second.frames.end() )
@@ -687,9 +597,9 @@ KittyGraphicsHandleResult KittyGraphicsProtocol::finishTransfer( PendingTransfer
 			createdFrame
 				? Rect( 0, 0, image->second.size.getWidth(), image->second.size.getHeight() )
 				: Rect( destinationX, destinationY, destinationX + width, destinationY + height );
-		update.rgba = createdFrame
-						  ? std::make_shared<const std::vector<Uint8>>( *destinationPixels )
-						  : std::make_shared<const std::vector<Uint8>>( std::move( finalPatch ) );
+		update.pixels = createdFrame
+							? std::make_shared<const std::vector<Uint8>>( *destinationPixels )
+							: std::make_shared<const std::vector<Uint8>>( std::move( finalPatch ) );
 		if ( !createdFrame && !mUpdates.empty() && mUpdates.back().type == update.type &&
 			 mUpdates.back().imageId == update.imageId &&
 			 mUpdates.back().frameNumber == update.frameNumber &&
@@ -730,7 +640,7 @@ KittyGraphicsHandleResult KittyGraphicsProtocol::finishTransfer( PendingTransfer
 				 false };
 
 	auto existing = mImages.find( imageId );
-	const size_t oldBytes = existing == mImages.end() ? 0 : existing->second.rgba->size();
+	const size_t oldBytes = existing == mImages.end() ? 0 : existing->second.pixels->size();
 	if ( !ensureCapacity( pixels.size(), imageId, existing == mImages.end() ) )
 		return { response( data, KittyGraphicsError::NoSpace ), KittyGraphicsError::NoSpace,
 				 false };
@@ -751,11 +661,11 @@ KittyGraphicsHandleResult KittyGraphicsProtocol::finishTransfer( PendingTransfer
 								  mPrimaryPlacements.end() );
 	}
 	std::shared_ptr<std::vector<Uint8>> pixelStorage;
-	if ( existing != mImages.end() && existing->second.rgba.use_count() == 1 ) {
-		pixelStorage = existing->second.rgba;
+	if ( existing != mImages.end() && existing->second.pixels.use_count() == 1 ) {
+		pixelStorage = existing->second.pixels;
 		pixelStorage->assign( pixels.begin(), pixels.end() );
 		if ( format == 24 )
-			recycleBuffer( pixels, mPixelScratch );
+			recycleBuffer( pixels, mDecodedScratch );
 		else if ( format == 32 )
 			recycleBuffer( pixels, mDecodedScratch );
 	} else {
@@ -767,8 +677,9 @@ KittyGraphicsHandleResult KittyGraphicsProtocol::finishTransfer( PendingTransfer
 	image.usageHint = data.usageHint.value_or( 0 );
 	image.anonymous = anonymous;
 	image.creationSerial = ++mCreationSerial;
-	image.rgba = std::move( pixelStorage );
-	mStorageBytes = mStorageBytes - oldBytes + image.rgba->size();
+	image.channels = format == 24 ? 3 : 4;
+	image.pixels = std::move( pixelStorage );
+	mStorageBytes = mStorageBytes - oldBytes + image.pixels->size();
 	const bool replaced = existing != mImages.end();
 	auto inserted = mImages.insert_or_assign( imageId, std::move( image ) ).first;
 
@@ -777,9 +688,10 @@ KittyGraphicsHandleResult KittyGraphicsProtocol::finishTransfer( PendingTransfer
 						   : TerminalGraphicsUpdateType::CreateImage;
 	update.imageId = imageId;
 	update.imageSize = inserted->second.size;
+	update.channels = inserted->second.channels;
 	update.region =
 		Rect( 0, 0, inserted->second.size.getWidth(), inserted->second.size.getHeight() );
-	update.rgba = inserted->second.rgba;
+	update.pixels = inserted->second.pixels;
 	mUpdates.emplace_back( std::move( update ) );
 	++mStats.fullImageUpdates;
 	++mPresentationGeneration;
@@ -1241,9 +1153,12 @@ KittyGraphicsProtocol::composeFrames( const KittyGraphicsCommandData& data ) {
 				 image == mImages.end() ? KittyGraphicsError::NotFound
 										: KittyGraphicsError::InvalidArgument,
 				 false };
+	if ( !ensureRootRGBA( imageId, image->second ) )
+		return { response( data, KittyGraphicsError::NoSpace ), KittyGraphicsError::NoSpace,
+				 false };
 	auto pixelsFor = [&]( Uint32 frameNumber ) -> std::vector<Uint8>* {
 		if ( frameNumber == 1 )
-			return image->second.rgba.get();
+			return image->second.pixels.get();
 		auto frame = image->second.frames.find( frameNumber );
 		return frame == image->second.frames.end() ? nullptr : &frame->second.rgba;
 	};
@@ -1284,9 +1199,9 @@ KittyGraphicsProtocol::composeFrames( const KittyGraphicsCommandData& data ) {
 					 source->data() + static_cast<size_t>( sourceY + row ) * imageStride +
 						 static_cast<size_t>( sourceX ) * 4,
 					 rowBytes );
-	if ( destinationFrame == 1 && image->second.rgba.use_count() != 1 ) {
-		image->second.rgba = std::make_shared<std::vector<Uint8>>( *image->second.rgba );
-		destination = image->second.rgba.get();
+	if ( destinationFrame == 1 && image->second.pixels.use_count() != 1 ) {
+		image->second.pixels = std::make_shared<std::vector<Uint8>>( *image->second.pixels );
+		destination = image->second.pixels.get();
 	}
 	const bool replace = data.cursorMovement.value_or( 0 ) == 1;
 	for ( Uint32 row = 0; row < height; ++row ) {
@@ -1320,7 +1235,7 @@ KittyGraphicsProtocol::composeFrames( const KittyGraphicsCommandData& data ) {
 	update.frameNumber = destinationFrame;
 	update.imageSize = image->second.size;
 	update.region = Rect( destinationX, destinationY, destinationX + width, destinationY + height );
-	update.rgba = std::make_shared<const std::vector<Uint8>>( std::move( result ) );
+	update.pixels = std::make_shared<const std::vector<Uint8>>( std::move( result ) );
 	mUpdates.emplace_back( std::move( update ) );
 	mPresentationDirty = true;
 	return { response( data, KittyGraphicsError::None, imageId ), KittyGraphicsError::None, true };
@@ -1366,7 +1281,7 @@ void KittyGraphicsProtocol::eraseImage( KittyImageId imageId ) {
 	auto image = mImages.find( imageId );
 	if ( image == mImages.end() )
 		return;
-	mStorageBytes -= image->second.rgba->size();
+	mStorageBytes -= image->second.pixels->size();
 	for ( const auto& frame : image->second.frames )
 		mFrameStorageBytes -= frame.second.rgba.size();
 	mImages.erase( image );
@@ -1379,7 +1294,7 @@ void KittyGraphicsProtocol::eraseImage( KittyImageId imageId ) {
 bool KittyGraphicsProtocol::ensureCapacity( size_t bytes, KittyImageId replacingId,
 											bool addingImage ) {
 	auto replaced = mImages.find( replacingId );
-	const size_t replacedBytes = replaced == mImages.end() ? 0 : replaced->second.rgba->size();
+	const size_t replacedBytes = replaced == mImages.end() ? 0 : replaced->second.pixels->size();
 	auto hasCapacity = [&] {
 		return bytes <= mMaxStorageBytes &&
 			   mStorageBytes - replacedBytes <= mMaxStorageBytes - bytes &&
@@ -1404,6 +1319,37 @@ bool KittyGraphicsProtocol::ensureCapacity( size_t bytes, KittyImageId replacing
 	return true;
 }
 
+bool KittyGraphicsProtocol::ensureRootRGBA( KittyImageId imageId, Image& image ) {
+	if ( image.channels == 4 )
+		return true;
+	const size_t pixelCount = static_cast<size_t>( image.size.getWidth() ) * image.size.getHeight();
+	const size_t rgbaBytes = pixelCount * 4;
+	if ( !ensureCapacity( rgbaBytes, imageId, false ) )
+		return false;
+	std::vector<Uint8> rgba;
+	rgba.resize( rgbaBytes );
+	for ( size_t pixel = 0; pixel < pixelCount; ++pixel ) {
+		const size_t source = pixel * 3;
+		const size_t destination = pixel * 4;
+		rgba[destination] = ( *image.pixels )[source];
+		rgba[destination + 1] = ( *image.pixels )[source + 1];
+		rgba[destination + 2] = ( *image.pixels )[source + 2];
+		rgba[destination + 3] = 255;
+	}
+	const size_t oldBytes = image.pixels->size();
+	image.pixels = std::make_shared<std::vector<Uint8>>( std::move( rgba ) );
+	image.channels = 4;
+	mStorageBytes = mStorageBytes - oldBytes + rgbaBytes;
+	TerminalGraphicsUpdate update;
+	update.type = TerminalGraphicsUpdateType::ReplaceImage;
+	update.imageId = imageId;
+	update.imageSize = image.size;
+	update.region = Rect( 0, 0, image.size.getWidth(), image.size.getHeight() );
+	update.pixels = image.pixels;
+	mUpdates.emplace_back( std::move( update ) );
+	return true;
+}
+
 std::string KittyGraphicsProtocol::response( const KittyGraphicsCommandData& data,
 											 KittyGraphicsError error,
 											 KittyImageId imageId ) const {
@@ -1528,7 +1474,7 @@ std::shared_ptr<TerminalGraphicsPresentation> KittyGraphicsProtocol::takePresent
 
 const std::vector<Uint8>* KittyGraphicsProtocol::imagePixels( KittyImageId imageId ) const {
 	auto image = mImages.find( imageId );
-	return image == mImages.end() ? nullptr : image->second.rgba.get();
+	return image == mImages.end() ? nullptr : image->second.pixels.get();
 }
 
 bool KittyGraphicsProtocol::hasVirtualPlacements() const {
@@ -1665,8 +1611,9 @@ void KittyGraphicsProtocol::resync() {
 		create.type = TerminalGraphicsUpdateType::CreateImage;
 		create.imageId = image.first;
 		create.imageSize = image.second.size;
+		create.channels = image.second.channels;
 		create.region = Rect( 0, 0, image.second.size.getWidth(), image.second.size.getHeight() );
-		create.rgba = image.second.rgba;
+		create.pixels = image.second.pixels;
 		mUpdates.emplace_back( std::move( create ) );
 		for ( const auto& frame : image.second.frames ) {
 			TerminalGraphicsUpdate createFrame;
@@ -1676,7 +1623,7 @@ void KittyGraphicsProtocol::resync() {
 			createFrame.imageSize = image.second.size;
 			createFrame.region =
 				Rect( 0, 0, image.second.size.getWidth(), image.second.size.getHeight() );
-			createFrame.rgba = std::make_shared<const std::vector<Uint8>>( frame.second.rgba );
+			createFrame.pixels = std::make_shared<const std::vector<Uint8>>( frame.second.rgba );
 			mUpdates.emplace_back( std::move( createFrame ) );
 		}
 	}
diff --git a/src/modules/eterm/src/eterm/terminal/kittygraphicsrenderer.cpp b/src/modules/eterm/src/eterm/terminal/kittygraphicsrenderer.cpp
index 7b8302d33..75e4e44cd 100644
--- a/src/modules/eterm/src/eterm/terminal/kittygraphicsrenderer.cpp
+++ b/src/modules/eterm/src/eterm/terminal/kittygraphicsrenderer.cpp
@@ -27,58 +27,68 @@ bool KittyGraphicsRenderer::applyUpdates( std::vector<TerminalGraphicsUpdate>&&
 
 		switch ( update.type ) {
 			case TerminalGraphicsUpdateType::ReplaceImage: {
-				if ( !update.rgba || update.imageSize.getWidth() <= 0 ||
-					 update.imageSize.getHeight() <= 0 )
+				if ( !update.pixels || update.imageSize.getWidth() <= 0 ||
+					 update.imageSize.getHeight() <= 0 ||
+					 ( update.channels != 3 && update.channels != 4 ) ||
+					 update.pixels->size() != static_cast<size_t>( update.imageSize.getWidth() ) *
+												  update.imageSize.getHeight() * update.channels )
 					return false;
 				auto existing = mImages.find( update.imageId );
 				if ( existing != mImages.end() && existing->second.texture &&
-					 existing->second.size == update.imageSize ) {
-					existing->second.texture->update( update.rgba->data(),
-													  update.imageSize.getWidth(),
-													  update.imageSize.getHeight(), 0, 0 );
+					 existing->second.size == update.imageSize &&
+					 existing->second.channels == update.channels ) {
+					existing->second.texture->update(
+						update.pixels->data(), update.imageSize.getWidth(),
+						update.imageSize.getHeight(), 0, 0,
+						Image::channelsToPixelFormat( update.channels ) );
 					existing->second.frames.clear();
 					break;
 				}
 				[[fallthrough]];
 			}
 			case TerminalGraphicsUpdateType::CreateImage: {
-				if ( !update.rgba || update.imageSize.getWidth() <= 0 ||
-					 update.imageSize.getHeight() <= 0 )
+				if ( !update.pixels || update.imageSize.getWidth() <= 0 ||
+					 update.imageSize.getHeight() <= 0 ||
+					 ( update.channels != 3 && update.channels != 4 ) ||
+					 update.pixels->size() != static_cast<size_t>( update.imageSize.getWidth() ) *
+												  update.imageSize.getHeight() * update.channels )
 					return false;
 				auto texture = TextureFactory::instance()->loadFromPixels(
-					update.rgba->data(), update.imageSize.getWidth(), update.imageSize.getHeight(),
-					4, false, Texture::ClampMode::ClampToEdge, false, false );
+					update.pixels->data(), update.imageSize.getWidth(),
+					update.imageSize.getHeight(), update.channels, false,
+					Texture::ClampMode::ClampToEdge, false, false );
 				if ( !texture )
 					return false;
 				texture->setFilter( Texture::Filter::Linear );
-				mImages.insert_or_assign( update.imageId,
-										  GPUImage{ std::move( texture ), {}, update.imageSize } );
+				mImages.insert_or_assign(
+					update.imageId,
+					GPUImage{ std::move( texture ), {}, update.imageSize, update.channels } );
 				break;
 			}
 			case TerminalGraphicsUpdateType::UpdateRegion: {
 				auto image = mImages.find( update.imageId );
-				if ( image == mImages.end() || !update.rgba )
+				if ( image == mImages.end() || !update.pixels )
 					return false;
 				const int width = update.region.Right - update.region.Left;
 				const int height = update.region.Bottom - update.region.Top;
 				if ( width <= 0 || height <= 0 || update.region.Left < 0 || update.region.Top < 0 ||
 					 update.region.Right > image->second.size.getWidth() ||
 					 update.region.Bottom > image->second.size.getHeight() ||
-					 update.rgba->size() != static_cast<size_t>( width ) * height * 4 )
+					 update.pixels->size() != static_cast<size_t>( width ) * height * 4 )
 					return false;
-				image->second.texture->update( update.rgba->data(), width, height,
+				image->second.texture->update( update.pixels->data(), width, height,
 											   update.region.Left, update.region.Top );
 				break;
 			}
 			case TerminalGraphicsUpdateType::CreateFrame:
 			case TerminalGraphicsUpdateType::ReplaceFrame: {
 				auto image = mImages.find( update.imageId );
-				if ( image == mImages.end() || !update.rgba || update.frameNumber <= 1 ||
-					 update.rgba->size() != static_cast<size_t>( image->second.size.getWidth() ) *
-												image->second.size.getHeight() * 4 )
+				if ( image == mImages.end() || !update.pixels || update.frameNumber <= 1 ||
+					 update.pixels->size() != static_cast<size_t>( image->second.size.getWidth() ) *
+												  image->second.size.getHeight() * 4 )
 					return false;
 				auto texture = TextureFactory::instance()->loadFromPixels(
-					update.rgba->data(), image->second.size.getWidth(),
+					update.pixels->data(), image->second.size.getWidth(),
 					image->second.size.getHeight(), 4, false, Texture::ClampMode::ClampToEdge,
 					false, false );
 				if ( !texture )
@@ -89,7 +99,7 @@ bool KittyGraphicsRenderer::applyUpdates( std::vector<TerminalGraphicsUpdate>&&
 			}
 			case TerminalGraphicsUpdateType::UpdateFrameRegion: {
 				auto image = mImages.find( update.imageId );
-				if ( image == mImages.end() || !update.rgba )
+				if ( image == mImages.end() || !update.pixels )
 					return false;
 				auto frame = image->second.frames.find( update.frameNumber );
 				if ( frame == image->second.frames.end() )
@@ -99,9 +109,9 @@ bool KittyGraphicsRenderer::applyUpdates( std::vector<TerminalGraphicsUpdate>&&
 				if ( width <= 0 || height <= 0 || update.region.Left < 0 || update.region.Top < 0 ||
 					 update.region.Right > image->second.size.getWidth() ||
 					 update.region.Bottom > image->second.size.getHeight() ||
-					 update.rgba->size() != static_cast<size_t>( width ) * height * 4 )
+					 update.pixels->size() != static_cast<size_t>( width ) * height * 4 )
 					return false;
-				frame->second->update( update.rgba->data(), width, height, update.region.Left,
+				frame->second->update( update.pixels->data(), width, height, update.region.Left,
 									   update.region.Top );
 				break;
 			}
diff --git a/src/modules/eterm/src/eterm/terminal/terminalemulator.cpp b/src/modules/eterm/src/eterm/terminal/terminalemulator.cpp
index 18c03c529..0e80b02de 100644
--- a/src/modules/eterm/src/eterm/terminal/terminalemulator.cpp
+++ b/src/modules/eterm/src/eterm/terminal/terminalemulator.cpp
@@ -61,6 +61,7 @@ using namespace EE::System;
 #include <stdlib.h>
 #include <string.h>
 #include <sys/types.h>
+#include <thread>
 
 #if defined( EE_ARCH_X86_64 )
 #include <immintrin.h>
@@ -124,6 +125,39 @@ TerminalCursorMode blinkingCursorVariant( TerminalCursorMode mode ) {
 } // namespace
 
 #if defined( EE_ARCH_X86_64 )
+#if defined( __GNUC__ ) || defined( __clang__ )
+__attribute__( ( target( "avx2" ) ) )
+#endif
+static int findKittyAPCControlAVX2( const char* data, int begin, int end ) {
+	const __m256i highThreeBitMask = _mm256_set1_epi8( static_cast<char>( 0xE0 ) );
+	const __m256i c1Prefix = _mm256_set1_epi8( static_cast<char>( 0x80 ) );
+	const __m256i bell = _mm256_set1_epi8( 0x07 );
+	const __m256i cancel = _mm256_set1_epi8( 0x18 );
+	const __m256i substitute = _mm256_set1_epi8( 0x1A );
+	const __m256i escape = _mm256_set1_epi8( 0x1B );
+	int offset = begin;
+	for ( ; offset + 32 <= end; offset += 32 ) {
+		const __m256i bytes =
+			_mm256_loadu_si256( reinterpret_cast<const __m256i*>( data + offset ) );
+		const __m256i c1 =
+			_mm256_cmpeq_epi8( _mm256_and_si256( bytes, highThreeBitMask ), c1Prefix );
+		const __m256i explicitControls = _mm256_or_si256(
+			_mm256_or_si256( _mm256_cmpeq_epi8( bytes, bell ), _mm256_cmpeq_epi8( bytes, cancel ) ),
+			_mm256_or_si256( _mm256_cmpeq_epi8( bytes, substitute ),
+							 _mm256_cmpeq_epi8( bytes, escape ) ) );
+		if ( _mm256_movemask_epi8( _mm256_or_si256( c1, explicitControls ) ) != 0 )
+			break;
+	}
+	while ( offset < end ) {
+		const unsigned char byte = static_cast<unsigned char>( data[offset] );
+		if ( byte == '\a' || byte == 030 || byte == 032 || byte == 033 ||
+			 ( byte >= 0x80 && byte <= 0x9F ) )
+			break;
+		++offset;
+	}
+	return offset;
+}
+
 #if defined( __GNUC__ ) || defined( __clang__ )
 __attribute__( ( target( "avx2" ) ) )
 #endif
@@ -165,6 +199,21 @@ static int trailingNonSpaceWidthNEON( Line line, int width, int minimumWidth ) {
 }
 #endif
 
+static int findKittyAPCControl( const char* data, int begin, int end ) {
+#if defined( EE_ARCH_X86_64 )
+	if ( CPU::hasAVX2() )
+		return findKittyAPCControlAVX2( data, begin, end );
+#endif
+	while ( begin < end ) {
+		const unsigned char byte = static_cast<unsigned char>( data[begin] );
+		if ( byte == '\a' || byte == 030 || byte == 032 || byte == 033 ||
+			 ( byte >= 0x80 && byte <= 0x9F ) )
+			break;
+		++begin;
+	}
+	return begin;
+}
+
 /* identification sequence returned in DA and DECID */
 static const char* vtiden = "\033[?6c";
 
@@ -3243,14 +3292,7 @@ int TerminalEmulator::twrite( const char* buf, int buflen, int show_ctrl ) {
 		 * behavior. */
 		if ( !show_ctrl && ( mTerm.esc & ESC_STR ) && mStrescseq.type == '_' &&
 			 mStrescseq.len > 0 && mStrescseq.buf[0] == 'G' ) {
-			int end = n;
-			while ( end < buflen ) {
-				const unsigned char byte = static_cast<unsigned char>( buf[end] );
-				if ( byte == '\a' || byte == 030 || byte == 032 || byte == 033 ||
-					 ( byte >= 0x80 && byte <= 0x9F ) )
-					break;
-				++end;
-			}
+			const int end = findKittyAPCControl( buf, n, buflen );
 			const size_t bytes = static_cast<size_t>( end - n );
 			if ( bytes != 0 ) {
 				if ( !mStrescseq.discarded ) {
@@ -4108,6 +4150,16 @@ bool TerminalEmulator::update() {
 		onProcessExit( mExitCode );
 	}
 
+	/* A non-blocking read can temporarily catch up with a producer that was blocked writing to the
+	 * PTY. Do not enter the worker's 8 ms idle wait immediately after consuming data: give the
+	 * producer a scheduling opportunity and probe the PTY once more. This matters for high-volume
+	 * protocols such as Kitty graphics, where otherwise every transport chunk can pay one idle
+	 * interval after the receiver becomes faster than the sender's wakeup latency. */
+	if ( reads > 0 && !readBudgetSaturated ) {
+		std::this_thread::yield();
+		return false;
+	}
+
 	return !readBudgetSaturated;
 }
 
diff --git a/src/tests/unit_tests/eterm_tests.cpp b/src/tests/unit_tests/eterm_tests.cpp
index c7493652e..d92f82456 100644
--- a/src/tests/unit_tests/eterm_tests.cpp
+++ b/src/tests/unit_tests/eterm_tests.cpp
@@ -14,13 +14,6 @@
 #include <limits>
 #include <thread>
 
-#if EE_PLATFORM != EE_PLATFORM_WIN && EE_PLATFORM != EE_PLATFORM_EMSCRIPTEN && \
-	EE_PLATFORM != EE_PLATFORM_ANDROID
-#include <fcntl.h>
-#include <sys/mman.h>
-#include <unistd.h>
-#endif
-
 using namespace eterm::Terminal;
 using namespace eterm::System;
 using namespace EE::System;
@@ -133,11 +126,11 @@ UTEST( eterm_session, graphics_update_queue_preserves_order_and_payloads ) {
 	TerminalGraphicsUpdate create;
 	create.type = TerminalGraphicsUpdateType::CreateImage;
 	create.imageId = 7;
-	create.rgba = pixels;
+	create.pixels = pixels;
 	TerminalGraphicsUpdate patch;
 	patch.type = TerminalGraphicsUpdateType::UpdateRegion;
 	patch.imageId = 7;
-	patch.rgba = pixels;
+	patch.pixels = pixels;
 
 	EXPECT_EQ( static_cast<Uint64>( 1 ), queue.enqueue( std::move( create ) ) );
 	EXPECT_EQ( static_cast<Uint64>( 2 ), queue.enqueue( std::move( patch ) ) );
@@ -148,14 +141,14 @@ UTEST( eterm_session, graphics_update_queue_preserves_order_and_payloads ) {
 	EXPECT_EQ( static_cast<Uint64>( 2 ), updates[1].sequence );
 	EXPECT_EQ( TerminalGraphicsUpdateType::CreateImage, updates[0].type );
 	EXPECT_EQ( TerminalGraphicsUpdateType::UpdateRegion, updates[1].type );
-	EXPECT_TRUE( pixels == updates[0].rgba );
+	EXPECT_TRUE( pixels == updates[0].pixels );
 }
 
 UTEST( eterm_session, graphics_update_queue_overflow_requires_resync ) {
 	TerminalGraphicsUpdateQueue queue( 2, 8 );
 	TerminalGraphicsUpdate update;
 	update.type = TerminalGraphicsUpdateType::UpdateRegion;
-	update.rgba = std::make_shared<const std::vector<Uint8>>( 8, 0xFF );
+	update.pixels = std::make_shared<const std::vector<Uint8>>( 8, 0xFF );
 	queue.enqueue( update );
 	queue.enqueue( std::move( update ) );
 
@@ -172,13 +165,13 @@ UTEST( eterm_session, graphics_update_queue_coalesces_superseded_video_frames )
 	TerminalGraphicsUpdate create;
 	create.type = TerminalGraphicsUpdateType::CreateImage;
 	create.imageId = 7;
-	create.rgba = std::make_shared<const std::vector<Uint8>>( 8, 1 );
+	create.pixels = std::make_shared<const std::vector<Uint8>>( 8, 1 );
 	EXPECT_EQ( static_cast<Uint64>( 1 ), queue.enqueue( std::move( create ) ) );
 	for ( Uint8 frame = 2; frame < 20; ++frame ) {
 		TerminalGraphicsUpdate replacement;
 		replacement.type = TerminalGraphicsUpdateType::ReplaceImage;
 		replacement.imageId = 7;
-		replacement.rgba = std::make_shared<const std::vector<Uint8>>( 8, frame );
+		replacement.pixels = std::make_shared<const std::vector<Uint8>>( 8, frame );
 		EXPECT_EQ( static_cast<Uint64>( 1 ), queue.enqueue( std::move( replacement ) ) );
 	}
 	EXPECT_FALSE( queue.needsResync() );
@@ -186,7 +179,7 @@ UTEST( eterm_session, graphics_update_queue_coalesces_superseded_video_frames )
 	auto updates = queue.drain();
 	ASSERT_EQ( static_cast<size_t>( 1 ), updates.size() );
 	EXPECT_EQ( TerminalGraphicsUpdateType::CreateImage, updates[0].type );
-	EXPECT_EQ( static_cast<Uint8>( 19 ), updates[0].rgba->front() );
+	EXPECT_EQ( static_cast<Uint8>( 19 ), updates[0].pixels->front() );
 }
 
 UTEST( eterm_session, ordered_selection_request ) {
@@ -561,7 +554,7 @@ UTEST( eterm, kitty_graphics_image_number_allocates_id_and_echoes_number ) {
 	ASSERT_EQ( static_cast<size_t>( 1 ), protocol.takePresentation()->placements.size() );
 }
 
-UTEST( eterm, kitty_graphics_rgb_and_zlib_normalize_to_rgba ) {
+UTEST( eterm, kitty_graphics_rgb_and_zlib_preserve_rgb24 ) {
 	const std::vector<Uint8> rgb{ 10, 20, 30, 40, 50, 60 };
 	std::vector<Uint8> compressed( Compression::getMaxCompressedBufferSize( rgb.size() ) );
 	IOStreamMemory source( reinterpret_cast<const char*>( rgb.data() ), rgb.size() );
@@ -578,8 +571,12 @@ UTEST( eterm, kitty_graphics_rgb_and_zlib_normalize_to_rgba ) {
 	EXPECT_TRUE( result.changed );
 	auto pixels = protocol.imagePixels( 9 );
 	ASSERT_TRUE( pixels != nullptr );
-	const std::vector<Uint8> expected{ 10, 20, 30, 255, 40, 50, 60, 255 };
+	const std::vector<Uint8> expected{ 10, 20, 30, 40, 50, 60 };
 	EXPECT_TRUE( expected == *pixels );
+	auto updates = protocol.takeUpdates();
+	ASSERT_EQ( static_cast<size_t>( 1 ), updates.size() );
+	EXPECT_EQ( static_cast<Uint8>( 3 ), updates.front().channels );
+	EXPECT_TRUE( updates.front().pixels && expected == *updates.front().pixels );
 }
 
 UTEST( eterm, kitty_graphics_png_decodes_to_rgba ) {
@@ -652,8 +649,8 @@ UTEST( eterm, kitty_graphics_resync_republishes_authoritative_images ) {
 	EXPECT_EQ( TerminalGraphicsUpdateType::ResetAll, updates[0].type );
 	EXPECT_EQ( TerminalGraphicsUpdateType::CreateImage, updates[1].type );
 	EXPECT_EQ( static_cast<KittyImageId>( 27 ), updates[1].imageId );
-	ASSERT_TRUE( updates[1].rgba != nullptr );
-	EXPECT_EQ( static_cast<size_t>( 4 ), updates[1].rgba->size() );
+	ASSERT_TRUE( updates[1].pixels != nullptr );
+	EXPECT_EQ( static_cast<size_t>( 4 ), updates[1].pixels->size() );
 }
 
 UTEST( eterm, kitty_graphics_root_frame_patch_publishes_only_changed_rectangle ) {
@@ -673,8 +670,28 @@ UTEST( eterm, kitty_graphics_root_frame_patch_publishes_only_changed_rectangle )
 	EXPECT_EQ( 0, updates[0].region.Top );
 	EXPECT_EQ( 2, updates[0].region.Right );
 	EXPECT_EQ( 1, updates[0].region.Bottom );
-	ASSERT_TRUE( updates[0].rgba != nullptr );
-	EXPECT_EQ( static_cast<size_t>( 4 ), updates[0].rgba->size() );
+	ASSERT_TRUE( updates[0].pixels != nullptr );
+	EXPECT_EQ( static_cast<size_t>( 4 ), updates[0].pixels->size() );
+}
+
+UTEST( eterm, kitty_graphics_rgb24_root_converts_only_when_mutated ) {
+	KittyGraphicsProtocol protocol;
+	ASSERT_EQ( KittyGraphicsError::None,
+			   protocol.handle( "a=t,f=24,s=2,v=1,i=92;AQIDBAUG" ).error );
+	auto updates = protocol.takeUpdates();
+	ASSERT_EQ( static_cast<size_t>( 1 ), updates.size() );
+	EXPECT_EQ( static_cast<Uint8>( 3 ), updates[0].channels );
+
+	ASSERT_EQ( KittyGraphicsError::None,
+			   protocol.handle( "a=f,i=92,r=1,f=32,s=1,v=1,x=1,y=0,X=1;BwgJCg==" ).error );
+	const std::vector<Uint8> expected{ 1, 2, 3, 255, 7, 8, 9, 10 };
+	ASSERT_TRUE( protocol.imagePixels( 92 ) != nullptr );
+	EXPECT_TRUE( expected == *protocol.imagePixels( 92 ) );
+	updates = protocol.takeUpdates();
+	ASSERT_EQ( static_cast<size_t>( 2 ), updates.size() );
+	EXPECT_EQ( TerminalGraphicsUpdateType::ReplaceImage, updates[0].type );
+	EXPECT_EQ( static_cast<Uint8>( 4 ), updates[0].channels );
+	EXPECT_EQ( TerminalGraphicsUpdateType::UpdateRegion, updates[1].type );
 }
 
 UTEST( eterm, kitty_graphics_animation_frame_create_control_and_compose ) {
@@ -688,9 +705,9 @@ UTEST( eterm, kitty_graphics_animation_frame_create_control_and_compose ) {
 	ASSERT_EQ( static_cast<size_t>( 1 ), updates.size() );
 	EXPECT_EQ( TerminalGraphicsUpdateType::CreateFrame, updates[0].type );
 	EXPECT_EQ( static_cast<Uint32>( 2 ), updates[0].frameNumber );
-	ASSERT_TRUE( updates[0].rgba != nullptr );
+	ASSERT_TRUE( updates[0].pixels != nullptr );
 	const std::vector<Uint8> expectedFrame{ 1, 2, 3, 4, 9, 10, 11, 12 };
-	EXPECT_TRUE( expectedFrame == *updates[0].rgba );
+	EXPECT_TRUE( expectedFrame == *updates[0].pixels );
 
 	auto control = protocol.handle( "a=a,i=30,c=2" );
 	EXPECT_TRUE( control.changed );
@@ -703,9 +720,9 @@ UTEST( eterm, kitty_graphics_animation_frame_create_control_and_compose ) {
 	updates = protocol.takeUpdates();
 	ASSERT_EQ( static_cast<size_t>( 1 ), updates.size() );
 	EXPECT_EQ( TerminalGraphicsUpdateType::UpdateFrameRegion, updates[0].type );
-	ASSERT_TRUE( updates[0].rgba != nullptr );
+	ASSERT_TRUE( updates[0].pixels != nullptr );
 	const std::vector<Uint8> expectedPatch{ 5, 6, 7, 8 };
-	EXPECT_TRUE( expectedPatch == *updates[0].rgba );
+	EXPECT_TRUE( expectedPatch == *updates[0].pixels );
 	EXPECT_TRUE( protocol.handle( "a=a,i=30,c=1,r=1,z=-1,s=3" ).changed );
 	EXPECT_TRUE( protocol.updateAnimations() );
 	presentation = protocol.takePresentation();
@@ -977,6 +994,25 @@ UTEST( eterm, kitty_graphics_strict_base64_rejects_invalid_payload_bytes ) {
 	EXPECT_EQ( KittyGraphicsError::InvalidData,
 			   protocol.handle( "a=t,f=32,s=1,v=1;AQ=DBA==" ).error );
 	EXPECT_EQ( KittyGraphicsError::InvalidData, protocol.handle( "a=t,f=32,s=1,v=1;AB==" ).error );
+	EXPECT_EQ( KittyGraphicsError::InvalidData,
+			   protocol.handle( "a=t,f=32,s=1,v=1,m=1;AQ==" ).error );
+	EXPECT_EQ( KittyGraphicsError::InvalidData,
+			   protocol.handle( "a=t,f=32,s=1,v=1,m=1;AQI!" ).error );
+}
+
+UTEST( eterm, kitty_graphics_strict_base64_validates_final_quantum ) {
+	auto decode = []( std::string_view encoded ) {
+		std::vector<Uint8> output( Base64::decodeSafeOutLen( encoded.size() ) );
+		return Base64::decode( encoded.size(), encoded.data(), output.size(), output.data(),
+							   Base64::DecodeMode::NoWhitespaceStrict );
+	};
+	EXPECT_EQ( static_cast<size_t>( 1 ), decode( "AA" ) );
+	EXPECT_EQ( static_cast<size_t>( 2 ), decode( "AAA" ) );
+	EXPECT_EQ( static_cast<size_t>( 1 ), decode( "AA==" ) );
+	EXPECT_EQ( static_cast<size_t>( 2 ), decode( "AAA=" ) );
+	for ( std::string_view invalid :
+		  { "A", "AB", "AAB", "AB==", "AAB=", "====", "A===", "AA=A", "AAAA=", "AAAA====" } )
+		EXPECT_EQ( static_cast<size_t>( -1 ), decode( invalid ) );
 }
 
 UTEST( eterm, kitty_graphics_strict_base64_decodes_boundaries_and_large_payloads ) {
@@ -1032,36 +1068,15 @@ UTEST( eterm, kitty_graphics_reuses_unreferenced_replacement_pixel_storage ) {
 	ASSERT_EQ( KittyGraphicsError::None,
 			   protocol.handle( "a=t,f=24,s=2,v=1,i=91,q=2;BwgJCgsM" ).error );
 	EXPECT_TRUE( storage == protocol.imagePixels( 91 ) );
-	const std::vector<Uint8> expected{ 7, 8, 9, 255, 10, 11, 12, 255 };
+	const std::vector<Uint8> expected{ 7, 8, 9, 10, 11, 12 };
 	EXPECT_TRUE( expected == *protocol.imagePixels( 91 ) );
 }
 
-#if EE_PLATFORM != EE_PLATFORM_WIN && EE_PLATFORM != EE_PLATFORM_EMSCRIPTEN && \
-	EE_PLATFORM != EE_PLATFORM_ANDROID
-UTEST( eterm, kitty_graphics_reads_and_unlinks_posix_shared_memory ) {
-	// mpv uses the Linux-compatible form without the optional leading slash.
-	const std::string name = "eterm-kitty-unit-" + std::to_string( getpid() );
-	shm_unlink( name.c_str() );
-	const int descriptor = shm_open( name.c_str(), O_CREAT | O_EXCL | O_RDWR, 0600 );
-	ASSERT_TRUE( descriptor >= 0 );
-	const Uint8 stored[] = { 99, 98, 1, 2, 3 };
-	ASSERT_EQ( static_cast<ssize_t>( sizeof( stored ) ),
-			   write( descriptor, stored, sizeof( stored ) ) );
-	close( descriptor );
-
-	std::string encodedName;
-	ASSERT_TRUE( Base64::encode( name, encodedName ) );
+UTEST( eterm, kitty_graphics_rejects_shared_memory_transmission ) {
 	KittyGraphicsProtocol protocol;
-	EXPECT_EQ( KittyGraphicsError::None,
-			   protocol.handle( "a=T,t=s,f=24,s=1,v=1,O=2,S=3,q=2,m=1;" + encodedName ).error );
-	auto updates = protocol.takeUpdates();
-	ASSERT_EQ( static_cast<size_t>( 1 ), updates.size() );
-	ASSERT_TRUE( updates[0].rgba != nullptr );
-	const std::vector<Uint8> expected{ 1, 2, 3, 255 };
-	EXPECT_TRUE( expected == *updates[0].rgba );
-	EXPECT_EQ( -1, shm_open( name.c_str(), O_RDONLY, 0 ) );
+	EXPECT_EQ( KittyGraphicsError::Unsupported,
+			   protocol.handle( "a=T,t=s,f=24,s=1,v=1,q=2;L2VlcHAtc2ht" ).error );
 }
-#endif
 
 UTEST( eterm, kitty_graphics_accepts_unchunked_direct_image_larger_than_eight_kibibytes ) {
 	std::vector<Uint8> rgb( 64 * 64 * 3 );
diff --git a/src/tests/unit_tests/gitconflict_tests.cpp b/src/tests/unit_tests/gitconflict_tests.cpp
index 9f574cee8..4ec037a0f 100644
--- a/src/tests/unit_tests/gitconflict_tests.cpp
+++ b/src/tests/unit_tests/gitconflict_tests.cpp
@@ -83,6 +83,7 @@ UTEST( GitConflict, KeepsMergeOperationAfterAllConflictsAreStaged ) {
 	ASSERT_EQ( EXIT_SUCCESS, run( { "init", "-b", "main" } ) );
 	ASSERT_EQ( EXIT_SUCCESS, run( { "config", "user.name", "eepp tests" } ) );
 	ASSERT_EQ( EXIT_SUCCESS, run( { "config", "user.email", "eepp-tests@example.invalid" } ) );
+	ASSERT_EQ( EXIT_SUCCESS, run( { "config", "rerere.enabled", "false" } ) );
 	const std::string file = ( temp.path / "conflict.txt" ).string();
 	ASSERT_TRUE( FileSystem::fileWrite( file, "base\n" ) );
 	ASSERT_EQ( EXIT_SUCCESS, run( { "add", "conflict.txt" } ) );