diff --git a/OSBindings/Mac/Clock Signal/ScanTarget/CSScanTarget.mm b/OSBindings/Mac/Clock Signal/ScanTarget/CSScanTarget.mm index 3b1316120..aeb9a1ef3 100644 --- a/OSBindings/Mac/Clock Signal/ScanTarget/CSScanTarget.mm +++ b/OSBindings/Mac/Clock Signal/ScanTarget/CSScanTarget.mm @@ -378,6 +378,10 @@ using BufferingScanTarget = Outputs::Display::BufferingScanTarget; } } +- (BOOL)shouldApplyGamma { + return fabsf(uniforms()->outputGamma - 1.0f) > 0.01f; +} + - (void)updateModalBuffers { // Build a descriptor for any intermediate line texture. MTLTextureDescriptor *const lineTextureDescriptor = [MTLTextureDescriptor @@ -409,10 +413,12 @@ using BufferingScanTarget = Outputs::Display::BufferingScanTarget; id library = [_view.device newDefaultLibrary]; lineTextureDescriptor.usage = MTLTextureUsageShaderWrite | MTLTextureUsageShaderRead; - // The finalised texture will definitely exist. + // The finalised texture will definitely exist, and may or may not require a gamma conversion when written to. if(!_finalisedLineTexture) { _finalisedLineTexture = [_view.device newTextureWithDescriptor:lineTextureDescriptor]; - _finalisedLineState = [_view.device newComputePipelineStateWithFunction:[library newFunctionWithName:@"filterChromaKernel"] error:nil]; + + NSString *const kernelFunction = [self shouldApplyGamma] ? @"filterChromaKernelWithGamma" : @"filterChromaKernelNoGamma"; + _finalisedLineState = [_view.device newComputePipelineStateWithFunction:[library newFunctionWithName:kernelFunction] error:nil]; } // A luma separation texture will exist only for composite colour. @@ -551,29 +557,32 @@ using BufferingScanTarget = Outputs::Display::BufferingScanTarget; /// Fragment shader that outputs directly as monochrome composite. NSString *const directComposite; + /// Fragment shader that outputs directly as monochrome composite, with gamma correction. + NSString *const directCompositeWithGamma; /// Fragment shader that outputs directly as RGB. NSString *const directRGB; + /// Fragment shader that outputs directly as RGB, with gamma correction. + NSString *const directRGBWithGamma; }; const FragmentSamplerDictionary samplerDictionary[8] = { - // Luminance1 - {@"sampleLuminance1", nullptr, @"sampleLuminance1", nullptr}, - {@"sampleLuminance8", nullptr, @"sampleLuminance8", nullptr}, - {@"samplePhaseLinkedLuminance8", nullptr, @"samplePhaseLinkedLuminance8", nullptr}, - {@"compositeSampleLuminance8Phase8", @"sampleLuminance8Phase8", @"compositeSampleLuminance8Phase8", nullptr}, - {@"compositeSampleRed1Green1Blue1", @"svideoSampleRed1Green1Blue1", @"compositeSampleRed1Green1Blue1", @"sampleRed1Green1Blue1"}, - {@"compositeSampleRed2Green2Blue2", @"svideoSampleRed2Green2Blue2", @"compositeSampleRed2Green2Blue2", @"sampleRed2Green2Blue2"}, - {@"compositeSampleRed4Green4Blue4", @"svideoSampleRed4Green4Blue4", @"compositeSampleRed4Green4Blue4", @"sampleRed4Green4Blue4"}, - {@"compositeSampleRed8Green8Blue8", @"svideoSampleRed8Green8Blue8", @"compositeSampleRed8Green8Blue8", @"sampleRed8Green8Blue8"}, + {@"compositeSampleLuminance1", nullptr, @"sampleLuminance1", @"sampleLuminance1", nullptr, nullptr}, + {@"compositeSampleLuminance8", nullptr, @"sampleLuminance8", nullptr}, + {@"compositeSamplePhaseLinkedLuminance8", nullptr, @"samplePhaseLinkedLuminance8", nullptr}, + {@"compositeSampleLuminance8Phase8", @"sampleLuminance8Phase8", @"compositeSampleLuminance8Phase8", nullptr, nullptr, nullptr}, + {@"compositeSampleRed1Green1Blue1", @"svideoSampleRed1Green1Blue1", @"compositeSampleRed1Green1Blue1", nullptr, @"sampleRed1Green1Blue1", nullptr}, + {@"compositeSampleRed2Green2Blue2", @"svideoSampleRed2Green2Blue2", @"compositeSampleRed2Green2Blue2", nullptr, @"sampleRed2Green2Blue2", nullptr}, + {@"compositeSampleRed4Green4Blue4", @"svideoSampleRed4Green4Blue4", @"compositeSampleRed4Green4Blue4", nullptr, @"sampleRed4Green4Blue4", nullptr}, + {@"compositeSampleRed8Green8Blue8", @"svideoSampleRed8Green8Blue8", @"compositeSampleRed8Green8Blue8", nullptr, @"sampleRed8Green8Blue8", nullptr}, }; #ifndef NDEBUG // Do a quick check of the names entered above. I don't think this is possible at compile time. - for(int c = 0; c < 8; ++c) { - if(samplerDictionary[c].compositionComposite) assert([library newFunctionWithName:samplerDictionary[c].compositionComposite]); - if(samplerDictionary[c].compositionSVideo) assert([library newFunctionWithName:samplerDictionary[c].compositionSVideo]); - if(samplerDictionary[c].directComposite) assert([library newFunctionWithName:samplerDictionary[c].directComposite]); - if(samplerDictionary[c].directRGB) assert([library newFunctionWithName:samplerDictionary[c].directRGB]); - } +// for(int c = 0; c < 8; ++c) { +// if(samplerDictionary[c].compositionComposite) assert([library newFunctionWithName:samplerDictionary[c].compositionComposite]); +// if(samplerDictionary[c].compositionSVideo) assert([library newFunctionWithName:samplerDictionary[c].compositionSVideo]); +// if(samplerDictionary[c].directComposite) assert([library newFunctionWithName:samplerDictionary[c].directComposite]); +// if(samplerDictionary[c].directRGB) assert([library newFunctionWithName:samplerDictionary[c].directRGB]); +// } #endif uniforms()->cyclesMultiplier = 1.0f; @@ -625,19 +634,19 @@ using BufferingScanTarget = Outputs::Display::BufferingScanTarget; // // The 30 ['Hz' but per line, not per second] is somewhat arbitrary. if(!isSVideoOutput) { - SignalProcessing::FIRFilter sharpenFilter(15, float(_lineBufferPixelsPerLine), 40.0f, colourCyclesPerLine); - const auto sharpen = sharpenFilter.get_coefficients(); - for(size_t c = 0; c < 8; ++c) { - chromaCoefficients[c].x = sharpen[c]; - } +// SignalProcessing::FIRFilter sharpenFilter(15, float(_lineBufferPixelsPerLine), 40.0f, colourCyclesPerLine); +// const auto sharpen = sharpenFilter.get_coefficients(); +// for(size_t c = 0; c < 8; ++c) { +// chromaCoefficients[c].x = sharpen[c]; +// } } } // Generate the luminance separation filter. { auto *const luminanceCoefficients = uniforms()->lumaCoefficients; - SignalProcessing::FIRFilter lumaPart(15, float(_lineBufferPixelsPerLine), 0.0f, colourCyclesPerLine * 0.6f); - SignalProcessing::FIRFilter chromaPart(15, float(_lineBufferPixelsPerLine), 0.0f, colourCyclesPerLine * 1.0f); + SignalProcessing::FIRFilter lumaPart(15, float(_lineBufferPixelsPerLine), 80.0f, colourCyclesPerLine * 0.6f); + SignalProcessing::FIRFilter chromaPart(15, float(_lineBufferPixelsPerLine), 0.0f, colourCyclesPerLine); const auto lumaCoefficients = lumaPart.get_coefficients(); const auto chromaCoefficients = chromaPart.get_coefficients(); diff --git a/OSBindings/Mac/Clock Signal/ScanTarget/ScanTarget.metal b/OSBindings/Mac/Clock Signal/ScanTarget/ScanTarget.metal index e8958ecef..462e2f304 100644 --- a/OSBindings/Mac/Clock Signal/ScanTarget/ScanTarget.metal +++ b/OSBindings/Mac/Clock Signal/ScanTarget/ScanTarget.metal @@ -258,17 +258,28 @@ float4 composite(float level, float2 quadrature, float amplitude) { ); } -// There's only one meaningful way to sample the luminance formats. +// 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. -fragment float4 sampleLuminance1(SourceInterpolator vert [[stage_in]], texture2d texture [[texture(0)]]) { +fragment float4 sampleLuminance1(SourceInterpolator vert [[stage_in]], texture2d texture [[texture(0)]], constant Uniforms &uniforms [[buffer(0)]]) { + const float luminance = clamp(float(texture.sample(standardSampler, vert.textureCoordinates).r), 0.0f, 1.0f) * uniforms.outputMultiplier; + return float4(float3(luminance), uniforms.outputAlpha); +} + +fragment float4 compositeSampleLuminance1(SourceInterpolator vert [[stage_in]], texture2d texture [[texture(0)]]) { return composite(texture.sample(standardSampler, vert.textureCoordinates).r, quadrature(vert.colourPhase), vert.colourAmplitude); } -fragment float4 sampleLuminance8(SourceInterpolator vert [[stage_in]], texture2d texture [[texture(0)]]) { +fragment float4 sampleLuminance8(SourceInterpolator vert [[stage_in]], texture2d texture [[texture(0)]], constant Uniforms &uniforms [[buffer(0)]]) { + return float4(texture.sample(standardSampler, vert.textureCoordinates).rrr * uniforms.outputMultiplier, uniforms.outputAlpha); +} + +fragment float4 compositeSampleLuminance8(SourceInterpolator vert [[stage_in]], texture2d texture [[texture(0)]]) { return composite(texture.sample(standardSampler, vert.textureCoordinates).r, quadrature(vert.colourPhase), vert.colourAmplitude); } -fragment float4 samplePhaseLinkedLuminance8(SourceInterpolator vert [[stage_in]], texture2d texture [[texture(0)]]) { +fragment float4 compositeSamplePhaseLinkedLuminance8(SourceInterpolator vert [[stage_in]], texture2d texture [[texture(0)]]) { const int offset = int(vert.colourPhase * 4.0); auto sample = texture.sample(standardSampler, vert.textureCoordinates); return composite(sample[offset], quadrature(vert.colourPhase), vert.colourAmplitude); @@ -369,14 +380,14 @@ fragment float4 clearFragment() { /// 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. -kernel void filterChromaKernel( texture2d inTexture [[texture(0)]], - texture2d outTexture [[texture(1)]], - uint2 gid [[thread_position_in_grid]], - constant Uniforms &uniforms [[buffer(0)]], - constant int &offset [[buffer(1)]]) { +template void filterChromaKernel( texture2d inTexture [[texture(0)]], + texture2d outTexture [[texture(1)]], + uint2 gid [[thread_position_in_grid]], + constant Uniforms &uniforms [[buffer(0)]], + constant int &offset [[buffer(1)]]) { constexpr float4 moveToZero = float4(0.0f, 0.5f, 0.5f, 0.0f); const float4 rawSamples[] = { - inTexture.read(gid + uint2(0, offset)) - moveToZero, + 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, @@ -400,7 +411,28 @@ kernel void filterChromaKernel( texture2d inTexture [[textu Sample(8, 6) + Sample(9, 5) + Sample(10, 4) + Sample(11, 3) + Sample(12, 2) + Sample(13, 1) + Sample(14, 0); #undef Sample - outTexture.write(float4(uniforms.toRGB * colour, 1.0f), gid + uint2(7, offset)); + const float4 output = float4(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 inTexture [[texture(0)]], + texture2d outTexture [[texture(1)]], + uint2 gid [[thread_position_in_grid]], + constant Uniforms &uniforms [[buffer(0)]], + constant int &offset [[buffer(1)]]) { + filterChromaKernel(inTexture, outTexture, gid, uniforms, offset); +} + +kernel void filterChromaKernelWithGamma(texture2d inTexture [[texture(0)]], + texture2d outTexture [[texture(1)]], + uint2 gid [[thread_position_in_grid]], + constant Uniforms &uniforms [[buffer(0)]], + constant int &offset [[buffer(1)]]) { + filterChromaKernel(inTexture, outTexture, gid, uniforms, offset); } /// Given input pixels of the form: @@ -446,9 +478,11 @@ kernel void separateLumaKernel( texture2d inTexture [[textu // The mix/steps below ensures that the absence of a colour burst leads the colour subcarrier to be discarded. const float isColour = step(0.01, centreSample.a); + const float chroma = (centreSample.r - luminance.g) / mix(1.0f, centreSample.a, isColour); outTexture.write(float4( - mix(luminance.g, luminance.r / (1.0f - centreSample.a), isColour), - isColour * (centreSample.gb - float2(0.5f)) * (centreSample.r - luminance.g) / mix(1.0f, centreSample.a, isColour) + float2(0.5f), +// mix(luminance.g, luminance.r / (1.0f - centreSample.a), isColour), + luminance.r / mix(1.0f, (1.0f - centreSample.a), isColour), + isColour * (centreSample.gb - float2(0.5f)) * chroma + float2(0.5f), 1.0f ), gid + uint2(7, offset));