diff --git a/Components/6560/6560.hpp b/Components/6560/6560.hpp
index 76f2c7d4d..fbc61bf94 100644
--- a/Components/6560/6560.hpp
+++ b/Components/6560/6560.hpp
@@ -445,20 +445,20 @@ template <class BusHandler> class MOS6560 {
 		// register state
 		struct {
 			bool interlaced = false, tall_characters = false;
-			uint8_t first_column_location, first_row_location;
-			uint8_t number_of_columns, number_of_rows;
-			uint16_t character_cell_start_address, video_matrix_start_address;
-			uint16_t backgroundColour, borderColour, auxiliary_colour;
+			uint8_t first_column_location = 0, first_row_location = 0;
+			uint8_t number_of_columns = 0, number_of_rows = 0;
+			uint16_t character_cell_start_address = 0, video_matrix_start_address = 0;
+			uint16_t backgroundColour = 0, borderColour = 0, auxiliary_colour = 0;
 			bool invertedCells = false;
 
-			uint8_t direct_values[16];
+			uint8_t direct_values[16] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 };
 		} registers_;
 
 		// output state
 		enum State {
 			Sync, ColourBurst, Border, Pixels
-		} this_state_, output_state_;
-		int cycles_in_state_;
+		} this_state_ = State::Sync, output_state_ = State::Sync;
+		int cycles_in_state_ = 0;
 
 		// counters that cover an entire field
 		int horizontal_counter_ = 0, vertical_counter_ = 0;
@@ -487,23 +487,23 @@ template <class BusHandler> class MOS6560 {
 
 		// latches dictating start and length of drawing
 		bool vertical_drawing_latch_ = false, horizontal_drawing_latch_ = false;
-		int rows_this_field_, columns_this_line_;
+		int rows_this_field_ = 0, columns_this_line_ = 0;
 
 		// current drawing position counter
-		int pixel_line_cycle_, column_counter_;
-		int current_row_;
-		uint16_t current_character_row_;
-		uint16_t video_matrix_address_counter_, base_video_matrix_address_counter_;
+		int pixel_line_cycle_ = 0, column_counter_ = 0;
+		int current_row_ = 0;
+		uint16_t current_character_row_ = 0;
+		uint16_t video_matrix_address_counter_ = 0, base_video_matrix_address_counter_ = 0;
 
 		// data latched from the bus
-		uint8_t character_code_, character_colour_, character_value_;
+		uint8_t character_code_ = 0, character_colour_ = 0, character_value_ = 0;
 
 		bool is_odd_frame_ = false, is_odd_line_ = false;
 
 		// lookup table from 6560 colour index to appropriate PAL/NTSC value
-		uint16_t colours_[16];
+		uint16_t colours_[16] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 };
 
-		uint16_t *pixel_pointer;
+		uint16_t *pixel_pointer = nullptr;
 		void output_border(int number_of_cycles) {
 			uint16_t *colour_pointer = reinterpret_cast<uint16_t *>(crt_.begin_data(1));
 			if(colour_pointer) *colour_pointer = registers_.borderColour;
@@ -511,13 +511,13 @@ template <class BusHandler> class MOS6560 {
 		}
 
 		struct {
-			int cycles_per_line;
-			int line_counter_increment_offset;
-			int final_line_increment_position;
-			int lines_per_progressive_field;
-			bool supports_interlacing;
+			int cycles_per_line = 0;
+			int line_counter_increment_offset = 0;
+			int final_line_increment_position = 0;
+			int lines_per_progressive_field = 0;
+			bool supports_interlacing = 0;
 		} timing_;
-		OutputMode output_mode_;
+		OutputMode output_mode_ = OutputMode::NTSC;
 };
 
 }
diff --git a/Components/9918/9918.cpp b/Components/9918/9918.cpp
index 11512c7ae..07ff15c8f 100644
--- a/Components/9918/9918.cpp
+++ b/Components/9918/9918.cpp
@@ -190,7 +190,9 @@ void TMS9918::run_for(const HalfCycles cycles) {
 	int read_cycles_pool = int_cycles;
 
 	while(write_cycles_pool || read_cycles_pool) {
+#ifndef NDEBUG
 		LineBufferPointer backup = read_pointer_;
+#endif
 
 		if(write_cycles_pool) {
 			// Determine how much writing to do.
@@ -329,8 +331,10 @@ void TMS9918::run_for(const HalfCycles cycles) {
 		}
 
 
+#ifndef NDEBUG
 		assert(backup.row == read_pointer_.row && backup.column == read_pointer_.column);
 		backup = write_pointer_;
+#endif
 
 
 		if(read_cycles_pool) {
diff --git a/Inputs/Keyboard.hpp b/Inputs/Keyboard.hpp
index 56105b279..0b4a4120c 100644
--- a/Inputs/Keyboard.hpp
+++ b/Inputs/Keyboard.hpp
@@ -46,6 +46,8 @@ class Keyboard {
 		/// Constructs a Keyboard that declares itself to observe only members of @c observed_keys.
 		Keyboard(const std::set<Key> &observed_keys, const std::set<Key> &essential_modifiers);
 
+		virtual ~Keyboard() {}
+
 		// Host interface.
 
 		/// @returns @c true if the key press affects the machine; @c false otherwise.
diff --git a/Machines/Apple/AppleII/Video.hpp b/Machines/Apple/AppleII/Video.hpp
index d75fd3f4e..67124fae4 100644
--- a/Machines/Apple/AppleII/Video.hpp
+++ b/Machines/Apple/AppleII/Video.hpp
@@ -541,7 +541,17 @@ template <class BusHandler, bool is_iie> class Video: public VideoBase {
 						const int colour_burst_start = std::max(first_sync_column + sync_length + 1, column_);
 						const int colour_burst_end = std::min(first_sync_column + sync_length + 4, ending_column);
 						if(colour_burst_end > colour_burst_start) {
-							crt_.output_colour_burst((colour_burst_end - colour_burst_start) * 14, 0);
+							// UGLY HACK AHOY!
+							// The OpenGL scan target introduces a phase error of 1/8th of a wave. The Metal one does not.
+							// Supply the real phase value if this is an Apple build.
+							// TODO: eliminate UGLY HACK.
+#ifdef __APPLE__
+							constexpr int phase = 224;
+#else
+							constexpr int phase = 0;
+#endif
+
+							crt_.output_colour_burst((colour_burst_end - colour_burst_start) * 14, phase);
 						}
 
 						second_blank_start = std::max(first_sync_column + sync_length + 3, column_);
diff --git a/Machines/Apple/Macintosh/Video.cpp b/Machines/Apple/Macintosh/Video.cpp
index 211401e36..3c98374f0 100644
--- a/Machines/Apple/Macintosh/Video.cpp
+++ b/Machines/Apple/Macintosh/Video.cpp
@@ -29,7 +29,17 @@ Video::Video(DeferredAudio &audio, DriveSpeedAccumulator &drive_speed_accumulato
  	crt_(704, 1, 370, 6, Outputs::Display::InputDataType::Luminance1) {
 
  	crt_.set_display_type(Outputs::Display::DisplayType::RGB);
+
+	// UGLY HACK. UGLY, UGLY HACK. UGLY!
+	// The OpenGL scan target fails properly to place visible areas which are not 4:3.
+	// The [newer] Metal scan target has no such issue. So assume that Apple => Metal,
+	// and set a visible area to work around the OpenGL issue if required.
+	// TODO: eliminate UGLY HACK.
+#ifdef __APPLE__
+	crt_.set_visible_area(Outputs::Display::Rect(0.08f, 10.0f / 368.0f, 0.82f, 344.0f / 368.0f));
+#else
 	crt_.set_visible_area(Outputs::Display::Rect(0.08f, -0.025f, 0.82f, 0.82f));
+#endif
 	crt_.set_aspect_ratio(1.73f);	// The Mac uses a non-standard scanning area.
 }
 
@@ -105,10 +115,13 @@ void Video::run_for(HalfCycles duration) {
 
 							pixel_buffer_ += 16;
 						}
+					} else {
+						video_address_ += size_t(final_pixel_word - first_word);
 					}
 
 					if(final_pixel_word == 32) {
 						crt_.output_data(512);
+						pixel_buffer_ = nullptr;
 					}
 				}
 
diff --git a/OSBindings/Mac/Clock Signal.xcodeproj/project.pbxproj b/OSBindings/Mac/Clock Signal.xcodeproj/project.pbxproj
index 9f491a899..c4c11eb2c 100644
--- a/OSBindings/Mac/Clock Signal.xcodeproj/project.pbxproj	
+++ b/OSBindings/Mac/Clock Signal.xcodeproj/project.pbxproj	
@@ -98,7 +98,6 @@
 		4B055AED1FAE9BA20060FFFF /* Z80Storage.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4B8334831F5DA0360097E338 /* Z80Storage.cpp */; };
 		4B055AEE1FAE9BBF0060FFFF /* Keyboard.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4B86E2591F8C628F006FAA45 /* Keyboard.cpp */; };
 		4B055AEF1FAE9BF00060FFFF /* Typer.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4B2B3A471F9B8FA70062DABF /* Typer.cpp */; };
-		4B055AF11FAE9C160060FFFF /* Accelerate.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 4BC76E6A1C98F43700E6EF73 /* Accelerate.framework */; };
 		4B055AF21FAE9C1C0060FFFF /* OpenGL.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 4B055AF01FAE9C080060FFFF /* OpenGL.framework */; };
 		4B08A2751EE35D56008B7065 /* Z80InterruptTests.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4B08A2741EE35D56008B7065 /* Z80InterruptTests.swift */; };
 		4B08A2781EE39306008B7065 /* TestMachine.mm in Sources */ = {isa = PBXBuildFile; fileRef = 4B08A2771EE39306008B7065 /* TestMachine.mm */; };
@@ -155,6 +154,9 @@
 		4B1D08061E0F7A1100763741 /* TimeTests.mm in Sources */ = {isa = PBXBuildFile; fileRef = 4B1D08051E0F7A1100763741 /* TimeTests.mm */; };
 		4B1E85811D176468001EF87D /* 6532Tests.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4B1E85801D176468001EF87D /* 6532Tests.swift */; };
 		4B1EDB451E39A0AC009D6819 /* chip.png in Resources */ = {isa = PBXBuildFile; fileRef = 4B1EDB431E39A0AC009D6819 /* chip.png */; };
+		4B228CD524D773B40077EF25 /* CSScanTarget.mm in Sources */ = {isa = PBXBuildFile; fileRef = 4B228CD424D773B30077EF25 /* CSScanTarget.mm */; };
+		4B228CD924DA12C60077EF25 /* CSScanTargetView.m in Sources */ = {isa = PBXBuildFile; fileRef = 4B228CD824DA12C60077EF25 /* CSScanTargetView.m */; };
+		4B228CDB24DA41890077EF25 /* ScanTarget.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4B228CDA24DA41880077EF25 /* ScanTarget.metal */; };
 		4B2530F4244E6774007980BF /* fm.json in Resources */ = {isa = PBXBuildFile; fileRef = 4B2530F3244E6773007980BF /* fm.json */; };
 		4B2A332D1DB86821002876E3 /* OricOptions.xib in Resources */ = {isa = PBXBuildFile; fileRef = 4B2A332B1DB86821002876E3 /* OricOptions.xib */; };
 		4B2A539F1D117D36003C6002 /* CSAudioQueue.m in Sources */ = {isa = PBXBuildFile; fileRef = 4B2A53911D117D36003C6002 /* CSAudioQueue.m */; };
@@ -215,7 +217,6 @@
 		4B54C0C51F8D91D90050900F /* Keyboard.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4B54C0C41F8D91D90050900F /* Keyboard.cpp */; };
 		4B54C0C81F8D91E50050900F /* Keyboard.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4B54C0C61F8D91E50050900F /* Keyboard.cpp */; };
 		4B54C0CB1F8D92590050900F /* Keyboard.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4B54C0CA1F8D92580050900F /* Keyboard.cpp */; };
-		4B55CE5D1C3B7D6F0093A61B /* CSOpenGLView.m in Sources */ = {isa = PBXBuildFile; fileRef = 4B55CE5C1C3B7D6F0093A61B /* CSOpenGLView.m */; };
 		4B55CE5F1C3B7D960093A61B /* MachineDocument.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4B55CE5E1C3B7D960093A61B /* MachineDocument.swift */; };
 		4B55DD8320DF06680043F2E5 /* MachinePicker.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4B55DD8020DF06680043F2E5 /* MachinePicker.swift */; };
 		4B55DD8420DF06680043F2E5 /* MachinePicker.xib in Resources */ = {isa = PBXBuildFile; fileRef = 4B55DD8120DF06680043F2E5 /* MachinePicker.xib */; };
@@ -370,7 +371,6 @@
 		4B778F6123A5F3560000D260 /* Disk.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4B8944FC201967B4007DE474 /* Disk.cpp */; };
 		4B778F6223A5F35F0000D260 /* File.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4B894500201967B4007DE474 /* File.cpp */; };
 		4B778F6323A5F3630000D260 /* Tape.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4B894501201967B4007DE474 /* Tape.cpp */; };
-		4B778F6423A5F3730000D260 /* Accelerate.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 4BC76E6A1C98F43700E6EF73 /* Accelerate.framework */; };
 		4B7913CC1DFCD80E00175A82 /* Video.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4B7913CA1DFCD80E00175A82 /* Video.cpp */; };
 		4B79A5011FC913C900EEDAD5 /* MSX.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4B79A4FF1FC913C900EEDAD5 /* MSX.cpp */; };
 		4B79E4441E3AF38600141F11 /* cassette.png in Resources */ = {isa = PBXBuildFile; fileRef = 4B79E4411E3AF38600141F11 /* cassette.png */; };
@@ -768,6 +768,10 @@
 		4BB73EAC1B587A5100552FC2 /* MainMenu.xib in Resources */ = {isa = PBXBuildFile; fileRef = 4BB73EAA1B587A5100552FC2 /* MainMenu.xib */; };
 		4BB73EB71B587A5100552FC2 /* AllSuiteATests.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4BB73EB61B587A5100552FC2 /* AllSuiteATests.swift */; };
 		4BB73EC21B587A5100552FC2 /* Clock_SignalUITests.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4BB73EC11B587A5100552FC2 /* Clock_SignalUITests.swift */; };
+		4BB8616E24E22DC500A00E03 /* BufferingScanTarget.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BB8616D24E22DC500A00E03 /* BufferingScanTarget.cpp */; };
+		4BB8616F24E22DC500A00E03 /* BufferingScanTarget.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BB8616D24E22DC500A00E03 /* BufferingScanTarget.cpp */; };
+		4BB8617124E22F5700A00E03 /* Accelerate.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 4BB8617024E22F4900A00E03 /* Accelerate.framework */; };
+		4BB8617224E22F5A00A00E03 /* Accelerate.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 4BB8617024E22F4900A00E03 /* Accelerate.framework */; };
 		4BBB70A4202011C2002FE009 /* MultiMediaTarget.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BBB70A3202011C2002FE009 /* MultiMediaTarget.cpp */; };
 		4BBB70A5202011C2002FE009 /* MultiMediaTarget.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BBB70A3202011C2002FE009 /* MultiMediaTarget.cpp */; };
 		4BBB70A8202014E2002FE009 /* MultiProducer.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BBB70A6202014E2002FE009 /* MultiProducer.cpp */; };
@@ -786,15 +790,12 @@
 		4BC1317B2346DF2B00E4FF3D /* MSA.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BC131782346DF2B00E4FF3D /* MSA.cpp */; };
 		4BC23A2C2467600F001A6030 /* OPLL.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BC23A2B2467600E001A6030 /* OPLL.cpp */; };
 		4BC23A2D2467600F001A6030 /* OPLL.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BC23A2B2467600E001A6030 /* OPLL.cpp */; };
-		4BC3C67C24C9230F0027BF76 /* BufferingScanTarget.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BC3C67A24C9230F0027BF76 /* BufferingScanTarget.cpp */; };
-		4BC3C67D24C9230F0027BF76 /* BufferingScanTarget.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BC3C67A24C9230F0027BF76 /* BufferingScanTarget.cpp */; };
 		4BC57CD92436A62900FBC404 /* State.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BC57CD82436A62900FBC404 /* State.cpp */; };
 		4BC57CDA2436A62900FBC404 /* State.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BC57CD82436A62900FBC404 /* State.cpp */; };
 		4BC5C3E022C994CD00795658 /* 68000MoveTests.mm in Sources */ = {isa = PBXBuildFile; fileRef = 4BC5C3DF22C994CC00795658 /* 68000MoveTests.mm */; };
 		4BC5FC3020CDDDEF00410AA0 /* AppleIIOptions.xib in Resources */ = {isa = PBXBuildFile; fileRef = 4BC5FC2E20CDDDEE00410AA0 /* AppleIIOptions.xib */; };
 		4BC751B21D157E61006C31D9 /* 6522Tests.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4BC751B11D157E61006C31D9 /* 6522Tests.swift */; };
 		4BC76E691C98E31700E6EF73 /* FIRFilter.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BC76E671C98E31700E6EF73 /* FIRFilter.cpp */; };
-		4BC76E6B1C98F43700E6EF73 /* Accelerate.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = 4BC76E6A1C98F43700E6EF73 /* Accelerate.framework */; };
 		4BC890D3230F86020025A55A /* DirectAccessDevice.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BC890D1230F86020025A55A /* DirectAccessDevice.cpp */; };
 		4BC890D4230F86020025A55A /* DirectAccessDevice.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BC890D1230F86020025A55A /* DirectAccessDevice.cpp */; };
 		4BC91B831D1F160E00884B76 /* CommodoreTAP.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BC91B811D1F160E00884B76 /* CommodoreTAP.cpp */; };
@@ -811,19 +812,14 @@
 		4BCE0060227D39AB000CA200 /* Video.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BCE005E227D39AB000CA200 /* Video.cpp */; };
 		4BCF1FA41DADC3DD0039D2E7 /* Oric.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BCF1FA21DADC3DD0039D2E7 /* Oric.cpp */; };
 		4BD0FBC3233706A200148981 /* CSApplication.m in Sources */ = {isa = PBXBuildFile; fileRef = 4BD0FBC2233706A200148981 /* CSApplication.m */; };
-		4BD191F42191180E0042E144 /* ScanTarget.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BD191F22191180E0042E144 /* ScanTarget.cpp */; };
 		4BD191F52191180E0042E144 /* ScanTarget.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BD191F22191180E0042E144 /* ScanTarget.cpp */; };
 		4BD388882239E198002D14B5 /* 68000Tests.mm in Sources */ = {isa = PBXBuildFile; fileRef = 4BD388872239E198002D14B5 /* 68000Tests.mm */; };
 		4BD3A30B1EE755C800B5B501 /* Video.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BD3A3091EE755C800B5B501 /* Video.cpp */; };
-		4BD424DF2193B5340097291A /* TextureTarget.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BD424DD2193B5340097291A /* TextureTarget.cpp */; };
 		4BD424E02193B5340097291A /* TextureTarget.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BD424DD2193B5340097291A /* TextureTarget.cpp */; };
-		4BD424E52193B5830097291A /* Shader.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BD424E12193B5820097291A /* Shader.cpp */; };
 		4BD424E62193B5830097291A /* Shader.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BD424E12193B5820097291A /* Shader.cpp */; };
-		4BD424E72193B5830097291A /* Rectangle.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BD424E22193B5820097291A /* Rectangle.cpp */; };
 		4BD424E82193B5830097291A /* Rectangle.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BD424E22193B5820097291A /* Rectangle.cpp */; };
 		4BD468F71D8DF41D0084958B /* 1770.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BD468F51D8DF41D0084958B /* 1770.cpp */; };
 		4BD4A8D01E077FD20020D856 /* PCMTrackTests.mm in Sources */ = {isa = PBXBuildFile; fileRef = 4BD4A8CF1E077FD20020D856 /* PCMTrackTests.mm */; };
-		4BD5D2682199148100DDF17D /* ScanTargetGLSLFragments.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BD5D2672199148100DDF17D /* ScanTargetGLSLFragments.cpp */; };
 		4BD5D2692199148100DDF17D /* ScanTargetGLSLFragments.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BD5D2672199148100DDF17D /* ScanTargetGLSLFragments.cpp */; };
 		4BD61664206B2AC800236112 /* QuickLoadOptions.xib in Resources */ = {isa = PBXBuildFile; fileRef = 4BD61662206B2AC700236112 /* QuickLoadOptions.xib */; };
 		4BD67DCB209BE4D700AB2146 /* StaticAnalyser.cpp in Sources */ = {isa = PBXBuildFile; fileRef = 4BD67DCA209BE4D600AB2146 /* StaticAnalyser.cpp */; };
@@ -1004,6 +1000,11 @@
 		4B1E857B1D174DEC001EF87D /* 6532.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; path = 6532.hpp; sourceTree = "<group>"; };
 		4B1E85801D176468001EF87D /* 6532Tests.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = 6532Tests.swift; sourceTree = "<group>"; };
 		4B1EDB431E39A0AC009D6819 /* chip.png */ = {isa = PBXFileReference; lastKnownFileType = image.png; path = chip.png; sourceTree = "<group>"; };
+		4B228CD424D773B30077EF25 /* CSScanTarget.mm */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.objcpp; path = CSScanTarget.mm; sourceTree = "<group>"; };
+		4B228CD624D773CA0077EF25 /* CSScanTarget.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = CSScanTarget.h; sourceTree = "<group>"; };
+		4B228CD724DA12C50077EF25 /* CSScanTargetView.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = CSScanTargetView.h; sourceTree = "<group>"; };
+		4B228CD824DA12C60077EF25 /* CSScanTargetView.m */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.objc; path = CSScanTargetView.m; sourceTree = "<group>"; };
+		4B228CDA24DA41880077EF25 /* ScanTarget.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ScanTarget.metal; sourceTree = "<group>"; xcLanguageSpecificationIdentifier = xcode.lang.metal; };
 		4B24095A1C45DF85004DA684 /* Stepper.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; path = Stepper.hpp; sourceTree = "<group>"; };
 		4B2530F3244E6773007980BF /* fm.json */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text.json; path = fm.json; sourceTree = "<group>"; };
 		4B2A332C1DB86821002876E3 /* Base */ = {isa = PBXFileReference; lastKnownFileType = file.xib; name = Base; path = "Clock Signal/Base.lproj/OricOptions.xib"; sourceTree = SOURCE_ROOT; };
@@ -1115,6 +1116,7 @@
 		4B4DC8271D2C2470003C5BF8 /* C1540.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; path = C1540.hpp; sourceTree = "<group>"; };
 		4B4DC8291D2C27A4003C5BF8 /* SerialBus.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = SerialBus.cpp; sourceTree = "<group>"; };
 		4B4DC82A1D2C27A4003C5BF8 /* SerialBus.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; path = SerialBus.hpp; sourceTree = "<group>"; };
+		4B4F2B7024DF99D4000DA6B0 /* CSScanTarget+CppScanTarget.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = "CSScanTarget+CppScanTarget.h"; sourceTree = "<group>"; };
 		4B50AF7F242817F40099BBD7 /* QuartzCore.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = QuartzCore.framework; path = System/Library/Frameworks/QuartzCore.framework; sourceTree = SDKROOT; };
 		4B51F70920A521D700AFA2C1 /* Source.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; path = Source.hpp; sourceTree = "<group>"; };
 		4B51F70A20A521D700AFA2C1 /* Observer.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; path = Observer.hpp; sourceTree = "<group>"; };
@@ -1129,8 +1131,6 @@
 		4B54C0C71F8D91E50050900F /* Keyboard.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; name = Keyboard.hpp; path = Electron/Keyboard.hpp; sourceTree = "<group>"; };
 		4B54C0C91F8D92580050900F /* Keyboard.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; name = Keyboard.hpp; path = ZX8081/Keyboard.hpp; sourceTree = "<group>"; };
 		4B54C0CA1F8D92580050900F /* Keyboard.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; name = Keyboard.cpp; path = ZX8081/Keyboard.cpp; sourceTree = "<group>"; };
-		4B55CE5B1C3B7D6F0093A61B /* CSOpenGLView.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; path = CSOpenGLView.h; sourceTree = "<group>"; };
-		4B55CE5C1C3B7D6F0093A61B /* CSOpenGLView.m */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.objc; path = CSOpenGLView.m; sourceTree = "<group>"; };
 		4B55CE5E1C3B7D960093A61B /* MachineDocument.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = MachineDocument.swift; sourceTree = "<group>"; };
 		4B55DD8020DF06680043F2E5 /* MachinePicker.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = MachinePicker.swift; sourceTree = "<group>"; };
 		4B55DD8220DF06680043F2E5 /* Base */ = {isa = PBXFileReference; lastKnownFileType = file.xib; name = Base; path = Base.lproj/MachinePicker.xib; sourceTree = "<group>"; };
@@ -1638,6 +1638,9 @@
 		4BB73EC11B587A5100552FC2 /* Clock_SignalUITests.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Clock_SignalUITests.swift; sourceTree = "<group>"; };
 		4BB73EC31B587A5100552FC2 /* Info.plist */ = {isa = PBXFileReference; lastKnownFileType = text.plist.xml; path = Info.plist; sourceTree = "<group>"; };
 		4BB73ECF1B587A6700552FC2 /* Clock Signal.entitlements */ = {isa = PBXFileReference; lastKnownFileType = text.xml; path = "Clock Signal.entitlements"; sourceTree = "<group>"; };
+		4BB8616C24E22DC500A00E03 /* BufferingScanTarget.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; path = BufferingScanTarget.hpp; sourceTree = "<group>"; };
+		4BB8616D24E22DC500A00E03 /* BufferingScanTarget.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = BufferingScanTarget.cpp; sourceTree = "<group>"; };
+		4BB8617024E22F4900A00E03 /* Accelerate.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = Accelerate.framework; path = System/Library/Frameworks/Accelerate.framework; sourceTree = SDKROOT; };
 		4BBB709C2020109C002FE009 /* DynamicMachine.hpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.h; path = DynamicMachine.hpp; sourceTree = "<group>"; };
 		4BBB70A2202011C2002FE009 /* MultiMediaTarget.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; path = MultiMediaTarget.hpp; sourceTree = "<group>"; };
 		4BBB70A3202011C2002FE009 /* MultiMediaTarget.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = MultiMediaTarget.cpp; sourceTree = "<group>"; };
@@ -1670,8 +1673,6 @@
 		4BC23A292467600E001A6030 /* OPLBase.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; path = OPLBase.hpp; sourceTree = "<group>"; };
 		4BC23A2A2467600E001A6030 /* EnvelopeGenerator.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; path = EnvelopeGenerator.hpp; sourceTree = "<group>"; };
 		4BC23A2B2467600E001A6030 /* OPLL.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = OPLL.cpp; sourceTree = "<group>"; };
-		4BC3C67A24C9230F0027BF76 /* BufferingScanTarget.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; name = BufferingScanTarget.cpp; path = ../../Outputs/ScanTargets/BufferingScanTarget.cpp; sourceTree = "<group>"; };
-		4BC3C67B24C9230F0027BF76 /* BufferingScanTarget.hpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.h; name = BufferingScanTarget.hpp; path = ../../Outputs/ScanTargets/BufferingScanTarget.hpp; sourceTree = "<group>"; };
 		4BC57CD2243427C700FBC404 /* AudioProducer.hpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.h; path = AudioProducer.hpp; sourceTree = "<group>"; };
 		4BC57CD32434282000FBC404 /* TimedMachine.hpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.h; path = TimedMachine.hpp; sourceTree = "<group>"; };
 		4BC57CD424342E0600FBC404 /* MachineTypes.hpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.h; path = MachineTypes.hpp; sourceTree = "<group>"; };
@@ -1682,7 +1683,6 @@
 		4BC751B11D157E61006C31D9 /* 6522Tests.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = 6522Tests.swift; sourceTree = "<group>"; };
 		4BC76E671C98E31700E6EF73 /* FIRFilter.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = FIRFilter.cpp; sourceTree = "<group>"; };
 		4BC76E681C98E31700E6EF73 /* FIRFilter.hpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.h; path = FIRFilter.hpp; sourceTree = "<group>"; };
-		4BC76E6A1C98F43700E6EF73 /* Accelerate.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = Accelerate.framework; path = System/Library/Frameworks/Accelerate.framework; sourceTree = SDKROOT; };
 		4BC890D1230F86020025A55A /* DirectAccessDevice.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = DirectAccessDevice.cpp; sourceTree = "<group>"; };
 		4BC890D2230F86020025A55A /* DirectAccessDevice.hpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.h; path = DirectAccessDevice.hpp; sourceTree = "<group>"; };
 		4BC91B811D1F160E00884B76 /* CommodoreTAP.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; path = CommodoreTAP.cpp; sourceTree = "<group>"; };
@@ -1809,7 +1809,7 @@
 			buildActionMask = 2147483647;
 			files = (
 				4B055AF21FAE9C1C0060FFFF /* OpenGL.framework in Frameworks */,
-				4B055AF11FAE9C160060FFFF /* Accelerate.framework in Frameworks */,
+				4BB8617224E22F5A00A00E03 /* Accelerate.framework in Frameworks */,
 				4B055ABD1FAE86530060FFFF /* libz.tbd in Frameworks */,
 				4B055A7A1FAE78A00060FFFF /* SDL2.framework in Frameworks */,
 			);
@@ -1819,8 +1819,8 @@
 			isa = PBXFrameworksBuildPhase;
 			buildActionMask = 2147483647;
 			files = (
+				4BB8617124E22F5700A00E03 /* Accelerate.framework in Frameworks */,
 				4B50AF80242817F40099BBD7 /* QuartzCore.framework in Frameworks */,
-				4BC76E6B1C98F43700E6EF73 /* Accelerate.framework in Frameworks */,
 				4B69FB461C4D950F00B5F0AA /* libz.tbd in Frameworks */,
 			);
 			runOnlyForDeploymentPostprocessing = 0;
