Merge pull request #834 from TomHarte/FloatingSpeaker

Resolves lowpass-speaker position aliasing
Since input and output are floating point, using an integer Stepper is not appropriate.
2025-10-25 09:27:01 +00:00 · 2020-09-16 19:05:54 -04:00 · 2020-09-16 18:53:44 -04:00 · 2020-09-16 18:19:58 -04:00 · 2020-09-16 18:15:57 -04:00 · 2020-09-16 17:24:18 -04:00
32 changed files with 2417 additions and 540 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/Outputs/Speaker/Implementation/LowpassSpeaker.hpp
+++ b/Outputs/Speaker/Implementation/LowpassSpeaker.hpp
@@ -10,7 +10,6 @@
 #define FilteringSpeaker_h

 #include "../Speaker.hpp"
-#include "../../../SignalProcessing/Stepper.hpp"
 #include "../../../SignalProcessing/FIRFilter.hpp"
 #include "../../../ClockReceiver/ClockReceiver.hpp"
 #include "../../../Concurrency/AsyncTaskQueue.hpp"
@@ -198,7 +197,8 @@ template <typename SampleSource> class LowpassSpeaker: public Speaker {
 		std::vector<int16_t> input_buffer_;
 		std::vector<int16_t> output_buffer_;

-		std::unique_ptr<SignalProcessing::Stepper> stepper_;
+		float step_rate_ = 0.0f;
+		float position_error_ = 0.0f;
 		std::unique_ptr<SignalProcessing::FIRFilter> filter_;

 		std::mutex filter_parameters_mutex_;
@@ -223,9 +223,8 @@ template <typename SampleSource> class LowpassSpeaker: public Speaker {
 			);
 			number_of_taps = (number_of_taps * 2) | 1;

-			stepper_ = std::make_unique<SignalProcessing::Stepper>(
-				uint64_t(filter_parameters.input_cycles_per_second),
-				uint64_t(filter_parameters.output_cycles_per_second));
+			step_rate_ = filter_parameters.input_cycles_per_second / filter_parameters.output_cycles_per_second;
+			position_error_ = 0.0f;

 			filter_ = std::make_unique<SignalProcessing::FIRFilter>(
 				unsigned(number_of_taps),
@@ -304,7 +303,8 @@ template <typename SampleSource> class LowpassSpeaker: public Speaker {
 			// If the next loop around is going to reuse some of the samples just collected, use a memmove to
 			// preserve them in the correct locations (TODO: use a longer buffer to fix that?) and don't skip
 			// anything. Otherwise skip as required to get to the next sample batch and don't expect to reuse.
-			const auto steps = stepper_->step() * (SampleSource::get_is_stereo() ? 2 : 1);
+			const size_t steps = size_t(step_rate_ + position_error_) * (SampleSource::get_is_stereo() ? 2 : 1);
+			position_error_ = fmodf(step_rate_ + position_error_, 1.0f);
 			if(steps < input_buffer_.size()) {
 				auto *const input_buffer = input_buffer_.data();
 				std::memmove(	input_buffer,
--- 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.
Author	SHA1	Message	Date
Thomas Harte	8b16da9695	Merge pull request #834 from TomHarte/FloatingSpeaker Resolves lowpass-speaker position aliasing	2020-09-16 19:05:54 -04:00
Thomas Harte	f783ec6269	Since input and output are floating point, using an integer Stepper is not appropriate.	2020-09-16 18:53:44 -04:00
Thomas Harte	22c9734874	Merge pull request #832 from TomHarte/MetalScanTarget Adds a Metal ScanTarget, for macOS.	2020-09-16 18:19:58 -04:00
Thomas Harte	a17d0e428f	Protects against some further uninitialised values.	2020-09-16 18:15:57 -04:00
Thomas Harte	bb57f0bcc7	Ensures all 6560 properties have a valid default value.	2020-09-16 17:24:18 -04:00
Thomas Harte	b1aefbfe85	Separates asserts.	2020-09-15 23:24:06 -04:00
Thomas Harte	061288f5a7	Add the Macintosh to the Mac kiosk mode informal test set.	2020-09-15 22:49:00 -04:00
Thomas Harte	5a53474536	Ensure MultiKeyboard deconstructs properly.	2020-09-15 22:48:44 -04:00
Thomas Harte	18d0fff8da	Graduates the Atari ST.	2020-09-15 22:46:38 -04:00
Thomas Harte	0ac2145740	Add Metal/OpenGL distinction.	2020-09-15 22:43:39 -04:00
Thomas Harte	bc8787ded6	Improves macro safety.	2020-09-15 22:26:33 -04:00
Thomas Harte	69d21daaa3	Improves commentary.	2020-09-15 22:21:05 -04:00
Thomas Harte	5651ef606d	Resolves failure to advance video address when output is blocked.	2020-09-15 22:20:06 -04:00
Thomas Harte	b831b31382	Adds a further sanity check.	2020-09-15 17:04:04 -04:00
Thomas Harte	2fd5cc056c	Adds std::atomic_thread_fences, but these seem not to be a magic bullet.	2020-09-15 16:34:34 -04:00
Thomas Harte	82dbdf7dfc	Switches to using regular linear interpolation for supersampling.	2020-09-14 22:36:00 -04:00
Thomas Harte	eb9903cd10	Defensively disables allocation of anything outside of visible lines.	2020-09-14 22:29:05 -04:00
Thomas Harte	227e98d6d7	Slightly simplifies control flow.	2020-09-14 22:27:25 -04:00
Thomas Harte	35476063b7	Resolves potential data races.	2020-09-14 21:07:50 -04:00
Thomas Harte	8557bb2136	Adds minor exposition.	2020-09-14 20:39:52 -04:00
Thomas Harte	c0c7818d5d	Reintroduces screenshots.	2020-09-14 20:33:05 -04:00
Thomas Harte	ceeadd6a33	Edges up towards reimplementing screenshots.	2020-09-13 22:30:17 -04:00
Thomas Harte	1a2545fdea	Excises dangling references to `OpenGLView`, reinstates display link.	2020-09-13 22:11:51 -04:00
Thomas Harte	c5e9a74c88	Uses DisplayMetrics to disable supersampling when too slow.	2020-09-13 21:07:59 -04:00
Thomas Harte	d7972a7b86	Enforces across-the-board supersampling. I'm damned if I can figure out how to talk an MTKView, or Metal in general, into supersampling so as a first effort this does it in software.	2020-09-13 19:30:26 -04:00
Thomas Harte	7dd4c67304	Corrects access to `data_type_size`, adds sanity check on output area return.	2020-09-13 18:59:27 -04:00
Thomas Harte	e113780fd1	Minor: ensures no possibility of a dangling(-ish) pointer within the Mac video.	2020-09-10 22:13:19 -04:00
Thomas Harte	e32ae6c191	Adds UGLY HACKs to workaround uncovered issues in the OpenGL scan target.	2020-09-10 22:10:24 -04:00
Thomas Harte	bcaceff378	Simplifies in-Metal transform logic, loading responsibility for setup onto the CPU. I think I've also finally excised whatever order-of-operations issue I was having with regard to non-4:3 displays.	2020-09-10 20:32:58 -04:00
Thomas Harte	d7b405c6f8	Ensures direct luminance -> 'RGB' video doesn't go down the composition pipeline.	2020-09-10 13:20:40 -04:00
Thomas Harte	edf8cf4dc6	Completes the set of with/without gamma, and ensures correct alpha selection. Also culls some other repetitive TODOs.	2020-09-09 19:28:38 -04:00
Thomas Harte	dfcc8e9822	Switches some of the interpolated fields to half precision.	2020-09-09 18:17:05 -04:00
Thomas Harte	016e96e6f8	Extends usage of `half`. Possibly towards its conclusion.	2020-09-09 15:10:19 -04:00
Thomas Harte	e7ce03c418	Attempts to ensure initial finalised line texture state. This was an attempt to remove the vertical line on the left of a composite display. Obviously the cause is not as suspected.	2020-09-09 13:15:21 -04:00
Thomas Harte	3d392dd81d	Completes conversion of composite & S-Video per-pixel processing to 16-bit floats.	2020-09-09 13:02:04 -04:00
Thomas Harte	42d810db7f	Switches per-pixel uniforms to halfs.	2020-09-09 10:53:09 -04:00
Thomas Harte	18571e8351	Also calculates a chroma kernel size, though it isn't used for anything yet.	2020-09-08 20:08:56 -04:00
Thomas Harte	dda1649ab7	Introduces smaller luma kernel functions where useable.	2020-09-08 19:55:37 -04:00
Thomas Harte	c82e0df071	Starts a transition towards half-precision arithmetic.	2020-09-08 19:37:36 -04:00
Thomas Harte	06b7ea5a6e	Strips the luma kernel back to 1d.	2020-09-08 19:15:19 -04:00
Thomas Harte	c49fcb9ec9	Based on further play: one box filter to separate luma/chroma, another to filter chroma, plus a FIR sharpen on luma.	2020-09-08 16:35:05 -04:00
Thomas Harte	0e44d6d214	Experiments with an all-box filter.	2020-09-08 16:19:08 -04:00
Thomas Harte	6adad7fbf5	Starts experimenting again with box filters.	2020-09-07 22:47:49 -04:00
Thomas Harte	de6ed7b615	Corrects phase-linked luminance support.	2020-09-07 20:53:28 -04:00
Thomas Harte	07dcb4dbb1	Starts reintroducing brightness, gamma and transparency for composite and S-Video pipelines.	2020-09-07 18:19:13 -04:00
Thomas Harte	e99896eadc	At least nominates alpha, gamma and brightness to metal.	2020-09-04 16:07:58 -04:00
Thomas Harte	489701afcb	Fixes window resizing.	2020-09-03 21:28:39 -04:00
Thomas Harte	55e576cc57	Ensures unpainted areas in composite displays have a non-asymptotic effect on luminance calculations.	2020-09-03 21:10:30 -04:00
Thomas Harte	6bd8ec9545	Alas, 1.0 seems to be the limit for proper artefact colour.	2020-09-03 20:53:45 -04:00
Thomas Harte	5cd8d86eef	Switches to dynamic generation of the 'sharpness' filter, correcting issues with the Apple II (amongst others).	2020-09-03 20:48:44 -04:00
Thomas Harte	74d0acdaec	Fixes non-RGB colour composite generation. The hard-coded sharpen filter proves to be a really bad fit for the Apple II though.	2020-09-03 19:04:17 -04:00
Thomas Harte	0288a1974b	Tries: separate filters for chroma and luma, plus a post-separation sharpen filter on the latter.	2020-09-03 13:18:21 -04:00
Thomas Harte	6efd8782fe	Tweaks coefficients some more; makes sure that data is never larger than the intermediate buffers.	2020-09-02 20:14:41 -04:00
Thomas Harte	8bab9d5d60	Corrects S-Video and composite generation for RGB[1/2/4] sources. Also toys with a double luminance filter in order to try to sharpen chrominance. But maybe I should be looking at other convolutions entirely?	2020-09-02 19:13:54 -04:00
Thomas Harte	6ef1dfd8be	Sets a more realistic colour subcarrier amplitude.	2020-09-02 15:52:05 -04:00
Thomas Harte	7e58648743	Corrects front-running bug, plays further with colour amplitude.	2020-09-02 15:51:48 -04:00
Thomas Harte	0f0c3e616d	Tweaks some numbers. I'm largely treading water here. Probably time to think about the race.	2020-09-02 08:17:01 -04:00
Thomas Harte	c7ce65ea4c	Attempts fully to restore composite video. Subject to some sort of nasty race condition for the time being.	2020-09-02 08:03:10 -04:00
Thomas Harte	c36247b609	Ensures reuse of offset buffers. There seems to be some sort of epic race condition as the drawing pipeline lags though. Will need to investigate.	2020-09-01 22:11:48 -04:00
Thomas Harte	15296e43a4	Attempts correctly to set up the CPU side of a composite video pipeline, at least. So: I think this is really close, but I'm out of time for the day.	2020-09-01 21:58:33 -04:00
Thomas Harte	f2929230a2	[Experimentally] introduces blending between computed S-Video fragments.	2020-09-01 21:37:36 -04:00
Thomas Harte	bf252b8061	Fixes sizing of buffers to the current output.	2020-09-01 21:33:54 -04:00
Thomas Harte	9e2bf2af7e	Restricts S-Video processing to updated lines.	2020-09-01 21:27:40 -04:00
Thomas Harte	245f2654f0	Shifts S-Video processing into the compute shader.	2020-09-01 20:37:11 -04:00
Thomas Harte	67ca298a72	Forces a no-op compute shader into the S-Video pipeline. The intention is to restrict the area acted over, and to do the S-Video filtering in there. Then I'll just need two such stages for composite.	2020-09-01 18:39:52 -04:00
Thomas Harte	67d4dbf91a	Starts girding for a third pipeline.	2020-08-31 20:01:59 -04:00
Thomas Harte	b344269140	I think I accept the need for intermediate steps now. This allocates storage.	2020-08-30 20:21:01 -04:00
Thomas Harte	bb547610f2	Adds commentary, shrinks some intermediate texture sizes.	2020-08-30 12:06:29 -04:00
Thomas Harte	1e1f007bb7	Possibly convinces myself that no-separation chroma/luma isn't practical. ... as appealing as it may be, were filters perfect.	2020-08-29 21:25:49 -04:00
Thomas Harte	c40d858f02	Switches back to angular stuff at input resolution; ensures all S-Video modes work. Now to roll back onto composite. Fingers crossed!	2020-08-29 20:54:46 -04:00
Thomas Harte	3d564d85fd	Proves that per-pixel sine/cos evaluation avoids phase issues. Even in PAL mode. But I'd rather not _require_ this as it kind of negates directly-sampled input.	2020-08-29 18:53:37 -04:00
Thomas Harte	02cea40ffa	Attempts to avoid introducing phase error in scanToComposition. Also brightens S-Video up to RGB levels.	2020-08-25 22:41:37 -04:00
Thomas Harte	e502d336db	Having decided that these things probably need to be separate, starts drilling down on S-Video.	2020-08-25 22:05:19 -04:00
Thomas Harte	807cb99f6d	Composite angles are signed.	2020-08-23 21:39:04 -04:00
Thomas Harte	8b6879a782	Brief detour: introduces myself to C++11 multiline string literals. A full cleaning coming soon, I imagine.	2020-08-23 21:18:38 -04:00
Thomas Harte	7ca0362f23	Treads water. Difficult current question: why does the Atari 2600's display change colours as the display tries to achieve horizontal lock? Phase should be unchanged. Ergo something is amiss.	2020-08-23 21:03:26 -04:00
Thomas Harte	56c7bd242a	Marginally tidies.	2020-08-22 16:38:36 -04:00
Thomas Harte	5c6112415a	Sets appropriate clear colour for composition render pass.	2020-08-21 22:41:54 -04:00
Thomas Harte	bf6a0c9fc4	Achieves a return of composite colour for RGB-producing machines.	2020-08-21 22:06:36 -04:00
Thomas Harte	d54b937ab6	Starts trying to do actual composite processing.	2020-08-21 21:11:25 -04:00
Thomas Harte	7c23c32e44	Corrects composition colour phase.	2020-08-20 20:45:45 -04:00
Thomas Harte	4e21d24b5f	Corrects composition colour amplitude.	2020-08-20 20:34:37 -04:00
Thomas Harte	ad6fb85fda	Corrects use of composition buffer. Something is still very obviously amiss in colour processing somewhere down the line, but the correct forms are once again visibly in evidence.	2020-08-20 20:21:28 -04:00
Thomas Harte	5dc39a5d24	Adds the composition render pass. Albeit that something here doesn't work at present.	2020-08-19 21:56:53 -04:00
Thomas Harte	3597f687de	Continues sidling towards composite & S-Video handling.	2020-08-19 21:20:06 -04:00
Thomas Harte	8811506adf	Starts towards building a compound[/composition?] buffer. I now need to discover whether I can use natural integer texture coordinates.	2020-08-17 22:10:02 -04:00
Thomas Harte	11dec6fc0f	Avoids a redundant clear.	2020-08-17 22:09:15 -04:00
Thomas Harte	59c4c8233f	Generalises existing `scanToDisplay` to add `lineToDisplay`.	2020-08-17 21:15:19 -04:00
Thomas Harte	9da79d2d81	Clarifies scaling logic.	2020-08-17 20:29:46 -04:00
Thomas Harte	246b474a25	Removes ONE_BIG_LOCK, having effectively neutered it anyway. Starts work on more explicit API usage validation. Maybe the issue isn't a race condition?	2020-08-16 22:09:25 -04:00
Thomas Harte	27e8a3a1b5	Obeys modals' zoom. Subject to an attempt at factoring aspect ratio differences.	2020-08-16 21:11:43 -04:00
Thomas Harte	745797b596	Introduces a stencil buffer plus the inter-frame clearing it allows.	2020-08-16 16:42:32 -04:00
Thomas Harte	940e9e037e	Adds `first_scan` to LineMetadata. Also reorders `Line` fields to match `Scan` fields, just for visual consistency.	2020-08-16 08:59:37 -04:00
Thomas Harte	512c0079a9	Makes thread safe.	2020-08-15 21:52:55 -04:00
Thomas Harte	645c29f853	Adds an intermediate buffer to correct inter-frame smoothing. Also goes someway back to the old scan output scheduling, albeit presently with limited thread safety.	2020-08-15 21:24:10 -04:00
Thomas Harte	e55945674d	Reduces main thread blocking.	2020-08-14 22:16:49 -04:00
Thomas Harte	7ac88536dd	Respects machine aspect ratio. To an extent. Doesn't currently deal with cropping of machines when the window aspect ratio is smaller.	2020-08-14 21:24:25 -04:00
Thomas Harte	230b9fc9e6	Permits multiple simultaneous scan reading ranges. Also updates the OpenGL scan target as per the latest movements of things.	2020-08-12 22:08:41 -04:00
Thomas Harte	27ca782cac	Enables blending; attempts to enable frame preservation. The latter seems to be evidencing a double buffer at play. More investigation required. On the plus side, the direct route is still well within GPU budget at 4k on my Core M. So a huge improvement there.	2020-08-12 19:34:07 -04:00
Thomas Harte	a136a00a2f	Takes a shot at adding RGB -> S-Video and composite conversion, for all RGB types.	2020-08-11 22:11:50 -04:00
Thomas Harte	637ec35d6a	Adds getters for standard colour-space conversion matrices. These are just more details on the meaning of the colour spaces, so I think they belong here.	2020-08-11 19:58:57 -04:00
Thomas Harte	4b55df1cb4	Starts upon a macro-oriented means of RGB input function generation.	2020-08-10 22:03:39 -04:00
Thomas Harte	b9309268ba	Possibly finally succeeds at moving Accelerate.framework to where it should be.	2020-08-10 21:46:11 -04:00
Thomas Harte	8fa89baf54	Slightly cleans up Xcode project; reenables kiosk-for-Mac builds.	2020-08-10 21:43:32 -04:00
Thomas Harte	8374a5e579	Adds superficially correct `compositeSampleLuminance8Phase8` function. Thereby uncovering a minor error in my decoding of colour phase.	2020-08-10 21:33:59 -04:00
Thomas Harte	525233e10b	Ensures all input data types are parseable in Metal. Though now I need to think a bit more about the best way to compose signal-type conversions, and whether output-type calculations (i.e. gamma, brightness) are applied.	2020-08-10 19:47:47 -04:00
Thomas Harte	eadda6a967	Further strips OpenGL from the macOS target.	2020-08-09 22:17:27 -04:00
Thomas Harte	3d6590af89	Throws out a little more OpenGL.	2020-08-09 22:11:31 -04:00
Thomas Harte	28d933d5d6	Does just enough to get 8-bit RGB and 1-bit luminance machines to display. Assuming an 'RGB' output.	2020-08-09 21:19:07 -04:00
Thomas Harte	c1dc42a094	Add comment on latent design aim. In the hope that I don't forget why I did this.	2020-08-09 21:18:23 -04:00
Thomas Harte	6384ff3ee7	Add fix for `data_type_size_` for owners that don't change texture pointer upon new modals.	2020-08-09 21:17:51 -04:00
Thomas Harte	a118594c8b	Hacks to make RGB1 visible (in a fashion).	2020-08-09 20:45:51 -04:00
Thomas Harte	93c6105442	Corrects calculation of dirty texture area.	2020-08-09 20:45:14 -04:00
Thomas Harte	ced4a75a1a	Adds note on the buffering scan target's minor adaptation of data_offset.	2020-08-09 20:44:46 -04:00
Thomas Harte	57fecdc09e	Ties everything together in an attempt to display RGB scans. I'm actually just getting a mess of pixels, but it's something!	2020-08-09 18:41:15 -04:00
Thomas Harte	cd491bb6e0	Cleans up project file; macOS 10.13 is definitely the deployment target.	2020-08-09 18:27:57 -04:00
Thomas Harte	f16ad8f71d	Takes a shot at submitting texture changes.	2020-08-09 17:59:52 -04:00
Thomas Harte	e340685a99	Seemingly proves that proper geometry is reaching Metal by drawing scans. No in-buffer accumulation yet, but this is progress. If I can add accumulation and stencil clearing, I'm not doing badly.	2020-08-08 23:11:44 -04:00
Thomas Harte	df89a8771c	Makes an attempt to have the emulator fill the actual GPU buffers. Not that they're drawn from correctly yet. I might first take a run at a new quick-path output route for emulated RGB displays, that just seeks to use the scans directly. No intermediate buffers. Besides probably being a good feature, it'll be a good way to ramp further up with Metal.	2020-08-08 22:49:02 -04:00
Thomas Harte	bdcf266e45	Having learnt a bit more: eliminates Metal `attribute` tags, switches to more natural expression of structs. Also thereby eliminates the need for a forced alignas(4) on various structs.	2020-08-08 17:27:32 -04:00
Thomas Harte	edf41b06fd	Eliminates the quad buffer. Vertices can be adduced from vertex ID.	2020-08-08 17:12:49 -04:00
Thomas Harte	38960a08d6	Adds adjustment for display aspect ratio. While also realising that I appear to be getting away without an MTLVertexDescriptor for Scans. Maybe OpenGL has prejudiced me, and they're actually optional for interleaved data?	2020-08-07 22:29:24 -04:00
Thomas Harte	fbda7aab23	Does just enough to get the correct (aspect ratio aside) output of scan outlines. So, up next, can I start streaming these things?	2020-08-07 22:20:01 -04:00
Thomas Harte	c575aa0640	Adds a buffer for scans, and posts two test instances.	2020-08-07 22:03:54 -04:00
Thomas Harte	583f6b1ba2	Modifies BufferingScanTarget to allow has-a relationship. I might switch fully to has-a. Further consideration required.	2020-08-07 22:03:27 -04:00
Thomas Harte	bb55ecc101	Disables --volume for kiosk mode testing.	2020-08-07 21:19:53 -04:00
Thomas Harte	4421acef34	Gets some uniforms in on the action. With some effort towards scans, but incompletely so.	2020-08-07 21:19:17 -04:00
Thomas Harte	4c9418f59a	Guarantees alignof(4) on all GPU-bound structures. Taken as given: Metal's requirement here is reasonable enough that it'll either be the same as other frameworks, or at least possibly help them down a fast path.	2020-08-07 21:18:08 -04:00
Thomas Harte	219923bd63	Reduces vertex size, draws a quad.	2020-08-05 21:33:25 -04:00
Thomas Harte	7551782a25	Switches to interleaved vertex data. This more closely relates to what I actually want to do.	2020-08-05 17:27:43 -04:00
Thomas Harte	5c836604c0	Reenable MaserSystem code. Accidental/poor branch management is evidenced here.	2020-08-04 21:50:54 -04:00
Thomas Harte	eff24a8726	My first baby steps in Metal continue; here's a triangle.	2020-08-04 21:49:01 -04:00
Thomas Harte	72df6e52cd	This is possibly at least dispatching an empty command buffer correctly.	2020-08-04 19:44:56 -04:00
Thomas Harte	e235a45abb	Breaks all output. ... by switching out NSOpenGLView for MKLView with no drawing infrastructure yet in place.	2020-08-04 18:22:14 -04:00