1// NTSC samples per color carrier cycle (assumed to be 15 or 12)
2override samples_per_color_cycle: i32;
3
4// FIR filter length (assumed to be between 1 and 84, inclusive)
5override fir_len: i32;
6
7// Number of NTSC samples to generate per frame buffer pixel (assumed to be at least 1)
8override upscale_factor: i32;
9
10// Phase offset to apply when demodulating U and V (should only be non-zero for NES NTSC output)
11override decode_hue_offset: f32 = 0.0;
12
13override decode_brightness: f32 = 1.0;
14override decode_saturation: f32 = 1.0;
15override decode_gamma: f32 = 2.2;
16
17// Used by rgb_to_ntsc
18@group(0) @binding(0) var<uniform> y_encode_lpf_coefficients: array<vec4f, 21>;
19@group(0) @binding(1) var<uniform> uv_encode_lpf_coefficients: array<vec4f, 21>;
20@group(0) @binding(2) var input_frame: texture_2d<f32>;
21@group(0) @binding(3) var ntsc_frame_w: texture_storage_2d<r32float, write>;
22
23// Used by separate_luma_chroma
24@group(0) @binding(4) var<uniform> luma_bsf_coefficients: array<vec4f, 21>;
25@group(0) @binding(5) var<uniform> chroma_bpf_coefficients: array<vec4f, 21>;
26@group(0) @binding(6) var ntsc_frame_r: texture_2d<f32>;
27@group(0) @binding(7) var ntsc_luma_w: texture_storage_2d<r32float, write>;
28@group(0) @binding(8) var ntsc_chroma_w: texture_storage_2d<r32float, write>;
29
30// Used by luma_chroma_to_rgb
31@group(0) @binding(9) var<uniform> y_decode_lpf_coefficients: array<vec4f, 21>;
32@group(0) @binding(10) var<uniform> uv_decode_lpf_coefficients: array<vec4f, 21>;
33@group(0) @binding(11) var ntsc_luma_r: texture_2d<f32>;
34@group(0) @binding(12) var ntsc_chroma_r: texture_2d<f32>;
35@group(0) @binding(13) var output_frame: texture_storage_2d<rgba8unorm, write>;
36
37struct ImmediateParams {
38 frame_phase_offset: i32,
39 per_line_phase_offset: i32,
40}
41
42// This is a uniform instead of an actual immediate because WebGPU does not support immediates (yet)
43@group(1) @binding(0) var<uniform> immediate_params: ImmediateParams;
44
45// https://www.nesdev.org/wiki/NTSC_video#Converting_YUV_to_signal_RGB
46// Y = 0.299*R + 0.587*G + 0.114*B
47// U = 0.492111 * (B - Y)
48// V = 0.877283 * (R - Y)
49const RGB_TO_YUV: mat3x3f = mat3x3f(
50 vec3f(0.299, 0.492111 * -0.299, 0.877283 * (1.0 - 0.299) ),
51 vec3f(0.587, 0.492111 * -0.587, 0.877283 * -0.587 ),
52 vec3f(0.114, 0.492111 * (1.0 - 0.114), 0.877283 * -0.114 ),
53);
54
55const YUV_TO_RGB: mat3x3f = mat3x3f(
56 vec3f(1.0, 1.0, 1.0),
57 vec3f(0.0, -0.394642, 2.032062),
58 vec3f(1.139883, -0.580622, 0.0),
59);
60
61const PI: f32 = radians(180.0);
62
63// Extra pixels to render at the horizontal edges, to avoid the NTSC signal sharply cutting off at the borders
64const BACKDROP_PIXELS: i32 = 6;
65
66// Convert from RGB to YUV, apply LPF to each YUV component, encode from YUV to NTSC
67// LPF to Y instead of BSF because I think it looks slightly better, and it's going to get LPFed during decoding anyway
68@compute @workgroup_size(16, 16, 1)
69fn rgb_to_ntsc(@builtin(global_invocation_id) invocation: vec3u) {
70 let ntsc_size = vec2i(textureDimensions(ntsc_frame_w));
71 let position = vec2i(invocation.xy);
72 if position.x >= ntsc_size.x || position.y >= ntsc_size.y {
73 return;
74 }
75
76 let input_size = vec2i(textureDimensions(input_frame));
77
78 let start_x = position.x + fir_len / 2 - upscale_factor * BACKDROP_PIXELS;
79 let input_divisor = vec2i(upscale_factor, 1);
80
81 var filtered_y = vec4f(0.0);
82 var filtered_u = vec4f(0.0);
83 var filtered_v = vec4f(0.0);
84 for (var i = 0; i < fir_len; i += 4) {
85 let rgb_pixels = array(
86 load_input_bounds_checked(vec2i(start_x - i, position.y) / input_divisor, input_size, vec3f(0.0)),
87 load_input_bounds_checked(vec2i(start_x - i - 1, position.y) / input_divisor, input_size, vec3f(0.0)),
88 load_input_bounds_checked(vec2i(start_x - i - 2, position.y) / input_divisor, input_size, vec3f(0.0)),
89 load_input_bounds_checked(vec2i(start_x - i - 3, position.y) / input_divisor, input_size, vec3f(0.0)),
90 );
91
92 let yuv_pixels = array(
93 RGB_TO_YUV * rgb_pixels[0],
94 RGB_TO_YUV * rgb_pixels[1],
95 RGB_TO_YUV * rgb_pixels[2],
96 RGB_TO_YUV * rgb_pixels[3],
97 );
98
99 let y_coefficients = y_encode_lpf_coefficients[i / 4];
100 let uv_coefficients = uv_encode_lpf_coefficients[i / 4];
101
102 filtered_y = fma(
103 y_coefficients,
104 vec4f(yuv_pixels[0].r, yuv_pixels[1].r, yuv_pixels[2].r, yuv_pixels[3].r),
105 filtered_y,
106 );
107 filtered_u = fma(
108 uv_coefficients,
109 vec4f(yuv_pixels[0].g, yuv_pixels[1].g, yuv_pixels[2].g, yuv_pixels[3].g),
110 filtered_u,
111 );
112 filtered_v = fma(
113 uv_coefficients,
114 vec4f(yuv_pixels[0].b, yuv_pixels[1].b, yuv_pixels[2].b, yuv_pixels[3].b),
115 filtered_v,
116 );
117 }
118
119 let yuv = vec3f(
120 dot(filtered_y, vec4f(1.0)),
121 dot(filtered_u, vec4f(1.0)),
122 dot(filtered_v, vec4f(1.0)),
123 );
124
125 let phase_x = position.x
126 + immediate_params.frame_phase_offset
127 + position.y * immediate_params.per_line_phase_offset;
128 let phase = f32(phase_x) / f32(samples_per_color_cycle) * 2.0 * PI;
129 let ntsc = yuv.r + yuv.g * sin(phase) + yuv.b * cos(phase);
130 textureStore(ntsc_frame_w, position, vec4f(ntsc, vec3f(0.0)));
131}
132
133fn load_input_bounds_checked(position: vec2i, input_size: vec2i, default_color: vec3f) -> vec3f {
134 if position.x < 0 || position.y < 0 || position.x >= input_size.x || position.y >= input_size.y {
135 // WGSL spec says implementations may return any texel within the texture if coordinates are out of bounds;
136 // guarantee that a black pixel gets loaded instead
137 return default_color;
138 }
139
140 return textureLoad(input_frame, position, 0).rgb;
141}
142
143// Apply BPF and BSF to NTSC signal
144@compute @workgroup_size(16, 16, 1)
145fn separate_luma_chroma(@builtin(global_invocation_id) invocation: vec3u) {
146 let frame_size = vec2i(textureDimensions(ntsc_frame_r));
147 let position = vec2i(invocation.xy);
148 if position.x >= frame_size.x || position.y >= frame_size.y {
149 return;
150 }
151
152 let start_x = position.x + fir_len / 2;
153
154 var pass_filtered = vec4f(0.0);
155 var stop_filtered = vec4f(0.0);
156 for (var i = 0; i < fir_len; i += 4) {
157 let ntsc_samples = vec4f(
158 textureLoad(ntsc_frame_r, vec2i(start_x - i, position.y), 0).r,
159 textureLoad(ntsc_frame_r, vec2i(start_x - i - 1, position.y), 0).r,
160 textureLoad(ntsc_frame_r, vec2i(start_x - i - 2, position.y), 0).r,
161 textureLoad(ntsc_frame_r, vec2i(start_x - i - 3, position.y), 0).r,
162 );
163
164 pass_filtered = fma(ntsc_samples, chroma_bpf_coefficients[i / 4], pass_filtered);
165 stop_filtered = fma(ntsc_samples, luma_bsf_coefficients[i / 4], stop_filtered);
166 }
167
168 let pass_sample = dot(pass_filtered, vec4f(1.0));
169 let stop_sample = dot(stop_filtered, vec4f(1.0));
170
171 textureStore(ntsc_chroma_w, position, vec4f(pass_sample, vec3f(0.0)));
172 textureStore(ntsc_luma_w, position, vec4f(stop_sample, vec3f(0.0)));
173}
174
175// Decode U and V from chroma, apply LPF to each YUV component, convert from YUV to RGB
176@compute @workgroup_size(16, 16, 1)
177fn luma_chroma_to_rgb(@builtin(global_invocation_id) invocation: vec3u) {
178 let output_size = vec2i(textureDimensions(output_frame));
179 let position = vec2i(invocation.xy);
180 if position.x >= output_size.x || position.y >= output_size.y {
181 return;
182 }
183
184 let start_x = position.x + fir_len / 2 + upscale_factor * BACKDROP_PIXELS;
185
186 var filtered_y = vec4f(0.0);
187 var filtered_u = vec4f(0.0);
188 var filtered_v = vec4f(0.0);
189 for (var i = 0; i < fir_len; i += 4) {
190 let luma_samples = vec4f(
191 textureLoad(ntsc_luma_r, vec2i(start_x - i, position.y), 0).r,
192 textureLoad(ntsc_luma_r, vec2i(start_x - i - 1, position.y), 0).r,
193 textureLoad(ntsc_luma_r, vec2i(start_x - i - 2, position.y), 0).r,
194 textureLoad(ntsc_luma_r, vec2i(start_x - i - 3, position.y), 0).r,
195 );
196
197 let chroma_samples = vec4f(
198 textureLoad(ntsc_chroma_r, vec2i(start_x - i, position.y), 0).r,
199 textureLoad(ntsc_chroma_r, vec2i(start_x - i - 1, position.y), 0).r,
200 textureLoad(ntsc_chroma_r, vec2i(start_x - i - 2, position.y), 0).r,
201 textureLoad(ntsc_chroma_r, vec2i(start_x - i - 3, position.y), 0).r,
202 );
203
204 let y_coefficients = y_decode_lpf_coefficients[i / 4];
205 let uv_coefficients = uv_decode_lpf_coefficients[i / 4];
206
207 let base_phases = vec4i(start_x - i, start_x - i - 1, start_x - i - 2, start_x - i - 3)
208 + immediate_params.frame_phase_offset
209 + position.y * immediate_params.per_line_phase_offset;
210 let phases = vec4f(base_phases) / f32(samples_per_color_cycle) * 2.0 * PI
211 + vec4f(decode_hue_offset);
212
213 // 2.0 multiplier in U/V for chroma saturation correction:
214 // https://www.nesdev.org/wiki/NTSC_video#Chroma_saturation_correction
215 filtered_y = fma(y_coefficients, luma_samples, filtered_y);
216 filtered_u = fma(uv_coefficients, chroma_samples * sin(phases) * 2.0, filtered_u);
217 filtered_v = fma(uv_coefficients, chroma_samples * cos(phases) * 2.0, filtered_v);
218 }
219
220 var yuv = vec3f(
221 dot(filtered_y, vec4f(1.0)),
222 dot(filtered_u, vec4f(1.0)),
223 dot(filtered_v, vec4f(1.0)),
224 );
225
226 yuv *= vec3f(decode_brightness) * vec3f(1.0, decode_saturation, decode_saturation);
227
228 var rgb = YUV_TO_RGB * yuv;
229 rgb = clamp(rgb, vec3f(0.0), vec3f(1.0));
230 rgb = pow(rgb, vec3f(2.2 / decode_gamma));
231
232 textureStore(output_frame, position, vec4f(rgb, 1.0));
233}
234
235// NES-to-NTSC based on https://www.nesdev.org/wiki/NTSC_video
236const NES_NTSC_BLACK: f32 = 0.312;
237const NES_NTSC_WHITE: f32 = 1.100;
238
239const NES_NTSC_LOW: vec4f = vec4f(0.228, 0.312, 0.552, 0.880);
240const NES_NTSC_HIGH: vec4f = vec4f(0.616, 0.840, 1.100, 1.100);
241const NES_NTSC_LOW_ATTENUATED: vec4f = vec4f(0.192, 0.256, 0.448, 0.712);
242const NES_NTSC_HIGH_ATTENUATED: vec4f = vec4f(0.500, 0.676, 0.896, 0.896);
243
244const NES_COLOR_BLACK: f32 = f32(0x1D) / 255.0;
245
246@compute @workgroup_size(16, 16, 1)
247fn nes_to_ntsc(@builtin(global_invocation_id) invocation: vec3u) {
248 let frame_size = vec2i(textureDimensions(ntsc_frame_w));
249 let position = vec2i(invocation.xy);
250 if position.x >= frame_size.x || position.y >= frame_size.y {
251 return;
252 }
253
254 let phase = immediate_params.frame_phase_offset
255 + position.y * immediate_params.per_line_phase_offset
256 + position.x;
257
258 let input_size = vec2i(textureDimensions(input_frame));
259
260 // Assume input frame buffer contains 6-bit NES colors (R) and 3-bit color emphasis (G) instead of RGB888 colors
261 let input_x = (position.x - upscale_factor * BACKDROP_PIXELS) / upscale_factor;
262 let input_texel = load_input_bounds_checked(vec2i(input_x, position.y), input_size, vec3f(NES_COLOR_BLACK, 0.0, 0.0));
263 let input_rg = vec2i(round(input_texel.rg * 255.0));
264 let nes_color = input_rg.r;
265 let color_emphasis = input_rg.g;
266
267 // Hue is lowest 4 bits of 6-bit color
268 let hue = nes_color & 0xF;
269
270 // Luma is forced to 1 when hue is 14 or 15
271 // Otherwise, highest 2 bits of 6-bit color
272 let luma = select(
273 1,
274 (nes_color >> 4) & 3,
275 hue < 0xE,
276 );
277
278 let emphasis_r = (color_emphasis & (1 << 0)) != 0;
279 let emphasis_g = (color_emphasis & (1 << 1)) != 0;
280 let emphasis_b = (color_emphasis & (1 << 2)) != 0;
281
282 // Color emphasis bits cause the PPU to attenuate half of the signal, or more than half if multiple bits are set
283 // Emphasis bits have no effect when hue is 14 or 15
284 let attenuate = hue < 0xE
285 && ((emphasis_r && nes_in_color_phase(0, phase))
286 || (emphasis_g && nes_in_color_phase(4, phase))
287 || (emphasis_b && nes_in_color_phase(8, phase)));
288
289 // Luma determines the two possible NTSC sample values
290 let low = select(NES_NTSC_LOW[luma], NES_NTSC_LOW_ATTENUATED[luma], attenuate);
291 let high = select(NES_NTSC_HIGH[luma], NES_NTSC_HIGH_ATTENUATED[luma], attenuate);
292
293 // NTSC signal is always high when hue is 0 and always low when hue is 13-15
294 // Otherwise the PPU outputs a square wave, phase shifted based on hue
295 let signal = select(
296 select(
297 select(
298 low,
299 high,
300 nes_in_color_phase(hue, phase),
301 ),
302 low,
303 hue >= 13,
304 ),
305 high,
306 hue == 0,
307 );
308
309 // Normalize so black=0 and white=1, with negative values (darker than black) possible and allowed
310 let normalized = (signal - NES_NTSC_BLACK) / (NES_NTSC_WHITE - NES_NTSC_BLACK);
311
312 textureStore(ntsc_frame_w, position, vec4f(normalized, vec3f(0.0)));
313}
314
315fn nes_in_color_phase(color: i32, phase: i32) -> bool {
316 return ((color + phase) % 12) < 6;
317}
ntsc.wgsl