diff --git a/grammar.py b/grammar.py index c03d74f..d0226a6 100755 --- a/grammar.py +++ b/grammar.py @@ -970,7 +970,7 @@ def parse_from_string(self, grammar_str): return 0 - def parse_from_file(self, filename): + def parse_from_file(self, filename, extra=None): """Parses grammar from file. Opens a text file, parses it and loads the grammar rules within. @@ -991,6 +991,10 @@ def parse_from_file(self, filename): print('Error reading ' + filename) return 1 self._definitions_dir = os.path.dirname(filename) + + if extra: + content = extra + content + return self.parse_from_string(content) def _compute_interesting_indices(self): diff --git a/webgpu/README.md b/webgpu/README.md new file mode 100644 index 0000000..532f77d --- /dev/null +++ b/webgpu/README.md @@ -0,0 +1,12 @@ +# Simple WebGPU fuzzer + +## Setup + +1. Populate the `wgsl/` directory with wgsl scripts. I would recomment copying wgsl test files from [tint's tests](https://source.chromium.org/chromium/chromium/src/+/main:third_party/dawn/test/tint/bug/). +2. From here, the usage is the same as [vanilla Domato's](https://github.com/googleprojectzero/domato). + +## Bugs +Chrome: [40063883](https://issues.chromium.org/u/0/issues/40063883), [40063356](https://issues.chromium.org/u/0/issues/40063356) + +## Building Grammars +This repo also contains a helper script that can be used to assist in generating Domato grammars using Chrome's [WebIDL compiler](https://source.chromium.org/chromium/chromium/src/+/main:third_party/blink/renderer/bindings/scripts/web_idl/README.md;l=1?q=f:md%20web_idl&sq=). It is far from complete but may help others generate grammars faster. diff --git a/webgpu/__init__.py b/webgpu/__init__.py deleted file mode 100644 index e69de29..0000000 diff --git a/webgpu/build_grammar.py b/webgpu/build_grammar.py new file mode 100644 index 0000000..eacf782 --- /dev/null +++ b/webgpu/build_grammar.py @@ -0,0 +1,150 @@ +# Domato - example generator script +# -------------------------------------- +# +# Written by Brendon Tiszka +# +# Copyright 2017 Google Inc. All Rights Reserved. +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +import web_idl + +# Path to your chromium build directory +PATH_TO_CHROME_BUILD_DIR = "" +web_idl_database_path = '/gen/third_party/blink/renderer/bindings/web_idl_database.pickle' +web_idl_database = web_idl.Database.read_from_file(PATH_TO_CHROME_BUILD_DIR + web_idl_database_path) + +def is_gpu(ident): + if "GPU" in ident: + return True + return False + +def is_promise(ident): + if "Promise" in ident: + return True + return False + +def remove_promise_info(s): + if "OrNullPromise" in s: + return s[:s.index("OrNullPromise")] + + if "Promise" in s: + return s[:s.index("Promise")] + + return s + +def parse_enums(): + builder = "" + + for enum in web_idl_database.enumerations: + if not is_gpu(enum.identifier): + continue + + for value in enum.values: + builder += "<{}> = \"{}\"\n".format(enum.identifier, value) + + builder += "\n" + + return builder + +def parse_namespaces(): + builder = "" + + for ns in web_idl_database.namespaces: + if not is_gpu(ns.identifier): + continue + + for const in ns.constants: + builder += "<{}> = {}.{}\n".format(ns.identifier, ns.identifier, const.identifier) + + builder += "\n" + + return builder + +def parse_dictionaries(): + builder = "" + + for dictionary in web_idl_database.dictionaries: + if not is_gpu(dictionary.identifier): + continue + + print(dictionary.identifier) + +def parse_interfaces(): + builder = "" + + for interface in web_idl_database.interfaces: + if not is_gpu(interface.identifier): + continue + + + builder += "#" + ("~"*16) + interface.identifier + ("~"*16) + "#\n" + for attribute in interface.attributes: + if attribute.is_readonly: + continue + + builder += "<{}>.{} = <{}>".format(interface.identifier, attribute.identifier, attribute.idl_type.type_name) + builder += "\n" + + for operation in interface.operations: + required_args = operation.num_of_required_arguments + num_args = len(operation.arguments) + + lifted_args = [] + for argument in operation.arguments: + lifted_args.append("<{}>".format(argument.idl_type.type_name)) + + for i in range(required_args, num_args + 1): + if operation.return_type.type_name != "Void": + return_type = remove_promise_info(operation.return_type.type_name) + builder += " = ".format(return_type) + + if is_promise(operation.return_type.type_name): + builder += "await " + + builder += "<{}>.{}(".format(interface.identifier, operation.identifier) + builder += "{}".format(",".join(lifted_args[:i])) + builder += ");\n" + + builder += "\n" + return builder + + +def parse_dictionaries(): + builder = "" + + for dictionary in web_idl_database.dictionaries: + if not is_gpu(dictionary.identifier): + continue + + + lifted_members = [] + for member in dictionary.members: + lifted_members.append("{}: <{}>".format(member.identifier, member.idl_type.type_name)) + + builder += "<{}> = ".format(dictionary.identifier) + builder += "{ " + builder += ", ".join(lifted_members) + builder += " };" + builder += "\n" + + return builder + +if __name__ == "__main__": + enums = parse_enums() + namespaces = parse_namespaces() + interfaces = parse_interfaces() + dictionaries = parse_dictionaries() + print(enums) + print(namespaces) + print(interfaces) + print(dictionaries) diff --git a/webgpu/generator.py b/webgpu/generator.py index fd14ab6..95657a3 100644 --- a/webgpu/generator.py +++ b/webgpu/generator.py @@ -1,9 +1,10 @@ # Domato - generator script for WebGPU # ------------------------------- # -# Written and maintained by Chrome Security Team +# Written and maintained by Ivan Fratric +# Modified by Brendon Tiszka to target webgpu # -# Copyright 2022 Google Inc. All Rights Reserved. +# Copyright 2017 Google Inc. All Rights Reserved. # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. # You may obtain a copy of the License at @@ -16,90 +17,113 @@ # See the License for the specific language governing permissions and # limitations under the License. - from __future__ import print_function +import glob import os -import re import random +import re import sys parent_dir = os.path.abspath(os.path.join(os.path.dirname(__file__), os.pardir)) sys.path.append(parent_dir) from grammar import Grammar -_N_MAIN_LINES = 100 -_N_EVENTHANDLER_LINES = 1 +_N_MAIN_LINES = 1000 +_N_SHADERS = 10 + +def extract_shader_stages_and_functions(code): + # Pattern to match both single-line and multiline stage attributes + pattern = r'@(?:compute|vertex|fragment)(?:\s+@[^(\n]+(?:\([^)]*\))?)*\s*(?:\n\s*)?fn\s+(\w+)' + matches = re.finditer(pattern, code) + + # Extract stage-function pairs + result = [] + for match in matches: + # Find the stage attribute within the matched attributes + stage_attr = re.search(r'@(compute|vertex|fragment)', match.group(0)).group(0) + function_name = match.group(1) + result.append((stage_attr, function_name)) + + return result + +def parse_entrypoints(shaders): + entrypoints = [] + for shader in shaders: + attr_fn_pairs = extract_shader_stages_and_functions(shader) + for attr, fn in attr_fn_pairs: + entrypoints.append(fn) -def generate_lines(grammar, num_lines): - output = '' - output += grammar._generate_code(num_lines) + result = "" + for entrypoint in entrypoints: + result += " = \"{}\"\n".format(entrypoint) - return output + return result -def GenerateNewSample(template, jsgrammar, wgslgrammar): - """Parses grammar rules from string. +def parse_bindings(shaders): + binding_pattern = r'@binding\((\d+)\)' + + binding_numbers = [] + for shader in shaders: + matches = re.finditer(binding_pattern, shader) + binding_numbers.extend(match.group(1) for match in matches) + + return "\n".join(f" = {binding}" for binding in binding_numbers) - Args: - template: A template string. - jsgrammar: Grammar for generating JS code. - wgslgrammar: Grammar for generating WGSL shader code. +def generate_function_body(webgpugrammar, num_lines): + js = '' + js += webgpugrammar._generate_code(num_lines) - Returns: - A string containing sample data. - """ + return js +def generate_new_sample(template, webgpugrammar): result = template - handlers = False - - # WebGPU JS while '' in result: - numlines = _N_MAIN_LINES - if handlers: - numlines = _N_EVENTHANDLER_LINES - else: - handlers = True result = result.replace( '', - generate_lines(jsgrammar, numlines), + generate_function_body(webgpugrammar, _N_MAIN_LINES), 1 ) return result -def generate_samples(grammar_dir, outfiles): - """Generates a set of samples and writes them to the output files. +def generate_samples(template, grammar_dir, outfiles): + extra = "" + shaders_dir = os.path.join(grammar_dir, "wgsl/*.wgsl") + shader_files = glob.glob(shaders_dir) - Args: - grammar_dir: directory to load grammar files from. - outfiles: A list of output filenames. - """ + shaders = [] + for i in range(_N_SHADERS): + shader_path = random.choice(shader_files) + with open(shader_path) as fp: + shader_src = fp.read() + shaders.append(shader_src) + + with open(os.path.join(grammar_dir, template), "r") as fp: + template_contents = fp.read() - f = open(os.path.join(grammar_dir, 'template.html')) - template = f.read() - f.close() + SHADER_CONST = "" + for i, shader in enumerate(shaders): + shader_template = SHADER_CONST % str(i) + template_contents = template_contents.replace(shader_template, shader) - jsgrammar = Grammar() - err = jsgrammar.parse_from_file(os.path.join(grammar_dir, 'webgpu.txt')) - if err > 0: - print('There were errors parsing WebGPU JS grammar') - return + extra += parse_entrypoints(shaders) + "\n" + extra += parse_bindings(shaders) + "\n" - wgslgrammar = Grammar() - err = jsgrammar.parse_from_file(os.path.join(grammar_dir, 'wgsl.txt')) + webgpugrammar = Grammar() + err = webgpugrammar.parse_from_file(os.path.join(grammar_dir, os.path.join(grammar_dir, 'webgpu.txt')), extra) if err > 0: - print('There were errors parsing WGSL grammar') + print('There were errors parsing grammar') return for outfile in outfiles: - result = GenerateNewSample(template, jsgrammar, wgslgrammar) + result = generate_new_sample(template_contents, webgpugrammar) if result is not None: print('Writing a sample to ' + outfile) try: - f = open(outfile, 'w') - f.write(result) - f.close() + with open(outfile, 'w') as f: + f.write(result) except IOError: print('Error writing to output') @@ -112,11 +136,11 @@ def get_option(option_name): return sys.argv[i][len(option_name) + 1:] return None - def main(): fuzzer_dir = os.path.dirname(__file__) multiple_samples = False + template = os.path.join(fuzzer_dir, "template.html") for a in sys.argv: if a.startswith('--output_dir='): @@ -125,29 +149,24 @@ def main(): multiple_samples = True if multiple_samples: + print('Running on ClusterFuzz') out_dir = get_option('--output_dir') nsamples = int(get_option('--no_of_files')) print('Output directory: ' + out_dir) print('Number of samples: ' + str(nsamples)) - if not os.path.exists(out_dir): os.mkdir(out_dir) - outfiles = [] for i in range(nsamples): outfiles.append(os.path.join(out_dir, 'fuzz-' + str(i).zfill(5) + '.html')) - - generate_samples(fuzzer_dir, outfiles) - - elif len(sys.argv) > 1: - outfile = sys.argv[1] - generate_samples(fuzzer_dir, [outfile]) - + generate_samples(template, fuzzer_dir, outfiles) else: print('Arguments missing') print("Usage:") - print("\tpython generator.py ") - print("\tpython generator.py --output_dir --no_of_files ") + print("""--input_dir . ** not used. ** + --output_dir . This is the output directory which the fuzzer should write to. + --no_of_files . This is the number of testcases which the fuzzer should write to the output directory. + """) if __name__ == '__main__': main() diff --git a/webgpu/template.html b/webgpu/template.html index ffb0c09..b34fa7a 100644 --- a/webgpu/template.html +++ b/webgpu/template.html @@ -1,65 +1,103 @@ + + - +setTimeout(function() { cleanup(); }, 12000); +function out_of_scope_gc_flush() { + return new Promise(resolve => setTimeout(() => { + // Let's ensure buffer is garbage collected + try { new ArrayBuffer(0x80000000); } catch(e) {} + + // Flush. + navigator.gpu.requestAdapter().then(adapter => { + adapter.requestDevice().then(device => {}) + }); + + }, 0)); +} + + +function cleanup() { + try { gc(); } catch(e) {} + try { new ArrayBuffer(0x80000000); } catch(e) {} + try { new ArrayBuffer(0x80000000); } catch(e) {} + try { new ArrayBuffer(0x80000000); } catch(e) {} + try { new ArrayBuffer(0x80000000); } catch(e) {} + + navigator.gpu.requestAdapter().then(adapter => { + adapter.requestDevice().then(device => {}) + }); +} + +async function trigger() { + const adapter = await navigator.gpu.requestAdapter(); + const device = await adapter.requestDevice(); + + const Shader0 = device.createShaderModule({ + code: ` + + ` + }); + + const Shader1 = device.createShaderModule({ + code: ` + + ` + }); + + const Shader2 = device.createShaderModule({ + code: ` + + ` + }); + + const Shader3 = device.createShaderModule({ + code: ` + + ` + }); + + const Shader4 = device.createShaderModule({ + code: ` + + ` + }); + + const Shader5 = device.createShaderModule({ + code: ` + + ` + }); + + const Shader6 = device.createShaderModule({ + code: ` + + ` + }); + + const Shader7 = device.createShaderModule({ + code: ` + + ` + }); + + const Shader8 = device.createShaderModule({ + code: ` + + ` + }); + + const Shader9 = device.createShaderModule({ + code: ` + + ` + }); + + + +} + + + \ No newline at end of file diff --git a/webgpu/webgpu.txt b/webgpu/webgpu.txt index 846fbb4..a7f5b5e 100644 --- a/webgpu/webgpu.txt +++ b/webgpu/webgpu.txt @@ -1,312 +1,716 @@ +!include webgpuhelpers.txt - = true - = false - - = 0 - = 1 - = 2 - = 4 - = 8 - = 16 - = 32 - = 64 - = 128 - - = 0 - = 1 - = 2 - = 4 - = 8 - = 16 - = 32 - = 64 - = 128 - - = 0x7fffff00 - = 0x64 - = 0x3e8 - = 0x4141 - = 0xefff - = 0xaa - = 0xaf43 - = -0x5a - = true - = false - - = 536870911 - = 536870912 - = 1073741823 - = 1073741824 - = 2147483647 - = 2147483648 - = 4294967295 - = 4294967296 - - = 100 - = 200 - = 500 - = 1000 - = 10000 - - = - = + - = + + - = "AAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAA" - = 17 - = 65 - = 257 - = 1025 - = 4097 - = 65537 - = String.fromCharCode() - = String.fromCharCode().repeat() - - = - = 256 - = 512 - = 1024 - = 4096 - = 16384 - = 65536 - - = - = - = - = - = - - = GPUBufferUsage.MAP_READ - = GPUBufferUsage.MAP_WRITE - = GPUBufferUsage.COPY_SRC - = GPUBufferUsage.COPY_DST - = GPUBufferUsage.STORAGE - = GPUBufferUsage.INDEX - = GPUBufferUsage.INDIRECT - = GPUBufferUsage.UNIFORM - = GPUBufferUsage.VERTEX - = GPUBufferUsage.QUERY_RESOLVE - = 0xffff - = 0x - = 0x003 - = 0x005 - = 0x006 - = 0x007 - = 0x1 - = 0x2 - = 0x3 - = 0x4 - = 0x5 - = 0x6 - = 0x7 - = 0x9 - = 0x11 - = 0x12 - = 0x081 - - = GPUMapMode.READ - = GPUMapMode.WRITE - - = 0 - = 1 - = 2 - - = GPUShaderStage.COMPUTE - = GPUShaderStage.FRAGMENT - = GPUShaderStage.VERTEX - - = "uniform" - = "storage" - = "read-only-storage" - = "filtering" - = "non-filtering" - = "comparison" - = "float" - = "unfilterable-float" - = "depth" - = "sint" - = "uint" - = "write-only" - - = - = - = - = - - = canvas0 - = canvas1 - = canvas2 - = canvas3 - = canvas4 - - = "2d" - = "webgl" - = "webgl2" - = "bitmaprenderer" - - = matrixA - = matrixB - = matrixC - - = "1d" - = "2d" - = "3d" - - = "r8unorm" - = "r8snorm" - = "r8uint" - = "r8sint" - = "r16uint" - = "r16sint" - = "r16float" - = "rg8unorm" - = "rg8snorm" - = "rg8uint" - = "rg8sint" - = "r32uint" - = "r32sint" - = "r32float" - = "rg16uint" - = "rg16sint" - = "rg16float" - = "rgba8unorm" - = "rgba8unorm-srgb" - = "rgba8snorm" - = "rgba8uint" - = "rgba8sint" - = "bgra8unorm" - = "bgra8unorm-srgb" - = "rgb9e5ufloat" - = "rgb10a2unorm" - = "rg11b10ufloat" - = "rg32uint" - = "rg32sint" - = "rg32float" - = "rgba16uint" - = "rgba16sint" - = "rgba16float" - = "rgba32uint" - = "rgba32sint" - = "rgba32float" - = "stencil8" - = "depth16unorm" - = "depth24plus" - = "depth24plus-stencil8" - = "depth32float" - = "depth32float-stencil8" - = "bc1-rgba-unorm" - = "bc1-rgba-unorm-srgb" - = "bc2-rgba-unorm" - = "bc2-rgba-unorm-srgb" - = "bc3-rgba-unorm" - = "bc3-rgba-unorm-srgb" - = "bc4-r-unorm" - = "bc4-r-snorm" - = "bc5-rg-unorm" - = "bc5-rg-snorm" - = "bc6h-rgb-ufloat" - = "bc6h-rgb-float" - = "bc7-rgba-unorm" - = "bc7-rgba-unorm-srgb" - = "etc2-rgb8unorm" - = "etc2-rgb8unorm-srgb" - = "etc2-rgb8a1unorm" - = "etc2-rgb8a1unorm-srgb" - = "etc2-rgba8unorm" - = "etc2-rgba8unorm-srgb" - = "eac-r11unorm" - = "eac-r11snorm" - = "eac-rg11unorm" - = "eac-rg11snorm" - = "astc-4x4-unorm" - = "astc-4x4-unorm-srgb" - = "astc-5x4-unorm" - = "astc-5x4-unorm-srgb" - = "astc-5x5-unorm" - = "astc-5x5-unorm-srgb" - = "astc-6x5-unorm" - = "astc-6x5-unorm-srgb" - = "astc-6x6-unorm" - = "astc-6x6-unorm-srgb" - = "astc-8x5-unorm" - = "astc-8x5-unorm-srgb" - = "astc-8x6-unorm" - = "astc-8x6-unorm-srgb" - = "astc-8x8-unorm" - = "astc-8x8-unorm-srgb" - = "astc-10x5-unorm" - = "astc-10x5-unorm-srgb" - = "astc-10x6-unorm" - = "astc-10x6-unorm-srgb" - = "astc-10x8-unorm" - = "astc-10x8-unorm-srgb" - = "astc-10x10-unorm" - = "astc-10x10-unorm-srgb", - = "astc-12x10-unorm" - = "astc-12x10-unorm-srgb" - = "astc-12x12-unorm" - = "astc-12x12-unorm-srgb" - - = "out-of-memory" - = "validation" - - = - -!include ../rules/common.txt -#include ../canvas/canvas.txt - -!lineguard try { } catch(e) { } + = Shader0 + = Shader1 + = Shader2 + = Shader3 + = Shader4 + = Shader5 + = Shader6 + = Shader7 + = Shader8 + = Shader9 + + = 1234 + = o + = wgsize + + = "srgb" + = "display-p3" + + = "uniform-buffer" + = "storage-buffer" + = "readonly-storage-buffer" + = "sampler" + = "comparison-sampler" + = "sampled-texture" + = "multisampled-texture" + = "writeonly-storage-texture" + + = "float" + = "sint" + = "uint" + = "depth-comparison" + + = "zero" + = "one" + = "src" + = "one-minus-src" + = "src-alpha" + = "one-minus-src-alpha" + = "dst" + = "one-minus-dst" + = "dst-alpha" + = "one-minus-dst-alpha" + = "src-alpha-saturated" + = "constant" + = "one-minus-constant" + + = "add" + = "subtract" + = "reverse-subtract" + = "min" + = "max" + + = "unmapped" + = "pending" + = "mapped" + + = "uniform" + = "storage" + = "read-only-storage" + + = "opaque" + = "premultiplied" + + = "error" + = "warning" + = "info" + + = "beginning" + = "end" + + = "out-of-memory" + = "validation" + = "internal" + + = "destroyed" + + = "auto" + + = "uint16" + = "uint32" + + = "occlusion" + = "pipeline-statistics" + = "timestamp" + + = "vertex-shader-invocations" + = "clipper-invocations" + = "clipper-primitives-out" + = "fragment-shader-invocations" + = "compute-shader-invocations" + + = "ccw" + = "cw" + + = "none" + = "front" + = "back" + + = "load" + = "clear" + + = "store" + = "discard" + + = "beginning" + = "end" + + = "point-list" + = "line-list" + = "line-strip" + = "triangle-list" + = "triangle-strip" + + = "low-power" + = "high-performance" + + = "filtering" + = "non-filtering" + = "comparison" + + = "clamp-to-edge" + = "repeat" + = "mirror-repeat" + + = "nearest" + = "linear" + + = "never" + = "less" + = "equal" + = "less-equal" + = "greater" + = "not-equal" + = "greater-equal" + = "always" + + = "keep" + = "zero" + = "replace" + = "invert" + = "increment-clamp" + = "decrement-clamp" + = "increment-wrap" + = "decrement-wrap" + + = "write-only" + + = "pipeline-statistics-query" + = "texture-compression-bc" + = "texture-compression-etc2" + = "texture-compression-astc" + = "timestamp-query" + = "timestamp-query-inside-passes" + = "shader-float16" + = "depth-clip-control" + = "depth32float-stencil8" + = "indirect-first-instance" + = "chromium-experimental-dp4a" + = "rg11b10ufloat-renderable" + + = "float" + = "unfilterable-float" + = "depth" + = "sint" + = "uint" + + = "1d" + = "2d" + = "3d" + + = "r8unorm" + = "r8snorm" + = "r8uint" + = "r8sint" + = "r16uint" + = "r16sint" + = "r16float" + = "rg8unorm" + = "rg8snorm" + = "rg8uint" + = "rg8sint" + = "r32uint" + = "r32sint" + = "r32float" + = "rg16uint" + = "rg16sint" + = "rg16float" + = "rgba8unorm" + = "rgba8unorm-srgb" + = "rgba8snorm" + = "rgba8uint" + = "rgba8sint" + = "bgra8unorm" + = "bgra8unorm-srgb" + = "rgb9e5ufloat" + = "rgb10a2unorm" + = "rg11b10ufloat" + = "rg32uint" + = "rg32sint" + = "rg32float" + = "rgba16uint" + = "rgba16sint" + = "rgba16float" + = "rgba32uint" + = "rgba32sint" + = "rgba32float" + = "depth32float" + = "depth32float-stencil8" + = "depth24plus" + = "depth24plus-stencil8" + = "depth16unorm" + = "stencil8" + = "bc1-rgba-unorm" + = "bc1-rgba-unorm-srgb" + = "bc2-rgba-unorm" + = "bc2-rgba-unorm-srgb" + = "bc3-rgba-unorm" + = "bc3-rgba-unorm-srgb" + = "bc4-r-unorm" + = "bc4-r-snorm" + = "bc5-rg-unorm" + = "bc5-rg-snorm" + = "bc6h-rgb-ufloat" + = "bc6h-rgb-float" + = "bc7-rgba-unorm" + = "bc7-rgba-unorm-srgb" + = "etc2-rgb8unorm" + = "etc2-rgb8unorm-srgb" + = "etc2-rgb8a1unorm" + = "etc2-rgb8a1unorm-srgb" + = "etc2-rgba8unorm" + = "etc2-rgba8unorm-srgb" + = "eac-r11unorm" + = "eac-r11snorm" + = "eac-rg11unorm" + = "eac-rg11snorm" + = "astc-4x4-unorm" + = "astc-4x4-unorm-srgb" + = "astc-5x4-unorm" + = "astc-5x4-unorm-srgb" + = "astc-5x5-unorm" + = "astc-5x5-unorm-srgb" + = "astc-6x5-unorm" + = "astc-6x5-unorm-srgb" + = "astc-6x6-unorm" + = "astc-6x6-unorm-srgb" + = "astc-8x5-unorm" + = "astc-8x5-unorm-srgb" + = "astc-8x6-unorm" + = "astc-8x6-unorm-srgb" + = "astc-8x8-unorm" + = "astc-8x8-unorm-srgb" + = "astc-10x5-unorm" + = "astc-10x5-unorm-srgb" + = "astc-10x6-unorm" + = "astc-10x6-unorm-srgb" + = "astc-10x8-unorm" + = "astc-10x8-unorm-srgb" + = "astc-10x10-unorm" + = "astc-10x10-unorm-srgb" + = "astc-12x10-unorm" + = "astc-12x10-unorm-srgb" + = "astc-12x12-unorm" + = "astc-12x12-unorm-srgb" + + = "1d" + = "2d" + = "2d-array" + = "cube" + = "cube-array" + = "3d" + + = "all" + = "stencil-only" + = "depth-only" + + = "uint8x2" + = "uint8x4" + = "sint8x2" + = "sint8x4" + = "unorm8x2" + = "unorm8x4" + = "snorm8x2" + = "snorm8x4" + = "uint16x2" + = "uint16x4" + = "sint16x2" + = "sint16x4" + = "unorm16x2" + = "unorm16x4" + = "snorm16x2" + = "snorm16x4" + = "float16x2" + = "float16x4" + = "float32" + = "float32x2" + = "float32x3" + = "float32x4" + = "uint32" + = "uint32x2" + = "uint32x3" + = "uint32x4" + = "sint32" + = "sint32x2" + = "sint32x3" + = "sint32x4" + + = "vertex" + = "instance" + +# Namespaces + = GPUBufferUsage.MAP_READ + = GPUBufferUsage.MAP_WRITE + = GPUBufferUsage.COPY_SRC + = GPUBufferUsage.COPY_DST + = GPUBufferUsage.INDEX + = GPUBufferUsage.VERTEX + = GPUBufferUsage.UNIFORM + = GPUBufferUsage.STORAGE + = GPUBufferUsage.INDIRECT + = GPUBufferUsage.QUERY_RESOLVE + + = GPUColorWrite.RED + = GPUColorWrite.GREEN + = GPUColorWrite.BLUE + = GPUColorWrite.ALPHA + = GPUColorWrite.ALL + + = GPUMapMode.READ + = GPUMapMode.WRITE + + = GPUShaderStage.VERTEX + = GPUShaderStage.FRAGMENT + = GPUShaderStage.COMPUTE + + = GPUTextureUsage.COPY_SRC + = GPUTextureUsage.COPY_DST + = GPUTextureUsage.TEXTURE_BINDING + = GPUTextureUsage.STORAGE_BINDING + = GPUTextureUsage.RENDER_ATTACHMENT + +#~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~# +#~~~~~~~~~DICTIONARIES~~~~~~~~~~~# +#~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~# + = { label: , entries: , layout: } + = { binding: , resource: } + = { binding: , resource: } + = { binding: , resource: } + = { binding: , resource: } + = { label: , entries: } + = { binding: , buffer: , visibility: } + = { binding: , externalTexture: , visibility: } + = { binding: , sampler: , visibility: } + = { binding: , texture: , visibility: } + = { dstFactor: , operation: , srcFactor: } + + = { alpha: , color: } + = { buffer: , offset: , size: } + = { hasDynamicOffset: , minBindingSize: , type: } + = { label: , mappedAtCreation: , size: , usage: } + = { alphaMode: , colorSpace: , device: , format: , usage: , viewFormats: } + = { a: , b: , g: , r: } + = { alphaBlend: , colorBlend: , format: , writeMask: } + = { blend: , format: , writeMask: } + = { label: } + = { label: } + = { label: , timestampWrites: } + = { location: , queryIndex: , querySet: } + = { label: , layout: , compute: } + = { label: , layout: , compute: } + = { depthBias: , depthBiasClamp: , depthBiasSlopeScale: , depthCompare: , depthWriteEnabled: , format: , stencilBack: , stencilFront: , stencilReadMask: , stencilWriteMask: } + = { depthCompare: , depthWriteEnabled: , format: , stencilBack: , stencilFront: , stencilReadMask: , stencilWriteMask: } + = { label: , defaultQueue: , requiredFeatures: , requiredLimits: } + = { depthOrArrayLayers: , height: , width: } + = { } + = { label: , colorSpace: , source: } + = { label: , colorSpace: , source: } + = { constants: , entryPoint: , module: , targets: } + = { constants: , entryPoint: , module: , targets: } + = { bytesPerRow: , offset: , rowsPerImage: , buffer: } + = { flipY: , origin: , source: } + = { flipY: , origin: , source: } + = { flipY: , origin: , source: } + = { flipY: , origin: , source: } + = { imageBitmap: , origin: } + = { aspect: , mipLevel: , origin: , texture: } + = { aspect: , mipLevel: , origin: , texture: , colorSpace: , premultipliedAlpha: } + = { bytesPerRow: , offset: , rowsPerImage: } + = { alphaToCoverageEnabled: , count: , mask: } + = { label: } + = { x: , y: } + = { x: , y: , z: } + = { label: , layout: } + = { label: , layout: } + = { label: , bindGroupLayouts: } + = { cullMode: , frontFace: , stripIndexFormat: , topology: , unclippedDepth: } + = { constants: , entryPoint: , module: } + = { label: , count: , pipelineStatistics: , type: } + = { label: } + = { cullMode: , depthBias: , depthBiasClamp: , depthBiasSlopeScale: , frontFace: } + = { label: } + = { label: , colorFormats: , depthStencilFormat: , sampleCount: , depthReadOnly: , stencilReadOnly: } + = { clearValue: , loadOp: , resolveTarget: , storeOp: , view: } + = { depthClearValue: , depthLoadOp: , depthReadOnly: , depthStoreOp: , stencilClearValue: , stencilLoadOp: , stencilReadOnly: , stencilStoreOp: , view: } + = { label: , colorAttachments: , depthStencilAttachment: , maxDrawCount: , occlusionQuerySet: , timestampWrites: } + = { label: , colorFormats: , depthStencilFormat: , sampleCount: } + = { location: , queryIndex: , querySet: } + = { label: , layout: , depthStencil: , fragment: , multisample: , primitive: , vertex: } + = { label: , layout: , depthStencil: , fragment: , multisample: , primitive: , vertex: } + = { powerPreference: } + = { type: } + = { label: , addressModeU: , addressModeV: , addressModeW: , compare: , lodMaxClamp: , lodMinClamp: , magFilter: , maxAnisotropy: , minFilter: , mipmapFilter: } + = { layout: } + = { layout: } + = { compare: , depthFailOp: , failOp: , passOp: } + = { access: , format: , viewDimension: } + = { multisampled: , sampleType: , viewDimension: } + = { label: , dimension: , format: , mipLevelCount: , sampleCount: , size: , usage: , viewFormats: } + = { label: , arrayLayerCount: , aspect: , baseArrayLayer: , baseMipLevel: , dimension: , format: , mipLevelCount: } + = { bubbles: , cancelable: , composed: , error: } + = { format: , offset: , shaderLocation: } + = { arrayStride: , attributes: , stepMode: } + = { constants: , entryPoint: , module: , buffers: } + + + = [] + = [, ] + = [, , ] + = [] + = [, ] + = [, , ] + = [] + = [, ] + = [, , ] + = [] + = [, ] + = [, , ] + = [] + = [, ] + = [, , ] + = [] + = [, ] + = [, , ] + = [] + = [, ] + = [, , ] + = [] + = [, ] + = [, , ] + = [] + = [, ] + = [, , ] + = [] + = [, ] + = [, , ] + = [] + = [, ] + = [, , ] + = [] + = [, ] + = [, , ] + = [] + = [, ] + = [, , ] + = [] + = [, ] + = [, , ] + = [, , , ] + = [] + = [, ] + = [, , ] + + = {} + = {} + = {} + = {} + = {} + = {} + = {} + = {} + = {} + = {} + = {} + = {} + = {} + = {} + = {} + = { : } + = { : } + = { : } + + = adapter + = device + +!lineguard try { } catch(e) {} !varformat fuzzvar%05d !begin lines -setTimeout(function(){location.reload();},); - = await navigator.gpu.requestAdapter(); - = await .requestDevice(); - = .queue; - = .createShaderModule({code: wgsl_shader_src}); - = .createBuffer({mappedAtCreation: , size: , usage: GPUBufferUsage.STORAGE}); - = .createBuffer({mappedAtCreation: , size: , usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC}); - = .createBuffer({mappedAtCreation: , size: , usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ}); - = ; - = ; - = ; -# = .createBuffer({mappedAtCreation: , size: , usage: }); - = .getMappedRange().fill(0x41); - = .getMappedRange(); - = new Uint8Array(.getMappedRange()); - = Float32Array(); - = Float32Array().set(); - = Uint8Array(); - = Uint8Array().set(); -for (var i = 0; i ; i++) {console.log([i]);} -for (var i = 0; i ; i++) {console.log([i]);} -.unmap(); -.map(); -.mapAsync(); -# = .createBindGroupLayout({entries: [{binding: 0, visibility: , buffer: {type: }}, {binding: 1, visibility: , buffer: {type: }}, {binding: 2, visibility: , buffer: {type: }}]}); - = .createBindGroupLayout({entries: [{binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: {type: "storage"}}, {binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: {type: "storage"}}, {binding: 2, visibility: GPUShaderStage.COMPUTE, buffer: {type: "storage"}}]}); -# = .createBindGroup({layout: , entries: [{binding: 0, resource: {buffer: }}, {binding: 1, resource: {buffer: }}, {binding: 2, resource: {buffer: }}]}); -activeBuffer = .createBuffer({size: , usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC}); - = .createBindGroup({layout: , entries: [{binding: 0, resource: {buffer: }}, {binding: 1, resource: {buffer: }}, {binding: 2, resource: {buffer: activeBuffer}}]}); - = .createComputePipeline({layout: .createPipelineLayout({bindGroupLayouts: []}), compute: {module: , entryPoint: "main"}}); - = .createCommandEncoder(); - = .beginComputePass(); -.setPipeline(); -.setBindGroup(, ) - = Math.ceil([] / ); -.dispatchWorkgroups(, ); -.end(); -.copyBufferToBuffer(, , , , ); - = .finish(); -.submit([]); - = .getContext(); -.configure({device: , format: , usage: }) -.transferControlToOffscreen(); -.pushErrorScope(); -.popErrorScope(); -.pushErrorScope(); -.popErrorScope().then().catch(e => {}); - = .getCurrentTexture(); - = .createView(); -.copyTextureToBuffer({texture: }, {buffer: , bytesPerRow: .width}, {width: , height: , depthOrArrayLayers: 1}); -for (var i = 0; i ; i++) {console.log([i]);} - = new OffscreenCanvas(original_image.width, original_image.height); -!end lines + + +#~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~# +#~~~~~~~~~~INTERFACES~~~~~~~~~~~~# +#~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~# + + +#~~~~~~~~~~~~~~~~GPU~~~~~~~~~~~~~~~~# + = .getPreferredCanvasFormat(); + +#~~~~~~~~~~~~~~~~GPUAdapter~~~~~~~~~~~~~~~~# + + = await .requestAdapterInfo(); + +#~~~~~~~~~~~~~~~~GPUBindGroup~~~~~~~~~~~~~~~~# +.label = + +#~~~~~~~~~~~~~~~~GPUBindGroupLayout~~~~~~~~~~~~~~~~# +.label = + +#~~~~~~~~~~~~~~~~GPUBuffer~~~~~~~~~~~~~~~~# +.label = +await .mapAsync(); +await .mapAsync(,); +await .mapAsync(,,); + = .getMappedRange(); + = .getMappedRange(); + = .getMappedRange(,); +.unmap(); +.destroy(); + +#~~~~~~~~~~~~~~~~GPUCanvasContext~~~~~~~~~~~~~~~~# +.configure(); +.unconfigure(); + = .getCurrentTexture(); + +#~~~~~~~~~~~~~~~~GPUCommandBuffer~~~~~~~~~~~~~~~~# +.label = + +#~~~~~~~~~~~~~~~~GPUCommandEncoder~~~~~~~~~~~~~~~~# +.label = + = .beginRenderPass(); + = .beginComputePass(); + = .beginComputePass(); +.copyBufferToBuffer(,,,,); +.copyBufferToTexture(,,); +.copyTextureToBuffer(,,); +.copyTextureToTexture(,,); +.pushDebugGroup(); +.popDebugGroup(); +.insertDebugMarker(); +.resolveQuerySet(,,,,); +.writeTimestamp(,); +.clearBuffer(); +.clearBuffer(,); +.clearBuffer(,,); + = .finish(); + = .finish(); + +#~~~~~~~~~~~~~~~~GPUComputePassEncoder~~~~~~~~~~~~~~~~# +.label = +.setPipeline(); +.dispatchWorkgroups(); +.dispatchWorkgroups(,); +.dispatchWorkgroups(,,); +.dispatchWorkgroupsIndirect(,); +.writeTimestamp(,); +.end(); +.setBindGroup(,); +.setBindGroup(,,); +.setBindGroup(,,,,); +.pushDebugGroup(); +.popDebugGroup(); +.insertDebugMarker(); + +#~~~~~~~~~~~~~~~~GPUComputePipeline~~~~~~~~~~~~~~~~# +.label = + = .getBindGroupLayout(); + +#~~~~~~~~~~~~~~~~GPUDevice~~~~~~~~~~~~~~~~# + = device.queu; +.onuncapturederror = null +.label = +.destroy(); + = .createBuffer(); + = .createTexture(); + = .experimentalImportTexture(,); + = .createSampler(); + = .createSampler(); + = .importExternalTexture(); + = .createBindGroup(); + = .createBindGroupLayout(); + = .createPipelineLayout(); + = .createRenderPipeline(); + = .createComputePipeline(); + = await .createRenderPipelineAsync(); + = await .createComputePipelineAsync(); + = .createCommandEncoder(); + = .createCommandEncoder(); + = .createRenderBundleEncoder(); + = .createQuerySet(); +.pushErrorScope(); + = await .popErrorScope(); + +#~~~~~~~~~~~~~~~~GPUExternalTexture~~~~~~~~~~~~~~~~# +.label = + +#~~~~~~~~~~~~~~~~GPUInternalError~~~~~~~~~~~~~~~~# + +#~~~~~~~~~~~~~~~~GPUPipelineLayout~~~~~~~~~~~~~~~~# +.label = + +#~~~~~~~~~~~~~~~~GPUQuerySet~~~~~~~~~~~~~~~~# +.label = +.destroy(); + +#~~~~~~~~~~~~~~~~GPUQueue~~~~~~~~~~~~~~~~# +.label = +.submit(); +await .onSubmittedWorkDone(); +.writeBuffer(,,); +.writeBuffer(,,,); +.writeBuffer(,,,,); +.writeBuffer(,,); +.writeBuffer(,,,); +.writeBuffer(,,,,); +.writeTexture(,,,); +.writeTexture(,,,); +.copyExternalImageToTexture(,,); + +#~~~~~~~~~~~~~~~~GPURenderBundle~~~~~~~~~~~~~~~~# +.label = + +#~~~~~~~~~~~~~~~~GPURenderBundleEncoder~~~~~~~~~~~~~~~~# +.label = + = .finish(); + = .finish(); +.setBindGroup(,); +.setBindGroup(,,); +.setBindGroup(,,,,); +.pushDebugGroup(); +.popDebugGroup(); +.insertDebugMarker(); +.setPipeline(); +.setIndexBuffer(,); +.setIndexBuffer(,,); +.setIndexBuffer(,,,); +.setVertexBuffer(,); +.setVertexBuffer(,,); +.setVertexBuffer(,,,); +.draw(); +.draw(,); +.draw(,,); +.draw(,,,); +.drawIndexed(); +.drawIndexed(,); +.drawIndexed(,,); +.drawIndexed(,,,); +.drawIndexed(,,,,); +.drawIndirect(,); +.drawIndexedIndirect(,); + +#~~~~~~~~~~~~~~~~GPURenderPassEncoder~~~~~~~~~~~~~~~~# +.label = +.setViewport(,,,,,); +.setScissorRect(,,,); +.setBlendConstant(); +.setStencilReference(); +.executeBundles(); +.beginOcclusionQuery(); +.endOcclusionQuery(); +.writeTimestamp(,); +.end(); +.setBindGroup(,); +.setBindGroup(,,); +.setBindGroup(,,,,); +.pushDebugGroup(); +.popDebugGroup(); +.insertDebugMarker(); +.setPipeline(); +.setIndexBuffer(,); +.setIndexBuffer(,,); +.setIndexBuffer(,,,); +.setVertexBuffer(,); +.setVertexBuffer(,,); +.setVertexBuffer(,,,); +.draw(); +.draw(,); +.draw(,,); +.draw(,,,); +.drawIndexed(); +.drawIndexed(,); +.drawIndexed(,,); +.drawIndexed(,,,); +.drawIndexed(,,,,); +.drawIndirect(,); +.drawIndexedIndirect(,); + +#~~~~~~~~~~~~~~~~GPURenderPipeline~~~~~~~~~~~~~~~~# +.label = + = .getBindGroupLayout(); + +#~~~~~~~~~~~~~~~~GPUSampler~~~~~~~~~~~~~~~~# +.label = + +#~~~~~~~~~~~~~~~~GPUShaderModule~~~~~~~~~~~~~~~~# +.label = + = await .compilationInfo(); + +#~~~~~~~~~~~~~~~~GPUTexture~~~~~~~~~~~~~~~~# +.label = + = .createView(); + = .createView(); +.destroy(); + +#~~~~~~~~~~~~~~~~GPUTextureView~~~~~~~~~~~~~~~~# +.label = diff --git a/webgpu/webgpuhelpers.txt b/webgpu/webgpuhelpers.txt new file mode 100644 index 0000000..e31b7b2 --- /dev/null +++ b/webgpu/webgpuhelpers.txt @@ -0,0 +1,136 @@ + + + = + = + = + = + = + = + = + = + = {} + + = self +=navigator +=self.navigator + = window + = document + = .contentWindow.document +=self + + = +=.gpu + = "fuzzstr_" + + = .transferControlToOffscreen() += + + = html_tag_to_id("canvas", ) + = + + = + = + = true + = false + + = + = + = . + = . + = + = + = + = + + = 0 + = 1 + = 2 + = 3 + = 4 + = 5 + = 6 + = 7 + = 8 + = 9 + = 10 + = 11 + = 12 + = 13 + = 14 + = 15 + = 16 + = + += + = 0 + = 1 + = 2 + = 4 + = 8 + = 16 + = 32 + = 64 + = 128 + + = 0x64 + = 0x3e8 + = 0x4141 + = 0xefff + = 0xaa + = 0xaf43 + = -0x9c + = true + = false + + = + = + = + = + + = 0x100 + = 65536 + = + = + + = 536870911 + = 536870912 + = 1073741823 + = 1073741824 + = 2147483647 + = 2147483648 + = 4294967295 + = 4294967296 + + +!lineguard try { } catch(e) { } +!varformat fuzzvar%05d +!begin lines + + = html_tag_to_id("iframe", ); + = html_tag_to_id("video", ); + = html_tag_to_id("video", ); + = html_tag_to_id("canvas", ); + + + = await createImageBitmap(,,,,); + = await createImageBitmap(,,,,); + = await createImageBitmap(,,,,); + = new VideoFrame(, { timestamp: }); + = new ImageData(, ); + = self.URL.createObjectURL(new Blob([worker_blob], { type: 'text/javascript' })); +=.getContext('webgpu'); + + + = new Int8Array(); + = new Int16Array(); + = new Int32Array(); + = new Uint8Array(); + = new Uint16Array(); + = new Uint32Array(); + = new Float32Array(); + = new Float64Array(); + = new Uint8ClampedArray(); + +out_of_scope_gc_flush(); +out_of_scope_gc_flush(); +out_of_scope_gc_flush(); diff --git a/webgpu/wgsl.txt b/webgpu/wgsl.txt deleted file mode 100644 index e69de29..0000000 diff --git a/webgpu/wgsl/domato-example.wgsl b/webgpu/wgsl/domato-example.wgsl new file mode 100644 index 0000000..61285b4 --- /dev/null +++ b/webgpu/wgsl/domato-example.wgsl @@ -0,0 +1,24 @@ +alias RTArr = array; + +struct S { + field0 : RTArr, +} + +var x_2 : vec3; + +@group(0) @binding(0) var x_5 : S; + +@group(0) @binding(1) var x_6 : S; + +fn main_1() { + let x_20 : u32 = x_2.x; + let x_22 : u32 = x_5.field0[x_20]; + x_6.field0[x_20] = bitcast(-(bitcast(x_22))); + return; +} + +@compute @workgroup_size(1, 1, 1) +fn main(@builtin(global_invocation_id) x_2_param : vec3) { + x_2 = x_2_param; + main_1(); +}