@@ -1830,7 +1830,6 @@
 			buildActionMask = 2147483647;
 			files = (
 				4B9F11CA2272433900701480 /* libz.tbd in Frameworks */,
-				4B778F6423A5F3730000D260 /* Accelerate.framework in Frameworks */,
 			);
 			runOnlyForDeploymentPostprocessing = 0;
 		};
@@ -1847,6 +1846,7 @@
 		4B055A761FAE78210060FFFF /* Frameworks */ = {
 			isa = PBXGroup;
 			children = (
+				4BB8617024E22F4900A00E03 /* Accelerate.framework */,
 				4B50AF7F242817F40099BBD7 /* QuartzCore.framework */,
 				4B055AF01FAE9C080060FFFF /* OpenGL.framework */,
 				4B055A771FAE78210060FFFF /* SDL2.framework */,
@@ -2060,6 +2060,17 @@
 			path = Icons;
 			sourceTree = "<group>";
 		};
+		4B228CD324D773B30077EF25 /* ScanTarget */ = {
+			isa = PBXGroup;
+			children = (
+				4B228CDA24DA41880077EF25 /* ScanTarget.metal */,
+				4B228CD424D773B30077EF25 /* CSScanTarget.mm */,
+				4B228CD624D773CA0077EF25 /* CSScanTarget.h */,
+				4B4F2B7024DF99D4000DA6B0 /* CSScanTarget+CppScanTarget.h */,
+			);
+			path = ScanTarget;
+			sourceTree = "<group>";
+		};
 		4B2409591C45DF85004DA684 /* SignalProcessing */ = {
 			isa = PBXGroup;
 			children = (
@@ -2223,6 +2234,7 @@
 				4BF52672218E752E00313227 /* ScanTarget.hpp */,
 				4B0CCC411C62D0B3001CAC5F /* CRT */,
 				4BD191D5219113B80042E144 /* OpenGL */,
+				4BB8616B24E22DC500A00E03 /* ScanTargets */,
 				4BD060A41FE49D3C006E14BE /* Speaker */,
 			);
 			name = Outputs;
@@ -2471,8 +2483,8 @@
 		4B55CE5A1C3B7D6F0093A61B /* Views */ = {
 			isa = PBXGroup;
 			children = (
-				4B55CE5B1C3B7D6F0093A61B /* CSOpenGLView.h */,
-				4B55CE5C1C3B7D6F0093A61B /* CSOpenGLView.m */,
+				4B228CD724DA12C50077EF25 /* CSScanTargetView.h */,
+				4B228CD824DA12C60077EF25 /* CSScanTargetView.m */,
 			);
 			path = Views;
 			sourceTree = "<group>";
@@ -3311,9 +3323,6 @@
 		4BB73E951B587A5100552FC2 = {
 			isa = PBXGroup;
 			children = (
-				4BC3C67A24C9230F0027BF76 /* BufferingScanTarget.cpp */,
-				4BC3C67B24C9230F0027BF76 /* BufferingScanTarget.hpp */,
-				4BC76E6A1C98F43700E6EF73 /* Accelerate.framework */,
 				4B51F70820A521D700AFA2C1 /* Activity */,
 				4B8944E2201967B4007DE474 /* Analyser */,
 				4BB73EA01B587A5100552FC2 /* Clock Signal */,
@@ -3371,6 +3380,7 @@
 				4BB73EAA1B587A5100552FC2 /* MainMenu.xib */,
 				4BE5F85A1C3E1C2500C43F01 /* Resources */,
 				4BDA00DB22E60EE900AC3CD0 /* ROMRequester */,
+				4B228CD324D773B30077EF25 /* ScanTarget */,
 				4B55CE5A1C3B7D6F0093A61B /* Views */,
 			);
 			path = "Clock Signal";
@@ -3480,6 +3490,16 @@
 			path = ../../Processors;
 			sourceTree = "<group>";
 		};
+		4BB8616B24E22DC500A00E03 /* ScanTargets */ = {
+			isa = PBXGroup;
+			children = (
+				4BB8616C24E22DC500A00E03 /* BufferingScanTarget.hpp */,
+				4BB8616D24E22DC500A00E03 /* BufferingScanTarget.cpp */,
+			);
+			name = ScanTargets;
+			path = ../../Outputs/ScanTargets;
+			sourceTree = "<group>";
+		};
 		4BBB70A1202011C2002FE009 /* Implementation */ = {
 			isa = PBXGroup;
 			children = (
@@ -4491,7 +4511,7 @@
 				4BEBFB522002DB30000708CC /* DiskROM.cpp in Sources */,
 				4BC23A2D2467600F001A6030 /* OPLL.cpp in Sources */,
 				4B055AA11FAE85DA0060FFFF /* OricMFMDSK.cpp in Sources */,
-				4BC3C67D24C9230F0027BF76 /* BufferingScanTarget.cpp in Sources */,
+				4BB8616F24E22DC500A00E03 /* BufferingScanTarget.cpp in Sources */,
 				4B0ACC2923775819008902D0 /* DMAController.cpp in Sources */,
 				4B055A951FAE85BB0060FFFF /* BitReverse.cpp in Sources */,
 				4B055ACE1FAE9B030060FFFF /* Plus3.cpp in Sources */,
@@ -4574,11 +4594,11 @@
 				4B0E04EA1FC9E5DA00F43484 /* CAS.cpp in Sources */,
 				4B7A90ED20410A85008514A2 /* StaticAnalyser.cpp in Sources */,
 				4B58601E1F806AB200AEE2E3 /* MFMSectorDump.cpp in Sources */,
+				4B228CD924DA12C60077EF25 /* CSScanTargetView.m in Sources */,
 				4B6AAEAD230E40250078E864 /* Target.cpp in Sources */,
 				4B448E841F1C4C480009ABD6 /* PulseQueuedTape.cpp in Sources */,
 				4B0E61071FF34737002A9DBD /* MSX.cpp in Sources */,
 				4B4518A01F75FD1C00926311 /* CPCDSK.cpp in Sources */,
-				4BD424DF2193B5340097291A /* TextureTarget.cpp in Sources */,
 				4B0CCC451C62D0B3001CAC5F /* CRT.cpp in Sources */,
 				4BC23A2C2467600F001A6030 /* OPLL.cpp in Sources */,
 				4B322E041F5A2E3C004EB04C /* Z80Base.cpp in Sources */,
@@ -4604,6 +4624,7 @@
 				4B1497921EE4B5A800CE2596 /* ZX8081.cpp in Sources */,
 				4B643F3F1D77B88000D431D6 /* DocumentController.swift in Sources */,
 				4BDA00E422E663B900AC3CD0 /* NSData+CRC32.m in Sources */,
+				4BB8616E24E22DC500A00E03 /* BufferingScanTarget.cpp in Sources */,
 				4BB4BFB022A42F290069048D /* MacintoshIMG.cpp in Sources */,
 				4B05401E219D1618001BF69C /* ScanTarget.cpp in Sources */,
 				4B4518861F75E91A00926311 /* MFMDiskController.cpp in Sources */,
@@ -4616,13 +4637,11 @@
 				4B4DC82B1D2C27A4003C5BF8 /* SerialBus.cpp in Sources */,
 				4BBFFEE61F7B27F1005F3FEB /* TrackSerialiser.cpp in Sources */,
 				4BAE49582032881E004BE78E /* CSZX8081.mm in Sources */,
-				4BD424E52193B5830097291A /* Shader.cpp in Sources */,
 				4B0333AF2094081A0050B93D /* AppleDSK.cpp in Sources */,
 				4B894518201967B4007DE474 /* ConfidenceCounter.cpp in Sources */,
 				4BCE005A227CFFCA000CA200 /* Macintosh.cpp in Sources */,
 				4B6AAEA4230E3E1D0078E864 /* MassStorageDevice.cpp in Sources */,
 				4B89452E201967B4007DE474 /* StaticAnalyser.cpp in Sources */,
-				4BD5D2682199148100DDF17D /* ScanTargetGLSLFragments.cpp in Sources */,
 				4BC890D3230F86020025A55A /* DirectAccessDevice.cpp in Sources */,
 				4B7BA03723CEB86000B98D9E /* BD500.cpp in Sources */,
 				4B38F3481F2EC11D00D9235D /* AmstradCPC.cpp in Sources */,
@@ -4638,7 +4657,6 @@
 				4B3BF5B01F146265005B6C36 /* CSW.cpp in Sources */,
 				4BCE0060227D39AB000CA200 /* Video.cpp in Sources */,
 				4B0ACC2E23775819008902D0 /* TIA.cpp in Sources */,
-				4BC3C67C24C9230F0027BF76 /* BufferingScanTarget.cpp in Sources */,
 				4B74CF85231370BC00500CE8 /* MacintoshVolume.cpp in Sources */,
 				4B4518A51F75FD1C00926311 /* SSD.cpp in Sources */,
 				4B55CE5F1C3B7D960093A61B /* MachineDocument.swift in Sources */,
@@ -4672,7 +4690,8 @@
 				4B8334841F5DA0360097E338 /* Z80Storage.cpp in Sources */,
 				4BA61EB01D91515900B3C876 /* NSData+StdVector.mm in Sources */,
 				4BDA00E022E644AF00AC3CD0 /* CSROMReceiverView.m in Sources */,
-				4BD191F42191180E0042E144 /* ScanTarget.cpp in Sources */,
+				4B228CDB24DA41890077EF25 /* ScanTarget.metal in Sources */,
+				4B228CD524D773B40077EF25 /* CSScanTarget.mm in Sources */,
 				4BCD634922D6756400F567F1 /* MacintoshDoubleDensityDrive.cpp in Sources */,
 				4B0F94FE208C1A1600FE41D9 /* NIB.cpp in Sources */,
 				4B89452A201967B4007DE474 /* File.cpp in Sources */,
@@ -4699,7 +4718,6 @@
 				4B0E04FA1FC9FA3100F43484 /* 9918.cpp in Sources */,
 				4B69FB3D1C4D908A00B5F0AA /* Tape.cpp in Sources */,
 				4B4518841F75E91A00926311 /* UnformattedTrack.cpp in Sources */,
-				4B55CE5D1C3B7D6F0093A61B /* CSOpenGLView.m in Sources */,
 				4B65086022F4CF8D009C1100 /* Keyboard.cpp in Sources */,
 				4B894528201967B4007DE474 /* Disk.cpp in Sources */,
 				4BBB70A4202011C2002FE009 /* MultiMediaTarget.cpp in Sources */,
@@ -4756,7 +4774,6 @@
 				4B37EE821D7345A6006A09A4 /* BinaryDump.cpp in Sources */,
 				4BCE0053227CE8CA000CA200 /* AppleII.cpp in Sources */,
 				4B8334821F5D9FF70097E338 /* PartialMachineCycle.cpp in Sources */,
-				4BD424E72193B5830097291A /* Rectangle.cpp in Sources */,
 				4B1B88C0202E3DB200B67DFF /* MultiConfigurable.cpp in Sources */,
 				4BFF1D3922337B0300838EA1 /* 68000Storage.cpp in Sources */,
 				4B54C0BC1F8D8E790050900F /* KeyboardMachine.cpp in Sources */,
@@ -5096,6 +5113,7 @@
 				);
 				GCC_C_LANGUAGE_STANDARD = gnu11;
 				GCC_OPTIMIZATION_LEVEL = 2;
+				GCC_PREPROCESSOR_DEFINITIONS = NDEBUG;
 				MACOSX_DEPLOYMENT_TARGET = 10.10;
 				PRODUCT_NAME = "$(TARGET_NAME)";
 			};
@@ -5149,8 +5167,9 @@
 				GCC_WARN_UNUSED_FUNCTION = YES;
 				GCC_WARN_UNUSED_PARAMETER = YES;
 				GCC_WARN_UNUSED_VARIABLE = YES;
-				MACOSX_DEPLOYMENT_TARGET = 10.10;
-				MTL_ENABLE_DEBUG_INFO = YES;
+				MACOSX_DEPLOYMENT_TARGET = 10.13;
+				MTL_ENABLE_DEBUG_INFO = INCLUDE_SOURCE;
+				MTL_FAST_MATH = YES;
 				ONLY_ACTIVE_ARCH = YES;
 				SDKROOT = macosx;
 				SWIFT_OPTIMIZATION_LEVEL = "-Onone";
@@ -5200,8 +5219,9 @@
 				GCC_WARN_UNUSED_FUNCTION = YES;
 				GCC_WARN_UNUSED_PARAMETER = YES;
 				GCC_WARN_UNUSED_VARIABLE = YES;
-				MACOSX_DEPLOYMENT_TARGET = 10.10;
+				MACOSX_DEPLOYMENT_TARGET = 10.13;
 				MTL_ENABLE_DEBUG_INFO = NO;
+				MTL_FAST_MATH = YES;
 				SDKROOT = macosx;
 				SWIFT_OPTIMIZATION_LEVEL = "-Owholemodule";
 				SWIFT_SWIFT3_OBJC_INFERENCE = Default;
@@ -5240,7 +5260,7 @@
 				GCC_WARN_UNUSED_LABEL = YES;
 				INFOPLIST_FILE = "Clock Signal/Info.plist";
 				LD_RUNPATH_SEARCH_PATHS = "$(inherited) @executable_path/../Frameworks";
-				MACOSX_DEPLOYMENT_TARGET = 10.12.2;
+				MTL_TREAT_WARNINGS_AS_ERRORS = YES;
 				OTHER_CPLUSPLUSFLAGS = (
 					"$(OTHER_CFLAGS)",
 					"-Wreorder",
@@ -5288,7 +5308,7 @@
 				GCC_WARN_UNUSED_LABEL = YES;
 				INFOPLIST_FILE = "Clock Signal/Info.plist";
 				LD_RUNPATH_SEARCH_PATHS = "$(inherited) @executable_path/../Frameworks";
-				MACOSX_DEPLOYMENT_TARGET = 10.12.2;
+				MTL_TREAT_WARNINGS_AS_ERRORS = YES;
 				OTHER_CPLUSPLUSFLAGS = (
 					"$(OTHER_CFLAGS)",
 					"-Wreorder",
diff --git a/OSBindings/Mac/Clock Signal.xcodeproj/xcshareddata/xcschemes/Clock Signal Kiosk.xcscheme b/OSBindings/Mac/Clock Signal.xcodeproj/xcshareddata/xcschemes/Clock Signal Kiosk.xcscheme
index 702ed8cf3..a98a6ccfd 100644
--- a/OSBindings/Mac/Clock Signal.xcodeproj/xcshareddata/xcschemes/Clock Signal Kiosk.xcscheme	
+++ b/OSBindings/Mac/Clock Signal.xcodeproj/xcshareddata/xcschemes/Clock Signal Kiosk.xcscheme	
@@ -31,7 +31,7 @@
       </Testables>
    </TestAction>
    <LaunchAction
-      buildConfiguration = "Debug"
+      buildConfiguration = "Release"
       selectedDebuggerIdentifier = "Xcode.DebuggerFoundation.Debugger.LLDB"
       selectedLauncherIdentifier = "Xcode.DebuggerFoundation.Launcher.LLDB"
       disableMainThreadChecker = "YES"
@@ -57,9 +57,13 @@
             isEnabled = "NO">
          </CommandLineArgument>
          <CommandLineArgument
-            argument = "--volume=0.001"
+            argument = "&quot;/Users/thomasharte/Library/Mobile Documents/com~apple~CloudDocs/Desktop/Soft/Macintosh/MusicWorks 0.42.image&quot;"
             isEnabled = "YES">
          </CommandLineArgument>
+         <CommandLineArgument
+            argument = "--volume=0.001"
+            isEnabled = "NO">
+         </CommandLineArgument>
          <CommandLineArgument
             argument = "--new=amstradcpc"
             isEnabled = "NO">
@@ -86,7 +90,7 @@
          </CommandLineArgument>
          <CommandLineArgument
             argument = "&quot;/Users/thomasharte/Library/Mobile Documents/com~apple~CloudDocs/Desktop/Soft/Amstrad CPC/Robocop.dsk&quot;"
-            isEnabled = "YES">
+            isEnabled = "NO">
          </CommandLineArgument>
          <CommandLineArgument
             argument = "--speed=5"
@@ -94,7 +98,7 @@
          </CommandLineArgument>
          <CommandLineArgument
             argument = "--rompath=/Users/thomasharte/Projects/CLK/ROMImages"
-            isEnabled = "NO">
+            isEnabled = "YES">
          </CommandLineArgument>
          <CommandLineArgument
             argument = "--help"
diff --git a/OSBindings/Mac/Clock Signal/Base.lproj/MachineDocument.xib b/OSBindings/Mac/Clock Signal/Base.lproj/MachineDocument.xib
index 8d955fcaf..b1eeadf5a 100644
--- a/OSBindings/Mac/Clock Signal/Base.lproj/MachineDocument.xib	
+++ b/OSBindings/Mac/Clock Signal/Base.lproj/MachineDocument.xib	
@@ -1,14 +1,14 @@
 <?xml version="1.0" encoding="UTF-8"?>
-<document type="com.apple.InterfaceBuilder3.Cocoa.XIB" version="3.0" toolsVersion="15705" targetRuntime="MacOSX.Cocoa" propertyAccessControl="none" useAutolayout="YES" customObjectInstantitationMethod="direct">
+<document type="com.apple.InterfaceBuilder3.Cocoa.XIB" version="3.0" toolsVersion="16097.2" targetRuntime="MacOSX.Cocoa" propertyAccessControl="none" useAutolayout="YES" customObjectInstantitationMethod="direct">
     <dependencies>
         <deployment identifier="macosx"/>
-        <plugIn identifier="com.apple.InterfaceBuilder.CocoaPlugin" version="15705"/>
+        <plugIn identifier="com.apple.InterfaceBuilder.CocoaPlugin" version="16097.2"/>
         <capability name="documents saved in the Xcode 8 format" minToolsVersion="8.0"/>
     </dependencies>
     <objects>
         <customObject id="-2" userLabel="File's Owner" customClass="MachineDocument" customModule="Clock_Signal" customModuleProvider="target">
             <connections>
-                <outlet property="openGLView" destination="DEG-fq-cjd" id="Gxs-2u-n7B"/>
+                <outlet property="scanTargetView" destination="DEG-fq-cjd" id="5aX-3R-eXQ"/>
                 <outlet property="volumeSlider" destination="zaz-lB-Iyt" id="flY-Th-oG4"/>
                 <outlet property="volumeView" destination="4ap-Gi-2AO" id="v4e-k6-Fqf"/>
                 <outlet property="window" destination="xOd-HO-29H" id="JIz-fz-R2o"/>
@@ -27,7 +27,7 @@
                 <rect key="frame" x="0.0" y="0.0" width="600" height="450"/>
                 <autoresizingMask key="autoresizingMask"/>
                 <subviews>
-                    <openGLView hidden="YES" wantsLayer="YES" useAuxiliaryDepthBufferStencil="NO" allowOffline="YES" wantsBestResolutionOpenGLSurface="YES" translatesAutoresizingMaskIntoConstraints="NO" id="DEG-fq-cjd" customClass="CSOpenGLView">
+                    <openGLView hidden="YES" wantsLayer="YES" useAuxiliaryDepthBufferStencil="NO" allowOffline="YES" wantsBestResolutionOpenGLSurface="YES" translatesAutoresizingMaskIntoConstraints="NO" id="DEG-fq-cjd" customClass="CSScanTargetView">
                         <rect key="frame" x="0.0" y="0.0" width="600" height="450"/>
                     </openGLView>
                     <box hidden="YES" boxType="custom" cornerRadius="4" title="Box" titlePosition="noTitle" translatesAutoresizingMaskIntoConstraints="NO" id="4ap-Gi-2AO">
diff --git a/OSBindings/Mac/Clock Signal/ClockSignal-Bridging-Header.h b/OSBindings/Mac/Clock Signal/ClockSignal-Bridging-Header.h
index ebbf066d9..dac72af09 100644
--- a/OSBindings/Mac/Clock Signal/ClockSignal-Bridging-Header.h	
+++ b/OSBindings/Mac/Clock Signal/ClockSignal-Bridging-Header.h	
@@ -10,7 +10,7 @@
 #import "CSStaticAnalyser.h"
 
 #import "CSAudioQueue.h"
-#import	"CSOpenGLView.h"
+#import	"CSScanTargetView.h"
 #import "CSROMReceiverView.h"
 
 #import "CSJoystickManager.h"
diff --git a/OSBindings/Mac/Clock Signal/Documents/MachineDocument.swift b/OSBindings/Mac/Clock Signal/Documents/MachineDocument.swift
index f5912a427..fffea8e3a 100644
--- a/OSBindings/Mac/Clock Signal/Documents/MachineDocument.swift	
+++ b/OSBindings/Mac/Clock Signal/Documents/MachineDocument.swift	
@@ -14,8 +14,7 @@ class MachineDocument:
 	NSDocument,
 	NSWindowDelegate,
 	CSMachineDelegate,
-	CSOpenGLViewDelegate,
-	CSOpenGLViewResponderDelegate,
+	CSScanTargetViewResponderDelegate,
 	CSAudioQueueDelegate,
 	CSROMReciverViewDelegate
 {
@@ -45,7 +44,7 @@ class MachineDocument:
 	// MARK: - Main NIB connections.
 
 	/// The OpenGL view to receive this machine's display.
-	@IBOutlet weak var openGLView: CSOpenGLView!
+	@IBOutlet weak var scanTargetView: CSScanTargetView!
 
 	/// The options panel, if any.
 	@IBOutlet var optionsPanel: MachinePanel!
@@ -100,8 +99,7 @@ class MachineDocument:
 		actionLock.lock()
 		drawLock.lock()
 		machine = nil
-		openGLView.delegate = nil
-		openGLView.invalidate()
+		scanTargetView.invalidate()
 		actionLock.unlock()
 		drawLock.unlock()
 
@@ -181,10 +179,10 @@ class MachineDocument:
 	// MARK: - Connections Between Machine and the Outside World
 
 	private func setupMachineOutput() {
-		if let machine = self.machine, let openGLView = self.openGLView, machine.view != openGLView {
+		if let machine = self.machine, let scanTargetView = self.scanTargetView, machine.view != scanTargetView {
 			// Establish the output aspect ratio and audio.
 			let aspectRatio = self.aspectRatio()
-			machine.setView(openGLView, aspectRatio: Float(aspectRatio.width / aspectRatio.height))
+			machine.setView(scanTargetView, aspectRatio: Float(aspectRatio.width / aspectRatio.height))
 
 			// Attach an options panel if one is available.
 			if let optionsPanelNibName = self.machineDescription?.optionsPanelNibName {
@@ -198,20 +196,20 @@ class MachineDocument:
 
 			// Callbacks from the OpenGL may come on a different thread, immediately following the .delegate set;
 			// hence the full setup of the best-effort updater prior to setting self as a delegate.
-			openGLView.delegate = self
-			openGLView.responderDelegate = self
+//			scanTargetView.delegate = self
+			scanTargetView.responderDelegate = self
 
 			// If this machine has a mouse, enable mouse capture; also indicate whether usurption
 			// of the command key is desired.
-			openGLView.shouldCaptureMouse = machine.hasMouse
-			openGLView.shouldUsurpCommand = machine.shouldUsurpCommand
+			scanTargetView.shouldCaptureMouse = machine.hasMouse
+			scanTargetView.shouldUsurpCommand = machine.shouldUsurpCommand
 
 			setupAudioQueueClockRate()
 
 			// Bring OpenGL view-holding window on top of the options panel and show the content.
-			openGLView.isHidden = false
-			openGLView.window!.makeKeyAndOrderFront(self)
-			openGLView.window!.makeFirstResponder(openGLView)
+			scanTargetView.isHidden = false
+			scanTargetView.window!.makeKeyAndOrderFront(self)
+			scanTargetView.window!.makeFirstResponder(scanTargetView)
 
 			// Start forwarding best-effort updates.
 			machine.start()
@@ -252,18 +250,6 @@ class MachineDocument:
 	final func audioQueueIsRunningDry(_ audioQueue: CSAudioQueue) {
 	}
 
-	/// Responds to the CSOpenGLViewDelegate redraw message by requesting a machine update if this is a timed
-	/// request, and ordering a redraw regardless of the motivation.
-	final func openGLViewRedraw(_ view: CSOpenGLView, event redrawEvent: CSOpenGLViewRedrawEvent) {
-		if drawLock.try() {
-			if redrawEvent == .timer {
-				machine.updateView(forPixelSize: view.backingSize)
-			}
-			machine.drawView(forPixelSize: view.backingSize)
-			drawLock.unlock()
-		}
-	}
-
 	// MARK: - Pasteboard Forwarding.
 
 	/// Forwards any text currently on the pasteboard into the active machine.
@@ -277,7 +263,7 @@ class MachineDocument:
 	// MARK: - Runtime Media Insertion.
 
 	/// Delegate message to receive drag and drop files.
-	final func openGLView(_ view: CSOpenGLView, didReceiveFileAt URL: URL) {
+	final func scanTargetView(_ view: CSScanTargetView, didReceiveFileAt URL: URL) {
 		let mediaSet = CSMediaSet(fileAt: URL)
 		if let mediaSet = mediaSet {
 			mediaSet.apply(to: self.machine)
@@ -310,7 +296,7 @@ class MachineDocument:
 			machine.clearAllKeys()
 			machine.joystickManager = nil
 		}
-		self.openGLView.releaseMouse()
+		self.scanTargetView.releaseMouse()
 	}
 
 	/// Upon becoming key, attaches joystick input to the machine.
@@ -608,23 +594,20 @@ class MachineDocument:
 		let url = pictursURL.appendingPathComponent(filename)
 
 		// Obtain the machine's current display.
-		var imageRepresentation: NSBitmapImageRep? = nil
-		self.openGLView.perform {
-			imageRepresentation = self.machine.imageRepresentation
-		}
+		let imageRepresentation = self.machine.imageRepresentation
 
 		// Encode as a PNG and save.
-		let pngData = imageRepresentation!.representation(using: .png, properties: [:])
+		let pngData = imageRepresentation.representation(using: .png, properties: [:])
 		try! pngData?.write(to: url)
 	}
 
 	// MARK: - Window Title Updates.
 	private var unadornedWindowTitle = ""
-	func openGLViewDidCaptureMouse(_ view: CSOpenGLView) {
+	internal func scanTargetViewDidCaptureMouse(_ view: CSScanTargetView) {
 		self.windowControllers[0].window?.title = self.unadornedWindowTitle + " (press ⌘+control to release mouse)"
 	}
 
-	func openGLViewDidReleaseMouse(_ view: CSOpenGLView) {
+	internal func scanTargetViewDidReleaseMouse(_ view: CSScanTargetView) {
 		self.windowControllers[0].window?.title = self.unadornedWindowTitle
 	}
 
@@ -750,7 +733,7 @@ class MachineDocument:
 	}
 	fileprivate var animationFader: ViewFader? = nil
 
-	func openGLViewDidShowOSMouseCursor(_ view: CSOpenGLView) {
+	internal func scanTargetViewDidShowOSMouseCursor(_ view: CSScanTargetView) {
 		// The OS mouse cursor became visible, so show the volume controls.
 		animationFader = nil
 		volumeView.layer?.removeAllAnimations()
@@ -758,7 +741,7 @@ class MachineDocument:
 		volumeView.layer?.opacity = 1.0
 	}
 
-	func openGLViewWillHideOSMouseCursor(_ view: CSOpenGLView) {
+	internal func scanTargetViewWillHideOSMouseCursor(_ view: CSScanTargetView) {
 		// The OS mouse cursor will be hidden, so hide the volume controls.
 		if !volumeView.isHidden && volumeView.layer?.animation(forKey: "opacity") == nil {
 			let fadeAnimation = CABasicAnimation(keyPath: "opacity")
diff --git a/OSBindings/Mac/Clock Signal/Machine/CSMachine.h b/OSBindings/Mac/Clock Signal/Machine/CSMachine.h
index d8595e06c..eea2c1f01 100644
--- a/OSBindings/Mac/Clock Signal/Machine/CSMachine.h	
+++ b/OSBindings/Mac/Clock Signal/Machine/CSMachine.h	
@@ -9,9 +9,9 @@
 #import <Foundation/Foundation.h>
 
 #import "CSAudioQueue.h"
-#import "CSOpenGLView.h"
-#import "CSStaticAnalyser.h"
 #import "CSJoystickManager.h"
+#import "CSScanTargetView.h"
+#import "CSStaticAnalyser.h"
 
 @class CSMachine;
 @protocol CSMachineDelegate
@@ -62,14 +62,11 @@ typedef NS_ENUM(NSInteger, CSMachineKeyboardInputMode) {
 - (BOOL)isStereo;
 - (void)setAudioSamplingRate:(float)samplingRate bufferSize:(NSUInteger)bufferSize stereo:(BOOL)stereo;
 
-- (void)setView:(nullable CSOpenGLView *)view aspectRatio:(float)aspectRatio;
+- (void)setView:(nullable CSScanTargetView *)view aspectRatio:(float)aspectRatio;
 
 - (void)start;
 - (void)stop;
 
-- (void)updateViewForPixelSize:(CGSize)pixelSize;
-- (void)drawViewForPixelSize:(CGSize)pixelSize;
-
 - (void)setKey:(uint16_t)key characters:(nullable NSString *)characters isPressed:(BOOL)isPressed;
 - (void)clearAllKeys;
 
@@ -77,7 +74,7 @@ typedef NS_ENUM(NSInteger, CSMachineKeyboardInputMode) {
 - (void)addMouseMotionX:(CGFloat)deltaX y:(CGFloat)deltaY;
 
 @property (atomic, strong, nullable) CSAudioQueue *audioQueue;
-@property (nonatomic, readonly, nonnull) CSOpenGLView *view;
+@property (nonatomic, readonly, nonnull) CSScanTargetView *view;
 @property (nonatomic, weak, nullable) id<CSMachineDelegate> delegate;
 
 @property (nonatomic, readonly, nonnull) NSString *userDefaultsPrefix;
diff --git a/OSBindings/Mac/Clock Signal/Machine/CSMachine.mm b/OSBindings/Mac/Clock Signal/Machine/CSMachine.mm
index cd5d8ba25..ff5e25265 100644
--- a/OSBindings/Mac/Clock Signal/Machine/CSMachine.mm	
+++ b/OSBindings/Mac/Clock Signal/Machine/CSMachine.mm	
@@ -9,8 +9,9 @@
 #import "CSMachine.h"
 #import "CSMachine+Target.h"
 
-#include "CSROMFetcher.hpp"
 #import "CSHighPrecisionTimer.h"
+#include "CSROMFetcher.hpp"
+#import "CSScanTarget+CppScanTarget.h"
 
 #include "MediaTarget.hpp"
 #include "JoystickMachine.hpp"
@@ -31,13 +32,7 @@
 #include <atomic>
 #include <bitset>
 
-#import <OpenGL/OpenGL.h>
-#include <OpenGL/gl3.h>
-
-#include "../../../../Outputs/OpenGL/ScanTarget.hpp"
-#include "../../../../Outputs/OpenGL/Screenshot.hpp"
-
-@interface CSMachine() <CSOpenGLViewDisplayLinkDelegate>
+@interface CSMachine() <CSScanTargetViewDisplayLinkDelegate>
 - (void)speaker:(Outputs::Speaker::Speaker *)speaker didCompleteSamples:(const int16_t *)samples length:(int)length;
 - (void)speakerDidChangeInputClock:(Outputs::Speaker::Speaker *)speaker;
 - (void)addLED:(NSString *)led;
@@ -154,7 +149,6 @@ struct ActivityObserver: public Activity::Observer {
 	NSMutableArray<NSString *> *_leds;
 
 	CSHighPrecisionTimer *_timer;
-	CGSize _pixelSize;
 	std::atomic_flag _isUpdating;
 	Time::Nanos _syncTime;
 	Time::Nanos _timeDiff;
@@ -165,7 +159,11 @@ struct ActivityObserver: public Activity::Observer {
 
 	NSTimer *_joystickTimer;
 
-	std::unique_ptr<Outputs::Display::OpenGL::ScanTarget> _scanTarget;
+	// This array exists to reduce blocking on the main queue; anything that would otherwise need
+	// to synchronise on self in order to post input to the machine can instead synchronise on
+	// _inputEvents and add a block to it. The main machine execution loop promises to synchronise
+	// on _inputEvents very briefly at the start of every tick and execute all enqueued blocks.
+	NSMutableArray<dispatch_block_t> *_inputEvents;
 }
 
 - (instancetype)initWithAnalyser:(CSStaticAnalyser *)result missingROMs:(inout NSMutableArray<CSMissingROM *> *)missingROMs {
@@ -217,6 +215,8 @@ struct ActivityObserver: public Activity::Observer {
 		_speakerDelegate.machine = self;
 		_speakerDelegate.machineAccessLock = _delegateMachineAccessLock;
 
+		_inputEvents = [[NSMutableArray alloc] init];
+
 		_joystickMachine = _machine->joystick_machine();
 		[self updateJoystickTimer];
 		_isUpdating.clear();
@@ -245,11 +245,11 @@ struct ActivityObserver: public Activity::Observer {
 	_speakerDelegate.machine = nil;
 	[_delegateMachineAccessLock unlock];
 
-	[_view performWithGLContext:^{
-		@synchronized(self) {
-			self->_scanTarget.reset();
-		}
-	}];
+//	[_view performWithGLContext:^{
+//		@synchronized(self) {
+//			self->_scanTarget.reset();
+//		}
+//	}];
 }
 
 - (float)idealSamplingRateFromRange:(NSRange)range {
@@ -351,30 +351,10 @@ struct ActivityObserver: public Activity::Observer {
 	}
 }
 
-- (void)setView:(CSOpenGLView *)view aspectRatio:(float)aspectRatio {
+- (void)setView:(CSScanTargetView *)view aspectRatio:(float)aspectRatio {
 	_view = view;
 	_view.displayLinkDelegate = self;
-	[view performWithGLContext:^{
-		[self setupOutputWithAspectRatio:aspectRatio];
-	} flushDrawable:NO];
-}
-
-- (void)setupOutputWithAspectRatio:(float)aspectRatio {
-	_scanTarget = std::make_unique<Outputs::Display::OpenGL::ScanTarget>();
-	_machine->scan_producer()->set_scan_target(_scanTarget.get());
-}
-
-- (void)updateViewForPixelSize:(CGSize)pixelSize {
-//	_pixelSize = pixelSize;
-
-//	@synchronized(self) {
-//		const auto scan_status = _machine->crt_machine()->get_scan_status();
-//		NSLog(@"FPS (hopefully): %0.2f [retrace: %0.4f]", 1.0f / scan_status.field_duration, scan_status.retrace_duration);
-//	}
-}
-
-- (void)drawViewForPixelSize:(CGSize)pixelSize {
-	_scanTarget->draw((int)pixelSize.width, (int)pixelSize.height);
+	_machine->scan_producer()->set_scan_target(_view.scanTarget.scanTarget);
 }
 
 - (void)paste:(NSString *)paste {
@@ -384,26 +364,7 @@ struct ActivityObserver: public Activity::Observer {
 }
 
 - (NSBitmapImageRep *)imageRepresentation {
-	// Grab a screenshot.
-	Outputs::Display::OpenGL::Screenshot screenshot(4, 3);
-
-	// Generate an NSBitmapImageRep containing the screenshot's data.
-	NSBitmapImageRep *const result =
-		[[NSBitmapImageRep alloc]
-			initWithBitmapDataPlanes:NULL
-			pixelsWide:screenshot.width
-			pixelsHigh:screenshot.height
-			bitsPerSample:8
-			samplesPerPixel:4
-			hasAlpha:YES
-			isPlanar:NO
-			colorSpaceName:NSDeviceRGBColorSpace
-			bytesPerRow:4 * screenshot.width
-			bitsPerPixel:0];
-
-	memcpy(result.bitmapData, screenshot.pixel_data.data(), size_t(screenshot.width*screenshot.height*4));
-
-	return result;
+	return self.view.imageRepresentation;
 }
 
 - (void)applyMedia:(const Analyser::Static::Media &)media {
@@ -428,84 +389,84 @@ struct ActivityObserver: public Activity::Observer {
 }
 
 - (void)setKey:(uint16_t)key characters:(NSString *)characters isPressed:(BOOL)isPressed {
-	auto keyboard_machine = _machine->keyboard_machine();
-	if(keyboard_machine && (self.inputMode != CSMachineKeyboardInputModeJoystick || !keyboard_machine->get_keyboard().is_exclusive())) {
-		Inputs::Keyboard::Key mapped_key = Inputs::Keyboard::Key::Help;	// Make an innocuous default guess.
+	[self applyInputEvent:^{
+		auto keyboard_machine = self->_machine->keyboard_machine();
+		if(keyboard_machine && (self.inputMode != CSMachineKeyboardInputModeJoystick || !keyboard_machine->get_keyboard().is_exclusive())) {
+			Inputs::Keyboard::Key mapped_key = Inputs::Keyboard::Key::Help;	// Make an innocuous default guess.
 #define BIND(source, dest) case source: mapped_key = Inputs::Keyboard::Key::dest; break;
-		// Connect the Carbon-era Mac keyboard scancodes to Clock Signal's 'universal' enumeration in order
-		// to pass into the platform-neutral realm.
-		switch(key) {
-			BIND(VK_ANSI_0, k0);	BIND(VK_ANSI_1, k1);	BIND(VK_ANSI_2, k2);	BIND(VK_ANSI_3, k3);	BIND(VK_ANSI_4, k4);
-			BIND(VK_ANSI_5, k5);	BIND(VK_ANSI_6, k6);	BIND(VK_ANSI_7, k7);	BIND(VK_ANSI_8, k8);	BIND(VK_ANSI_9, k9);
+			// Connect the Carbon-era Mac keyboard scancodes to Clock Signal's 'universal' enumeration in order
+			// to pass into the platform-neutral realm.
+			switch(key) {
+				BIND(VK_ANSI_0, k0);	BIND(VK_ANSI_1, k1);	BIND(VK_ANSI_2, k2);	BIND(VK_ANSI_3, k3);	BIND(VK_ANSI_4, k4);
+				BIND(VK_ANSI_5, k5);	BIND(VK_ANSI_6, k6);	BIND(VK_ANSI_7, k7);	BIND(VK_ANSI_8, k8);	BIND(VK_ANSI_9, k9);
 
-			BIND(VK_ANSI_Q, Q);		BIND(VK_ANSI_W, W);		BIND(VK_ANSI_E, E);		BIND(VK_ANSI_R, R);		BIND(VK_ANSI_T, T);
-			BIND(VK_ANSI_Y, Y);		BIND(VK_ANSI_U, U);		BIND(VK_ANSI_I, I);		BIND(VK_ANSI_O, O);		BIND(VK_ANSI_P, P);
+				BIND(VK_ANSI_Q, Q);		BIND(VK_ANSI_W, W);		BIND(VK_ANSI_E, E);		BIND(VK_ANSI_R, R);		BIND(VK_ANSI_T, T);
+				BIND(VK_ANSI_Y, Y);		BIND(VK_ANSI_U, U);		BIND(VK_ANSI_I, I);		BIND(VK_ANSI_O, O);		BIND(VK_ANSI_P, P);
 
-			BIND(VK_ANSI_A, A);		BIND(VK_ANSI_S, S);		BIND(VK_ANSI_D, D);		BIND(VK_ANSI_F, F);		BIND(VK_ANSI_G, G);
-			BIND(VK_ANSI_H, H);		BIND(VK_ANSI_J, J);		BIND(VK_ANSI_K, K);		BIND(VK_ANSI_L, L);
+				BIND(VK_ANSI_A, A);		BIND(VK_ANSI_S, S);		BIND(VK_ANSI_D, D);		BIND(VK_ANSI_F, F);		BIND(VK_ANSI_G, G);
+				BIND(VK_ANSI_H, H);		BIND(VK_ANSI_J, J);		BIND(VK_ANSI_K, K);		BIND(VK_ANSI_L, L);
 
-			BIND(VK_ANSI_Z, Z);		BIND(VK_ANSI_X, X);		BIND(VK_ANSI_C, C);		BIND(VK_ANSI_V, V);
-			BIND(VK_ANSI_B, B);		BIND(VK_ANSI_N, N);		BIND(VK_ANSI_M, M);
+				BIND(VK_ANSI_Z, Z);		BIND(VK_ANSI_X, X);		BIND(VK_ANSI_C, C);		BIND(VK_ANSI_V, V);
+				BIND(VK_ANSI_B, B);		BIND(VK_ANSI_N, N);		BIND(VK_ANSI_M, M);
 
-			BIND(VK_F1, F1);		BIND(VK_F2, F2);		BIND(VK_F3, F3);		BIND(VK_F4, F4);
-			BIND(VK_F5, F5);		BIND(VK_F6, F6);		BIND(VK_F7, F7);		BIND(VK_F8, F8);
-			BIND(VK_F9, F9);		BIND(VK_F10, F10);		BIND(VK_F11, F11);		BIND(VK_F12, F12);
+				BIND(VK_F1, F1);		BIND(VK_F2, F2);		BIND(VK_F3, F3);		BIND(VK_F4, F4);
+				BIND(VK_F5, F5);		BIND(VK_F6, F6);		BIND(VK_F7, F7);		BIND(VK_F8, F8);
+				BIND(VK_F9, F9);		BIND(VK_F10, F10);		BIND(VK_F11, F11);		BIND(VK_F12, F12);
 
-			BIND(VK_ANSI_Keypad0, Keypad0);		BIND(VK_ANSI_Keypad1, Keypad1);		BIND(VK_ANSI_Keypad2, Keypad2);
-			BIND(VK_ANSI_Keypad3, Keypad3);		BIND(VK_ANSI_Keypad4, Keypad4);		BIND(VK_ANSI_Keypad5, Keypad5);
-			BIND(VK_ANSI_Keypad6, Keypad6);		BIND(VK_ANSI_Keypad7, Keypad7);		BIND(VK_ANSI_Keypad8, Keypad8);
-			BIND(VK_ANSI_Keypad9, Keypad9);
+				BIND(VK_ANSI_Keypad0, Keypad0);		BIND(VK_ANSI_Keypad1, Keypad1);		BIND(VK_ANSI_Keypad2, Keypad2);
+				BIND(VK_ANSI_Keypad3, Keypad3);		BIND(VK_ANSI_Keypad4, Keypad4);		BIND(VK_ANSI_Keypad5, Keypad5);
+				BIND(VK_ANSI_Keypad6, Keypad6);		BIND(VK_ANSI_Keypad7, Keypad7);		BIND(VK_ANSI_Keypad8, Keypad8);
+				BIND(VK_ANSI_Keypad9, Keypad9);
 
-			BIND(VK_ANSI_Equal, Equals);						BIND(VK_ANSI_Minus, Hyphen);
-			BIND(VK_ANSI_RightBracket, CloseSquareBracket);		BIND(VK_ANSI_LeftBracket, OpenSquareBracket);
-			BIND(VK_ANSI_Quote, Quote);							BIND(VK_ANSI_Grave, BackTick);
+				BIND(VK_ANSI_Equal, Equals);						BIND(VK_ANSI_Minus, Hyphen);
+				BIND(VK_ANSI_RightBracket, CloseSquareBracket);		BIND(VK_ANSI_LeftBracket, OpenSquareBracket);
+				BIND(VK_ANSI_Quote, Quote);							BIND(VK_ANSI_Grave, BackTick);
 
-			BIND(VK_ANSI_Semicolon, Semicolon);
-			BIND(VK_ANSI_Backslash, Backslash);					BIND(VK_ANSI_Slash, ForwardSlash);
-			BIND(VK_ANSI_Comma, Comma);							BIND(VK_ANSI_Period, FullStop);
+				BIND(VK_ANSI_Semicolon, Semicolon);
+				BIND(VK_ANSI_Backslash, Backslash);					BIND(VK_ANSI_Slash, ForwardSlash);
+				BIND(VK_ANSI_Comma, Comma);							BIND(VK_ANSI_Period, FullStop);
 
-			BIND(VK_ANSI_KeypadDecimal, KeypadDecimalPoint);	BIND(VK_ANSI_KeypadEquals, KeypadEquals);
-			BIND(VK_ANSI_KeypadMultiply, KeypadAsterisk);		BIND(VK_ANSI_KeypadDivide, KeypadSlash);
-			BIND(VK_ANSI_KeypadPlus, KeypadPlus);				BIND(VK_ANSI_KeypadMinus, KeypadMinus);
-			BIND(VK_ANSI_KeypadClear, KeypadDelete);			BIND(VK_ANSI_KeypadEnter, KeypadEnter);
+				BIND(VK_ANSI_KeypadDecimal, KeypadDecimalPoint);	BIND(VK_ANSI_KeypadEquals, KeypadEquals);
+				BIND(VK_ANSI_KeypadMultiply, KeypadAsterisk);		BIND(VK_ANSI_KeypadDivide, KeypadSlash);
+				BIND(VK_ANSI_KeypadPlus, KeypadPlus);				BIND(VK_ANSI_KeypadMinus, KeypadMinus);
+				BIND(VK_ANSI_KeypadClear, KeypadDelete);			BIND(VK_ANSI_KeypadEnter, KeypadEnter);
 
-			BIND(VK_Return, Enter);					BIND(VK_Tab, Tab);
-			BIND(VK_Space, Space);					BIND(VK_Delete, Backspace);
-			BIND(VK_Control, LeftControl);			BIND(VK_Option, LeftOption);
-			BIND(VK_Command, LeftMeta);				BIND(VK_Shift, LeftShift);
-			BIND(VK_RightControl, RightControl);	BIND(VK_RightOption, RightOption);
-			BIND(VK_Escape, Escape);				BIND(VK_CapsLock, CapsLock);
-			BIND(VK_Home, Home);					BIND(VK_End, End);
-			BIND(VK_PageUp, PageUp);				BIND(VK_PageDown, PageDown);
+				BIND(VK_Return, Enter);					BIND(VK_Tab, Tab);
+				BIND(VK_Space, Space);					BIND(VK_Delete, Backspace);
+				BIND(VK_Control, LeftControl);			BIND(VK_Option, LeftOption);
+				BIND(VK_Command, LeftMeta);				BIND(VK_Shift, LeftShift);
+				BIND(VK_RightControl, RightControl);	BIND(VK_RightOption, RightOption);
+				BIND(VK_Escape, Escape);				BIND(VK_CapsLock, CapsLock);
+				BIND(VK_Home, Home);					BIND(VK_End, End);
+				BIND(VK_PageUp, PageUp);				BIND(VK_PageDown, PageDown);
 
-			BIND(VK_RightShift, RightShift);
-			BIND(VK_Help, Help);
-			BIND(VK_ForwardDelete, Delete);
+				BIND(VK_RightShift, RightShift);
+				BIND(VK_Help, Help);
+				BIND(VK_ForwardDelete, Delete);
 
-			BIND(VK_LeftArrow, Left);		BIND(VK_RightArrow, Right);
-			BIND(VK_DownArrow, Down);		BIND(VK_UpArrow, Up);
-		}
+				BIND(VK_LeftArrow, Left);		BIND(VK_RightArrow, Right);
+				BIND(VK_DownArrow, Down);		BIND(VK_UpArrow, Up);
+			}
 #undef BIND
 
-		// Pick an ASCII code, if any.
-		char pressedKey = '\0';
-		if(characters.length) {
-			unichar firstCharacter = [characters characterAtIndex:0];
-			if(firstCharacter < 128) {
-				pressedKey = (char)firstCharacter;
+			// Pick an ASCII code, if any.
+			char pressedKey = '\0';
+			if(characters.length) {
+				unichar firstCharacter = [characters characterAtIndex:0];
+				if(firstCharacter < 128) {
+					pressedKey = (char)firstCharacter;
+				}
+			}
+
+			@synchronized(self) {
+				if(keyboard_machine->apply_key(mapped_key, pressedKey, isPressed, self.inputMode == CSMachineKeyboardInputModeKeyboardLogical)) {
+					return;
+				}
 			}
 		}
 
-		@synchronized(self) {
-			if(keyboard_machine->apply_key(mapped_key, pressedKey, isPressed, self.inputMode == CSMachineKeyboardInputModeKeyboardLogical)) {
-				return;
-			}
-		}
-	}
-
-	auto joystick_machine = _machine->joystick_machine();
-	if(self.inputMode == CSMachineKeyboardInputModeJoystick && joystick_machine) {
-		@synchronized(self) {
+		auto joystick_machine = self->_machine->joystick_machine();
+		if(self.inputMode == CSMachineKeyboardInputModeJoystick && joystick_machine) {
 			auto &joysticks = joystick_machine->get_joysticks();
 			if(!joysticks.empty()) {
 				// Convert to a C++ bool so that the following calls are resolved correctly even if overloaded.
@@ -530,49 +491,55 @@ struct ActivityObserver: public Activity::Observer {
 				}
 			}
 		}
+	}];
+}
+
+- (void)applyInputEvent:(dispatch_block_t)event {
+	@synchronized(_inputEvents) {
+		[_inputEvents addObject:event];
 	}
 }
 
 - (void)clearAllKeys {
 	const auto keyboard_machine = _machine->keyboard_machine();
 	if(keyboard_machine) {
-		@synchronized(self) {
+		[self applyInputEvent:^{
 			keyboard_machine->get_keyboard().reset_all_keys();
-		}
+		}];
 	}
 
 	const auto joystick_machine = _machine->joystick_machine();
 	if(joystick_machine) {
-		@synchronized(self) {
+		[self applyInputEvent:^{
 			for(auto &joystick : joystick_machine->get_joysticks()) {
 				joystick->reset_all_inputs();
 			}
-		}
+		}];
 	}
 
 	const auto mouse_machine = _machine->mouse_machine();
 	if(mouse_machine) {
-		@synchronized(self) {
+		[self applyInputEvent:^{
 			mouse_machine->get_mouse().reset_all_buttons();
-		}
+		}];
 	}
 }
 
 - (void)setMouseButton:(int)button isPressed:(BOOL)isPressed {
 	auto mouse_machine = _machine->mouse_machine();
 	if(mouse_machine) {
-		@synchronized(self) {
+		[self applyInputEvent:^{
 			mouse_machine->get_mouse().set_button_pressed(button % mouse_machine->get_mouse().get_number_of_buttons(), isPressed);
-		}
+		}];
 	}
 }
 
 - (void)addMouseMotionX:(CGFloat)deltaX y:(CGFloat)deltaY {
 	auto mouse_machine = _machine->mouse_machine();
 	if(mouse_machine) {
-		@synchronized(self) {
+		[self applyInputEvent:^{
 			mouse_machine->get_mouse().move(int(deltaX), int(deltaY));
-		}
+		}];
 	}
 }
 
@@ -737,11 +704,10 @@ struct ActivityObserver: public Activity::Observer {
 
 #pragma mark - Timer
 
-- (void)openGLViewDisplayLinkDidFire:(CSOpenGLView *)view now:(const CVTimeStamp *)now outputTime:(const CVTimeStamp *)outputTime {
+- (void)scanTargetViewDisplayLinkDidFire:(CSScanTargetView *)view now:(const CVTimeStamp *)now outputTime:(const CVTimeStamp *)outputTime {
 	// First order of business: grab a timestamp.
 	const auto timeNow = Time::nanos_now();
 
-	CGSize pixelSize = view.backingSize;
 	BOOL isSyncLocking;
 	@synchronized(self) {
 		// Store a means to map from CVTimeStamp.hostTime to Time::Nanos;
@@ -753,9 +719,6 @@ struct ActivityObserver: public Activity::Observer {
 		// Store the next end-of-frame time. TODO: and start of next and implied visible duration, if raster racing?
 		_syncTime = int64_t(now->hostTime) + _timeDiff;
 
-		// Also crib the current view pixel size.
-		_pixelSize = pixelSize;
-
 		// Set the current refresh period.
 		_refreshPeriod = double(now->videoRefreshPeriod) / double(now->videoTimeScale);
 
@@ -765,9 +728,7 @@ struct ActivityObserver: public Activity::Observer {
 
 	// Draw the current output. (TODO: do this within the timer if either raster racing or, at least, sync matching).
 	if(!isSyncLocking) {
-		[self.view performWithGLContext:^{
-			self->_scanTarget->draw((int)pixelSize.width, (int)pixelSize.height);
-		} flushDrawable:YES];
+		[self.view draw];
 	}
 }
 
@@ -783,9 +744,16 @@ struct ActivityObserver: public Activity::Observer {
 		lastTime = std::max(timeNow - Time::Nanos(10'000'000'000 / TICKS), lastTime);
 		const auto duration = timeNow - lastTime;
 
-		CGSize pixelSize;
 		BOOL splitAndSync = NO;
 		@synchronized(self) {
+			// Post on input events.
+			@synchronized(self->_inputEvents) {
+				for(dispatch_block_t action: self->_inputEvents) {
+					action();
+				}
+				[self->_inputEvents removeAllObjects];
+			}
+
 			// If this tick includes vsync then inspect the machine.
 			if(timeNow >= self->_syncTime && lastTime < self->_syncTime) {
 				splitAndSync = self->_isSyncLocking = self->_scanSynchroniser.can_synchronise(self->_machine->scan_producer()->get_scan_status(), self->_refreshPeriod);
@@ -806,7 +774,6 @@ struct ActivityObserver: public Activity::Observer {
 			if(!splitAndSync) {
 				self->_machine->timed_machine()->run_for((double)duration / 1e9);
 			}
-			pixelSize = self->_pixelSize;
 		}
 
 		// If this was not a split-and-sync then dispatch the update request asynchronously, unless
@@ -822,13 +789,10 @@ struct ActivityObserver: public Activity::Observer {
 		}
 		if(!wasUpdating) {
 			dispatch_async(dispatch_get_global_queue(QOS_CLASS_USER_INTERACTIVE, 0), ^{
-				[self.view performWithGLContext:^{
-					self->_scanTarget->update((int)pixelSize.width, (int)pixelSize.height);
-
-					if(splitAndSync) {
-						self->_scanTarget->draw((int)pixelSize.width, (int)pixelSize.height);
-					}
-				} flushDrawable:splitAndSync];
+				[self.view updateBacking];
+				if(splitAndSync) {
+					[self.view draw];
+				}
 				self->_isUpdating.clear();
 			});
 		}
diff --git a/OSBindings/Mac/Clock Signal/ScanTarget/CSScanTarget+CppScanTarget.h b/OSBindings/Mac/Clock Signal/ScanTarget/CSScanTarget+CppScanTarget.h
new file mode 100644
index 000000000..d0dbc0a69
--- /dev/null
+++ b/OSBindings/Mac/Clock Signal/ScanTarget/CSScanTarget+CppScanTarget.h	
@@ -0,0 +1,16 @@
+//
+//  CSScanTarget+C__ScanTarget.h
+//  Clock Signal
+//
+//  Created by Thomas Harte on 08/08/2020.
+//  Copyright © 2020 Thomas Harte. All rights reserved.
+//
+
+#import "CSScanTarget.h"
+#include "ScanTarget.hpp"
+
+@interface CSScanTarget (CppScanTarget)
+
+@property (nonatomic, readonly, nonnull) Outputs::Display::ScanTarget *scanTarget;
+
+@end
diff --git a/OSBindings/Mac/Clock Signal/ScanTarget/CSScanTarget.h b/OSBindings/Mac/Clock Signal/ScanTarget/CSScanTarget.h
new file mode 100644
index 000000000..993e1fd04
--- /dev/null
+++ b/OSBindings/Mac/Clock Signal/ScanTarget/CSScanTarget.h	
@@ -0,0 +1,25 @@
+//
+//  ScanTarget.h
+//  Clock Signal
+//
+//  Created by Thomas Harte on 02/08/2020.
+//  Copyright © 2020 Thomas Harte. All rights reserved.
+//
+
+#import <Foundation/Foundation.h>
+#import <MetalKit/MetalKit.h>
+
+/*!
+	Provides a ScanTarget that uses Metal as its back-end.
+*/
+@interface CSScanTarget : NSObject <MTKViewDelegate>
+
+- (nonnull instancetype)initWithView:(nonnull MTKView *)view;
+
+// Draws all scans currently residing at the scan target to the backing store,
+// ready for output when next requested.
+- (void)updateFrameBuffer;
+
+- (nonnull NSBitmapImageRep *)imageRepresentation;
+
+@end
diff --git a/OSBindings/Mac/Clock Signal/ScanTarget/CSScanTarget.mm b/OSBindings/Mac/Clock Signal/ScanTarget/CSScanTarget.mm
new file mode 100644
index 000000000..7076afa80
--- /dev/null
+++ b/OSBindings/Mac/Clock Signal/ScanTarget/CSScanTarget.mm	
@@ -0,0 +1,1185 @@
+//
+//  ScanTarget.m
+//  Clock Signal
+//
+//  Created by Thomas Harte on 02/08/2020.
+//  Copyright © 2020 Thomas Harte. All rights reserved.
+//
+
+#import "CSScanTarget.h"
+
+#import <Metal/Metal.h>
+
+#include <algorithm>
+#include <atomic>
+#include <cmath>
+
+#include "BufferingScanTarget.hpp"
+#include "FIRFilter.hpp"
+
+/*
+
+	RGB and composite monochrome
+	----------------------------
+
+	Source data is converted to 32bpp RGB or to composite directly from its input, at output resolution.
+	Gamma correction is applied unless the inputs are 1bpp (e.g. Macintosh-style black/white, TTL-style RGB).
+
+	S-Video
+	-------
+
+	Source data is pasted together with a common clock in the composition buffer. Colour phase is baked in
+	at this point. Format within the composition buffer is:
+
+		.r = luminance
+		.g = 0.5 + 0.5 * chrominance * cos(phase)
+		.b = 0.5 + 0.5 * chrominance * sin(phase)
+
+	Contents of the composition buffer are then drawn into the finalised line texture; at this point a suitable
+	low-filter is applied to the two chrominance channels, colours are converted to RGB and gamma corrected.
+
+	Contents from the finalised line texture are then painted to the display.
+
+	Composite colour
+	----------------
+
+	Source data is pasted together with a common clock in the composition buffer. Colour phase and amplitude are
+	recorded at this point. Format within the composition buffer is:
+
+		.r = composite value
+		.g = 0.5 + 0.5 * cos(phase)
+		.b = 0.5 + 0.5 * sin(phase)
+		.a = amplitude
+
+	[aside: upfront calculation of cos/sin is just because it'll need to be calculated at this precision anyway,
+	and doing it here avoids having to do unit<->radian conversions on phase alone]
+
+	Contents of the composition buffer are transferred to the separated-luma buffer, subject to a low-paass filter
+	that has sought to separate luminance and chrominance, and with phase and amplitude now baked into the latter:
+
+		.r = luminance
+		.g = 0.5 + 0.5 * chrominance * cos(phase)
+		.b = 0.5 + 0.5 * chrominance * sin(phase)
+
+	The process now continues as per the corresponding S-Video steps.
+
+	NOTES
+	-----
+
+		1)	for many of the input pixel formats it would be possible to do the trigonometric side of things at
+			arbitrary precision. Since it would always be necessary to support fixed-precision processing because
+			of the directly-sampled input formats, I've used fixed throughout to reduce the number of permutations
+			and combinations of code I need to support. The precision is always selected to be at least four times
+			the colour clock.
+
+		2)	I experimented with skipping the separated-luma buffer for composite colour based on the observation that
+			just multiplying the raw signal by sin and cos and then filtering well below the colour subcarrier frequency
+			should be sufficient. It wasn't in practice because the bits of luminance that don't quite separate are then
+			of such massive amplitude that you get huge bands of bright colour in place of the usual chroma dots.
+
+		3)	I also initially didn't want to have a finalied-line texture, but processing costs changed my mind on that.
+			If you accept that output will be fixed precision, anyway. In that case, processing for a typical NTSC frame
+			in its original resolution means applying filtering (i.e. at least 15 samples per pixel) likely between
+			218,400 and 273,000 times per output frame, then upscaling from there at 1 sample per pixel. Count the second
+			sample twice for the original store and you're talking between 16*218,400 = 3,494,400 to 16*273,000 = 4,368,000
+			total pixel accesses. Though that's not a perfect way to measure cost, roll with it.
+
+			On my 4k monitor, doing it at actual output resolution would instead cost 3840*2160*15 = 124,416,000 total
+			accesses. Which doesn't necessarily mean "more than 28 times as much", but does mean "a lot more".
+
+			(going direct-to-display for composite monochrome means evaluating sin/cos a lot more often than it might
+			with more buffering in between, but that doesn't provisionally seem to be as much of a bottleneck)
+*/
+
+namespace {
+
+/// Provides a container for __fp16 versions of tightly-packed single-precision plain old data with a copy assignment constructor.
+template <typename NaturalType> struct HalfConverter {
+	__fp16 elements[sizeof(NaturalType) / sizeof(float)];
+
+	void operator =(const NaturalType &rhs) {
+		const float *floatRHS = reinterpret_cast<const float *>(&rhs);
+		for(size_t c = 0; c < sizeof(elements) / sizeof(*elements); ++c) {
+			elements[c] = __fp16(floatRHS[c]);
+		}
+	}
+};
+
+// Tracks the Uniforms struct declared in ScanTarget.metal; see there for field definitions.
+//
+// __fp16 is a Clang-specific type which I'm using as equivalent to a Metal half, i.e. an IEEE 754 binary16.
+struct Uniforms {
+	int32_t scale[2];
+	float cyclesMultiplier;
+	float lineWidth;
+
+	simd::float3x3 sourcetoDisplay;
+
+	HalfConverter<simd::float3x3> toRGB;
+	HalfConverter<simd::float3x3> fromRGB;
+
+	HalfConverter<simd::float3> chromaKernel[8];
+	__fp16 lumaKernel[8];
+
+	__fp16 outputAlpha;
+	__fp16 outputGamma;
+	__fp16 outputMultiplier;
+};
+
+constexpr size_t NumBufferedLines = 500;
+constexpr size_t NumBufferedScans = NumBufferedLines * 4;
+
+/// The shared resource options this app would most favour; applied as widely as possible.
+constexpr MTLResourceOptions SharedResourceOptionsStandard = MTLResourceCPUCacheModeWriteCombined | MTLResourceStorageModeShared;
+
+/// The shared resource options used for the write-area texture; on macOS it can't be MTLResourceStorageModeShared so this is a carve-out.
+constexpr MTLResourceOptions SharedResourceOptionsTexture = MTLResourceCPUCacheModeWriteCombined | MTLResourceStorageModeManaged;
+
+#define uniforms() reinterpret_cast<Uniforms *>(_uniformsBuffer.contents)
+
+#define RangePerform(start, end, size, func)	\
+	if((start) != (end)) {	\
+		if((start) < (end)) {	\
+			func((start), (end) - (start));	\
+		} else {	\
+			func((start), (size) - (start));	\
+			if(end) {	\
+				func(0, (end));	\
+			}	\
+		}	\
+	}
+
+/// @returns the proper 1d kernel to apply a box filter around a certain point a pixel density of @c radiansPerPixel and applying an
+///		angular limit of @c cutoff. The values returned will be the first eight of a fifteen-point filter that is symmetrical around its centre.
+std::array<float, 8> boxCoefficients(float radiansPerPixel, float cutoff) {
+	std::array<float, 8> filter;
+	float total = 0.0f;
+
+	for(size_t c = 0; c < 8; ++c) {
+		// This coefficient occupies the angular window [6.5-c, 7.5-c]*radiansPerPixel.
+		const float startAngle = (6.5f - float(c)) * radiansPerPixel;
+		const float endAngle = (7.5f - float(c)) * radiansPerPixel;
+
+		float coefficient = 0.0f;
+		if(endAngle < cutoff) {
+			coefficient = 1.0f;
+		} else if(startAngle >= cutoff) {
+			coefficient = 0.0f;
+		} else {
+			coefficient = (cutoff - startAngle) / radiansPerPixel;
+		}
+		total += 2.0f * coefficient;	// All but the centre coefficient will be used twice.
+		filter[c] = coefficient;
+	}
+	total = total - filter[7];			// As per above; ensure the centre coefficient is counted only once.
+
+	for(size_t c = 0; c < 8; ++c) {
+		filter[c] /= total;
+	}
+
+	return filter;
+}
+
+}
+
+using BufferingScanTarget = Outputs::Display::BufferingScanTarget;
+
+@implementation CSScanTarget {
+	// The command queue for the device in use.
+	id<MTLCommandQueue> _commandQueue;
+
+	// Pipelines.
+	id<MTLRenderPipelineState> _composePipeline;		// For rendering to the composition texture.
+	id<MTLRenderPipelineState> _outputPipeline;			// For drawing to the frame buffer.
+	id<MTLRenderPipelineState> _copyPipeline;			// For copying from one texture to another.
+	id<MTLRenderPipelineState> _supersamplePipeline;	// For resampling from one texture to one that is 1/4 as large.
+	id<MTLRenderPipelineState> _clearPipeline;			// For applying additional inter-frame clearing (cf. the stencil).
+
+	// Buffers.
+	id<MTLBuffer> _uniformsBuffer;	// A static buffer, containing a copy of the Uniforms struct.
+	id<MTLBuffer> _scansBuffer;		// A dynamic buffer, into which the CPU writes Scans for later display.
+	id<MTLBuffer> _linesBuffer;		// A dynamic buffer, into which the CPU writes Lines for later display.
+
+	// Textures: the write area.
+	//
+	// The write area receives fragments of output from the emulated machine.
+	// So it is written by the CPU and read by the GPU.
+	id<MTLTexture> _writeAreaTexture;
+	id<MTLBuffer> _writeAreaBuffer;		// The storage underlying the write-area texture.
+	size_t _bytesPerInputPixel;			// Determines per-pixel sizing within the write-area texture.
+	size_t _totalTextureBytes;			// Holds the total size of the write-area texture.
+
+	// Textures: the frame buffer.
+	//
+	// When inter-frame blending is in use, the frame buffer contains the most recent output.
+	// Metal isn't really set up for single-buffered output, so this acts as if it were that
+	// single buffer. This texture is complete 2d data, copied directly to the display.
+	id<MTLTexture> _frameBuffer;
+	MTLRenderPassDescriptor *_frameBufferRenderPass;	// The render pass for _drawing to_ the frame buffer.
+	BOOL _dontClearFrameBuffer;
+
+	// Textures: the stencil.
+	//
+	// Scan targets recceive scans, not full frames. Those scans may not cover the entire display,
+	// either because unlit areas have been omitted or because a sync discrepancy means that the full
+	// potential vertical or horizontal width of the display isn't used momentarily.
+	//
+	// In order to manage inter-frame blending correctly in those cases, a stencil is attached to the
+	// frame buffer so that a clearing step can darken any pixels that weren't naturally painted during
+	// any frame.
+	id<MTLTexture> _frameBufferStencil;
+	id<MTLDepthStencilState> _drawStencilState;		// Always draws, sets stencil to 1.
+	id<MTLDepthStencilState> _clearStencilState;	// Draws only where stencil is 0, clears all to 0.
+
+	// Textures: the composition texture.
+	//
+	// If additional temporal processing is required (i.e. for S-Video and colour composite output),
+	// fragments from the write-area texture are assembled into the composition texture, where they
+	// properly adjoin their neighbours and everything is converted to a common clock.
+	id<MTLTexture> _compositionTexture;
+	MTLRenderPassDescriptor *_compositionRenderPass;	// The render pass for _drawing to_ the composition buffer.
+
+	enum class Pipeline {
+		/// Scans are painted directly to the frame buffer.
+		DirectToDisplay,
+		/// Scans are painted to the composition buffer, which is processed to the finalised line buffer,
+		/// from which lines are painted to the frame buffer.
+		SVideo,
+		/// Scans are painted to the composition buffer, which is processed to the separated luma buffer and then the finalised line buffer,
+		/// from which lines are painted to the frame buffer.
+		CompositeColour
+
+		// TODO: decide what to do for downard-scaled direct-to-display. Obvious options are to include lowpass
+		// filtering into the scan outputter and contine hoping that the vertical takes care of itself, or maybe
+		// to stick with DirectToDisplay but with a minimum size for the frame buffer and apply filtering from
+		// there to the screen.
+	};
+	Pipeline _pipeline;
+
+	// Textures: additional storage used when processing S-Video and composite colour input.
+	id<MTLTexture> _finalisedLineTexture;
+	id<MTLComputePipelineState> _finalisedLineState;
+	id<MTLTexture> _separatedLumaTexture;
+	id<MTLComputePipelineState> _separatedLumaState;
+	NSUInteger _lineBufferPixelsPerLine;
+
+	size_t _lineOffsetBuffer;
+	id<MTLBuffer> _lineOffsetBuffers[NumBufferedLines];	// Allocating NumBufferedLines buffers ensures these can't possibly be exhausted;
+														// for this list to be exhausted there'd have to be more draw calls in flight than
+														// there are lines for them to operate upon.
+
+	// The scan target in C++-world terms and the non-GPU storage for it.
+	BufferingScanTarget _scanTarget;
+	BufferingScanTarget::LineMetadata _lineMetadataBuffer[NumBufferedLines];
+	std::atomic_flag _isDrawing;
+
+	// Additional pipeline information.
+	size_t _lumaKernelSize;
+	size_t _chromaKernelSize;
+	std::atomic<bool> _isUsingSupersampling;
+
+	// The output view.
+	__weak MTKView *_view;
+}
+
+- (nonnull instancetype)initWithView:(nonnull MTKView *)view {
+	self = [super init];
+	if(self) {
+		_view = view;
+		_commandQueue = [view.device newCommandQueue];
+
+		// Allocate space for uniforms.
+		_uniformsBuffer = [view.device
+			newBufferWithLength:sizeof(Uniforms)
+			options:MTLResourceCPUCacheModeWriteCombined | MTLResourceStorageModeShared];
+
+		// Allocate buffers for scans and lines and for the write area texture.
+		_scansBuffer = [view.device
+			newBufferWithLength:sizeof(Outputs::Display::BufferingScanTarget::Scan)*NumBufferedScans
+			options:SharedResourceOptionsStandard];
+		_linesBuffer = [view.device
+			newBufferWithLength:sizeof(Outputs::Display::BufferingScanTarget::Line)*NumBufferedLines
+			options:SharedResourceOptionsStandard];
+		_writeAreaBuffer = [view.device
+			newBufferWithLength:BufferingScanTarget::WriteAreaWidth*BufferingScanTarget::WriteAreaHeight*4
+			options:SharedResourceOptionsTexture];
+
+		// Install all that storage in the buffering scan target.
+		_scanTarget.set_write_area(reinterpret_cast<uint8_t *>(_writeAreaBuffer.contents));
+		_scanTarget.set_line_buffer(reinterpret_cast<BufferingScanTarget::Line *>(_linesBuffer.contents), _lineMetadataBuffer, NumBufferedLines);
+		_scanTarget.set_scan_buffer(reinterpret_cast<BufferingScanTarget::Scan *>(_scansBuffer.contents), NumBufferedScans);
+
+		// Generate copy and clear pipelines.
+		id<MTLLibrary> library = [_view.device newDefaultLibrary];
+		MTLRenderPipelineDescriptor *const pipelineDescriptor = [[MTLRenderPipelineDescriptor alloc] init];
+		pipelineDescriptor.colorAttachments[0].pixelFormat = _view.colorPixelFormat;
+		pipelineDescriptor.vertexFunction = [library newFunctionWithName:@"copyVertex"];
+		pipelineDescriptor.fragmentFunction = [library newFunctionWithName:@"copyFragment"];
+		_copyPipeline = [_view.device newRenderPipelineStateWithDescriptor:pipelineDescriptor error:nil];
+
+		pipelineDescriptor.fragmentFunction = [library newFunctionWithName:@"interpolateFragment"];
+		_supersamplePipeline = [_view.device newRenderPipelineStateWithDescriptor:pipelineDescriptor error:nil];
+
+		pipelineDescriptor.fragmentFunction = [library newFunctionWithName:@"clearFragment"];
+		pipelineDescriptor.stencilAttachmentPixelFormat = MTLPixelFormatStencil8;
+		_clearPipeline = [_view.device newRenderPipelineStateWithDescriptor:pipelineDescriptor error:nil];
+
+		// Clear stencil: always write the reference value (of 0), but draw only where the stencil already
+		// had that value.
+		MTLDepthStencilDescriptor *depthStencilDescriptor = [[MTLDepthStencilDescriptor alloc] init];
+		depthStencilDescriptor.frontFaceStencil.stencilCompareFunction = MTLCompareFunctionEqual;
+		depthStencilDescriptor.frontFaceStencil.depthStencilPassOperation = MTLStencilOperationReplace;
+		depthStencilDescriptor.frontFaceStencil.stencilFailureOperation = MTLStencilOperationReplace;
+		_clearStencilState = [view.device newDepthStencilStateWithDescriptor:depthStencilDescriptor];
+
+		// Allocate a large number of single-int buffers, for supplying offsets to the compute shaders.
+		// There's a ridiculous amount of overhead in this, but it avoids allocations during drawing,
+		// and a single int per instance is all I need.
+		for(size_t c = 0; c < NumBufferedLines; ++c) {
+			_lineOffsetBuffers[c] = [_view.device newBufferWithLength:sizeof(int) options:SharedResourceOptionsStandard];
+		}
+
+		// Ensure the is-drawing flag is initially clear.
+		_isDrawing.clear();
+
+		// Set initial aspect-ratio multiplier and generate buffers.
+		[self mtkView:view drawableSizeWillChange:view.drawableSize];
+	}
+
+	return self;
+}
+
+/*!
+ @method mtkView:drawableSizeWillChange:
+ @abstract Called whenever the drawableSize of the view will change
+ @discussion Delegate can recompute view and projection matricies or regenerate any buffers to be compatible with the new view size or resolution
+ @param view MTKView which called this method
+ @param size New drawable size in pixels
+ */
+- (void)mtkView:(nonnull MTKView *)view drawableSizeWillChange:(CGSize)size {
+	[self setAspectRatio];
+
+	@synchronized(self) {
+		// Always [re]try multisampling upon a resize.
+		_scanTarget.display_metrics_.announce_did_resize();
+		_isUsingSupersampling = true;
+		[self updateSizeBuffersToSize:size];
+	}
+}
+
+- (void)updateSizeBuffers {
+	@synchronized(self) {
+		[self updateSizeBuffersToSize:_view.drawableSize];
+	}
+}
+
+- (id<MTLCommandBuffer>)copyTexture:(id<MTLTexture>)source to:(id<MTLTexture>)destination {
+	MTLRenderPassDescriptor *const copyTextureDescriptor = [[MTLRenderPassDescriptor alloc] init];
+	copyTextureDescriptor.colorAttachments[0].texture = destination;
+	copyTextureDescriptor.colorAttachments[0].loadAction = MTLLoadActionDontCare;
+	copyTextureDescriptor.colorAttachments[0].storeAction = MTLStoreActionStore;
+
+	id<MTLCommandBuffer> commandBuffer = [_commandQueue commandBuffer];
+	id<MTLRenderCommandEncoder> encoder = [commandBuffer renderCommandEncoderWithDescriptor:copyTextureDescriptor];
+
+	[encoder setRenderPipelineState:_copyPipeline];
+	[encoder setVertexTexture:source atIndex:0];
+	[encoder setFragmentTexture:source atIndex:0];
+
+	[encoder drawPrimitives:MTLPrimitiveTypeTriangleStrip vertexStart:0 vertexCount:4];
+	[encoder endEncoding];
+	[commandBuffer commit];
+
+	return commandBuffer;
+}
+
+- (void)updateSizeBuffersToSize:(CGSize)size {
+	const NSUInteger frameBufferWidth = NSUInteger(size.width * _view.layer.contentsScale) * (_isUsingSupersampling ? 2 : 1);
+	const NSUInteger frameBufferHeight = NSUInteger(size.height * _view.layer.contentsScale) * (_isUsingSupersampling ? 2 : 1);
+
+	// Generate a framebuffer and a stencil.
+	MTLTextureDescriptor *const textureDescriptor = [MTLTextureDescriptor
+		texture2DDescriptorWithPixelFormat:_view.colorPixelFormat
+		width:frameBufferWidth
+		height:frameBufferHeight
+		mipmapped:NO];
+	textureDescriptor.usage = MTLTextureUsageRenderTarget | MTLTextureUsageShaderRead;
+	textureDescriptor.resourceOptions = MTLResourceStorageModePrivate;
+	id<MTLTexture> _oldFrameBuffer = _frameBuffer;
+	_frameBuffer = [_view.device newTextureWithDescriptor:textureDescriptor];
+
+	MTLTextureDescriptor *const stencilTextureDescriptor = [MTLTextureDescriptor
+		texture2DDescriptorWithPixelFormat:MTLPixelFormatStencil8
+		width:frameBufferWidth
+		height:frameBufferHeight
+		mipmapped:NO];
+	stencilTextureDescriptor.usage = MTLTextureUsageRenderTarget;
+	stencilTextureDescriptor.resourceOptions = MTLResourceStorageModePrivate;
+	_frameBufferStencil = [_view.device newTextureWithDescriptor:stencilTextureDescriptor];
+
+	// Generate a render pass with that framebuffer and stencil.
+	_frameBufferRenderPass = [[MTLRenderPassDescriptor alloc] init];
+	_frameBufferRenderPass.colorAttachments[0].texture = _frameBuffer;
+	_frameBufferRenderPass.colorAttachments[0].loadAction = MTLLoadActionLoad;
+	_frameBufferRenderPass.colorAttachments[0].storeAction = MTLStoreActionStore;
+
+	_frameBufferRenderPass.stencilAttachment.clearStencil = 0;
+	_frameBufferRenderPass.stencilAttachment.texture = _frameBufferStencil;
+	_frameBufferRenderPass.stencilAttachment.loadAction = MTLLoadActionLoad;
+	_frameBufferRenderPass.stencilAttachment.storeAction = MTLStoreActionStore;
+
+	// Establish intended stencil useage; it's only to track which pixels haven't been painted
+	// at all at the end of every frame. So: always paint, and replace the stored stencil value
+	// (which is seeded as 0) with the nominated one (a 1).
+	MTLDepthStencilDescriptor *depthStencilDescriptor = [[MTLDepthStencilDescriptor alloc] init];
+	depthStencilDescriptor.frontFaceStencil.stencilCompareFunction = MTLCompareFunctionAlways;
+	depthStencilDescriptor.frontFaceStencil.depthStencilPassOperation = MTLStencilOperationReplace;
+	_drawStencilState = [_view.device newDepthStencilStateWithDescriptor:depthStencilDescriptor];
+
+	// Draw from _oldFrameBuffer to _frameBuffer.
+	if(_oldFrameBuffer) {
+		[self copyTexture:_oldFrameBuffer to:_frameBuffer];
+
+		// Don't clear the framebuffer at the end of this frame.
+		_dontClearFrameBuffer = YES;
+	}
+}
+
+- (BOOL)shouldApplyGamma {
+	return fabsf(float(uniforms()->outputGamma) - 1.0f) > 0.01f;
+}
+
+- (void)updateModalBuffers {
+	// Build a descriptor for any intermediate line texture.
+	MTLTextureDescriptor *const lineTextureDescriptor = [MTLTextureDescriptor
+		texture2DDescriptorWithPixelFormat:MTLPixelFormatBGRA8Unorm
+		width:2048		// This 'should do'.
+		height:NumBufferedLines
+		mipmapped:NO];
+	lineTextureDescriptor.resourceOptions = MTLResourceStorageModePrivate;
+
+	if(_pipeline == Pipeline::DirectToDisplay) {
+		// Buffers are not required when outputting direct to display; so if this isn't that then release anything
+		// currently being held and return.
+		_finalisedLineTexture = nil;
+		_finalisedLineState = nil;
+		_separatedLumaTexture = nil;
+		_separatedLumaState = nil;
+		_compositionTexture = nil;
+		_compositionRenderPass = nil;
+		return;
+	}
+
+	// Create a composition texture if one does not yet exist.
+	if(!_compositionTexture) {
+		lineTextureDescriptor.usage = MTLTextureUsageRenderTarget | MTLTextureUsageShaderRead;
+		_compositionTexture = [_view.device newTextureWithDescriptor:lineTextureDescriptor];
+	}
+
+	// Grab the shader library.
+	id<MTLLibrary> library = [_view.device newDefaultLibrary];
+	lineTextureDescriptor.usage = MTLTextureUsageShaderWrite | MTLTextureUsageShaderRead;
+
+	// The finalised texture will definitely exist, and may or may not require a gamma conversion when written to.
+	if(!_finalisedLineTexture) {
+		_finalisedLineTexture = [_view.device newTextureWithDescriptor:lineTextureDescriptor];
+
+		NSString *const kernelFunction = [self shouldApplyGamma] ? @"filterChromaKernelWithGamma" : @"filterChromaKernelNoGamma";
+		_finalisedLineState = [_view.device newComputePipelineStateWithFunction:[library newFunctionWithName:kernelFunction] error:nil];
+
+		// Ensure finalised line texture is initially clear.
+		id<MTLComputePipelineState> clearPipeline = [_view.device newComputePipelineStateWithFunction:[library newFunctionWithName:@"clearKernel"] error:nil];
+		id<MTLCommandBuffer> commandBuffer = [_commandQueue commandBuffer];
+		id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];
+
+		[self dispatchComputeCommandEncoder:computeEncoder pipelineState:clearPipeline width:lineTextureDescriptor.width height:lineTextureDescriptor.height offsetBuffer:[self bufferForOffset:0]];
+
+		[computeEncoder endEncoding];
+		[commandBuffer commit];
+	}
+
+	// A luma separation texture will exist only for composite colour.
+	if(_pipeline == Pipeline::CompositeColour) {
+		if(!_separatedLumaTexture) {
+			_separatedLumaTexture = [_view.device newTextureWithDescriptor:lineTextureDescriptor];
+
+			NSString *kernelFunction;
+			switch(_lumaKernelSize) {
+				default:	kernelFunction = @"separateLumaKernel15";	break;
+				case 9:		kernelFunction = @"separateLumaKernel9";	break;
+				case 7:		kernelFunction = @"separateLumaKernel7";	break;
+				case 1:
+				case 3:
+				case 5:		kernelFunction = @"separateLumaKernel5";	break;
+			}
+
+			_separatedLumaState = [_view.device newComputePipelineStateWithFunction:[library newFunctionWithName:kernelFunction] error:nil];
+		}
+	} else {
+		_separatedLumaTexture = nil;
+	}
+}
+
+- (void)setAspectRatio {
+	const auto modals = _scanTarget.modals();
+	const auto viewAspectRatio = (_view.bounds.size.width / _view.bounds.size.height);
+	simd::float3x3 sourceToDisplay{1.0f};
+
+	// The starting coordinate space is [0, 1].
+
+	// Move the centre of the cropping rectangle to the centre of the display.
+	{
+		simd::float3x3 recentre{1.0f};
+		recentre.columns[2][0] = 0.5f - (modals.visible_area.origin.x + modals.visible_area.size.width * 0.5f);
+		recentre.columns[2][1] = 0.5f - (modals.visible_area.origin.y + modals.visible_area.size.height * 0.5f);
+		sourceToDisplay = recentre * sourceToDisplay;
+	}
+
+	// Convert from the internal [0, 1] to centred [-1, 1] (i.e. Metal's eye coordinates, though also appropriate
+	// for the zooming step that follows).
+	{
+		simd::float3x3 convertToEye;
+		convertToEye.columns[0][0] = 2.0f;
+		convertToEye.columns[1][1] = -2.0f;
+		convertToEye.columns[2][0] = -1.0f;
+		convertToEye.columns[2][1] = 1.0f;
+		convertToEye.columns[2][2] = 1.0f;
+		sourceToDisplay = convertToEye * sourceToDisplay;
+	}
+
+	// Determine the correct zoom level. This is a combination of (i) the necessary horizontal stretch to produce a proper
+	// aspect ratio; and (ii) the necessary zoom from there to either fit the visible area width or height as per a decision
+	// on letterboxing or pillarboxing.
+	const float aspectRatioStretch = float(modals.aspect_ratio / viewAspectRatio);
+	const float fitWidthZoom = 1.0f / (float(modals.visible_area.size.width) * aspectRatioStretch);
+	const float fitHeightZoom = 1.0f / float(modals.visible_area.size.height);
+	const float zoom = std::min(fitWidthZoom, fitHeightZoom);
+
+	// Convert from there to the proper aspect ratio by stretching or compressing width.
+	// After this the output is exactly centred, filling the vertical space and being as wide or slender as it likes.
+	{
+		simd::float3x3 applyAspectRatio{1.0f};
+		applyAspectRatio.columns[0][0] = aspectRatioStretch * zoom;
+		applyAspectRatio.columns[1][1] = zoom;
+		sourceToDisplay = applyAspectRatio * sourceToDisplay;
+	}
+
+	// Store.
+	uniforms()->sourcetoDisplay = sourceToDisplay;
+}
+
+- (void)setModals:(const Outputs::Display::ScanTarget::Modals &)modals {
+	//
+	// Populate uniforms.
+	//
+	uniforms()->scale[0] = modals.output_scale.x;
+	uniforms()->scale[1] = modals.output_scale.y;
+	uniforms()->lineWidth = 1.05f / modals.expected_vertical_lines;
+	[self setAspectRatio];
+
+	const auto toRGB = to_rgb_matrix(modals.composite_colour_space);
+	uniforms()->toRGB = simd::float3x3(
+		simd::float3{toRGB[0], toRGB[1], toRGB[2]},
+		simd::float3{toRGB[3], toRGB[4], toRGB[5]},
+		simd::float3{toRGB[6], toRGB[7], toRGB[8]}
+	);
+
+	const auto fromRGB = from_rgb_matrix(modals.composite_colour_space);
+	uniforms()->fromRGB = simd::float3x3(
+		simd::float3{fromRGB[0], fromRGB[1], fromRGB[2]},
+		simd::float3{fromRGB[3], fromRGB[4], fromRGB[5]},
+		simd::float3{fromRGB[6], fromRGB[7], fromRGB[8]}
+	);
+
+	// This is fixed for now; consider making it a function of frame rate and/or of whether frame syncing
+	// is ongoing (which would require a way to signal that to this scan target).
+	uniforms()->outputAlpha = __fp16(0.64f);
+	uniforms()->outputMultiplier = __fp16(modals.brightness);
+
+	const float displayGamma = 2.2f;	// This is assumed.
+	uniforms()->outputGamma = __fp16(displayGamma / modals.intended_gamma);
+
+
+
+	//
+	// Generate input texture.
+	//
+	MTLPixelFormat pixelFormat;
+	_bytesPerInputPixel = size_for_data_type(modals.input_data_type);
+	if(data_type_is_normalised(modals.input_data_type)) {
+		switch(_bytesPerInputPixel) {
+			default:
+			case 1: pixelFormat = MTLPixelFormatR8Unorm;	break;
+			case 2: pixelFormat = MTLPixelFormatRG8Unorm;	break;
+			case 4: pixelFormat = MTLPixelFormatRGBA8Unorm;	break;
+		}
+	} else {
+		switch(_bytesPerInputPixel) {
+			default:
+			case 1: pixelFormat = MTLPixelFormatR8Uint;		break;
+			case 2: pixelFormat = MTLPixelFormatRG8Uint;	break;
+			case 4: pixelFormat = MTLPixelFormatRGBA8Uint;	break;
+		}
+	}
+	MTLTextureDescriptor *const textureDescriptor = [MTLTextureDescriptor
+		texture2DDescriptorWithPixelFormat:pixelFormat
+		width:BufferingScanTarget::WriteAreaWidth
+		height:BufferingScanTarget::WriteAreaHeight
+		mipmapped:NO];
+	textureDescriptor.resourceOptions = SharedResourceOptionsTexture;
+	if(@available(macOS 10.14, *)) {
+		textureDescriptor.allowGPUOptimizedContents = NO;
+	}
+
+	// TODO: the call below is the only reason why this project now requires macOS 10.13; is it all that helpful versus just uploading each frame?
+	const NSUInteger bytesPerRow = BufferingScanTarget::WriteAreaWidth * _bytesPerInputPixel;
+	_writeAreaTexture = [_writeAreaBuffer
+		newTextureWithDescriptor:textureDescriptor
+		offset:0
+		bytesPerRow:bytesPerRow];
+	_totalTextureBytes = bytesPerRow * BufferingScanTarget::WriteAreaHeight;
+
+
+
+	//
+	// Generate scan pipeline.
+	//
+	id<MTLLibrary> library = [_view.device newDefaultLibrary];
+	MTLRenderPipelineDescriptor *pipelineDescriptor = [[MTLRenderPipelineDescriptor alloc] init];
+
+	// Occasions when the composition buffer isn't required are slender: the output must be neither RGB nor composite monochrome.
+	const bool isComposition =
+		modals.display_type != Outputs::Display::DisplayType::RGB &&
+		modals.display_type != Outputs::Display::DisplayType::CompositeMonochrome;
+	const bool isSVideoOutput = modals.display_type == Outputs::Display::DisplayType::SVideo;
+
+	if(!isComposition) {
+		_pipeline = Pipeline::DirectToDisplay;
+	} else {
+		_pipeline = isSVideoOutput ? Pipeline::SVideo : Pipeline::CompositeColour;
+	}
+
+	struct FragmentSamplerDictionary {
+		/// Fragment shader that outputs to the composition buffer for composite processing.
+		NSString *const compositionComposite;
+		/// Fragment shader that outputs to the composition buffer for S-Video processing.
+		NSString *const compositionSVideo;
+
+		/// Fragment shader that outputs directly as monochrome composite.
+		NSString *const directComposite;
+		/// Fragment shader that outputs directly as monochrome composite, with gamma correction.
+		NSString *const directCompositeWithGamma;
+		/// Fragment shader that outputs directly as RGB.
+		NSString *const directRGB;
+		/// Fragment shader that outputs directly as RGB, with gamma correction.
+		NSString *const directRGBWithGamma;
+	};
+	const FragmentSamplerDictionary samplerDictionary[8] = {
+		// Composite formats.
+		{@"compositeSampleLuminance1", 				nil,	@"sampleLuminance1",				@"sampleLuminance1",						@"sampleLuminance1",				@"sampleLuminance1"},
+		{@"compositeSampleLuminance8", 				nil,	@"sampleLuminance8", 				@"sampleLuminance8WithGamma",				@"sampleLuminance8", 				@"sampleLuminance8WithGamma"},
+		{@"compositeSamplePhaseLinkedLuminance8", 	nil,	@"samplePhaseLinkedLuminance8",		@"samplePhaseLinkedLuminance8WithGamma",	@"samplePhaseLinkedLuminance8",		@"samplePhaseLinkedLuminance8WithGamma"},
+
+		// S-Video formats.
+		{@"compositeSampleLuminance8Phase8", @"sampleLuminance8Phase8", @"directCompositeSampleLuminance8Phase8", @"directCompositeSampleLuminance8Phase8WithGamma", @"directCompositeSampleLuminance8Phase8", @"directCompositeSampleLuminance8Phase8WithGamma"},
+
+		// RGB formats.
+		{@"compositeSampleRed1Green1Blue1", @"svideoSampleRed1Green1Blue1", @"directCompositeSampleRed1Green1Blue1", @"directCompositeSampleRed1Green1Blue1WithGamma", @"sampleRed1Green1Blue1", @"sampleRed1Green1Blue1"},
+		{@"compositeSampleRed2Green2Blue2", @"svideoSampleRed2Green2Blue2", @"directCompositeSampleRed2Green2Blue2", @"directCompositeSampleRed2Green2Blue2WithGamma", @"sampleRed2Green2Blue2", @"sampleRed2Green2Blue2WithGamma"},
+		{@"compositeSampleRed4Green4Blue4", @"svideoSampleRed4Green4Blue4", @"directCompositeSampleRed4Green4Blue4", @"directCompositeSampleRed4Green4Blue4WithGamma", @"sampleRed4Green4Blue4", @"sampleRed4Green4Blue4WithGamma"},
+		{@"compositeSampleRed8Green8Blue8", @"svideoSampleRed8Green8Blue8", @"directCompositeSampleRed8Green8Blue8", @"directCompositeSampleRed8Green8Blue8WithGamma", @"sampleRed8Green8Blue8", @"sampleRed8Green8Blue8WithGamma"},
+	};
+
+#ifndef NDEBUG
+	// Do a quick check that all the shaders named above are defined in the Metal code. I don't think this is possible at compile time.
+	for(int c = 0; c < 8; ++c) {
+#define Test(x)	if(samplerDictionary[c].x)	assert([library newFunctionWithName:samplerDictionary[c].x]);
+		Test(compositionComposite);
+		Test(compositionSVideo);
+		Test(directComposite);
+		Test(directCompositeWithGamma);
+		Test(directRGB);
+		Test(directRGBWithGamma);
+#undef Test
+	}
+#endif
+
+	uniforms()->cyclesMultiplier = 1.0f;
+	if(_pipeline != Pipeline::DirectToDisplay) {
+		// Pick a suitable cycle multiplier.
+		const float minimumSize = 4.0f * float(modals.colour_cycle_numerator) / float(modals.colour_cycle_denominator);
+		while(uniforms()->cyclesMultiplier * modals.cycles_per_line < minimumSize) {
+			uniforms()->cyclesMultiplier += 1.0f;
+
+			if(uniforms()->cyclesMultiplier * modals.cycles_per_line > 2048) {
+				uniforms()->cyclesMultiplier -= 1.0f;
+				break;
+			}
+		}
+
+		// Create suitable filters.
+		_lineBufferPixelsPerLine = NSUInteger(modals.cycles_per_line) * NSUInteger(uniforms()->cyclesMultiplier);
+		const float colourCyclesPerLine = float(modals.colour_cycle_numerator) / float(modals.colour_cycle_denominator);
+
+		// Compute radians per pixel.
+		const float radiansPerPixel = (colourCyclesPerLine * 3.141592654f * 2.0f) / float(_lineBufferPixelsPerLine);
+
+		// Generate the chrominance filter.
+		{
+			simd::float3 firCoefficients[8];
+			const auto chromaCoefficients = boxCoefficients(radiansPerPixel, 3.141592654f);
+			_chromaKernelSize = 15;
+			for(size_t c = 0; c < 8; ++c) {
+				firCoefficients[c].y = firCoefficients[c].z = (isSVideoOutput ? 2.0f : 1.0f) * chromaCoefficients[c];
+				firCoefficients[c].x = 0.0f;
+				if(fabsf(chromaCoefficients[c]) < 0.01f) {
+					_chromaKernelSize -= 2;
+				}
+			}
+			firCoefficients[7].x = 1.0f;
+
+			// Luminance will be very soft as a result of the separation phase; apply a sharpen filter to try to undo that.
+			// This is applied separately because the first composite processing step is going to select between the nominal
+			// chroma and luma parts to take the place of luminance depending on whether a colour burst was found, and high-pass
+			// filtering the chrominance channel would be visually detrimental.
+			//
+			// The low cut off ['Hz' but per line, not per second] is somewhat arbitrary.
+			if(!isSVideoOutput) {
+				SignalProcessing::FIRFilter sharpenFilter(15, float(_lineBufferPixelsPerLine), 40.0f, colourCyclesPerLine);
+				const auto sharpen = sharpenFilter.get_coefficients();
+				size_t sharpenFilterSize = 15;
+				bool isStart = true;
+				for(size_t c = 0; c < 8; ++c) {
+					firCoefficients[c].x = sharpen[c];
+					if(fabsf(sharpen[c]) > 0.01f) isStart = false;
+					if(isStart) sharpenFilterSize -= 2;
+				}
+				_chromaKernelSize = std::max(_chromaKernelSize, sharpenFilterSize);
+			}
+
+			// Convert to half-size floats.
+			for(size_t c = 0; c < 8; ++c) {
+				uniforms()->chromaKernel[c] = firCoefficients[c];
+			}
+		}
+
+		// Generate the luminance separation filter and determine its required size.
+		{
+			auto *const filter = uniforms()->lumaKernel;
+			const auto coefficients = boxCoefficients(radiansPerPixel, 3.141592654f);
+			_lumaKernelSize = 15;
+			for(size_t c = 0; c < 8; ++c) {
+				filter[c] = __fp16(coefficients[c]);
+				if(fabsf(coefficients[c]) < 0.01f) {
+					_lumaKernelSize -= 2;
+				}
+			}
+		}
+	}
+
+	// Update intermediate storage.
+	[self updateModalBuffers];
+
+	if(_pipeline != Pipeline::DirectToDisplay) {
+		// Create the composition render pass.
+ 		pipelineDescriptor.colorAttachments[0].pixelFormat = _compositionTexture.pixelFormat;
+		pipelineDescriptor.vertexFunction = [library newFunctionWithName:@"scanToComposition"];
+		pipelineDescriptor.fragmentFunction =
+			[library newFunctionWithName:isSVideoOutput ? samplerDictionary[int(modals.input_data_type)].compositionSVideo : samplerDictionary[int(modals.input_data_type)].compositionComposite];
+
+		_composePipeline = [_view.device newRenderPipelineStateWithDescriptor:pipelineDescriptor error:nil];
+
+		_compositionRenderPass = [[MTLRenderPassDescriptor alloc] init];
+		_compositionRenderPass.colorAttachments[0].texture = _compositionTexture;
+		_compositionRenderPass.colorAttachments[0].loadAction = MTLLoadActionClear;
+		_compositionRenderPass.colorAttachments[0].storeAction = MTLStoreActionStore;
+		_compositionRenderPass.colorAttachments[0].clearColor = MTLClearColorMake(0.0, 0.5, 0.5, 0.3);
+	}
+
+	// Build the output pipeline.
+	pipelineDescriptor.colorAttachments[0].pixelFormat = _view.colorPixelFormat;
+	pipelineDescriptor.vertexFunction = [library newFunctionWithName:_pipeline == Pipeline::DirectToDisplay ? @"scanToDisplay" : @"lineToDisplay"];
+
+	if(_pipeline != Pipeline::DirectToDisplay) {
+		pipelineDescriptor.fragmentFunction = [library newFunctionWithName:@"interpolateFragment"];
+	} else {
+		const bool isRGBOutput = modals.display_type == Outputs::Display::DisplayType::RGB;
+
+		NSString *shaderName;
+		if(isRGBOutput) {
+			shaderName = [self shouldApplyGamma] ? samplerDictionary[int(modals.input_data_type)].directRGBWithGamma : samplerDictionary[int(modals.input_data_type)].directRGB;
+		} else {
+			shaderName = [self shouldApplyGamma] ? samplerDictionary[int(modals.input_data_type)].directCompositeWithGamma : samplerDictionary[int(modals.input_data_type)].directComposite;
+		}
+		pipelineDescriptor.fragmentFunction = [library newFunctionWithName:shaderName];
+	}
+
+	// Enable blending.
+	pipelineDescriptor.colorAttachments[0].blendingEnabled = YES;
+	pipelineDescriptor.colorAttachments[0].sourceRGBBlendFactor = MTLBlendFactorSourceAlpha;
+	pipelineDescriptor.colorAttachments[0].destinationRGBBlendFactor = MTLBlendFactorOneMinusSourceAlpha;
+
+	// Set stencil format.
+	pipelineDescriptor.stencilAttachmentPixelFormat = MTLPixelFormatStencil8;
+
+	// Finish.
+	_outputPipeline = [_view.device newRenderPipelineStateWithDescriptor:pipelineDescriptor error:nil];
+}
+
+- (void)outputFrom:(size_t)start to:(size_t)end commandBuffer:(id<MTLCommandBuffer>)commandBuffer {
+	if(start == end) return;
+
+	// Generate a command encoder for the view.
+	id<MTLRenderCommandEncoder> encoder = [commandBuffer renderCommandEncoderWithDescriptor:_frameBufferRenderPass];
+
+	// Final output. Could be scans or lines.
+	[encoder setRenderPipelineState:_outputPipeline];
+
+	if(_pipeline != Pipeline::DirectToDisplay) {
+		[encoder setFragmentTexture:_finalisedLineTexture atIndex:0];
+		[encoder setVertexBuffer:_linesBuffer offset:0 atIndex:0];
+	} else {
+		[encoder setFragmentTexture:_writeAreaTexture atIndex:0];
+		[encoder setVertexBuffer:_scansBuffer offset:0 atIndex:0];
+	}
+	[encoder setVertexBuffer:_uniformsBuffer offset:0 atIndex:1];
+	[encoder setFragmentBuffer:_uniformsBuffer offset:0 atIndex:0];
+
+	[encoder setDepthStencilState:_drawStencilState];
+	[encoder setStencilReferenceValue:1];
+#ifndef NDEBUG
+	// Quick aid for debugging: the stencil test is predicated on front-facing pixels, so make sure they're
+	// being generated.
+	[encoder setCullMode:MTLCullModeBack];
+#endif
+
+#define OutputStrips(start, size)	[encoder drawPrimitives:MTLPrimitiveTypeTriangleStrip vertexStart:0 vertexCount:4 instanceCount:size baseInstance:start]
+	RangePerform(start, end, _pipeline != Pipeline::DirectToDisplay ? NumBufferedLines : NumBufferedScans, OutputStrips);
+#undef OutputStrips
+
+	// Complete encoding.
+	[encoder endEncoding];
+}
+
+- (void)outputFrameCleanerToCommandBuffer:(id<MTLCommandBuffer>)commandBuffer {
+	// Generate a command encoder for the view.
+	id<MTLRenderCommandEncoder> encoder = [commandBuffer renderCommandEncoderWithDescriptor:_frameBufferRenderPass];
+
+	[encoder setRenderPipelineState:_clearPipeline];
+	[encoder setDepthStencilState:_clearStencilState];
+	[encoder setStencilReferenceValue:0];
+
+	[encoder setVertexTexture:_frameBuffer atIndex:0];
+	[encoder setFragmentTexture:_frameBuffer atIndex:0];
+	[encoder setFragmentBuffer:_uniformsBuffer offset:0 atIndex:0];
+
+	[encoder drawPrimitives:MTLPrimitiveTypeTriangleStrip vertexStart:0 vertexCount:4];
+	[encoder endEncoding];
+}
+
+- (void)composeOutputArea:(const BufferingScanTarget::OutputArea &)outputArea commandBuffer:(id<MTLCommandBuffer>)commandBuffer {
+	// Output all scans to the composition buffer.
+	const id<MTLRenderCommandEncoder> encoder = [commandBuffer renderCommandEncoderWithDescriptor:_compositionRenderPass];
+	[encoder setRenderPipelineState:_composePipeline];
+
+	[encoder setVertexBuffer:_scansBuffer offset:0 atIndex:0];
+	[encoder setVertexBuffer:_uniformsBuffer offset:0 atIndex:1];
+	[encoder setVertexTexture:_compositionTexture atIndex:0];
+
+	[encoder setFragmentBuffer:_uniformsBuffer offset:0 atIndex:0];
+	[encoder setFragmentTexture:_writeAreaTexture atIndex:0];
+
+#define OutputScans(start, size)	[encoder drawPrimitives:MTLPrimitiveTypeLine vertexStart:0 vertexCount:2 instanceCount:size baseInstance:start]
+	RangePerform(outputArea.start.scan, outputArea.end.scan, NumBufferedScans, OutputScans);
+#undef OutputScans
+	[encoder endEncoding];
+}
+
+- (id<MTLBuffer>)bufferForOffset:(size_t)offset {
+	// Store and apply the offset.
+	const auto buffer = _lineOffsetBuffers[_lineOffsetBuffer];
+	*(reinterpret_cast<int *>(_lineOffsetBuffers[_lineOffsetBuffer].contents)) = int(offset);
+	_lineOffsetBuffer = (_lineOffsetBuffer + 1) % NumBufferedLines;
+	return buffer;
+}
+
+- (void)dispatchComputeCommandEncoder:(id<MTLComputeCommandEncoder>)encoder pipelineState:(id<MTLComputePipelineState>)pipelineState width:(NSUInteger)width height:(NSUInteger)height offsetBuffer:(id<MTLBuffer>)offsetBuffer {
+	[encoder setBuffer:offsetBuffer offset:0 atIndex:1];
+
+	// This follows the recommendations at https://developer.apple.com/documentation/metal/calculating_threadgroup_and_grid_sizes ;
+	// I currently have no independent opinion whatsoever.
+	const MTLSize threadsPerThreadgroup = MTLSizeMake(
+		pipelineState.threadExecutionWidth,
+		pipelineState.maxTotalThreadsPerThreadgroup / pipelineState.threadExecutionWidth,
+		1
+	);
+	const MTLSize threadsPerGrid = MTLSizeMake(width, height, 1);
+
+	// Set the pipeline state and dispatch the drawing. Which may slightly overdraw.
+	[encoder setComputePipelineState:pipelineState];
+	[encoder dispatchThreads:threadsPerGrid threadsPerThreadgroup:threadsPerThreadgroup];
+}
+
+- (void)updateFrameBuffer {
+	// TODO: rethink BufferingScanTarget::perform. Is it now really just for guarding the modals?
+	_scanTarget.perform([=] {
+		const Outputs::Display::ScanTarget::Modals *const newModals = _scanTarget.new_modals();
+		if(newModals) {
+			[self setModals:*newModals];
+		}
+	});
+
+	@synchronized(self) {
+		if(!_frameBufferRenderPass) return;
+
+		const auto outputArea = _scanTarget.get_output_area();
+
+		if(outputArea.end.line != outputArea.start.line) {
+
+			// Ensure texture changes are noted.
+			const auto writeAreaModificationStart = size_t(outputArea.start.write_area_x + outputArea.start.write_area_y * 2048) * _bytesPerInputPixel;
+			const auto writeAreaModificationEnd = size_t(outputArea.end.write_area_x + outputArea.end.write_area_y * 2048) * _bytesPerInputPixel;
+#define FlushRegion(start, size)	[_writeAreaBuffer didModifyRange:NSMakeRange(start, size)]
+			RangePerform(writeAreaModificationStart, writeAreaModificationEnd, _totalTextureBytes, FlushRegion);
+#undef FlushRegion
+
+			// Obtain a source for render command encoders.
+			id<MTLCommandBuffer> commandBuffer = [_commandQueue commandBuffer];
+
+			//
+			// Drawing algorithm used below, in broad terms:
+			//
+			// Maintain a persistent buffer of current CRT state.
+			//
+			// During each frame, paint to the persistent buffer anything new. Update a stencil buffer to track
+			// every pixel so-far touched.
+			//
+			// At the end of the frame, draw a 'frame cleaner', which is a whole-screen rect that paints over
+			// only those areas that the stencil buffer indicates weren't painted this frame.
+			//
+			// Hence every pixel is touched every frame, regardless of the machine's output.
+			//
+
+			switch(_pipeline) {
+				case Pipeline::DirectToDisplay: {
+					// Output scans directly, broken up by frame.
+					size_t line = outputArea.start.line;
+					size_t scan = outputArea.start.scan;
+					while(line != outputArea.end.line) {
+						if(_lineMetadataBuffer[line].is_first_in_frame) {
+							[self outputFrom:scan to:_lineMetadataBuffer[line].first_scan commandBuffer:commandBuffer];
+							scan = _lineMetadataBuffer[line].first_scan;
+
+							if(_lineMetadataBuffer[line].previous_frame_was_complete && !_dontClearFrameBuffer) {
+								[self outputFrameCleanerToCommandBuffer:commandBuffer];
+							}
+							_dontClearFrameBuffer = NO;
+						}
+						line = (line + 1) % NumBufferedLines;
+					}
+					[self outputFrom:scan to:outputArea.end.scan commandBuffer:commandBuffer];
+				} break;
+
+				case Pipeline::CompositeColour:
+				case Pipeline::SVideo: {
+					// Build the composition buffer.
+					[self composeOutputArea:outputArea commandBuffer:commandBuffer];
+
+					if(_pipeline == Pipeline::SVideo) {
+						// Filter from composition to the finalised line texture.
+						id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];
+						[computeEncoder setTexture:_compositionTexture atIndex:0];
+						[computeEncoder setTexture:_finalisedLineTexture atIndex:1];
+						[computeEncoder setBuffer:_uniformsBuffer offset:0 atIndex:0];
+
+						if(outputArea.end.line > outputArea.start.line) {
+							[self dispatchComputeCommandEncoder:computeEncoder pipelineState:_finalisedLineState width:_lineBufferPixelsPerLine height:outputArea.end.line - outputArea.start.line offsetBuffer:[self bufferForOffset:outputArea.start.line]];
+						} else {
+							[self dispatchComputeCommandEncoder:computeEncoder pipelineState:_finalisedLineState width:_lineBufferPixelsPerLine height:NumBufferedLines - outputArea.start.line offsetBuffer:[self bufferForOffset:outputArea.start.line]];
+							if(outputArea.end.line) {
+								[self dispatchComputeCommandEncoder:computeEncoder pipelineState:_finalisedLineState width:_lineBufferPixelsPerLine height:outputArea.end.line offsetBuffer:[self bufferForOffset:0]];
+							}
+						}
+
+						[computeEncoder endEncoding];
+					} else {
+						// Separate luminance.
+						id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];
+						[computeEncoder setTexture:_compositionTexture atIndex:0];
+						[computeEncoder setTexture:_separatedLumaTexture atIndex:1];
+						[computeEncoder setBuffer:_uniformsBuffer offset:0 atIndex:0];
+
+						__unsafe_unretained id<MTLBuffer> offsetBuffers[2] = {nil, nil};
+						offsetBuffers[0] = [self bufferForOffset:outputArea.start.line];
+
+						if(outputArea.end.line > outputArea.start.line) {
+							[self dispatchComputeCommandEncoder:computeEncoder pipelineState:_separatedLumaState width:_lineBufferPixelsPerLine height:outputArea.end.line - outputArea.start.line offsetBuffer:offsetBuffers[0]];
+						} else {
+							[self dispatchComputeCommandEncoder:computeEncoder pipelineState:_separatedLumaState width:_lineBufferPixelsPerLine height:NumBufferedLines - outputArea.start.line offsetBuffer:offsetBuffers[0]];
+							if(outputArea.end.line) {
+								offsetBuffers[1] = [self bufferForOffset:0];
+								[self dispatchComputeCommandEncoder:computeEncoder pipelineState:_separatedLumaState width:_lineBufferPixelsPerLine height:outputArea.end.line offsetBuffer:offsetBuffers[1]];
+							}
+						}
+
+						// Filter resulting chrominance.
+						[computeEncoder setTexture:_separatedLumaTexture atIndex:0];
+						[computeEncoder setTexture:_finalisedLineTexture atIndex:1];
+						[computeEncoder setBuffer:_uniformsBuffer offset:0 atIndex:0];
+
+						if(outputArea.end.line > outputArea.start.line) {
+							[self dispatchComputeCommandEncoder:computeEncoder pipelineState:_finalisedLineState width:_lineBufferPixelsPerLine height:outputArea.end.line - outputArea.start.line offsetBuffer:offsetBuffers[0]];
+						} else {
+							[self dispatchComputeCommandEncoder:computeEncoder pipelineState:_finalisedLineState width:_lineBufferPixelsPerLine height:NumBufferedLines - outputArea.start.line offsetBuffer:offsetBuffers[0]];
+							if(outputArea.end.line) {
+								[self dispatchComputeCommandEncoder:computeEncoder pipelineState:_finalisedLineState width:_lineBufferPixelsPerLine height:outputArea.end.line offsetBuffer:offsetBuffers[1]];
+							}
+						}
+
+						[computeEncoder endEncoding];
+					}
+
+					// Output lines, broken up by frame.
+					size_t startLine = outputArea.start.line;
+					size_t line = outputArea.start.line;
+					while(line != outputArea.end.line) {
+						if(_lineMetadataBuffer[line].is_first_in_frame) {
+							[self outputFrom:startLine to:line commandBuffer:commandBuffer];
+							startLine = line;
+
+							if(_lineMetadataBuffer[line].previous_frame_was_complete && !_dontClearFrameBuffer) {
+								[self outputFrameCleanerToCommandBuffer:commandBuffer];
+							}
+							_dontClearFrameBuffer = NO;
+						}
+						line = (line + 1) % NumBufferedLines;
+					}
+					[self outputFrom:startLine to:outputArea.end.line commandBuffer:commandBuffer];
+				} break;
+			}
+
+			// Add a callback to update the scan target buffer and commit the drawing.
+			[commandBuffer addCompletedHandler:^(id<MTLCommandBuffer> _Nonnull) {
+				self->_scanTarget.complete_output_area(outputArea);
+			}];
+			[commandBuffer commit];
+		} else {
+			// There was no work, but to be contractually correct, remember to announce completion,
+			// and do it after finishing an empty command queue, as a cheap way to ensure this doen't
+			// front run any actual processing. TODO: can I do a better job of that?
+			id<MTLCommandBuffer> commandBuffer = [_commandQueue commandBuffer];
+			[commandBuffer addCompletedHandler:^(id<MTLCommandBuffer> _Nonnull) {
+				self->_scanTarget.complete_output_area(outputArea);
+			}];
+			[commandBuffer commit];
+
+			// TODO: reenable these and work out how on earth the Master System + Alex Kidd (US) is managing
+			// to provide write_area_y = 0, start_x = 0, end_x = 1.
+//			assert(outputArea.end.line == outputArea.start.line);
+//			assert(outputArea.end.scan == outputArea.start.scan);
+//			assert(outputArea.end.write_area_y == outputArea.start.write_area_y);
+//			assert(outputArea.end.write_area_x == outputArea.start.write_area_x);
+		}
+	}
+}
+
+/*!
+ @method drawInMTKView:
+ @abstract Called on the delegate when it is asked to render into the view
+ @discussion Called on the delegate when it is asked to render into the view
+ */
+- (void)drawInMTKView:(nonnull MTKView *)view {
+	if(_isDrawing.test_and_set()) {
+		_scanTarget.display_metrics_.announce_draw_status(false);
+		return;
+	}
+
+	// Disable supersampling if performance requires it.
+	if(_isUsingSupersampling && _scanTarget.display_metrics_.should_lower_resolution()) {
+		_isUsingSupersampling = false;
+		[self updateSizeBuffers];
+	}
+
+	// Schedule a copy from the current framebuffer to the view; blitting is unavailable as the target is a framebuffer texture.
+	id<MTLCommandBuffer> commandBuffer = [_commandQueue commandBuffer];
+
+	// Every pixel will be drawn, so don't clear or reload.
+	view.currentRenderPassDescriptor.colorAttachments[0].loadAction = MTLLoadActionDontCare;
+	id<MTLRenderCommandEncoder> encoder = [commandBuffer renderCommandEncoderWithDescriptor:view.currentRenderPassDescriptor];
+
+	[encoder setRenderPipelineState:_isUsingSupersampling ? _supersamplePipeline : _copyPipeline];
+	[encoder setVertexTexture:_frameBuffer atIndex:0];
+	[encoder setFragmentTexture:_frameBuffer atIndex:0];
+
+	[encoder drawPrimitives:MTLPrimitiveTypeTriangleStrip vertexStart:0 vertexCount:4];
+	[encoder endEncoding];
+
+	[commandBuffer presentDrawable:view.currentDrawable];
+	[commandBuffer addCompletedHandler:^(id<MTLCommandBuffer> _Nonnull) {
+		self->_isDrawing.clear();
+		self->_scanTarget.display_metrics_.announce_draw_status(true);
+	}];
+	[commandBuffer commit];
+}
+
+-  (Outputs::Display::ScanTarget *)scanTarget {
+	return &_scanTarget;
+}
+
+- (NSBitmapImageRep *)imageRepresentation {
+	// Create an NSBitmapRep as somewhere to copy pixel data to.
+	NSBitmapImageRep *const result =
+		[[NSBitmapImageRep alloc]
+			initWithBitmapDataPlanes:NULL
+			pixelsWide:(NSInteger)_frameBuffer.width
+			pixelsHigh:(NSInteger)_frameBuffer.height
+			bitsPerSample:8
+			samplesPerPixel:4
+			hasAlpha:YES
+			isPlanar:NO
+			colorSpaceName:NSDeviceRGBColorSpace
+			bytesPerRow:4 * (NSInteger)_frameBuffer.width
+			bitsPerPixel:0];
+
+	// Create a CPU-accessible texture and copy the current contents of the _frameBuffer to it.
+	// TODO: supersample rather than directly copy if appropriate?
+	id<MTLTexture> cpuTexture;
+	MTLTextureDescriptor *const textureDescriptor = [MTLTextureDescriptor
+		texture2DDescriptorWithPixelFormat:_view.colorPixelFormat
+		width:_frameBuffer.width
+		height:_frameBuffer.height
+		mipmapped:NO];
+	textureDescriptor.usage = MTLTextureUsageRenderTarget | MTLTextureUsageShaderRead;
+	textureDescriptor.resourceOptions = MTLResourceStorageModeManaged;
+	cpuTexture = [_view.device newTextureWithDescriptor:textureDescriptor];
+	[[self copyTexture:_frameBuffer to:cpuTexture] waitUntilCompleted];
+
+	// Copy from the CPU-visible texture to the bitmap image representation.
+	uint8_t *const bitmapData = result.bitmapData;
+	[cpuTexture
+		getBytes:bitmapData
+		bytesPerRow:_frameBuffer.width*4
+		fromRegion:MTLRegionMake2D(0, 0, _frameBuffer.width, _frameBuffer.height)
+		mipmapLevel:0];
+
+	// Set alpha to fully opaque and do some byte shuffling if necessary;
+	// Apple likes BGR for output but RGB is the best I can specify to NSBitmapImageRep.
+	//
+	// I'm not putting my foot down and having the GPU do the conversion I want
+	// because this lets me reuse _copyPipeline and thereby cut down on boilerplate,
+	// especially given that screenshots are not a bottleneck.
+	const NSUInteger totalBytes = _frameBuffer.width * _frameBuffer.height * 4;
+	const bool flipRedBlue = _view.colorPixelFormat == MTLPixelFormatBGRA8Unorm;
+	for(NSUInteger offset = 0; offset < totalBytes; offset += 4) {
+		if(flipRedBlue) {
+			const uint8_t red = bitmapData[offset];
+			bitmapData[offset] = bitmapData[offset+2];
+			bitmapData[offset+2] = red;
+		}
+		bitmapData[offset+3] = 0xff;
+	}
+
+	return result;
+}
+
+@end
diff --git a/OSBindings/Mac/Clock Signal/ScanTarget/ScanTarget.metal b/OSBindings/Mac/Clock Signal/ScanTarget/ScanTarget.metal
new file mode 100644
index 000000000..88b0b5e99
--- /dev/null
+++ b/OSBindings/Mac/Clock Signal/ScanTarget/ScanTarget.metal	
@@ -0,0 +1,588 @@
+//
+//  ScanTarget.metal
+//  Clock Signal
+//
+//  Created by Thomas Harte on 04/08/2020.
+//  Copyright © 2020 Thomas Harte. All rights reserved.
+//
+
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct Uniforms {
+	// This is used to scale scan positions, i.e. it provides the range
+	// for mapping from scan-style integer positions into eye space.
+	int2 scale;
+
+	// Applies a multiplication to all cyclesSinceRetrace values.
+	float cycleMultiplier;
+
+	// This provides the intended height of a scan, in eye-coordinate terms.
+	float lineWidth;
+
+	// Provides zoom and offset to scale the source data.
+	float3x3 sourceToDisplay;
+
+	// Provides conversions to and from RGB for the active colour space.
+	half3x3 toRGB;
+	half3x3 fromRGB;
+
+	// Describes the filter in use for chroma filtering; it'll be
+	// 15 coefficients but they're symmetrical around the centre.
+	half3 chromaKernel[8];
+
+	// Describes the filter in use for luma filtering; 15 coefficients
+	// symmetrical around the centre.
+	half lumaKernel[8];
+
+	// Sets the opacity at which output strips are drawn.
+	half outputAlpha;
+
+	// Sets the gamma power to which output colours are raised.
+	half outputGamma;
+
+	// Sets a brightness multiplier for output colours.
+	half outputMultiplier;
+};
+
+namespace {
+
+constexpr sampler standardSampler(	coord::pixel,
+									address::clamp_to_edge,	// Although arbitrary, stick with this address mode for compatibility all the way to MTLFeatureSet_iOS_GPUFamily1_v1.
+									filter::nearest);
+
+constexpr sampler linearSampler(	coord::pixel,
+									address::clamp_to_edge,	// Although arbitrary, stick with this address mode for compatibility all the way to MTLFeatureSet_iOS_GPUFamily1_v1.
+									filter::linear);
+
+}
+
+// MARK: - Structs used for receiving data from the emulation.
+
+// This is intended to match the net effect of `Scan` as defined by the BufferingScanTarget.
+struct Scan {
+	struct EndPoint {
+		uint16_t position[2];
+		uint16_t dataOffset;
+		int16_t compositeAngle;
+		uint16_t cyclesSinceRetrace;
+	} endPoints[2];
+
+	uint8_t compositeAmplitude;
+	uint16_t dataY;
+	uint16_t line;
+};
+
+// This matches the BufferingScanTarget's `Line`.
+struct Line {
+	struct EndPoint {
+		uint16_t position[2];
+		int16_t compositeAngle;
+		uint16_t cyclesSinceRetrace;
+	} endPoints[2];
+
+	uint8_t compositeAmplitude;
+	uint16_t line;
+};
+
+// MARK: - Intermediate structs.
+
+struct SourceInterpolator {
+	float4 position [[position]];
+	float2 textureCoordinates;
+	half unitColourPhase;		// i.e. one unit per circle.
+	half colourPhase;			// i.e. 2*pi units per circle, just regular radians.
+	half colourAmplitude [[flat]];
+};
+
+struct CopyInterpolator {
+	float4 position [[position]];
+	float2 textureCoordinates;
+};
+
+// MARK: - Vertex shaders.
+
+float2 textureLocation(constant Line *line, float offset, constant Uniforms &uniforms) {
+	return float2(
+		uniforms.cycleMultiplier * mix(line->endPoints[0].cyclesSinceRetrace, line->endPoints[1].cyclesSinceRetrace, offset),
+		line->line + 0.5f);
+}
+
+float2 textureLocation(constant Scan *scan, float offset, constant Uniforms &) {
+	return float2(
+		mix(scan->endPoints[0].dataOffset, scan->endPoints[1].dataOffset, offset),
+		scan->dataY + 0.5f);
+}
+
+template <typename Input> SourceInterpolator toDisplay(
+	constant Uniforms &uniforms [[buffer(1)]],
+	constant Input *inputs [[buffer(0)]],
+	uint instanceID [[instance_id]],
+	uint vertexID [[vertex_id]]) {
+	SourceInterpolator output;
+
+	// Get start and end vertices in regular float2 form.
+	const float2 start = float2(
+		float(inputs[instanceID].endPoints[0].position[0]) / float(uniforms.scale.x),
+		float(inputs[instanceID].endPoints[0].position[1]) / float(uniforms.scale.y)
+	);
+	const float2 end = float2(
+		float(inputs[instanceID].endPoints[1].position[0]) / float(uniforms.scale.x),
+		float(inputs[instanceID].endPoints[1].position[1]) / float(uniforms.scale.y)
+	);
+
+	// Calculate the tangent and normal.
+	const float2 tangent = (end - start);
+	const float2 normal = float2(tangent.y, -tangent.x) / length(tangent);
+
+	// Load up the colour details.
+	output.colourAmplitude = float(inputs[instanceID].compositeAmplitude) / 255.0f;
+	output.unitColourPhase = mix(
+		float(inputs[instanceID].endPoints[0].compositeAngle),
+		float(inputs[instanceID].endPoints[1].compositeAngle),
+		float((vertexID&2) >> 1)
+	) / 64.0f;
+	output.colourPhase = 2.0f * 3.141592654f * output.unitColourPhase;
+
+	// Hence determine this quad's real shape, using vertexID to pick a corner.
+
+	// position2d is now in the range [0, 1].
+	const float2 sourcePosition = start + (float(vertexID&2) * 0.5f) * tangent + (float(vertexID&1) - 0.5f) * normal * uniforms.lineWidth;
+	const float2 position2d = (uniforms.sourceToDisplay * float3(sourcePosition, 1.0f)).xy;
+
+	output.position = float4(
+		position2d,
+		0.0f,
+		1.0f
+	);
+	output.textureCoordinates = textureLocation(&inputs[instanceID], float((vertexID&2) >> 1), uniforms);
+
+	return output;
+}
+
+// These next two assume the incoming geometry to be a four-vertex triangle strip; each instance will therefore
+// produce a quad.
+
+vertex SourceInterpolator scanToDisplay(	constant Uniforms &uniforms [[buffer(1)]],
+											constant Scan *scans [[buffer(0)]],
+											uint instanceID [[instance_id]],
+											uint vertexID [[vertex_id]]) {
+	return toDisplay(uniforms, scans, instanceID, vertexID);
+}
+
+vertex SourceInterpolator lineToDisplay(	constant Uniforms &uniforms [[buffer(1)]],
+											constant Line *lines [[buffer(0)]],
+											uint instanceID [[instance_id]],
+											uint vertexID [[vertex_id]]) {
+	return toDisplay(uniforms, lines, instanceID, vertexID);
+}
+
+// This assumes that it needs to generate endpoints for a line segment.
+
+vertex SourceInterpolator scanToComposition(	constant Uniforms &uniforms [[buffer(1)]],
+												constant Scan *scans [[buffer(0)]],
+												uint instanceID [[instance_id]],
+												uint vertexID [[vertex_id]],
+												texture2d<float> texture [[texture(0)]]) {
+	SourceInterpolator result;
+
+	// Populate result as if direct texture access were available.
+	result.position.x = uniforms.cycleMultiplier * mix(scans[instanceID].endPoints[0].cyclesSinceRetrace, scans[instanceID].endPoints[1].cyclesSinceRetrace, float(vertexID));
+	result.position.y = scans[instanceID].line;
+	result.position.zw = float2(0.0f, 1.0f);
+
+	result.textureCoordinates.x = mix(scans[instanceID].endPoints[0].dataOffset, scans[instanceID].endPoints[1].dataOffset, float(vertexID));
+	result.textureCoordinates.y = scans[instanceID].dataY;
+
+	result.unitColourPhase = mix(
+		float(scans[instanceID].endPoints[0].compositeAngle),
+		float(scans[instanceID].endPoints[1].compositeAngle),
+		float(vertexID)
+	) / 64.0f;
+	result.colourPhase = 2.0f * 3.141592654f * result.unitColourPhase;
+	result.colourAmplitude = float(scans[instanceID].compositeAmplitude) / 255.0f;
+
+	// Map position into eye space, allowing for target texture dimensions.
+	const float2 textureSize = float2(texture.get_width(), texture.get_height());
+	result.position.xy =
+		((result.position.xy + float2(0.0f, 0.5f)) / textureSize)
+		* float2(2.0f, -2.0f) + float2(-1.0f, 1.0f);
+
+	return result;
+}
+
+vertex CopyInterpolator copyVertex(uint vertexID [[vertex_id]], texture2d<float> texture [[texture(0)]]) {
+	CopyInterpolator vert;
+
+	const uint x = vertexID & 1;
+	const uint y = (vertexID >> 1) & 1;
+
+	vert.textureCoordinates = float2(
+		x * texture.get_width(),
+		y * texture.get_height()
+	);
+	vert.position = float4(
+		float(x) * 2.0 - 1.0,
+		1.0 - float(y) * 2.0,
+		0.0,
+		1.0
+	);
+
+	return vert;
+}
+
+// MARK: - Various input format conversion samplers.
+
+half2 quadrature(float phase) {
+	return half2(cos(phase), sin(phase));
+}
+
+half4 composite(half level, half2 quadrature, half amplitude) {
+	return half4(
+		level,
+		half2(0.5f) + quadrature*half(0.5f),
+		amplitude
+	);
+}
+
+// The luminance formats can be sampled either in their natural format, or to the intermediate
+// composite format used for composition. Direct sampling is always for final output, so the two
+// 8-bit formats also provide a gamma option.
+
+half convertLuminance1(SourceInterpolator vert [[stage_in]], texture2d<ushort> texture [[texture(0)]]) {
+	return clamp(half(texture.sample(standardSampler, vert.textureCoordinates).r), half(0.0f), half(1.0f));
+}
+
+half convertLuminance8(SourceInterpolator vert [[stage_in]], texture2d<half> texture [[texture(0)]]) {
+	return texture.sample(standardSampler, vert.textureCoordinates).r;
+}
+
+half convertPhaseLinkedLuminance8(SourceInterpolator vert [[stage_in]], texture2d<half> texture [[texture(0)]]) {
+	const int offset = int(vert.unitColourPhase * 4.0f) & 3;
+	auto sample = texture.sample(standardSampler, vert.textureCoordinates);
+	return sample[offset];
+}
+
+
+#define CompositeSet(name, type)	\
+	fragment half4 sample##name(SourceInterpolator vert [[stage_in]], texture2d<type> texture [[texture(0)]], constant Uniforms &uniforms [[buffer(0)]]) {	\
+		const half luminance = convert##name(vert, texture) * uniforms.outputMultiplier;	\
+		return half4(half3(luminance), uniforms.outputAlpha);	\
+	}	\
+	\
+	fragment half4 sample##name##WithGamma(SourceInterpolator vert [[stage_in]], texture2d<type> texture [[texture(0)]], constant Uniforms &uniforms [[buffer(0)]]) {	\
+		const half luminance = pow(convert##name(vert, texture) * uniforms.outputMultiplier, uniforms.outputGamma);	\
+		return half4(half3(luminance), uniforms.outputAlpha);	\
+	}	\
+	\
+	fragment half4 compositeSample##name(SourceInterpolator vert [[stage_in]], texture2d<type> texture [[texture(0)]], constant Uniforms &uniforms [[buffer(0)]]) {	\
+		const half luminance = convert##name(vert, texture) * uniforms.outputMultiplier;	\
+		return composite(luminance, quadrature(vert.colourPhase), vert.colourAmplitude);	\
+	}
+
+CompositeSet(Luminance1, ushort);
+CompositeSet(Luminance8, half);
+CompositeSet(PhaseLinkedLuminance8, half);
+
+#undef CompositeSet
+
+// The luminance/phase format can produce either composite or S-Video.
+
+/// @returns A 2d vector comprised where .x = luminance; .y = chroma.
+half2 convertLuminance8Phase8(SourceInterpolator vert [[stage_in]], texture2d<half> texture [[texture(0)]]) {
+	const auto luminancePhase = texture.sample(standardSampler, vert.textureCoordinates).rg;
+	const half phaseOffset = 3.141592654 * 4.0 * luminancePhase.g;
+	const half rawChroma = step(luminancePhase.g, half(0.75f)) * cos(vert.colourPhase + phaseOffset);
+	return half2(luminancePhase.r, rawChroma);
+}
+
+fragment half4 compositeSampleLuminance8Phase8(SourceInterpolator vert [[stage_in]], texture2d<half> texture [[texture(0)]]) {
+	const half2 luminanceChroma = convertLuminance8Phase8(vert, texture);
+	const half luminance = mix(luminanceChroma.r, luminanceChroma.g, vert.colourAmplitude);
+	return composite(luminance, quadrature(vert.colourPhase), vert.colourAmplitude);
+}
+
+fragment half4 sampleLuminance8Phase8(SourceInterpolator vert [[stage_in]], texture2d<half> texture [[texture(0)]]) {
+	const half2 luminanceChroma = convertLuminance8Phase8(vert, texture);
+	const half2 qam = quadrature(vert.colourPhase) * half(0.5f);
+	return half4(luminanceChroma.r,
+			half2(0.5f) + luminanceChroma.g*qam,
+			half(1.0f));
+}
+
+fragment half4 directCompositeSampleLuminance8Phase8(SourceInterpolator vert [[stage_in]], texture2d<half> texture [[texture(0)]], constant Uniforms &uniforms [[buffer(0)]]) {
+	const half2 luminanceChroma = convertLuminance8Phase8(vert, texture);
+	const half luminance = mix(luminanceChroma.r * uniforms.outputMultiplier, luminanceChroma.g, vert.colourAmplitude);
+	return half4(half3(luminance), uniforms.outputAlpha);
+}
+
+fragment half4 directCompositeSampleLuminance8Phase8WithGamma(SourceInterpolator vert [[stage_in]], texture2d<half> texture [[texture(0)]], constant Uniforms &uniforms [[buffer(0)]]) {
+	const half2 luminanceChroma = convertLuminance8Phase8(vert, texture);
+	const half luminance = mix(pow(luminanceChroma.r * uniforms.outputMultiplier, uniforms.outputGamma), luminanceChroma.g, vert.colourAmplitude);
+	return half4(half3(luminance), uniforms.outputAlpha);
+}
+
+
+// All the RGB formats can produce RGB, composite or S-Video.
+
+half3 convertRed8Green8Blue8(SourceInterpolator vert, texture2d<half> texture) {
+	return texture.sample(standardSampler, vert.textureCoordinates).rgb;
+}
+
+half3 convertRed4Green4Blue4(SourceInterpolator vert, texture2d<ushort> texture) {
+	const auto sample = texture.sample(standardSampler, vert.textureCoordinates).rg;
+	return clamp(half3(sample.r&15, (sample.g >> 4)&15, sample.g&15), half(0.0f), half(1.0f));
+}
+
+half3 convertRed2Green2Blue2(SourceInterpolator vert, texture2d<ushort> texture) {
+	const auto sample = texture.sample(standardSampler, vert.textureCoordinates).r;
+	return clamp(half3((sample >> 4)&3, (sample >> 2)&3, sample&3), half(0.0f), half(1.0f));
+}
+
+half3 convertRed1Green1Blue1(SourceInterpolator vert, texture2d<ushort> texture) {
+	const auto sample = texture.sample(standardSampler, vert.textureCoordinates).r;
+	return clamp(half3(sample&4, sample&2, sample&1), half(0.0f), half(1.0f));
+}
+
+#define DeclareShaders(name, pixelType)	\
+	fragment half4 sample##name(SourceInterpolator vert [[stage_in]], texture2d<pixelType> texture [[texture(0)]], constant Uniforms &uniforms [[buffer(0)]]) {	\
+		return half4(convert##name(vert, texture), uniforms.outputAlpha);	\
+	}	\
+	\
+	fragment half4 sample##name##WithGamma(SourceInterpolator vert [[stage_in]], texture2d<pixelType> texture [[texture(0)]], constant Uniforms &uniforms [[buffer(0)]]) {	\
+		return half4(pow(convert##name(vert, texture), uniforms.outputGamma), uniforms.outputAlpha);	\
+	}	\
+	\
+	fragment half4 svideoSample##name(SourceInterpolator vert [[stage_in]], texture2d<pixelType> texture [[texture(0)]], constant Uniforms &uniforms [[buffer(0)]]) {	\
+		const auto colour = uniforms.fromRGB * convert##name(vert, texture);	\
+		const half2 qam = quadrature(vert.colourPhase);	\
+		const half chroma = dot(colour.gb, qam);	\
+		return half4(	\
+			colour.r,	\
+			half2(0.5f) + chroma*qam*half(0.5f),	\
+			half(1.0f)		\
+		);	\
+	}	\
+	\
+	half composite##name(SourceInterpolator vert, texture2d<pixelType> texture, constant Uniforms &uniforms, half2 colourSubcarrier) {	\
+		const auto colour = uniforms.fromRGB * convert##name(vert, texture);	\
+		return mix(colour.r, dot(colour.gb, colourSubcarrier), half(vert.colourAmplitude));	\
+	}	\
+	\
+	fragment half4 compositeSample##name(SourceInterpolator vert [[stage_in]], texture2d<pixelType> texture [[texture(0)]], constant Uniforms &uniforms [[buffer(0)]]) {	\
+		const half2 colourSubcarrier = quadrature(vert.colourPhase);	\
+		return composite(composite##name(vert, texture, uniforms, colourSubcarrier), colourSubcarrier, vert.colourAmplitude);	\
+	}	\
+	\
+	fragment half4 directCompositeSample##name(SourceInterpolator vert [[stage_in]], texture2d<pixelType> texture [[texture(0)]], constant Uniforms &uniforms [[buffer(0)]]) {	\
+		const half level = composite##name(vert, texture, uniforms, quadrature(vert.colourPhase)); 	\
+		return half4(half3(level), uniforms.outputAlpha);	\
+	}	\
+	\
+	fragment half4 directCompositeSample##name##WithGamma(SourceInterpolator vert [[stage_in]], texture2d<pixelType> texture [[texture(0)]], constant Uniforms &uniforms [[buffer(0)]]) {	\
+		const half level = pow(composite##name(vert, texture, uniforms, quadrature(vert.colourPhase)), uniforms.outputGamma); 	\
+		return half4(half3(level), uniforms.outputAlpha);	\
+	}
+
+DeclareShaders(Red8Green8Blue8, half)
+DeclareShaders(Red4Green4Blue4, ushort)
+DeclareShaders(Red2Green2Blue2, ushort)
+DeclareShaders(Red1Green1Blue1, ushort)
+
+fragment half4 copyFragment(CopyInterpolator vert [[stage_in]], texture2d<half> texture [[texture(0)]]) {
+	return texture.sample(standardSampler, vert.textureCoordinates);
+}
+
+fragment half4 interpolateFragment(CopyInterpolator vert [[stage_in]], texture2d<half> texture [[texture(0)]]) {
+	return texture.sample(linearSampler, vert.textureCoordinates);
+}
+
+fragment half4 clearFragment(constant Uniforms &uniforms [[buffer(0)]]) {
+	return half4(0.0, 0.0, 0.0, uniforms.outputAlpha);
+}
+
+// MARK: - Compute kernels
+
+/// Given input pixels of the form (luminance, 0.5 + 0.5*chrominance*cos(phase), 0.5 + 0.5*chrominance*sin(phase)), applies a lowpass
+/// filter to the two chrominance parts, then uses the toRGB matrix to convert to RGB and stores.
+template <bool applyGamma> void filterChromaKernel(	texture2d<half, access::read> inTexture [[texture(0)]],
+													texture2d<half, access::write> outTexture [[texture(1)]],
+													uint2 gid [[thread_position_in_grid]],
+													constant Uniforms &uniforms [[buffer(0)]],
+													constant int &offset [[buffer(1)]]) {
+	constexpr half4 moveToZero(0.0f, 0.5f, 0.5f, 0.0f);
+	const half4 rawSamples[] = {
+		inTexture.read(gid + uint2(0, offset)) - moveToZero,
+		inTexture.read(gid + uint2(1, offset)) - moveToZero,
+		inTexture.read(gid + uint2(2, offset)) - moveToZero,
+		inTexture.read(gid + uint2(3, offset)) - moveToZero,
+		inTexture.read(gid + uint2(4, offset)) - moveToZero,
+		inTexture.read(gid + uint2(5, offset)) - moveToZero,
+		inTexture.read(gid + uint2(6, offset)) - moveToZero,
+		inTexture.read(gid + uint2(7, offset)) - moveToZero,
+		inTexture.read(gid + uint2(8, offset)) - moveToZero,
+		inTexture.read(gid + uint2(9, offset)) - moveToZero,
+		inTexture.read(gid + uint2(10, offset)) - moveToZero,
+		inTexture.read(gid + uint2(11, offset)) - moveToZero,
+		inTexture.read(gid + uint2(12, offset)) - moveToZero,
+		inTexture.read(gid + uint2(13, offset)) - moveToZero,
+		inTexture.read(gid + uint2(14, offset)) - moveToZero,
+	};
+
+#define Sample(x, y) uniforms.chromaKernel[y] * rawSamples[x].rgb
+	const half3 colour =
+		Sample(0, 0) + Sample(1, 1) + Sample(2, 2) + Sample(3, 3) + Sample(4, 4) + Sample(5, 5) + Sample(6, 6) +
+		Sample(7, 7) +
+		Sample(8, 6) + Sample(9, 5) + Sample(10, 4) + Sample(11, 3) + Sample(12, 2) + Sample(13, 1) + Sample(14, 0);
+#undef Sample
+
+	const half4 output = half4(uniforms.toRGB * colour * uniforms.outputMultiplier, uniforms.outputAlpha);
+	if(applyGamma) {
+		outTexture.write(pow(output, uniforms.outputGamma), gid + uint2(7, offset));
+	} else {
+		outTexture.write(output, gid + uint2(7, offset));
+	}
+}
+
+kernel void filterChromaKernelNoGamma(	texture2d<half, access::read> inTexture [[texture(0)]],
+										texture2d<half, access::write> outTexture [[texture(1)]],
+										uint2 gid [[thread_position_in_grid]],
+										constant Uniforms &uniforms [[buffer(0)]],
+										constant int &offset [[buffer(1)]]) {
+	filterChromaKernel<false>(inTexture, outTexture, gid, uniforms, offset);
+}
+
+kernel void filterChromaKernelWithGamma(	texture2d<half, access::read> inTexture [[texture(0)]],
+											texture2d<half, access::write> outTexture [[texture(1)]],
+											uint2 gid [[thread_position_in_grid]],
+											constant Uniforms &uniforms [[buffer(0)]],
+											constant int &offset [[buffer(1)]]) {
+	filterChromaKernel<true>(inTexture, outTexture, gid, uniforms, offset);
+}
+
+void setSeparatedLumaChroma(half luminance, half4 centreSample, texture2d<half, access::write> outTexture, uint2 gid, int offset) {
+	// The mix/steps below ensures that the absence of a colour burst leads the colour subcarrier to be discarded.
+	const half isColour = step(half(0.01f), centreSample.a);
+	const half chroma = (centreSample.r - luminance) / mix(half(1.0f), centreSample.a, isColour);
+	outTexture.write(half4(
+			luminance / mix(half(1.0f), (half(1.0f) - centreSample.a), isColour),
+			isColour * (centreSample.gb - half2(0.5f)) * chroma + half2(0.5f),
+			1.0f
+		),
+		gid + uint2(7, offset));
+}
+
+
+/// Given input pixels of the form:
+///
+///	(composite sample, cos(phase), sin(phase), colour amplitude), applies a lowpass
+///
+/// Filters to separate luminance, subtracts that and scales and maps the remaining chrominance in order to output
+/// pixels in the form:
+///
+///	(luminance, 0.5 + 0.5*chrominance*cos(phase), 0.5 + 0.5*chrominance*sin(phase))
+///
+/// i.e. the input form for the filterChromaKernel, above].
+kernel void separateLumaKernel15(	texture2d<half, access::read> inTexture [[texture(0)]],
+									texture2d<half, access::write> outTexture [[texture(1)]],
+									uint2 gid [[thread_position_in_grid]],
+									constant Uniforms &uniforms [[buffer(0)]],
+									constant int &offset [[buffer(1)]]) {
+	const half4 centreSample = inTexture.read(gid + uint2(7, offset));
+	const half rawSamples[] = {
+		inTexture.read(gid + uint2(0, offset)).r,	inTexture.read(gid + uint2(1, offset)).r,
+		inTexture.read(gid + uint2(2, offset)).r,	inTexture.read(gid + uint2(3, offset)).r,
+		inTexture.read(gid + uint2(4, offset)).r,	inTexture.read(gid + uint2(5, offset)).r,
+		inTexture.read(gid + uint2(6, offset)).r,
+		centreSample.r,
+		inTexture.read(gid + uint2(8, offset)).r,
+		inTexture.read(gid + uint2(9, offset)).r,	inTexture.read(gid + uint2(10, offset)).r,
+		inTexture.read(gid + uint2(11, offset)).r,	inTexture.read(gid + uint2(12, offset)).r,
+		inTexture.read(gid + uint2(13, offset)).r,	inTexture.read(gid + uint2(14, offset)).r,
+	};
+
+#define Sample(x, y) uniforms.lumaKernel[y] * rawSamples[x]
+	const half luminance =
+		Sample(0, 0) + Sample(1, 1) + Sample(2, 2) + Sample(3, 3) + Sample(4, 4) + Sample(5, 5) + Sample(6, 6) +
+		Sample(7, 7) +
+		Sample(8, 6) + Sample(9, 5) + Sample(10, 4) + Sample(11, 3) + Sample(12, 2) + Sample(13, 1) + Sample(14, 0);
+#undef Sample
+
+	return setSeparatedLumaChroma(luminance, centreSample, outTexture, gid, offset);
+}
+
+kernel void separateLumaKernel9(	texture2d<half, access::read> inTexture [[texture(0)]],
+									texture2d<half, access::write> outTexture [[texture(1)]],
+									uint2 gid [[thread_position_in_grid]],
+									constant Uniforms &uniforms [[buffer(0)]],
+									constant int &offset [[buffer(1)]]) {
+	const half4 centreSample = inTexture.read(gid + uint2(7, offset));
+	const half rawSamples[] = {
+		inTexture.read(gid + uint2(3, offset)).r,	inTexture.read(gid + uint2(4, offset)).r,
+		inTexture.read(gid + uint2(5, offset)).r,	inTexture.read(gid + uint2(6, offset)).r,
+		centreSample.r,
+		inTexture.read(gid + uint2(8, offset)).r,	inTexture.read(gid + uint2(9, offset)).r,
+		inTexture.read(gid + uint2(10, offset)).r,	inTexture.read(gid + uint2(11, offset)).r
+	};
+
+#define Sample(x, y) uniforms.lumaKernel[y] * rawSamples[x]
+	const half luminance =
+		Sample(0, 3) + Sample(1, 4) + Sample(2, 5) + Sample(3, 6) +
+		Sample(4, 7) +
+		Sample(5, 6) + Sample(6, 5) + Sample(7, 4) + Sample(8, 3);
+#undef Sample
+
+	return setSeparatedLumaChroma(luminance, centreSample, outTexture, gid, offset);
+}
+
+kernel void separateLumaKernel7(	texture2d<half, access::read> inTexture [[texture(0)]],
+									texture2d<half, access::write> outTexture [[texture(1)]],
+									uint2 gid [[thread_position_in_grid]],
+									constant Uniforms &uniforms [[buffer(0)]],
+									constant int &offset [[buffer(1)]]) {
+	const half4 centreSample = inTexture.read(gid + uint2(7, offset));
+	const half rawSamples[] = {
+		inTexture.read(gid + uint2(4, offset)).r,
+		inTexture.read(gid + uint2(5, offset)).r,	inTexture.read(gid + uint2(6, offset)).r,
+		centreSample.r,
+		inTexture.read(gid + uint2(8, offset)).r,	inTexture.read(gid + uint2(9, offset)).r,
+		inTexture.read(gid + uint2(10, offset)).r
+	};
+
+#define Sample(x, y) uniforms.lumaKernel[y] * rawSamples[x]
+	const half luminance =
+		Sample(0, 4) + Sample(1, 5) + Sample(2, 6) +
+		Sample(3, 7) +
+		Sample(4, 6) + Sample(5, 5) + Sample(6, 4);
+#undef Sample
+
+	return setSeparatedLumaChroma(luminance, centreSample, outTexture, gid, offset);
+}
+
+kernel void separateLumaKernel5(	texture2d<half, access::read> inTexture [[texture(0)]],
+									texture2d<half, access::write> outTexture [[texture(1)]],
+									uint2 gid [[thread_position_in_grid]],
+									constant Uniforms &uniforms [[buffer(0)]],
+									constant int &offset [[buffer(1)]]) {
+	const half4 centreSample = inTexture.read(gid + uint2(7, offset));
+	const half rawSamples[] = {
+		inTexture.read(gid + uint2(5, offset)).r,	inTexture.read(gid + uint2(6, offset)).r,
+		centreSample.r,
+		inTexture.read(gid + uint2(8, offset)).r,	inTexture.read(gid + uint2(9, offset)).r,
+	};
+
+#define Sample(x, y) uniforms.lumaKernel[y] * rawSamples[x]
+	const half luminance =
+		Sample(0, 5) + Sample(1, 6) +
+		Sample(2, 7) +
+		Sample(3, 6) + Sample(4, 5);
+#undef Sample
+
+	return setSeparatedLumaChroma(luminance, centreSample, outTexture, gid, offset);
+}
+
+kernel void clearKernel(	texture2d<half, access::write> outTexture [[texture(1)]],
+							uint2 gid [[thread_position_in_grid]]) {
+	outTexture.write(half4(0.0f, 0.0f, 0.0f, 1.0f), gid);
+}
diff --git a/OSBindings/Mac/Clock Signal/Views/CSOpenGLView.h b/OSBindings/Mac/Clock Signal/Views/CSScanTargetView.h
similarity index 60%
rename from OSBindings/Mac/Clock Signal/Views/CSOpenGLView.h
rename to OSBindings/Mac/Clock Signal/Views/CSScanTargetView.h
index f4090fd7a..58917f785 100644
--- a/OSBindings/Mac/Clock Signal/Views/CSOpenGLView.h	
+++ b/OSBindings/Mac/Clock Signal/Views/CSScanTargetView.h	
@@ -1,5 +1,5 @@
 //
-//  CSOpenGLView.h
+//  CSScanTargetView.h
 //  Clock Signal
 //
 //  Created by Thomas Harte on 16/07/2015.
@@ -8,63 +8,12 @@
 
 #import <Foundation/Foundation.h>
 #import <AppKit/AppKit.h>
+#import <MetalKit/MetalKit.h>
 
-@class CSOpenGLView;
+@class CSScanTargetView;
+@class CSScanTarget;
 
-typedef NS_ENUM(NSInteger, CSOpenGLViewRedrawEvent) {
-	/// Indicates that AppKit requested a redraw for some reason (mostly likely, the window is being resized). So,
-	/// if the delegate doesn't redraw the view, the user is likely to see a graphical flaw.
-	CSOpenGLViewRedrawEventAppKit,
-	/// Indicates that the view's display-linked timer has triggered a redraw request. So, if the delegate doesn't
-	/// redraw the view, the user will just see the previous drawing without interruption.
-	CSOpenGLViewRedrawEventTimer
-};
-
-@protocol CSOpenGLViewDelegate
-/*!
-	Requests that the delegate produce an image of its current output state. May be called on
-	any queue or thread.
-	@param view The view making the request.
-	@param redrawEvent If @c YES then the delegate may decline to redraw if its output would be
-	identical to the previous frame. If @c NO then the delegate must draw.
-*/
-- (void)openGLViewRedraw:(nonnull CSOpenGLView *)view event:(CSOpenGLViewRedrawEvent)redrawEvent;
-
-/*!
-	Announces receipt of a file by drag and drop to the delegate.
-	@param view The view making the request.
-	@param URL The file URL of the received file.
-*/
-- (void)openGLView:(nonnull CSOpenGLView *)view didReceiveFileAtURL:(nonnull NSURL *)URL;
-
-/*!
-	Announces 'capture' of the mouse — i.e. that the view is now preventing the mouse from exiting
-	the window, in order to forward continuous mouse motion.
-	@param view The view making the announcement.
-*/
-- (void)openGLViewDidCaptureMouse:(nonnull CSOpenGLView *)view;
-
-/*!
-	Announces that the mouse is no longer captured.
-	@param view The view making the announcement.
-*/
-- (void)openGLViewDidReleaseMouse:(nonnull CSOpenGLView *)view;
-
-/*!
-	Announces that the OS mouse cursor is now being displayed again, after having been invisible.
-	@param view The view making the announcement.
-*/
-- (void)openGLViewDidShowOSMouseCursor:(nonnull CSOpenGLView *)view;
-
-/*!
-	Announces that the OS mouse cursor will now be hidden.
-	@param view The view making the announcement.
-*/
-- (void)openGLViewWillHideOSMouseCursor:(nonnull CSOpenGLView *)view;
-
-@end
-
-@protocol CSOpenGLViewResponderDelegate <NSObject>
+@protocol CSScanTargetViewResponderDelegate <NSObject>
 /*!
 	Supplies a keyDown event to the delegate.
 	@param event The @c NSEvent describing the keyDown.
@@ -111,41 +60,72 @@ typedef NS_ENUM(NSInteger, CSOpenGLViewRedrawEvent) {
 */
 - (void)mouseUp:(nonnull NSEvent *)event;
 
+/*!
+	Announces 'capture' of the mouse — i.e. that the view is now preventing the mouse from exiting
+	the window, in order to forward continuous mouse motion.
+	@param view The view making the announcement.
+*/
+- (void)scanTargetViewDidCaptureMouse:(nonnull CSScanTargetView *)view;
+
+/*!
+	Announces that the mouse is no longer captured.
+	@param view The view making the announcement.
+*/
+- (void)scanTargetViewDidReleaseMouse:(nonnull CSScanTargetView *)view;
+
+/*!
+	Announces that the OS mouse cursor is now being displayed again, after having been invisible.
+	@param view The view making the announcement.
+*/
+- (void)scanTargetViewDidShowOSMouseCursor:(nonnull CSScanTargetView *)view;
+
+/*!
+	Announces that the OS mouse cursor will now be hidden.
+	@param view The view making the announcement.
+*/
+- (void)scanTargetViewWillHideOSMouseCursor:(nonnull CSScanTargetView *)view;
+
+/*!
+	Announces receipt of a file by drag and drop to the delegate.
+	@param view The view making the request.
+	@param URL The file URL of the received file.
+*/
+- (void)scanTargetView:(nonnull CSScanTargetView *)view didReceiveFileAtURL:(nonnull NSURL *)URL;
+
 @end
 
 /*!
-	Although I'm still on the fence about this as a design decision, CSOpenGLView  is itself responsible
+	Although I'm still on the fence about this as a design decision, CSScanTargetView  is itself responsible
 	for creating and destroying a CVDisplayLink. There's a practical reason for this: you'll get real synchronisation
-	only if a link is explicitly tied to a particular display, and the CSOpenGLView therefore owns the knowledge
+	only if a link is explicitly tied to a particular display, and the CSScanTargetView therefore owns the knowledge
 	necessary to decide when to create and modify them. It doesn't currently just propagate "did change screen"-type
 	messages because I haven't yet found a way to track that other than polling, in which case I might as well put
 	that into the display link callback.
 */
-@protocol CSOpenGLViewDisplayLinkDelegate
+@protocol CSScanTargetViewDisplayLinkDelegate
 
 /*!
 	Informs the delegate that the display link has fired.
 */
-- (void)openGLViewDisplayLinkDidFire:(nonnull CSOpenGLView *)view now:(nonnull const CVTimeStamp *)now outputTime:(nonnull const CVTimeStamp *)outputTime;
+- (void)scanTargetViewDisplayLinkDidFire:(nonnull CSScanTargetView *)view now:(nonnull const CVTimeStamp *)now outputTime:(nonnull const CVTimeStamp *)outputTime;
 
 @end
 
 /*!
-	Provides an OpenGL canvas with a refresh-linked update timer that can forward a subset
+	Provides a visible scan target with a refresh-linked update timer that can forward a subset
 	of typical first-responder actions.
 */
-@interface CSOpenGLView : NSOpenGLView
+@interface CSScanTargetView : MTKView
 
-@property (atomic, weak, nullable) id <CSOpenGLViewDelegate> delegate;
-@property (nonatomic, weak, nullable) id <CSOpenGLViewResponderDelegate> responderDelegate;
-@property (atomic, weak, nullable) id <CSOpenGLViewDisplayLinkDelegate> displayLinkDelegate;
+@property (nonatomic, weak, nullable) id <CSScanTargetViewResponderDelegate> responderDelegate;
+@property (atomic, weak, nullable) id <CSScanTargetViewDisplayLinkDelegate> displayLinkDelegate;
 
 /// Determines whether the view offers mouse capturing — i.e. if the user clicks on the view then
-/// then the system cursor is disabled and the mouse events defined by CSOpenGLViewResponderDelegate
+/// then the system cursor is disabled and the mouse events defined by CSScanTargetViewResponderDelegate
 /// are forwarded, unless and until the user releases the mouse using the control+command shortcut.
 @property (nonatomic, assign) BOOL shouldCaptureMouse;
 
-/// Determines whether the CSOpenGLViewResponderDelegate of this window expects to use the command
+/// Determines whether the CSScanTargetViewResponderDelegate of this window expects to use the command
 /// key as though it were any other key — i.e. all command combinations should be forwarded to the delegate,
 /// not being allowed to trigger regular application shortcuts such as command+q or command+h.
 ///
@@ -162,19 +142,24 @@ typedef NS_ENUM(NSInteger, CSOpenGLViewRedrawEvent) {
 */
 - (void)invalidate;
 
-/// The size in pixels of the OpenGL canvas, factoring in screen pixel density and view size in points.
-@property (nonatomic, readonly) CGSize backingSize;
-
 /*!
-	Locks this view's OpenGL context and makes it current, performs @c action and then unlocks
-	the context. @c action is performed on the calling queue.
+	Ensures output begins on all pending scans.
 */
-- (void)performWithGLContext:(nonnull dispatch_block_t)action flushDrawable:(BOOL)flushDrawable;
-- (void)performWithGLContext:(nonnull dispatch_block_t)action;
+- (void)updateBacking;
 
 /*!
 	Instructs that the mouse cursor, if currently captured, should be released.
 */
 - (void)releaseMouse;
 
+/*!
+	@returns An image of the view's current contents.
+*/
+- (nonnull NSBitmapImageRep *)imageRepresentation;
+
+/*!
+	@returns The CSScanTarget being used for this display.
+*/
+@property(nonatomic, readonly, nonnull) CSScanTarget *scanTarget;
+
 @end
diff --git a/OSBindings/Mac/Clock Signal/Views/CSOpenGLView.m b/OSBindings/Mac/Clock Signal/Views/CSScanTargetView.m
similarity index 71%
rename from OSBindings/Mac/Clock Signal/Views/CSOpenGLView.m
rename to OSBindings/Mac/Clock Signal/Views/CSScanTargetView.m
index 2b6759170..65ed41fb1 100644
--- a/OSBindings/Mac/Clock Signal/Views/CSOpenGLView.m	
+++ b/OSBindings/Mac/Clock Signal/Views/CSScanTargetView.m	
@@ -1,24 +1,24 @@
 //
-//  CSOpenGLView
+//  CSScanTargetView
 //  CLK
 //
 //  Created by Thomas Harte on 16/07/2015.
 //  Copyright 2015 Thomas Harte. All rights reserved.
 //
 
-#import "CSOpenGLView.h"
+#import "CSScanTargetView.h"
 #import "CSApplication.h"
+#import "CSScanTarget.h"
 @import CoreVideo;
 @import GLKit;
 
 #include <stdatomic.h>
 
-@interface CSOpenGLView () <NSDraggingDestination, CSApplicationEventDelegate>
+@interface CSScanTargetView () <NSDraggingDestination, CSApplicationEventDelegate>
 @end
 
-@implementation CSOpenGLView {
+@implementation CSScanTargetView {
 	CVDisplayLinkRef _displayLink;
-	CGSize _backingSize;
 	NSNumber *_currentScreenNumber;
 
 	NSTrackingArea *_mouseTrackingArea;
@@ -27,20 +27,8 @@
 
 	atomic_int _isDrawingFlag;
 	BOOL _isInvalid;
-}
 
-- (void)prepareOpenGL {
-	[super prepareOpenGL];
-
-	// Prepare the atomic int.
-	atomic_init(&_isDrawingFlag, 0);
-
-	// Set the clear colour.
-	[self.openGLContext makeCurrentContext];
-	glClearColor(0.0, 0.0, 0.0, 1.0);
-
-	// Setup the [initial] display link.
-	[self setupDisplayLink];
+	CSScanTarget *_scanTarget;
 }
 
 - (void)setupDisplayLink {
@@ -58,17 +46,12 @@
 	// Set the renderer output callback function.
 	CVDisplayLinkSetOutputCallback(_displayLink, DisplayLinkCallback, (__bridge void * __nullable)(self));
 
-	// Set the display link for the current renderer.
-	CGLContextObj cglContext = [[self openGLContext] CGLContextObj];
-	CGLPixelFormatObj cglPixelFormat = [[self pixelFormat] CGLPixelFormatObj];
-	CVDisplayLinkSetCurrentCGDisplayFromOpenGLContext(_displayLink, cglContext, cglPixelFormat);
-
 	// Activate the display link.
 	CVDisplayLinkStart(_displayLink);
 }
 
 static CVReturn DisplayLinkCallback(__unused CVDisplayLinkRef displayLink, const CVTimeStamp *now, const CVTimeStamp *outputTime, __unused CVOptionFlags flagsIn, __unused CVOptionFlags *flagsOut, void *displayLinkContext) {
-	CSOpenGLView *const view = (__bridge CSOpenGLView *)displayLinkContext;
+	CSScanTargetView *const view = (__bridge CSScanTargetView *)displayLinkContext;
 
 	// Schedule an opportunity to check that the display link is still linked to the correct display.
 	dispatch_async(dispatch_get_main_queue(), ^{
@@ -78,7 +61,7 @@ static CVReturn DisplayLinkCallback(__unused CVDisplayLinkRef displayLink, const
 	// Ensure _isDrawingFlag has value 1 when drawing, 0 otherwise.
 	atomic_store(&view->_isDrawingFlag, 1);
 
-	[view.displayLinkDelegate openGLViewDisplayLinkDidFire:view now:now outputTime:outputTime];
+	[view.displayLinkDelegate scanTargetViewDisplayLinkDidFire:view now:now outputTime:outputTime];
 	/*
 		Do not touch the display link from after this call; there's a bit of a race condition with setupDisplayLink.
 		Specifically: Apple provides CVDisplayLinkStop but a call to that merely prevents future calls to the callback,
@@ -106,30 +89,12 @@ static CVReturn DisplayLinkCallback(__unused CVDisplayLinkRef displayLink, const
 	// feels fine.
 	NSNumber *const screenNumber = self.window.screen.deviceDescription[@"NSScreenNumber"];
 	if(![_currentScreenNumber isEqual:screenNumber]) {
-		// Issue a reshape, in case a switch to/from a Retina display has
-		// happened, changing the results of -convertSizeToBacking:, etc.
-		[self reshape];
-
 		// Also switch display links, to make sure synchronisation is with the display
 		// the window is actually on, and at its rate.
 		[self setupDisplayLink];
 	}
 }
 
-- (void)drawAtTime:(const CVTimeStamp *)now frequency:(double)frequency {
-	[self redrawWithEvent:CSOpenGLViewRedrawEventTimer];
-}
-
-- (void)drawRect:(NSRect)dirtyRect {
-	[self redrawWithEvent:CSOpenGLViewRedrawEventAppKit];
-}
-
-- (void)redrawWithEvent:(CSOpenGLViewRedrawEvent)event  {
-	[self performWithGLContext:^{
-		[self.delegate openGLViewRedraw:self event:event];
-	} flushDrawable:YES];
-}
-
 - (void)invalidate {
 	_isInvalid = YES;
 	[self stopDisplayLink];
@@ -160,65 +125,35 @@ static CVReturn DisplayLinkCallback(__unused CVDisplayLinkRef displayLink, const
 	CVDisplayLinkRelease(_displayLink);
 }
 
-- (CGSize)backingSize {
-	@synchronized(self) {
-		return _backingSize;
-	}
+- (CSScanTarget *)scanTarget {
+	return _scanTarget;
 }
 
-- (void)reshape {
-	[super reshape];
-	@synchronized(self) {
-		_backingSize = [self convertSizeToBacking:self.bounds.size];
-	}
-
-	[self performWithGLContext:^{
-		CGSize viewSize = [self backingSize];
-		glViewport(0, 0, (GLsizei)viewSize.width, (GLsizei)viewSize.height);
-	} flushDrawable:NO];
+- (void)updateBacking {
+	[_scanTarget updateFrameBuffer];
 }
 
 - (void)awakeFromNib {
-	NSOpenGLPixelFormatAttribute attributes[] = {
-		NSOpenGLPFADoubleBuffer,
-		NSOpenGLPFAOpenGLProfile,	NSOpenGLProfileVersion3_2Core,
-//		NSOpenGLPFAMultisample,
-//		NSOpenGLPFASampleBuffers,	1,
-//		NSOpenGLPFASamples,			2,
-		0
-	};
+	// Use the preferred device if available.
+	if(@available(macOS 10.15, *)) {
+		self.device = self.preferredDevice;
+	} else {
+		self.device = MTLCreateSystemDefaultDevice();
+	}
 
-	NSOpenGLPixelFormat *pixelFormat = [[NSOpenGLPixelFormat alloc] initWithAttributes:attributes];
-	NSOpenGLContext *context = [[NSOpenGLContext alloc] initWithFormat:pixelFormat shareContext:nil];
+	// Configure for explicit drawing.
+	self.paused = YES;
+	self.enableSetNeedsDisplay = NO;
 
-#ifdef DEBUG
-	// When we're using a CoreProfile context, crash if we call a legacy OpenGL function
-	// This will make it much more obvious where and when such a function call is made so
-	// that we can remove such calls.
-	// Without this we'd simply get GL_INVALID_OPERATION error for calling legacy functions
-	// but it would be more difficult to see where that function was called.
-	CGLEnable([context CGLContextObj], kCGLCECrashOnRemovedFunctions);
-#endif
-
-	self.pixelFormat = pixelFormat;
-	self.openGLContext = context;
-	self.wantsBestResolutionOpenGLSurface = YES;
+	// Create the scan target.
+	_scanTarget = [[CSScanTarget alloc] initWithView:self];
+	self.delegate = _scanTarget;
 
 	// Register to receive dragged and dropped file URLs.
 	[self registerForDraggedTypes:@[(__bridge NSString *)kUTTypeFileURL]];
-}
 
-- (void)performWithGLContext:(dispatch_block_t)action flushDrawable:(BOOL)flushDrawable {
-	CGLLockContext([[self openGLContext] CGLContextObj]);
-	[self.openGLContext makeCurrentContext];
-	action();
-	CGLUnlockContext([[self openGLContext] CGLContextObj]);
-
-	if(flushDrawable) CGLFlushDrawable([[self openGLContext] CGLContextObj]);
-}
-
-- (void)performWithGLContext:(nonnull dispatch_block_t)action {
-	[self performWithGLContext:action flushDrawable:NO];
+	// Setup the [initial] display link.
+	[self setupDisplayLink];
 }
 
 #pragma mark - NSResponder
@@ -259,12 +194,16 @@ static CVReturn DisplayLinkCallback(__unused CVDisplayLinkRef displayLink, const
 	[self.responderDelegate paste:sender];
 }
 
+- (NSBitmapImageRep *)imageRepresentation {
+	return self.scanTarget.imageRepresentation;
+}
+
 #pragma mark - NSDraggingDestination
 
 - (BOOL)performDragOperation:(id <NSDraggingInfo>)sender {
 	for(NSPasteboardItem *item in [[sender draggingPasteboard] pasteboardItems]) {
 		NSURL *URL = [NSURL URLWithString:[item stringForType:(__bridge NSString *)kUTTypeFileURL]];
-		[self.delegate openGLView:self didReceiveFileAtURL:URL];
+		[self.responderDelegate scanTargetView:self didReceiveFileAtURL:URL];
 	}
 	return YES;
 }
@@ -300,13 +239,13 @@ static CVReturn DisplayLinkCallback(__unused CVDisplayLinkRef displayLink, const
 
 		_mouseHideTimer = [NSTimer scheduledTimerWithTimeInterval:3.0 repeats:NO block:^(__unused NSTimer * _Nonnull timer) {
 			[NSCursor setHiddenUntilMouseMoves:YES];
-			[self.delegate openGLViewWillHideOSMouseCursor:self];
+			[self.responderDelegate scanTargetViewWillHideOSMouseCursor:self];
 		}];
 	}
 }
 
 - (void)mouseEntered:(NSEvent *)event {
-	[self.delegate openGLViewDidShowOSMouseCursor:self];
+	[self.responderDelegate scanTargetViewDidShowOSMouseCursor:self];
 	[super mouseEntered:event];
 	[self scheduleMouseHide];
 }
@@ -315,7 +254,7 @@ static CVReturn DisplayLinkCallback(__unused CVDisplayLinkRef displayLink, const
 	[super mouseExited:event];
 	[_mouseHideTimer invalidate];
 	_mouseHideTimer = nil;
-	[self.delegate openGLViewWillHideOSMouseCursor:self];
+	[self.responderDelegate scanTargetViewWillHideOSMouseCursor:self];
 }
 
 - (void)releaseMouse {
@@ -323,8 +262,8 @@ static CVReturn DisplayLinkCallback(__unused CVDisplayLinkRef displayLink, const
 		_mouseIsCaptured = NO;
 		CGAssociateMouseAndMouseCursorPosition(true);
 		[NSCursor unhide];
-		[self.delegate openGLViewDidReleaseMouse:self];
-		[self.delegate openGLViewDidShowOSMouseCursor:self];
+		[self.responderDelegate scanTargetViewDidReleaseMouse:self];
+		[self.responderDelegate scanTargetViewDidShowOSMouseCursor:self];
 		((CSApplication *)[NSApplication sharedApplication]).eventDelegate = nil;
 	}
 }
@@ -336,7 +275,7 @@ static CVReturn DisplayLinkCallback(__unused CVDisplayLinkRef displayLink, const
 		// Mouse capture is off, so don't play games with the cursor, just schedule it to
 		// hide in the near future.
 		[self scheduleMouseHide];
-		[self.delegate openGLViewDidShowOSMouseCursor:self];
+		[self.responderDelegate scanTargetViewDidShowOSMouseCursor:self];
 	} else {
 		if(_mouseIsCaptured) {
 			// Mouse capture is on, so move the cursor back to the middle of the window, and
@@ -354,7 +293,7 @@ static CVReturn DisplayLinkCallback(__unused CVDisplayLinkRef displayLink, const
 
 			[self.responderDelegate mouseMoved:event];
 		} else {
-			[self.delegate openGLViewDidShowOSMouseCursor:self];
+			[self.responderDelegate scanTargetViewDidShowOSMouseCursor:self];
 		}
 	}
 }
@@ -387,8 +326,8 @@ static CVReturn DisplayLinkCallback(__unused CVDisplayLinkRef displayLink, const
 			_mouseIsCaptured = YES;
 			[NSCursor hide];
 			CGAssociateMouseAndMouseCursorPosition(false);
-			[self.delegate openGLViewWillHideOSMouseCursor:self];
-			[self.delegate openGLViewDidCaptureMouse:self];
+			[self.responderDelegate scanTargetViewWillHideOSMouseCursor:self];
+			[self.responderDelegate scanTargetViewDidCaptureMouse:self];
 			if(self.shouldUsurpCommand) {
 				((CSApplication *)[NSApplication sharedApplication]).eventDelegate = self;
 			}
diff --git a/OSBindings/Mac/Clock SignalTests/MasterSystemVDPTests.mm b/OSBindings/Mac/Clock SignalTests/MasterSystemVDPTests.mm
index 57e9a7488..5a6167bd4 100644
--- a/OSBindings/Mac/Clock SignalTests/MasterSystemVDPTests.mm	
+++ b/OSBindings/Mac/Clock SignalTests/MasterSystemVDPTests.mm	
@@ -7,7 +7,6 @@
 //
 
 #import <XCTest/XCTest.h>
-#import <OpenGL/OpenGL.h>
 
 #include "9918.hpp"
 
diff --git a/Outputs/CRT/CRT.cpp b/Outputs/CRT/CRT.cpp
index 637be7021..7d9acee7d 100644
--- a/Outputs/CRT/CRT.cpp
+++ b/Outputs/CRT/CRT.cpp
@@ -194,8 +194,8 @@ Outputs::Display::ScanTarget::Scan::EndPoint CRT::end_point(uint16_t data_offset
 	end_point.y = uint16_t(vertical_flywheel_->get_current_output_position() / vertical_flywheel_output_divider_);
 	end_point.data_offset = data_offset;
 
-	// TODO: this is a workaround for the limited precision that can be posted onwards;
-	// it'd be better to make time_multiplier_ an explicit modal and just not divide by it.
+	// Ensure .composite_angle is sampled at the location indicated by .cycles_since_end_of_horizontal_retrace.
+	// TODO: I could supply time_multiplier_ as a modal and just not round .cycles_since_end_of_horizontal_retrace. Would that be better?
 	const auto lost_precision = cycles_since_horizontal_sync_ % time_multiplier_;
 	end_point.composite_angle = int16_t(((phase_numerator_ - lost_precision * colour_cycle_numerator_) << 6) / phase_denominator_) * (is_alernate_line_ ? -1 : 1);
 	end_point.cycles_since_end_of_horizontal_retrace = uint16_t(cycles_since_horizontal_sync_ / time_multiplier_);
@@ -427,7 +427,8 @@ void CRT::set_immediate_default_phase(float phase) {
 
 void CRT::output_data(int number_of_cycles, size_t number_of_samples) {
 #ifndef NDEBUG
-	assert(number_of_samples > 0 && number_of_samples <= allocated_data_length_);
+	assert(number_of_samples > 0);
+	assert(number_of_samples <= allocated_data_length_);
 	allocated_data_length_ = std::numeric_limits<size_t>::min();
 #endif
 	scan_target_->end_data(number_of_samples);
diff --git a/Outputs/CRT/CRT.hpp b/Outputs/CRT/CRT.hpp
index 8c9a79f5d..7ab9dbcf5 100644
--- a/Outputs/CRT/CRT.hpp
+++ b/Outputs/CRT/CRT.hpp
@@ -81,7 +81,7 @@ class CRT {
 
 		Outputs::Display::ScanTarget *scan_target_ = &Outputs::Display::NullScanTarget::singleton;
 		Outputs::Display::ScanTarget::Modals scan_target_modals_;
-		static constexpr uint8_t DefaultAmplitude = 80;
+		static constexpr uint8_t DefaultAmplitude = 41;	// Based upon a black level to maximum excursion and positive burst peak of: NTSC: 882 & 143; PAL: 933 & 150.
 
 #ifndef NDEBUG
 		size_t allocated_data_length_ = std::numeric_limits<size_t>::min();
diff --git a/Outputs/DisplayMetrics.cpp b/Outputs/DisplayMetrics.cpp
index 05632b4b1..f238f5065 100644
--- a/Outputs/DisplayMetrics.cpp
+++ b/Outputs/DisplayMetrics.cpp
@@ -50,7 +50,7 @@ void Metrics::announce_did_resize() {
 	frames_missed_ = frames_hit_ = 0;
 }
 
-void Metrics::announce_draw_status(size_t, std::chrono::high_resolution_clock::duration, bool complete) {
+void Metrics::announce_draw_status(bool complete) {
 	if(!complete) {
 		++frames_missed_;
 	} else {
@@ -79,6 +79,10 @@ void Metrics::announce_draw_status(size_t, std::chrono::high_resolution_clock::d
 	}
 }
 
+void Metrics::announce_draw_status(size_t, std::chrono::high_resolution_clock::duration, bool complete) {
+	announce_draw_status(complete);
+}
+
 bool Metrics::should_lower_resolution() const {
 	// If less than 100 frames are on record, return no opinion; otherwise
 	// suggest a lower resolution if more than 10 frames in the last 100-200
diff --git a/Outputs/DisplayMetrics.hpp b/Outputs/DisplayMetrics.hpp
index 9cd4f7889..1983bdac1 100644
--- a/Outputs/DisplayMetrics.hpp
+++ b/Outputs/DisplayMetrics.hpp
@@ -12,6 +12,7 @@
 #include "ScanTarget.hpp"
 
 #include <array>
+#include <atomic>
 #include <chrono>
 
 namespace Outputs {
@@ -33,6 +34,9 @@ class Metrics {
 		/// Provides Metrics with a new data point for output speed estimation.
 		void announce_draw_status(size_t lines, std::chrono::high_resolution_clock::duration duration, bool complete);
 
+		/// Provides Metrics with a new data point for output speed estimation, albeit without line-specific information.
+		void announce_draw_status(bool complete);
+
 		/// @returns @c true if Metrics thinks a lower output buffer resolution is desirable in the abstract; @c false otherwise.
 		bool should_lower_resolution() const;
 
@@ -48,8 +52,8 @@ class Metrics {
 		size_t line_total_history_pointer_ = 0;
 		void add_line_total(int);
 
-		int frames_hit_ = 0;
-		int frames_missed_ = 0;
+		std::atomic<int> frames_hit_ = 0;
+		std::atomic<int> frames_missed_ = 0;
 };
 
 }
diff --git a/Outputs/OpenGL/OpenGL.hpp b/Outputs/OpenGL/OpenGL.hpp
index 1cef48dda..d8e86ef90 100644
--- a/Outputs/OpenGL/OpenGL.hpp
+++ b/Outputs/OpenGL/OpenGL.hpp
@@ -16,6 +16,8 @@
 #ifdef __APPLE__
 	#if TARGET_OS_IPHONE
 	#else
+		// These remain so that I can, at least for now, build the kiosk version under macOS.
+		// They can be eliminated if and when Apple fully withdraws OpenGL support.
 		#include <OpenGL/OpenGL.h>
 		#include <OpenGL/gl3.h>
 		#include <OpenGL/gl3ext.h>
diff --git a/Outputs/OpenGL/ScanTarget.cpp b/Outputs/OpenGL/ScanTarget.cpp
index 3a1078422..55c27a216 100644
--- a/Outputs/OpenGL/ScanTarget.cpp
+++ b/Outputs/OpenGL/ScanTarget.cpp
@@ -118,8 +118,9 @@ void ScanTarget::setup_pipeline() {
 	const auto data_type_size = Outputs::Display::size_for_data_type(modals.input_data_type);
 
 	// Resize the texture only if required.
-	if(data_type_size != write_area_data_size()) {
-		write_area_texture_.resize(WriteAreaWidth*WriteAreaHeight*data_type_size);
+	const size_t required_size = WriteAreaWidth*WriteAreaHeight*data_type_size;
+	if(required_size != write_area_data_size()) {
+		write_area_texture_.resize(required_size);
 		set_write_area(write_area_texture_.data());
 	}
 
@@ -186,7 +187,9 @@ void ScanTarget::update(int, int output_height) {
 		true);
 
 	// Grab the new output list.
-	perform([=] (const OutputArea &area) {
+	perform([=] {
+		OutputArea area = get_output_area();
+
 		// Establish the pipeline if necessary.
 		const auto new_modals = BufferingScanTarget::new_modals();
 		const bool did_setup_pipeline = bool(new_modals);
@@ -478,6 +481,7 @@ void ScanTarget::update(int, int output_height) {
 		// Grab a fence sync object to avoid busy waiting upon the next extry into this
 		// function, and reset the is_updating_ flag.
 		fence_ = glFenceSync(GL_SYNC_GPU_COMMANDS_COMPLETE, 0);
+		complete_output_area(area);
 	});
 }
 
diff --git a/Outputs/OpenGL/ScanTarget.hpp b/Outputs/OpenGL/ScanTarget.hpp
index 21e21962f..fb7a75db5 100644
--- a/Outputs/OpenGL/ScanTarget.hpp
+++ b/Outputs/OpenGL/ScanTarget.hpp
@@ -38,7 +38,7 @@ namespace OpenGL {
 	this uses various internal buffers so that the only geometry
 	drawn to the target framebuffer is a quad.
 */
-class ScanTarget: public Outputs::Display::BufferingScanTarget {
+class ScanTarget: public Outputs::Display::BufferingScanTarget {	// TODO: use private inheritance and expose only display_metrics() and a custom cast?
 	public:
 		ScanTarget(GLuint target_framebuffer = 0, float output_gamma = 2.2f);
 		~ScanTarget();
diff --git a/Outputs/OpenGL/ScanTargetGLSLFragments.cpp b/Outputs/OpenGL/ScanTargetGLSLFragments.cpp
index 9e5b83e53..7a55329a0 100644
--- a/Outputs/OpenGL/ScanTargetGLSLFragments.cpp
+++ b/Outputs/OpenGL/ScanTargetGLSLFragments.cpp
@@ -489,38 +489,40 @@ std::unique_ptr<Shader> ScanTarget::conversion_shader() const {
 std::unique_ptr<Shader> ScanTarget::composition_shader() const {
 	const auto modals = BufferingScanTarget::modals();
 	const std::string vertex_shader =
-		"#version 150\n"
+	R"x(#version 150
 
-		"in float startDataX;"
-		"in float startClock;"
+		in float startDataX;
+		in float startClock;
 
-		"in float endDataX;"
-		"in float endClock;"
+		in float endDataX;
+		in float endClock;
 
-		"in float dataY;"
-		"in float lineY;"
+		in float dataY;
+		in float lineY;
 
-		"out vec2 textureCoordinate;"
-		"uniform usampler2D textureName;"
+		out vec2 textureCoordinate;
+		uniform usampler2D textureName;
 
-		"void main(void) {"
-			"float lateral = float(gl_VertexID & 1);"
-			"float longitudinal = float((gl_VertexID & 2) >> 1);"
+		void main(void) {
+			float lateral = float(gl_VertexID & 1);
+			float longitudinal = float((gl_VertexID & 2) >> 1);
 
-			"textureCoordinate = vec2(mix(startDataX, endDataX, lateral), dataY + 0.5) / textureSize(textureName, 0);"
-			"vec2 eyePosition = vec2(mix(startClock, endClock, lateral), lineY + longitudinal) / vec2(2048.0, 2048.0);"
-			"gl_Position = vec4(eyePosition*2.0 - vec2(1.0), 0.0, 1.0);"
-		"}";
+			textureCoordinate = vec2(mix(startDataX, endDataX, lateral), dataY + 0.5) / textureSize(textureName, 0);
+			vec2 eyePosition = vec2(mix(startClock, endClock, lateral), lineY + longitudinal) / vec2(2048.0, 2048.0);
+			gl_Position = vec4(eyePosition*2.0 - vec2(1.0), 0.0, 1.0);
+		}
+	)x";
 
 	std::string fragment_shader =
-		"#version 150\n"
+	R"x(#version 150
 
-		"out vec4 fragColour;"
-		"in vec2 textureCoordinate;"
+		out vec4 fragColour;
+		in vec2 textureCoordinate;
 
-		"uniform usampler2D textureName;"
+		uniform usampler2D textureName;
 
-		"void main(void) {";
+		void main(void) {
+	)x";
 
 	switch(modals.input_data_type) {
 		case InputDataType::Luminance1:
diff --git a/Outputs/ScanTarget.hpp b/Outputs/ScanTarget.hpp
index 63be212d1..0860af010 100644
--- a/Outputs/ScanTarget.hpp
+++ b/Outputs/ScanTarget.hpp
@@ -9,6 +9,7 @@
 #ifndef Outputs_Display_ScanTarget_h
 #define Outputs_Display_ScanTarget_h
 
+#include <array>
 #include <cstddef>
 #include <cstdint>
 #include "../ClockReceiver/TimeTypes.hpp"
@@ -53,6 +54,9 @@ enum class DisplayType {
 
 /*!
 	Enumerates the potential formats of input data.
+
+	All types are designed to be 1, 2 or 4 bytes per pixel; this hopefully creates appropriate alignment
+	on all formats.
 */
 enum class InputDataType {
 
@@ -72,8 +76,10 @@ enum class InputDataType {
 	// of a colour subcarrier. So they can be used to generate a luminance signal,
 	// or an s-video pipeline.
 
-	Luminance8Phase8,		// 2 bytes/pixel; first is luminance, second is phase.
-							// Phase is encoded on a 192-unit circle; anything
+	Luminance8Phase8,		// 2 bytes/pixel; first is luminance, second is phase
+							// of a cosine wave.
+							//
+							// Phase is encoded on a 128-unit circle; anything
 							// greater than 192 implies that the colour part of
 							// the signal should be omitted.
 
@@ -86,7 +92,8 @@ enum class InputDataType {
 	Red8Green8Blue8,		// 4 bytes/pixel; first is red, second is green, third is blue, fourth is vacant.
 };
 
-inline size_t size_for_data_type(InputDataType data_type) {
+/// @returns the number of bytes per sample for data of type @c data_type.
+constexpr inline size_t size_for_data_type(InputDataType data_type) {
 	switch(data_type) {
 		case InputDataType::Luminance1:
 		case InputDataType::Luminance8:
@@ -107,7 +114,28 @@ inline size_t size_for_data_type(InputDataType data_type) {
 	}
 }
 
-inline DisplayType natural_display_type_for_data_type(InputDataType data_type) {
+/// @returns @c true if this data type presents normalised data, i.e. each byte holds a
+/// value in the range [0, 255] representing a real number in the range [0.0, 1.0]; @c false otherwise.
+constexpr inline size_t data_type_is_normalised(InputDataType data_type) {
+	switch(data_type) {
+		case InputDataType::Luminance8:
+		case InputDataType::Luminance8Phase8:
+		case InputDataType::Red8Green8Blue8:
+		case InputDataType::PhaseLinkedLuminance8:
+			return true;
+
+		default:
+		case InputDataType::Luminance1:
+		case InputDataType::Red1Green1Blue1:
+		case InputDataType::Red2Green2Blue2:
+		case InputDataType::Red4Green4Blue4:
+			return false;
+	}
+}
+
+/// @returns The 'natural' display type for data of type @c data_type. The natural display is whichever would
+/// display it with the least number of conversions. Caveat: a colour display is assumed for pure-composite data types.
+constexpr inline DisplayType natural_display_type_for_data_type(InputDataType data_type) {
 	switch(data_type) {
 		default:
 		case InputDataType::Luminance1:
@@ -126,6 +154,34 @@ inline DisplayType natural_display_type_for_data_type(InputDataType data_type) {
 	}
 }
 
+/// @returns A 3x3 matrix in row-major order to convert from @c colour_space to RGB.
+inline std::array<float, 9> to_rgb_matrix(ColourSpace colour_space) {
+	const std::array<float, 9> yiq_to_rgb = {1.0f, 1.0f, 1.0f, 0.956f, -0.272f, -1.106f, 0.621f, -0.647f, 1.703f};
+	const std::array<float, 9> yuv_to_rgb = {1.0f, 1.0f, 1.0f, 0.0f, -0.39465f, 2.03211f, 1.13983f, -0.58060f, 0.0f};
+
+	switch(colour_space) {
+		case ColourSpace::YIQ:	return yiq_to_rgb;
+		case ColourSpace::YUV:	return yuv_to_rgb;
+	}
+
+	// Should be unreachable.
+	return std::array<float, 9>{};
+}
+
+/// @returns A 3x3 matrix in row-major order to convert to @c colour_space to RGB.
+inline std::array<float, 9> from_rgb_matrix(ColourSpace colour_space) {
+	const std::array<float, 9> rgb_to_yiq = {0.299f, 0.596f, 0.211f, 0.587f, -0.274f, -0.523f, 0.114f, -0.322f, 0.312f};
+	const std::array<float, 9> rgb_to_yuv = {0.299f, -0.14713f, 0.615f, 0.587f, -0.28886f, -0.51499f, 0.114f, 0.436f, -0.10001f};
+
+	switch(colour_space) {
+		case ColourSpace::YIQ:	return rgb_to_yiq;
+		case ColourSpace::YUV:	return rgb_to_yuv;
+	}
+
+	// Should be unreachable.
+	return std::array<float, 9>{};
+}
+
 /*!
 	Provides an abstract target for 'scans' i.e. continuous sweeps of output data,
 	which are identified by 2d start and end coordinates, and the PCM-sampled data
@@ -325,22 +381,22 @@ struct ScanTarget {
 
 struct ScanStatus {
 	/// The current (prediced) length of a field (including retrace).
-	Time::Seconds field_duration;
+	Time::Seconds field_duration = 0.0;
 	/// The difference applied to the field_duration estimate during the last field.
-	Time::Seconds field_duration_gradient;
+	Time::Seconds field_duration_gradient = 0.0;
 	/// The amount of time this device spends in retrace.
-	Time::Seconds retrace_duration;
+	Time::Seconds retrace_duration = 0.0;
 	/// The distance into the current field, from a small negative amount (in retrace) through
 	/// 0 (start of visible area field) to 1 (end of field).
 	///
 	/// This will increase monotonically, being a measure
 	/// of the current vertical position — i.e. if current_position = 0.8 then a caller can
 	/// conclude that the top 80% of the visible part of the display has been painted.
-	float current_position;
+	float current_position = 0.0f;
 	/// The total number of hsyncs so far encountered;
-	int hsync_count;
+	int hsync_count = 0;
 	/// @c true if retrace is currently going on; @c false otherwise.
-	bool is_in_retrace;
+	bool is_in_retrace = false;
 
 	/*!
 		@returns this ScanStatus, with time-relative fields scaled by dividing them by @c dividend.
diff --git a/Outputs/ScanTargets/BufferingScanTarget.cpp b/Outputs/ScanTargets/BufferingScanTarget.cpp
index 96b9f5bf1..c2d0286aa 100644
--- a/Outputs/ScanTargets/BufferingScanTarget.cpp
+++ b/Outputs/ScanTargets/BufferingScanTarget.cpp
@@ -11,12 +11,6 @@
 #include <cassert>
 #include <cstring>
 
-// If enabled, this uses the producer lock to cover both production and consumption
-// rather than attempting to proceed lockfree. This is primarily for diagnostic purposes;
-// it allows empirical exploration of whether the logical and memory barriers that are
-// meant to mediate things between the read pointers and the submit pointers are functioning.
-#define ONE_BIG_LOCK
-
 #define TextureAddressGetY(v)	uint16_t((v) >> 11)
 #define TextureAddressGetX(v)	uint16_t((v) & 0x7ff)
 #define TextureSub(a, b)		(((a) - (b)) & 0x3fffff)
@@ -64,13 +58,18 @@ uint8_t *BufferingScanTarget::begin_data(size_t required_length, size_t required
 		end_x = aligned_start_x + uint16_t(1 + required_length);
 	}
 
-	// Check whether that steps over the read pointer.
+	// Check whether that steps over the read pointer; if so then the final address will be closer
+	// to the write pointer than the old.
 	const auto end_address = TextureAddress(end_x, output_y);
 	const auto read_pointers = read_pointers_.load(std::memory_order::memory_order_relaxed);
 
 	const auto end_distance = TextureSub(end_address, read_pointers.write_area);
 	const auto previous_distance = TextureSub(write_pointers_.write_area, read_pointers.write_area);
 
+	// Perform a quick sanity check.
+	assert(end_distance >= 0);
+	assert(previous_distance >= 0);
+
 	// If allocating this would somehow make the write pointer back away from the read pointer,
 	// there must not be enough space left.
 	if(end_distance < previous_distance) {
@@ -79,6 +78,7 @@ uint8_t *BufferingScanTarget::begin_data(size_t required_length, size_t required
 	}
 
 	// Everything checks out, note expectation of a future end_data and return the pointer.
+	assert(!data_is_allocated_);
 	data_is_allocated_ = true;
 	vended_write_area_pointer_ = write_pointers_.write_area = TextureAddress(aligned_start_x, output_y);
 
@@ -131,31 +131,42 @@ Outputs::Display::ScanTarget::Scan *BufferingScanTarget::begin_scan() {
 		return nullptr;
 	}
 
-	const auto result = &scan_buffer_[write_pointers_.scan_buffer];
+	const auto result = &scan_buffer_[write_pointers_.scan];
 	const auto read_pointers = read_pointers_.load(std::memory_order::memory_order_relaxed);
 
 	// Advance the pointer.
-	const auto next_write_pointer = decltype(write_pointers_.scan_buffer)((write_pointers_.scan_buffer + 1) % scan_buffer_size_);
+	const auto next_write_pointer = decltype(write_pointers_.scan)((write_pointers_.scan + 1) % scan_buffer_size_);
 
 	// Check whether that's too many.
-	if(next_write_pointer == read_pointers.scan_buffer) {
+	if(next_write_pointer == read_pointers.scan) {
 		allocation_has_failed_ = true;
 		vended_scan_ = nullptr;
 		return nullptr;
 	}
-	write_pointers_.scan_buffer = next_write_pointer;
+	write_pointers_.scan = next_write_pointer;
 	++provided_scans_;
 
 	// Fill in extra OpenGL-specific details.
 	result->line = write_pointers_.line;
 
 	vended_scan_ = result;
+
+#ifndef NDEBUG
+	assert(!scan_is_ongoing_);
+	scan_is_ongoing_ = true;
+#endif
+
 	return &result->scan;
 }
 
 void BufferingScanTarget::end_scan() {
 	std::lock_guard lock_guard(producer_mutex_);
 
+#ifndef NDEBUG
+	assert(scan_is_ongoing_);
+	scan_is_ongoing_ = false;
+#endif
+
 	// Complete the scan only if one is afoot.
 	if(vended_scan_) {
 		vended_scan_->data_y = TextureAddressGetY(vended_write_area_pointer_);
@@ -178,7 +189,7 @@ void BufferingScanTarget::announce(Event event, bool is_visible, const Outputs::
 		// The previous-frame-is-complete flag is subject to a two-slot queue because
 		// measurement for *this* frame needs to begin now, meaning that the previous
 		// result needs to be put somewhere — it'll be attached to the first successful
-		// line output.
+		// line output, whenever that comes.
 		is_first_in_frame_ = true;
 		previous_frame_was_complete_ = frame_is_complete_;
 		frame_is_complete_ = true;
@@ -188,18 +199,18 @@ void BufferingScanTarget::announce(Event event, bool is_visible, const Outputs::
 	if(output_is_visible_ == is_visible) return;
 	output_is_visible_ = is_visible;
 
+#ifndef NDEBUG
+	assert(!scan_is_ongoing_);
+#endif
+
 	if(is_visible) {
 		const auto read_pointers = read_pointers_.load(std::memory_order::memory_order_relaxed);
 
-		// Attempt to allocate a new line, noting allocation failure if necessary.
+		// Attempt to allocate a new line, noting allocation success or failure.
 		const auto next_line = uint16_t((write_pointers_.line + 1) % line_buffer_size_);
-		if(next_line == read_pointers.line) {
-			allocation_has_failed_ = true;
-		}
-		provided_scans_ = 0;
-
-		// If there was space for a new line, establish its start.
+		allocation_has_failed_ = next_line == read_pointers.line;
 		if(!allocation_has_failed_) {
+			// If there was space for a new line, establish its start and reset the count of provided scans.
 			Line &active_line = line_buffer_[size_t(write_pointers_.line)];
 			active_line.end_points[0].x = location.x;
 			active_line.end_points[0].y = location.y;
@@ -207,16 +218,24 @@ void BufferingScanTarget::announce(Event event, bool is_visible, const Outputs::
 			active_line.end_points[0].composite_angle = location.composite_angle;
 			active_line.line = write_pointers_.line;
 			active_line.composite_amplitude = composite_amplitude;
+
+			provided_scans_ = 0;
 		}
 	} else {
 		// Commit the most recent line only if any scans fell on it and all allocation was successful.
 		if(!allocation_has_failed_ && provided_scans_) {
+			const auto submit_pointers = submit_pointers_.load(std::memory_order::memory_order_relaxed);
+
 			// Store metadata.
 			LineMetadata &metadata = line_metadata_buffer_[size_t(write_pointers_.line)];
 			metadata.is_first_in_frame = is_first_in_frame_;
 			metadata.previous_frame_was_complete = previous_frame_was_complete_;
+			metadata.first_scan = submit_pointers.scan;
 			is_first_in_frame_ = false;
 
+			// Sanity check.
+			assert(((metadata.first_scan + size_t(provided_scans_)) % scan_buffer_size_) == write_pointers_.scan);
+
 			// Store actual line data.
 			Line &active_line = line_buffer_[size_t(write_pointers_.line)];
 			active_line.end_points[1].x = location.x;
@@ -228,6 +247,7 @@ void BufferingScanTarget::announce(Event event, bool is_visible, const Outputs::
 			write_pointers_.line = uint16_t((write_pointers_.line + 1) % line_buffer_size_);
 
 			// Update the submit pointers with all lines, scans and data written during this line.
+			std::atomic_thread_fence(std::memory_order::memory_order_release);
 			submit_pointers_.store(write_pointers_, std::memory_order::memory_order_release);
 		} else {
 			// Something failed, or there was nothing on the line anyway, so reset all pointers to where they
@@ -236,9 +256,8 @@ void BufferingScanTarget::announce(Event event, bool is_visible, const Outputs::
 			frame_is_complete_ &= !allocation_has_failed_;
 		}
 
-		// Reset the allocation-has-failed flag for the next line
-		// and mark no line as active.
-		allocation_has_failed_ = false;
+		// Don't permit anything to be allocated on invisible areas.
+		allocation_has_failed_ = true;
 	}
 }
 
@@ -248,6 +267,9 @@ void BufferingScanTarget::will_change_owner() {
 	std::lock_guard lock_guard(producer_mutex_);
 	allocation_has_failed_ = true;
 	vended_scan_ = nullptr;
+#ifdef DEBUG
+	data_is_allocated_ = false;
+#endif
 }
 
 const Outputs::Display::Metrics &BufferingScanTarget::display_metrics() {
@@ -255,16 +277,8 @@ const Outputs::Display::Metrics &BufferingScanTarget::display_metrics() {
 }
 
 void BufferingScanTarget::set_write_area(uint8_t *base) {
-	// This is a bit of a hack. This call needs the producer mutex and should be
-	// safe to call from a @c perform block in order to support all potential consumers.
-	// But the temporary hack of ONE_BIG_LOCK then implies that either I need a recursive
-	// mutex, or I have to make a coupling assumption about my caller. I've done the latter,
-	// because ONE_BIG_LOCK is really really meant to be temporary. I hope.
-#ifndef ONE_BIG_LOCK
 	std::lock_guard lock_guard(producer_mutex_);
-#endif
 	write_area_ = base;
-	data_type_size_ = Outputs::Display::size_for_data_type(modals_.input_data_type);
 	write_pointers_ = submit_pointers_ = read_pointers_ = PointerSet();
 	allocation_has_failed_ = true;
 	vended_scan_ = nullptr;
@@ -285,37 +299,52 @@ void BufferingScanTarget::set_modals(Modals modals) {
 
 // MARK: - Consumer.
 
-void BufferingScanTarget::perform(const std::function<void(const OutputArea &)> &function) {
-#ifdef ONE_BIG_LOCK
-	std::lock_guard lock_guard(producer_mutex_);
-#endif
-
+BufferingScanTarget::OutputArea BufferingScanTarget::get_output_area() {
 	// The area to draw is that between the read pointers, representing wherever reading
 	// last stopped, and the submit pointers, representing all the new data that has been
 	// cleared for submission.
 	const auto submit_pointers = submit_pointers_.load(std::memory_order::memory_order_acquire);
-	const auto read_pointers = read_pointers_.load(std::memory_order::memory_order_relaxed);
+	const auto read_ahead_pointers = read_ahead_pointers_.load(std::memory_order::memory_order_relaxed);
+	std::atomic_thread_fence(std::memory_order::memory_order_acquire);
 
 	OutputArea area;
 
-	area.start.line = read_pointers.line;
+	area.start.line = read_ahead_pointers.line;
 	area.end.line = submit_pointers.line;
 
-	area.start.scan = read_pointers.scan_buffer;
-	area.end.scan = submit_pointers.scan_buffer;
+	area.start.scan = read_ahead_pointers.scan;
+	area.end.scan = submit_pointers.scan;
 
-	area.start.write_area_x = TextureAddressGetX(read_pointers.write_area);
-	area.start.write_area_y = TextureAddressGetY(read_pointers.write_area);
+	area.start.write_area_x = TextureAddressGetX(read_ahead_pointers.write_area);
+	area.start.write_area_y = TextureAddressGetY(read_ahead_pointers.write_area);
 	area.end.write_area_x = TextureAddressGetX(submit_pointers.write_area);
 	area.end.write_area_y = TextureAddressGetY(submit_pointers.write_area);
 
-	// Perform only while holding the is_updating lock.
-	while(is_updating_.test_and_set(std::memory_order_acquire));
-	function(area);
-	is_updating_.clear(std::memory_order_release);
+	// Update the read-ahead pointers.
+	read_ahead_pointers_.store(submit_pointers, std::memory_order::memory_order_relaxed);
 
-	// Update the read pointers.
-	read_pointers_.store(submit_pointers, std::memory_order::memory_order_relaxed);
+#ifndef NDEBUG
+	area.counter = output_area_counter_;
+	++output_area_counter_;
+#endif
+
+	return area;
+}
+
+void BufferingScanTarget::complete_output_area(const OutputArea &area) {
+	// TODO: check that this is the expected next area if in DEBUG mode.
+
+	PointerSet new_read_pointers;
+	new_read_pointers.line = uint16_t(area.end.line);
+	new_read_pointers.scan = uint16_t(area.end.scan);
+	new_read_pointers.write_area = TextureAddress(area.end.write_area_x, area.end.write_area_y);
+	read_pointers_.store(new_read_pointers, std::memory_order::memory_order_relaxed);
+
+#ifndef NDEBUG
+	// This will fire if the caller is announcing completed output areas out of order.
+	assert(area.counter == output_area_next_returned_);
+	++output_area_next_returned_;
+#endif
 }
 
 void BufferingScanTarget::perform(const std::function<void(void)> &function) {
@@ -340,6 +369,13 @@ const Outputs::Display::ScanTarget::Modals *BufferingScanTarget::new_modals() {
 		return nullptr;
 	}
 	modals_are_dirty_ = false;
+
+	// MAJOR SHARP EDGE HERE: assume that because the new_modals have been fetched then the caller will
+	// now ensure their texture buffer is appropriate. They might provide a new pointer and might now.
+	// But either way it's now appropriate to start treating the data size as implied by the data type.
+	std::lock_guard lock_guard(producer_mutex_);
+	data_type_size_ = Outputs::Display::size_for_data_type(modals_.input_data_type);
+
 	return &modals_;
 }
 
diff --git a/Outputs/ScanTargets/BufferingScanTarget.hpp b/Outputs/ScanTargets/BufferingScanTarget.hpp
index 1f6b40ae0..aa74e6579 100644
--- a/Outputs/ScanTargets/BufferingScanTarget.hpp
+++ b/Outputs/ScanTargets/BufferingScanTarget.hpp
@@ -39,7 +39,6 @@ class BufferingScanTarget: public Outputs::Display::ScanTarget {
 		/*! @returns The DisplayMetrics object that this ScanTarget has been providing with announcements and draw overages. */
 		const Metrics &display_metrics();
 
-	protected:
 		static constexpr int WriteAreaWidth = 2048;
 		static constexpr int WriteAreaHeight = 2048;
 
@@ -49,13 +48,15 @@ class BufferingScanTarget: public Outputs::Display::ScanTarget {
 		// It is the subclass's responsibility to post timings.
 		Metrics display_metrics_;
 
-		// Extends the definition of a Scan to include two extra fields,
-		// completing this scan's source data and destination locations.
+		/// Extends the definition of a Scan to include two extra fields,
+		/// completing this scan's source data and destination locations.
 		struct Scan {
 			Outputs::Display::ScanTarget::Scan scan;
 
 			/// Stores the y coordinate for this scan's data within the write area texture.
 			/// Use this plus the scan's endpoints' data_offsets to locate this data in 2d.
+			/// Note that the data_offsets will have been adjusted to be relative to the line
+			/// they fall within, not the data allocation.
 			uint16_t data_y;
 			/// Stores the y coordinate assigned to this scan within the intermediate buffers.
 			/// Use this plus this scan's endpoints' x locations to determine where to composite
@@ -69,11 +70,12 @@ class BufferingScanTarget: public Outputs::Display::ScanTarget {
 		struct Line {
 			struct EndPoint {
 				uint16_t x, y;
-				uint16_t cycles_since_end_of_horizontal_retrace;
 				int16_t composite_angle;
+				uint16_t cycles_since_end_of_horizontal_retrace;
 			} end_points[2];
-			uint16_t line;
+
 			uint8_t composite_amplitude;
+			uint16_t line;
 		};
 
 		/// Provides additional metadata about lines; this is separate because it's unlikely to be of
@@ -86,6 +88,8 @@ class BufferingScanTarget: public Outputs::Display::ScanTarget {
 			/// from a frame if performance problems mean that the emulated machine is running
 			/// more quickly than complete frames can be generated.
 			bool previous_frame_was_complete;
+			/// The index of the first scan that will appear on this line.
+			size_t first_scan;
 		};
 
 		/// Sets the area of memory to use as a scan buffer.
@@ -112,6 +116,10 @@ class BufferingScanTarget: public Outputs::Display::ScanTarget {
 		/// (iii) the number of lines that have been completed.
 		///
 		/// New write areas and scans are exposed only upon completion of the corresponding lines.
+		/// The values indicated by the start point are the first that should be drawn. Those indicated
+		/// by the end point are one after the final that should be drawn.
+		///
+		/// So e.g. start.scan = 23, end.scan = 24 means draw a single scan, index 23.
 		struct OutputArea {
 			struct Endpoint {
 				int write_area_x, write_area_y;
@@ -120,14 +128,32 @@ class BufferingScanTarget: public Outputs::Display::ScanTarget {
 			};
 
 			Endpoint start, end;
+
+#ifndef NDEBUG
+			size_t counter;
+#endif
 		};
+
+		/// Gets the current range of content that has been posted but not yet returned by
+		/// a previous call to get_output_area().
+		///
+		/// Does not require the caller to be within a @c perform block.
+		OutputArea get_output_area();
+
+		/// Announces that the output area has now completed output, freeing up its memory for
+		/// further modification.
+		///
+		/// It is the caller's responsibility to ensure that the areas passed to complete_output_area
+		/// are those from get_output_area and are marked as completed in the same order that
+		/// they were originally provided.
+		///
+		/// Does not require the caller to be within a @c perform block.
+		void complete_output_area(const OutputArea &);
+
 		/// Performs @c action ensuring that no other @c perform actions, or any
 		/// change to modals, occurs simultaneously.
 		void perform(const std::function<void(void)> &action);
 
-		/// Acts as per void(void) @c perform but also dequeues all latest available video output.
-		void perform(const std::function<void(const OutputArea &)> &);
-
 		/// @returns new Modals if any have been set since the last call to get_new_modals().
 		///		The caller must be within a @c perform block.
 		const Modals *new_modals();
@@ -185,7 +211,7 @@ class BufferingScanTarget: public Outputs::Display::ScanTarget {
 			int32_t write_area = 1;
 
 			// Points into the scan buffer.
-			uint16_t scan_buffer = 0;
+			uint16_t scan = 0;
 
 			// Points into the line buffer.
 			uint16_t line = 0;
@@ -199,6 +225,8 @@ class BufferingScanTarget: public Outputs::Display::ScanTarget {
 		/// may run and is therefore used by both producer and consumer.
 		std::atomic<PointerSet> read_pointers_;
 
+		std::atomic<PointerSet> read_ahead_pointers_;
+
 		/// This is used as a spinlock to guard `perform` calls.
 		std::atomic_flag is_updating_;
 
@@ -226,6 +254,13 @@ class BufferingScanTarget: public Outputs::Display::ScanTarget {
 		// from a call to @c get_new_modals.
 		Modals modals_;
 		bool modals_are_dirty_ = false;
+
+#ifndef NDEBUG
+		// Debug features; these amount to API validation.
+		bool scan_is_ongoing_ = false;
+		size_t output_area_counter_ = 0;
+		size_t output_area_next_returned_ = 0;
+#endif
 };
 
 
diff --git a/README.md b/README.md
index 8f27e6a83..df088a029 100644
--- a/README.md
+++ b/README.md
@@ -4,7 +4,7 @@ Clock Signal ('CLK') is an emulator for tourists that seeks to be invisible. Use
 
 macOS and source releases are [hosted on GitHub](https://github.com/TomHarte/CLK/releases). For desktop Linux it is also available as a [Snap](https://snapcraft.io/clock-signal).
 
-On the Mac it is a native Cocoa application; under Linux, BSD and other UNIXes and UNIX-alikes it can be built either with Qt or with SDL; the Qt build should be considered preliminary and is currently closely bound to X11 as Qt doesn't abstract game-like keyboard handling.
+On the Mac it is a native Cocoa and Metal application; under Linux, BSD and other UNIXes and UNIX-alikes it uses OpenGL and can be built either with Qt or with SDL.
 
 So its aims are:
 * single-click load of any piece of source media for any supported platform;
@@ -17,6 +17,7 @@ It currently contains emulations of the:
 * Amstrad CPC;
 * Apple II/II+ and IIe;
 * Atari 2600;
+* Atari ST;
 * ColecoVision;
 * Commodore Vic-20 (and Commodore 1540/1);
 * Macintosh 512ke and Plus;
@@ -25,8 +26,6 @@ It currently contains emulations of the:
 * Sega Master System; and
 * Sinclair ZX80/81.
 
-In addition, emulation of the Atari ST is experimental.
-
 ## Single-click Loading
 
 Through the combination of static analysis and runtime analysis, CLK seeks to be able automatically to select and configure the appropriate machine to run any provided disk, tape or ROM; to issue any commands necessary to run the software contained on the disk, tape or ROM; and to provide accelerated loading where feasible.