Merge pull request #832 from TomHarte/MetalScanTarget

Adds a Metal ScanTarget, for macOS.
2024-07-04 18:29:40 +00:00 · 2020-09-16 18:19:58 -04:00 · 2020-09-16 18:19:58 -04:00 · 22c9734874
commit 22c9734874
parent d20c11e401 a17d0e428f
31 changed files with 2411 additions and 534 deletions
--- 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;
 };

 }
--- 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) {
--- 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.
--- 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_);
--- 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;
 					}
 				}

--- a/Signal.xcodeproj/project.pbxproj
+++ b/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",
--- a/Signal.xcodeproj/xcshareddata/xcschemes/Clock
+++ b/Signal.xcodeproj/xcshareddata/xcschemes/Clock
@ -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"
--- a/Signal/Base.lproj/MachineDocument.xib
+++ b/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">
--- a/Signal/ClockSignal-Bridging-Header.h
+++ b/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"
--- a/Signal/Documents/MachineDocument.swift
+++ b/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")
--- a/Signal/Machine/CSMachine.h
+++ b/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;
--- a/Signal/Machine/CSMachine.mm
+++ b/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();
 			});
 		}
--- a/Signal/ScanTarget/CSScanTarget+CppScanTarget.h
+++ b/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
--- a/Signal/ScanTarget/CSScanTarget.h
+++ b/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
--- a/Signal/ScanTarget/CSScanTarget.mm
+++ b/Signal/ScanTarget/CSScanTarget.mm
--- a/Signal/ScanTarget/ScanTarget.metal
+++ b/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);
+}
--- a/Signal/Views/CSScanTargetView.h
+++ b/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
--- a/Signal/Views/CSScanTargetView.m
+++ b/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;
 			}
--- a/SignalTests/MasterSystemVDPTests.mm
+++ b/SignalTests/MasterSystemVDPTests.mm
@ -7,7 +7,6 @@
 //

 #import <XCTest/XCTest.h>
-#import <OpenGL/OpenGL.h>

 #include "9918.hpp"

--- 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);
--- 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();
--- 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
--- 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;
 };

 }
--- 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>
--- 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);
 	});
 }

--- 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();
--- 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:
--- 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.
--- 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_;
 }

--- 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
 };


--- 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.