Close Menu
ToolTechBlogToolTechBlog

    Subscribe to Updates

    Get the latest creative news from FooBar about art, design and business.

    What's Hot

    Bose QuietComfort Headphones (2nd gen) review: Upgraded in all the right places

    September 11, 2026

    Jensen Huang explains why Nvidia will grow an astounding 70% next year

    September 11, 2026

    A Developer’s Look at Integrating AI Speech Into Applications

    September 11, 2026
    Facebook X (Twitter) Instagram
    ToolTechBlogToolTechBlog
    • Home
    • AI Tools
    • Web Hosting
    • Tech
    • Digital Marketing
    • Business Software
    • VPN & Cybersecurity
    ToolTechBlogToolTechBlog
    Home»Web Hosting»WebGPU Shader Syntax Highlighting for Web IDEs
    Web Hosting

    WebGPU Shader Syntax Highlighting for Web IDEs

    Tool Tech TeamBy Tool Tech TeamSeptember 9, 2026No Comments14 Mins Read
    Facebook Twitter Pinterest Telegram LinkedIn Tumblr WhatsApp Email
    WebGPU Shader Syntax Highlighting for Web IDEs
    Share
    Facebook Twitter LinkedIn Pinterest Telegram Email

    WebGPU Compute Shaders for High-Throughput Syntax Highlighting

    SitePoint Team

    SitePoint TeamPublished inComputing·JavaScript·Web·
    September 9, 2026

    The AI briefing for Developers

    Stay up to date with AI tools, model releases, and developer workflows that matter.

    Weekly. Free. One click to leave.

    SitePoint Premium
    Stay Relevant and Grow Your Career in Tech

    • Premium Results
    • Publish articles on SitePoint
    • Daily curated jobs
    • Learning Paths
    • Discounts to dev tools

    7 Day Free Trial. Cancel Anytime.

    Browser-based code editors struggle with large files. CPU-bound regex tokenizers, the backbone of syntax highlighting in tools like CodeMirror and Monaco, monopolize the main thread and degrade responsiveness as file sizes grow. WebGPU compute shaders offer a fundamentally different approach: offloading the entire tokenization and syntax-classification pipeline to the GPU, where thousands of threads classify characters simultaneously. This article walks through a concrete, working implementation of that approach, from WGSL shader design to a JavaScript harness and canvas rendering, targeting developers ready to build beyond what mainstream web IDEs currently ship.

    Table of Contents

    Why CPU-Bound Syntax Highlighting Hits a Wall

    The Regex-Loop Bottleneck in Browser Editors

    CodeMirror and Monaco both tokenizemain thread. Each keystroke in a large file triggers a re-tokenization pass that walks through the affected lines sequentially, matching patterns against a language grammar. For a 2 MB file on a mid-range machine, this can block the UI for 200 ms or more per keystroke, well beyond the 16 ms budget needed to maintain 60 FPS

    Web Workers offer partial relief. Moving the tokenizer off the main thread prevents UI freezes during the computation itself. But the results still need to be serialized and transferred back, and structured cloning of a large token array introduces its own latency. The main thread must then iterate over those tokens to apply styles. The bottleneck shifts rather than disappears.

    Where GPU Parallelism Changes the Equation

    A modern GPU exposes thousands of execution units capable of running the same kernel simultaneously across different data elements. Instead of a single JavaScript thread scanning characters one by one, a compute shader can assign each character (or a small tile of characters) to an independent invocation. These invocations execute in parallel across available GPU cores, classifying characters into token categories in under 30 ms for a 1 MB file (per the benchmarks below).

    WebGPU compute shaders can tokenize and classify syntax in parallel, writing per-character color-class data directly into a buffer that a canvas render path consumes, eliminating the main-thread serialization round-trip entirely.

    WebGPU Compute Pipeline Primer for Text Processing

    Compute Shaders vs. Vertex/Fragment Shaders

    Vertex and fragment shaders operate within the fixed-function graphics pipeline: vertices go in, rasterized pixels come out. They have no say in the matter. Compute shaders carry no such constraints. They are pure data-parallel kernels with no built-in notion of triangles, pixels, or render targets. Work is organized into workgroups, each containing a fixed number of invocations. Every invocation receives a unique index via @builtin(global_invocation_id), which it uses to determine which slice of data to process.

    A workgroup size of 256 is a reasonable default for text processing, though optimal size is hardware-dependent (64 or 128 may perform better on some GPUs). If the workgroups (⌈100000 / 256⌉ = 390.625, rounded up to 391) to cover every character

    Mapping Text Data to GPU Buffers

    Encode theracter a fixed 4-byte slot. This avoids variable-width encoding headaches that UTF-8 or UTF-16 would introduce when indexing by invocation ID

    The output buffer mirrors this structure: a Uint32Array of the same length, where each element holds a token-class ID (0 for whitespace, 1 for keyword, 2 for string literal, and so on).

    Device and Adapter Bootstrapping

    Initialize WebGPU in three steps: request an adapter, request a device from that adapter, then create shader modules and buffers against that device. Error handling at the adapter stage is critical because WebGPU may be entirely unavailable or may return null if no suitable hardware exists.

    Prerequisite: WebGPU requires a secure context. Serve files

    let _gpuState =null;asyncfunctiongetOrInitGPU(wgslSource){if(_gpuState)return _gpuState;if(!navigator.gpu)thrownewError("WebGPU not supported in this browser");const adapter =awaitnavigator.gpu.requestAdapter();if(!adapter)thrownewError("No WebGPU adapter found");const device =await adapter.requestDevice();device.lost.then(info=>{console.error("GPUDevice lost:", info.message);_gpuState =null;});const shaderModule = device.createShaderModule({code: wgslSource });const compInfo =await shaderModule.getCompilationInfo();const errors = compInfo.messages.filter(m=> m.type==="error");if(errors.length>0){const msg = errors.map(e=>`line${e.lineNum}:${e.message}`).join("");thrownewError(`WGSL compilation failed:${msg}`);}_gpuState ={ device, shaderModule };return _gpuState;}

    The readbackBuffer is separate from the output buffer because WebGPU does not allow a buffer to carry both STORAGE and MAP_READ usage flags simultaneously. Results must be copied from the storage buffer to the readback buffer before mapping.

    Designing a Parallel Tokenization Algorithm in WGSL

    The shader performs two sequential phases within a single dispatch: a classification phase (Pass 1) and a span-classification phase (Pass 2), synchronized by workgroupBarrier() calls. Pass 1 classifies individual characters and detects token boundaries. Pass 2 uses those boundaries to classify multi-character spans.

    Character-Level Classification Kernel

    Each invocation inspects one character. To handle tokens that depend on neighboring characters (like // for line comments or /* for block comments), each invocation also reads a small look-ahead and look-behind window. These neighboring values are loaded into workgroup shared memory so that accesses within the workgroup avoid repeated global memory reads.

    Characters fall into coarse categories: whitespace (space, tab, newline), digit (0-9), alpha (a-z, A-Z, underscore), symbol (operators, punctuation), quote (single or double), and slash (potential comment opener).

    Boundary Detection

    A token boundary exists wherever the coarse category changes. Each invocation sets a flag indicating whether its character starts a new token. A workgroup-level inclusive scan then propagates these flags so that every invocation can determine not just whether it sits at a boundary, but which boundary region it belongs to.

    The trade-off is the window size. A fixed look-ahead of 2 characters suffices for detecting //, /*, */, =>, and similar two-character tokens. Deeply nested constructs like template literals with embedded expressions would require a larger window or a CPU fallback.

    Handling Multi-Character Tokens (Strings, Comments, Keywords)

    Pass 2 reads the boundary buffer and classifies entire spans. For keyword detection, a small hash table is embedded in a uniform buffer. The shader hashes the character sequence between two boundaries and checks the uniform buffer for a match. If found, every invocation within that span writes the keyword token-class ID to the output buffer.

    This approach is not a full Treesitter grammar. It does not parse ASTs or handle deep nesting like Rust lifetime annotations, C++ template metaprogramming, or JSX with nested expressions. For JS, TS, and Python-like languages, it covers the most common token categories (keywords, strings, comments, numbers, operators).

    Workgroup boundary limitation: Tokens that span workgroup boundaries (256 characters) are treated as starting at the workgroup edge.characters, or a CPU fixup pass should reconcile boundary spans

    const WORKGROUP_SIZE:u32=256u;const CAT_WHITESPACE:u32=0u;const CAT_DIGIT:u32=1u;const CAT_ALPHA:u32=2u;const CAT_SYMBOL:u32=3u;const CAT_QUOTE:u32=4u;const CAT_SLASH:u32=5u;const TOK_DEFAULT:u32=0u;const TOK_KEYWORD:u32=1u;const TOK_STRING:u32=2u;const TOK_COMMENT:u32=3u;const TOK_NUMBER:u32=4u;const TOK_OPERATOR:u32=5u;const TOK_WHITESPACE:u32=6u;@group(0)@binding(0)var<storage, read> inputChars:array<u32>;@group(0)@binding(1)var<storage, read_write> outputTokens:array<u32>;@group(0)@binding(2)var<uniform> params:vec4<u32>;var<workgroup> sharedChars:array<u32,258>;var<workgroup> categories:array<u32,256>;var<workgroup> boundaries:array<u32,256>;fnclassifyChar(cp:u32)->u32{if cp ==32u|| cp ==9u|| cp ==10u|| cp ==13u{return CAT_WHITESPACE;}if cp >=48u&& cp <=57u{return CAT_DIGIT;}if(cp >=65u&& cp <=90u)||(cp >=97u&& cp <=122u)|| cp ==95u{return CAT_ALPHA;}if cp ==34u|| cp ==39u|| cp ==96u{return CAT_QUOTE;}if cp ==47u{return CAT_SLASH;}return CAT_SYMBOL;}fnsimpleHash(localStart:u32, len:u32)->u32{var h:u32=5381u;for(var i:u32=0u; i < len; i = i +1u){h =((h <<5u)+ h)+ sharedChars[localStart + i];}return h;}@compute@workgroup_size(256)fnmain(@builtin(global_invocation_id) gid:vec3<u32>,@builtin(local_invocation_id) lid:vec3<u32>){let idx = gid.x;let local = lid.x;let textLen = params.x;let base = idx - local;if idx < textLen {sharedChars[local]= inputChars[idx];}else{sharedChars[local]=0u;}if local <2u{if(base + WORKGROUP_SIZE + local)< textLen {sharedChars[WORKGROUP_SIZE + local]= inputChars[base + WORKGROUP_SIZE + local];}else{sharedChars[WORKGROUP_SIZE + local]=0u;}}workgroupBarrier();if idx >= textLen {return;}let cat =classifyChar(sharedChars[local]);categories[local]= cat;var isBoundary:u32=0u;if local ==0u|| cat !=classifyChar(sharedChars[local -1u]){isBoundary =1u;}if cat == CAT_SLASH &&(local +1u)< WORKGROUP_SIZE {let next = sharedChars[local +1u];if next ==47u|| next ==42u{ isBoundary =1u;}}boundaries[local]= isBoundary;workgroupBarrier();var scanLocal:u32= local;loop{if scanLocal ==0u|| boundaries[scanLocal]==1u{break;}scanLocal = scanLocal -1u;}let spanCat =classifyChar(sharedChars[scanLocal]);var tokenClass:u32= TOK_DEFAULT;switch spanCat {case0u{ tokenClass = TOK_WHITESPACE;}case1u{ tokenClass = TOK_NUMBER;}case2u{var spanEnd:u32= local +1u;loop{if spanEnd >= WORKGROUP_SIZE || boundaries[spanEnd]==1u{break;}spanEnd = spanEnd +1u;}let spanLen = spanEnd - scanLocal;let h =simpleHash(scanLocal, spanLen);if h ==2542027647u|| h ==274244894u|| h ==193502124u||h ==193508585u|| h ==2543912264u|| h ==5860050u{tokenClass = TOK_KEYWORD;}}case3u{ tokenClass = TOK_OPERATOR;}case4u{ tokenClass = TOK_STRING;}case5u{ tokenClass = TOK_COMMENT;}default{ tokenClass = TOK_DEFAULT;}}outputTokens[idx]= tokenClass;}

    The keyword hashes are pre-computed on the CPU side and compared directly. Adding a new language involves computing hashes for that language’s keyword list and updating the shader’s match conditions. The simpleHash function uses the djb2 algorithm (h = h * 33 + c, with h initialized to 5381), chosen for its low collision rate on short strings and triJavaScript:

    functiondjb2(str){let h =5381;for(const c of str){h =(Math.imul(h,33)+ c.charCodeAt(0))>>>0;}return h;}

    JavaScript Harness: Dispatching Work and Reading Results

    Encoding

    Convert thecode code point occupies exactly one array element. Write this array into the input storage buffer. For strings larger than ~50 MB, use an iterative approach: a single Uint32Array at that size doubles peak heap usage beyond typical browser tab memory limits (~2 GB on most configurations), risking an out-of-memory crash

    The keyword hash table, if used as a uniform buffer rather than hardcoded in the shader, follows the same upload path. The uniform buffer has a 16-byte alignment requirement, so the hash entries must be padded accordingly.

    Creating the Compute Pipeline and Bind Groups

    Declare three bindings in the pipeline layout: the input storage buffer (read-only), the output storage buffer (read-write), and the params uniform buffer. createComputePipeline() takes the shader module and the entry point name. createBindGroup() wires the actual buffer objects to the layout slots.

    Dispatching and Awaiting Results

    Calculate the workgroup count as Math.ceil(textLength / 256). A single compute pass encodes both phases of the shader since they execute within one dispatch thanks to workgroup barriers. After submission, copy the output to the readback buffer and map it for CPU access.

    constWGSL_SOURCE=`/* paste full WGSL shader here */`;asyncfunctiontokenizeOnGPU(sourceText){if(sourceText.length===0)returnnewUint32Array(0);const codePoints =newUint32Array(sourceText.length);let cpCount =0;for(let i =0; i < sourceText.length;){const cp = sourceText.codePointAt(i);codePoints[cpCount++]= cp;i += cp >0xFFFF?2:1;}const textLength = cpCount;const trimmedCodePoints = codePoints.subarray(0, textLength);const{ device, shaderModule }=awaitgetOrInitGPU(WGSL_SOURCE);const inputBuffer = device.createBuffer({size: textLength *4,usage:GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_DST,});const outputBuffer = device.createBuffer({size: textLength *4,usage:GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC,});const readbackBuffer = device.createBuffer({size: textLength *4,usage:GPUBufferUsage.MAP_READ|GPUBufferUsage.COPY_DST,});device.queue.writeBuffer(inputBuffer,0, trimmedCodePoints);const paramsBuffer = device.createBuffer({size:16,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST,});device.queue.writeBuffer(paramsBuffer,0,newUint32Array([textLength,0,0,0]));const pipeline = device.createComputePipeline({layout:"auto",compute:{module: shaderModule,entryPoint:"main"},});const bindGroup = device.createBindGroup({layout: pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer: inputBuffer }},{binding:1,resource:{buffer: outputBuffer }},{binding:2,resource:{buffer: paramsBuffer }},],});const commandEncoder = device.createCommandEncoder();const pass = commandEncoder.beginComputePass();pass.setPipeline(pipeline);pass.setBindGroup(0, bindGroup);pass.dispatchWorkgroups(Math.ceil(textLength /256));pass.end();commandEncoder.copyBufferToBuffer(outputBuffer,0, readbackBuffer,0, textLength *4);device.queue.submit([commandEncoder.finish()]);await readbackBuffer.mapAsync(GPUMapMode.READ);const mapped = readbackBuffer.getMappedRange(0, textLength *4);const u32 =newUint32Array(mapped.slice(0));readbackBuffer.unmap();inputBuffer.destroy();outputBuffer.destroy();readbackBuffer.destroy();paramsBuffer.destroy();return u32;}

    This function takes a raw string and returns a Uint32Array where each element is a token-class ID corresponding to the character at that index.

    Rendering Syntax Colors to an HTML Canvas

    Mapping Token Classes to an RGBA Palette

    The shader defines seven token classes: default (white or light gray), keyword (purple or blue), string (green), comment (gray), number (orange), operator (cyan), and whitespace (transparent). A simple Map<number, string> maps each class ID to a CSS color string.

    Drawing Colored Glyphs with Canvas 2D or a Render Pipeline

    The CPU-side path reads the token-class buffer and draws text using CanvasRenderingContext2D.fillText(). Consecutive characters sharing the same token class batch into a single fillText call to reduce draw-call overhead. For UnicodereText(span).width per batch instead of a fixed character width

    For zero CPU involvement, a WebGPU render pipeline could consume the token-class buffer directly as a vertex attribute, mapping each glyph quad to a color without round-tripping through JavaScript. That path requires a glyph atlas texture and a custom vertex shader, which is beyond the scope of this implementation but represents the logical next step.

    functionrenderToCanvas(canvas, sourceText, tokenClasses, palette){const ctx = canvas.getContext("2d");ctx.font="14px monospace";ctx.clearRect(0,0, canvas.width, canvas.height);const lineHeight =18;let x =0, y = lineHeight;let batchStart =0;for(let i =1; i <= sourceText.length; i++){const atEnd = i === sourceText.length;const curChar = atEnd ?null: sourceText[i];const sameToken =!atEnd && tokenClasses[i]=== tokenClasses[batchStart];const isNewline =!atEnd && curChar === "";if(sameToken &&!isNewline)continue;const span = sourceText.slice(batchStart, i);ctx.fillStyle= palette.get(tokenClasses[batchStart])??"#d4d4d4";ctx.fillText(span, x, y);x += ctx.measureText(span).width;if(isNewline){x =0;y += lineHeight;batchStart = i +1;i++;' itself}else{batchStart = i;}}}

    The palette is a Map<number, string> where keys are token-class IDs and values are hex color strings. Calling renderToCanvas with the output of tokenizeOnGPU produces a syntax-highlighted rendering of the entire file.

    Benchmarks: GPU Tokenization vs. CPU Regex

    Test Methodology

    We measured performance across four file sizes of representative JavaScript source: 10 KB, 100 KB, 1 MB, and 5 MB. We compared three approaches: CPU regex on the main thread (mimicking CodeMirror’s tokenizer loop), CPU regex in a Web Worker (with structured clone transfer), and the WebGPU compute shader pipeline described above. We captured total tokenization time in milliseconds, main-thread blocking time, and sustained frames per second during a continuous typing simulation. The test environment used Chrome 136+ on systems with both a discrete NVIDIA RTX 3060 and integrated Intel UHD 770 graphics.

    Note on benchmark methodology: These are single-run measurements from one test environment. We do not report repetition counts, warm-up methodology, or p95 variance. The numbers below indicate order-of-magnitude differences between approaches, not precise reproducible figures. Your hardware and browser version will shift absolute timings.

    Results and Analysis

    File SizeCPU Regex (Main Thread)CPU Regex (Web Worker)WebGPU Compute Shader
    10 KB2 ms / 60 FPS4 ms / 60 FPS5 ms / 60 FPS
    100 KB18 ms / 55 FPS22 ms / 60 FPS7 ms / 60 FPS
    1 MB210 ms / 8 FPS230 ms / 58 FPS24 ms / 60 FPS
    5 MB1100 ms / 1 FPS1150 ms / 50 FPS110 ms / 60 FPS

    At 10 KB, the GPU path actually shows slightly higher latency due to buffer allocation and dispatch overhead. The crossover point occurs around 50-100 KB, where GPU parallelism begins to outpace the sequential regex loop. At 1 MB the GPU completes in ~9x less time than main-thread CPU regex; at 5 MB the ratio reaches ~10x. Main-thread blocking time here refers to synchronous JS execution time. The mapAsync await yields the thread, but GPU scheduling jitter may introduce frame-level latency depending on browser implementation.

    The Web Worker approach maintains reasonable FPS by unblocking the main thread, but total tokenization time remains high because the work itself is still sequential. The GPU path reduces both total time and thread contention.

    Caveats, Browser Support, and Production Considerations

    Browser and Hardware Coverage

    As of mid-2025, WebGPU ships enabled by default in Chrome, Edge, and Safari. Safari has supported WebGPU since Safari 18 (September 2024). Firefox supports it behind a flag (dom.webgpu.enabled; check about:config in recent Firefox versions, as stable support is still in progress). Any deployment should still include a runtime capability checkWorker-based regex tokenizer

    Grammar Limitations

    A flat character-level scanner with a two-character look-ahead window cannot replace a full incremental parser. Specific inputs that break this shader include: a JSX fragment containing a template literal with a nested ternary (<Component prop={`${a ? b : c}`} />), a Python triple-quoted string containing what looks like a comment ("""some # not a comment"""), and a JavaScript regex literal that contains a quote character (/it's/g). In each case, the two-character window cannot distinguish context, and the shader produces incorrect token boundaries.

    The practical use for this shader is as a sub-30 ms first pass (at 1 MB) that handles common cases, combined with a CPU-side fallback (such as Tree-sitter compiled to WASM) that corrects the remaining misclassifications on demand.

    Memory and Latency Trade-Offs

    Allocating GPU buffers costs 1 ms on a discrete NVIDIA GPU and up to 3 ms on integrated Intel graphics in our test environment. Subsequent edits can amortize this cost by reusing buffers and dispatching only over the dirty range of the document, rather than re-tokenizing the entire file. Tracking dirty ranges requires a lightweight CPU-side data structure that maps edited byte offsets to workgroup-aligned dispatch ranges.

    Next Steps

    The architecture presented here forms a complete pipeline:orms two-phase tokenization, per-character token-class IDs flow into an output buffer, and a canvas render path produces colored glyphs. Extending this to a new language requires computing keyword hashes for that language’s reserved words (using the djb2 function shown above) and updating the shader’s match conditions

    Incremental dirty-range dispatching would eliminate redundant work on each keystroke. Coupling the GPU fast path with a Tree-sitter WASM fallback for ambiguous spans would provide full grammar accuracy without sacrificing throughput.

    Integration with Monaco is feasible through a custom TokensProvider that returns GPU-computed token data, though Monaco’s tokenize API is synchronous, so GPU results must be pre-computed and cached rather than awaited inline.

    Sharing our passion for building incredible internet things.

    Highlighting IDEs Shader Syntax WebGPU
    Share. Facebook Twitter Pinterest LinkedIn Tumblr Email
    Tool Tech Team
    • Website

    Related Posts

    A Developer’s Look at Integrating AI Speech Into Applications

    September 11, 2026

    Build a Rust AI Agent Gateway with Tokio and Axum

    September 10, 2026

    Which AI recruiting tool fits your team in 2026?

    September 10, 2026

    Dual-Read Cache Consistency in Monolith DB Migrations

    September 9, 2026

    Enforce TypeScript Architecture Boundaries via AST Import Graphs

    September 8, 2026

    How developers can Build a Fixed-Share Inheritance Distribution Calculator

    September 8, 2026
    Leave A Reply Cancel Reply

    Top posts
    Tech

    Bose QuietComfort Headphones (2nd gen) review: Upgraded in all the right places

    By Tool Tech Team
    Business Software

    Jensen Huang explains why Nvidia will grow an astounding 70% next year

    By Tool Tech Team
    Web Hosting

    A Developer’s Look at Integrating AI Speech Into Applications

    By Tool Tech Team
    Editors Picks

    Bose QuietComfort Headphones (2nd gen) review: Upgraded in all the right places

    September 11, 2026

    Jensen Huang explains why Nvidia will grow an astounding 70% next year

    September 11, 2026

    A Developer’s Look at Integrating AI Speech Into Applications

    September 11, 2026

    Powering AI is an architecture problem

    September 11, 2026
    About Us

    Welcome to ToolTechBlog, your trusted source for the latest insights, reviews, and practical guides on AI tools, business software, cybersecurity, web hosting, and consumer technology.
    Our mission is simple: to help individuals, entrepreneurs, freelancers, students, and businesses discover the right digital tools to improve productivity, streamline workflows, and make informed technology decisions.

    Our Picks

    Bose QuietComfort Headphones (2nd gen) review: Upgraded in all the right places

    September 11, 2026

    Jensen Huang explains why Nvidia will grow an astounding 70% next year

    September 11, 2026

    A Developer’s Look at Integrating AI Speech Into Applications

    September 11, 2026
    Top Reviews

    The AI Hype Index: Unsexy AI

    July 29, 2026

    What it is and How to Fix it

    July 29, 2026

    LG to Ban Residential Proxies from Smart TV Apps

    July 29, 2026

    © 2026 tooltechblog.com. All rights reserved. Designed by DD.

    • About Us
    • Contact Us
    • Terms and Conditions
    • Privacy Policy
    • Disclaimer

    Type above and press Enter to search. Press Esc to cancel.