Merge remote-tracking branch 'tint/main' into HEAD

Integrates Tint repo into Dawn

KIs:
- Building docs for Tint is turned off, because it fails due to lack
  of annotations in Dawn source files.
- Dawn CQ needs to be updated to run Tint specific tests
- Significant post-merge cleanup needed

R=bclayton,cwallez
BUG=dawn:1339

Change-Id: I6c9714a0030934edd6c51f3cac4684dcd59d1ea3

.clang-format

@@ -1,3 +1,4 @@
+# http://clang.llvm.org/docs/ClangFormatStyleOptions.html
 BasedOnStyle: Chromium
 Standard: Cpp11
 

.gitignore

@@ -9,12 +9,15 @@
 /testing
 /third_party/abseil-cpp/
 /third_party/angle
+/third_party/benchmark
+/third_party/binutils
 /third_party/catapult
 /third_party/clang-format
+/third_party/cpplint
 /third_party/glfw
 /third_party/googletest
 /third_party/gpuweb
-/third_party/webgpu-cts
+/third_party/gpuweb-cts
 /third_party/jinja2
 /third_party/jsoncpp
 /third_party/llvm-build
@@ -22,10 +25,11 @@
 /third_party/node
 /third_party/node-addon-api
 /third_party/node-api-headers
+/third_party/protobuf
 /third_party/swiftshader
-/third_party/tint
 /third_party/vulkan-deps
 /third_party/vulkan_memory_allocator
+/third_party/webgpu-cts
 /third_party/zlib
 /tools/clang
 /tools/cmake
@@ -33,7 +37,7 @@
 /tools/memory
 /out
 
-# Modified from https://www.gitignore.io/api/vim,macos,linux,emacs,windows,sublimetext,visualstudio,visualstudiocode
+# Modified from https://www.gitignore.io/api/vim,macos,linux,emacs,windows,sublimetext,visualstudio,visualstudiocode,intellij
 
 ### Emacs ###
 *~
@@ -101,8 +105,19 @@ ehthumbs_vista.db
 Desktop.ini
 $RECYCLE.BIN/
 
+### Intellij ###
+.idea
+
 ### Dawn node tools binaries
 src/dawn/node/tools/bin/
 
 ### Cached node transpiled tools
 /.node_transpile_work_dir
+
+# Misc inherited from Tint
+/test.wgsl
+coverage.summary
+default.profraw
+lcov.info
+/cmake-build-*/
+/testing

AUTHORS

@@ -1,6 +1,7 @@
-# This is the list of Dawn authors for copyright purposes.
+# This is the list of Dawn & Tint authors for copyright purposes.
 #
 # This does not necessarily list everyone who has contributed code, since in
 # some cases, their employer may be the copyright holder.  To see the full list
 # of contributors, see the revision history in source control.
-Google Inc.
+Google LLC
+Vasyl Teliman

AUTHORS.dawn

@@ -0,0 +1,6 @@
+# This is the list of Dawn authors for copyright purposes.
+#
+# This does not necessarily list everyone who has contributed code, since in
+# some cases, their employer may be the copyright holder.  To see the full list
+# of contributors, see the revision history in source control.
+Google Inc.

AUTHORS.tint

@@ -0,0 +1,8 @@
+# This is the list of the Tint authors for copyright purposes.
+#
+# This does not necessarily list everyone who has contributed code, since in
+# some cases, their employer may be the copyright holder.  To see the full list
+# of contributors, see the revision history in source control.
+
+Google LLC
+Vasyl Teliman

BUILD.gn

@@ -21,9 +21,15 @@ group("all") {
     "src/dawn/native:webgpu_dawn",
     "src/dawn/tests",
     "src/fuzzers/dawn:dawn_fuzzers",
+    "src/tint/fuzzers",
+    "src/tint:libtint",
+    "test/tint:tint_unittests",
   ]
   if (dawn_standalone) {
-    deps += [ "samples/dawn:samples" ]
+    deps += [
+      "samples/dawn:samples",
+      "src/tint/cmd:tint",
+    ]
   }
 }
 

CMakeLists.txt

@@ -1,4 +1,4 @@
-# Copyright 2020 The Dawn Authors
+# Copyright 2022 The Dawn & Tint Authors
 #
 # Licensed under the Apache License, Version 2.0 (the "License");
 # you may not use this file except in compliance with the License.
@@ -12,7 +12,7 @@
 # See the License for the specific language governing permissions and
 # limitations under the License.
 
-cmake_minimum_required(VERSION 3.10)
+cmake_minimum_required(VERSION 3.10.2)
 
 # When upgrading to CMake 3.11 we can remove DAWN_DUMMY_FILE because source-less add_library
 # becomes available.
@@ -26,13 +26,18 @@ project(
     DESCRIPTION "Dawn, a WebGPU implementation"
     LANGUAGES C CXX
 )
+enable_testing()
 
 set_property(GLOBAL PROPERTY USE_FOLDERS ON)
 
-if(NOT CMAKE_BUILD_TYPE)
+set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR})
-    message(WARNING "CMAKE_BUILD_TYPE not set, forcing it to Debug")
+set(CMAKE_POSITION_INDEPENDENT_CODE ON)
-    set(CMAKE_BUILD_TYPE "Debug" CACHE STRING
+set(CMAKE_CXX_STANDARD 17)
-        "Build type (Debug, Release, RelWithDebInfo, MinSizeRel)" FORCE)
+set(CMAKE_DEBUG_POSTFIX "")
+
+if ("${CMAKE_BUILD_TYPE}" STREQUAL "")
+  message(STATUS "No build type selected, default to Debug")
+  set(CMAKE_BUILD_TYPE "Debug")
 endif()
 
 set(DAWN_BUILD_GEN_DIR "${Dawn_BINARY_DIR}/gen")
@@ -132,7 +137,7 @@ set_if_not_defined(DAWN_GLFW_DIR "${DAWN_THIRD_PARTY_DIR}/glfw" "Directory in wh
 set_if_not_defined(DAWN_JINJA2_DIR "${DAWN_THIRD_PARTY_DIR}/jinja2" "Directory in which to find Jinja2")
 set_if_not_defined(DAWN_SPIRV_HEADERS_DIR "${DAWN_THIRD_PARTY_DIR}/vulkan-deps/spirv-headers/src" "Directory in which to find SPIRV-Headers")
 set_if_not_defined(DAWN_SPIRV_TOOLS_DIR "${DAWN_THIRD_PARTY_DIR}/vulkan-deps/spirv-tools/src" "Directory in which to find SPIRV-Tools")
-set_if_not_defined(DAWN_TINT_DIR "${DAWN_THIRD_PARTY_DIR}/tint" "Directory in which to find Tint")
+set_if_not_defined(DAWN_TINT_DIR "${Dawn_SOURCE_DIR}" "Directory in which to find Tint")
 set_if_not_defined(DAWN_VULKAN_HEADERS_DIR "${DAWN_THIRD_PARTY_DIR}/vulkan-deps/vulkan-headers/src" "Directory in which to find Vulkan-Headers")
 
 # Dependencies for DAWN_BUILD_NODE_BINDINGS
@@ -203,11 +208,355 @@ endif()
 
 set(CMAKE_CXX_STANDARD "17")
 
+################################################################################
+# Tint
+################################################################################
+
+# TINT_IS_SUBPROJECT is 1 if added via add_subdirectory() from another project.
+get_directory_property(TINT_IS_SUBPROJECT PARENT_DIRECTORY)
+if(TINT_IS_SUBPROJECT)
+  set(TINT_IS_SUBPROJECT 1)
+
+  # If tint is used as a subproject, default to disabling the building of
+  # documentation and tests. These are unlikely to be desirable, but can be
+  # enabled.
+  set(TINT_BUILD_DOCS_DEFAULT OFF)
+  set(TINT_BUILD_TESTS_DEFAULT OFF)
+else()
+  set(TINT_BUILD_DOCS_DEFAULT ON)
+  set(TINT_BUILD_TESTS_DEFAULT ON)
+endif()
+
+# Forcing building docs off right now, since currently this will try to build docs for both Tint & Dawn, and Dawn isn't annotated yet.
+set(TINT_BUILD_DOCS_DEFAULT OFF)
+
+option_if_not_defined(TINT_BUILD_SAMPLES "Build samples" ON)
+option_if_not_defined(TINT_BUILD_DOCS "Build documentation" ${TINT_BUILD_DOCS_DEFAULT})
+option_if_not_defined(TINT_DOCS_WARN_AS_ERROR "When building documentation, treat warnings as errors" OFF)
+option_if_not_defined(TINT_BUILD_SPV_READER "Build the SPIR-V input reader" ON)
+option_if_not_defined(TINT_BUILD_WGSL_READER "Build the WGSL input reader" ON)
+option_if_not_defined(TINT_BUILD_GLSL_WRITER "Build the GLSL output writer" ON)
+option_if_not_defined(TINT_BUILD_HLSL_WRITER "Build the HLSL output writer" ON)
+option_if_not_defined(TINT_BUILD_MSL_WRITER "Build the MSL output writer" ON)
+option_if_not_defined(TINT_BUILD_SPV_WRITER "Build the SPIR-V output writer" ON)
+option_if_not_defined(TINT_BUILD_WGSL_WRITER "Build the WGSL output writer" ON)
+option_if_not_defined(TINT_BUILD_FUZZERS "Build fuzzers" OFF)
+option_if_not_defined(TINT_BUILD_SPIRV_TOOLS_FUZZER "Build SPIRV-Tools fuzzer" OFF)
+option_if_not_defined(TINT_BUILD_AST_FUZZER "Build AST fuzzer" OFF)
+option_if_not_defined(TINT_BUILD_REGEX_FUZZER "Build regex fuzzer" OFF)
+option_if_not_defined(TINT_BUILD_BENCHMARKS "Build benchmarks" OFF)
+option_if_not_defined(TINT_BUILD_TESTS "Build tests" ${TINT_BUILD_TESTS_DEFAULT})
+option_if_not_defined(TINT_BUILD_AS_OTHER_OS "Override OS detection to force building of *_other.cc files" OFF)
+option_if_not_defined(TINT_BUILD_REMOTE_COMPILE "Build the remote-compile tool for validating shaders on a remote machine" OFF)
+
+set(TINT_LIB_FUZZING_ENGINE_LINK_OPTIONS "" CACHE STRING "Used by OSS-Fuzz to control, via link options, which fuzzing engine should be used")
+
+option_if_not_defined(TINT_ENABLE_MSAN "Enable memory sanitizer" OFF)
+option_if_not_defined(TINT_ENABLE_ASAN "Enable address sanitizer" OFF)
+option_if_not_defined(TINT_ENABLE_UBSAN "Enable undefined behaviour sanitizer" OFF)
+
+option_if_not_defined(TINT_ENABLE_BREAK_IN_DEBUGGER "Enable tint::debugger::Break()" OFF)
+
+option_if_not_defined(TINT_EMIT_COVERAGE "Emit code coverage information" OFF)
+
+option_if_not_defined(TINT_CHECK_CHROMIUM_STYLE "Check for [chromium-style] issues during build" OFF)
+
+option_if_not_defined(TINT_SYMBOL_STORE_DEBUG_NAME "Enable storing of name in tint::ast::Symbol to help debugging the AST" OFF)
+
+message(STATUS "Tint build samples: ${TINT_BUILD_SAMPLES}")
+message(STATUS "Tint build docs: ${TINT_BUILD_DOCS}")
+message(STATUS "Tint build docs with warn as error: ${TINT_DOCS_WARN_AS_ERROR}")
+message(STATUS "Tint build SPIR-V reader: ${TINT_BUILD_SPV_READER}")
+message(STATUS "Tint build WGSL reader: ${TINT_BUILD_WGSL_READER}")
+message(STATUS "Tint build GLSL writer: ${TINT_BUILD_GLSL_WRITER}")
+message(STATUS "Tint build HLSL writer: ${TINT_BUILD_HLSL_WRITER}")
+message(STATUS "Tint build MSL writer: ${TINT_BUILD_MSL_WRITER}")
+message(STATUS "Tint build SPIR-V writer: ${TINT_BUILD_SPV_WRITER}")
+message(STATUS "Tint build WGSL writer: ${TINT_BUILD_WGSL_WRITER}")
+message(STATUS "Tint build fuzzers: ${TINT_BUILD_FUZZERS}")
+message(STATUS "Tint build SPIRV-Tools fuzzer: ${TINT_BUILD_SPIRV_TOOLS_FUZZER}")
+message(STATUS "Tint build AST fuzzer: ${TINT_BUILD_AST_FUZZER}")
+message(STATUS "Tint build regex fuzzer: ${TINT_BUILD_REGEX_FUZZER}")
+message(STATUS "Tint build benchmarks: ${TINT_BUILD_BENCHMARKS}")
+message(STATUS "Tint build tests: ${TINT_BUILD_TESTS}")
+message(STATUS "Tint build with ASAN: ${TINT_ENABLE_ASAN}")
+message(STATUS "Tint build with MSAN: ${TINT_ENABLE_MSAN}")
+message(STATUS "Tint build with UBSAN: ${TINT_ENABLE_UBSAN}")
+message(STATUS "Tint build checking [chromium-style]: ${TINT_CHECK_CHROMIUM_STYLE}")
+message(STATUS "Tint build remote-compile tool: ${TINT_BUILD_REMOTE_COMPILE}")
+
+if (NOT ${TINT_LIB_FUZZING_ENGINE_LINK_OPTIONS} STREQUAL "")
+  message(STATUS "Using provided LIB_FUZZING_ENGINE options: ${TINT_LIB_FUZZING_ENGINE_LINK_OPTIONS}")
+endif()
+
+message(STATUS "Using python3")
+find_package(PythonInterp 3 REQUIRED)
+
+if (${TINT_BUILD_SPIRV_TOOLS_FUZZER})
+  message(STATUS "TINT_BUILD_SPIRV_TOOLS_FUZZER is ON - setting
+      TINT_BUILD_FUZZERS
+      TINT_BUILD_SPV_READER
+      TINT_BUILD_SPV_WRITER
+      TINT_BUILD_WGSL_READER
+      TINT_BUILD_WGSL_WRITER
+      TINT_BUILD_GLSL_WRITER
+      TINT_BUILD_HLSL_WRITER
+      TINT_BUILD_MSL_WRITER to ON")
+  set(TINT_BUILD_FUZZERS ON CACHE BOOL "Build tint fuzzers" FORCE)
+  set(TINT_BUILD_SPV_READER ON CACHE BOOL "Build SPIR-V reader" FORCE)
+  set(TINT_BUILD_SPV_WRITER ON CACHE BOOL "Build SPIR-V writer" FORCE)
+  set(TINT_BUILD_WGSL_READER ON CACHE BOOL "Build WGSL reader" FORCE)
+  set(TINT_BUILD_WGSL_WRITER ON CACHE BOOL "Build WGSL writer" FORCE)
+  set(TINT_BUILD_GLSL_WRITER ON CACHE BOOL "Build HLSL writer" FORCE)
+  set(TINT_BUILD_HLSL_WRITER ON CACHE BOOL "Build HLSL writer" FORCE)
+  set(TINT_BUILD_MSL_WRITER ON CACHE BOOL "Build MSL writer" FORCE)
+endif()
+
+if (${TINT_BUILD_AST_FUZZER})
+  message(STATUS "TINT_BUILD_AST_FUZZER is ON - setting
+      TINT_BUILD_FUZZERS
+      TINT_BUILD_WGSL_READER
+      TINT_BUILD_WGSL_WRITER
+      TINT_BUILD_SPV_WRITER
+      TINT_BUILD_MSL_WRITER
+      TINT_BUILD_GLSL_WRITER
+      TINT_BUILD_HLSL_WRITER to ON")
+  set(TINT_BUILD_FUZZERS ON CACHE BOOL "Build tint fuzzers" FORCE)
+  set(TINT_BUILD_WGSL_READER ON CACHE BOOL "Build WGSL reader" FORCE)
+  set(TINT_BUILD_WGSL_WRITER ON CACHE BOOL "Build WGSL writer" FORCE)
+  set(TINT_BUILD_SPV_WRITER ON CACHE BOOL "Build SPIR-V writer" FORCE)
+  set(TINT_BUILD_MSL_WRITER ON CACHE BOOL "Build MSL writer" FORCE)
+  set(TINT_BUILD_GLSL_WRITER ON CACHE BOOL "Build GLSL writer" FORCE)
+  set(TINT_BUILD_HLSL_WRITER ON CACHE BOOL "Build HLSL writer" FORCE)
+endif()
+
+if (${TINT_BUILD_REGEX_FUZZER})
+  message(STATUS "TINT_BUILD_REGEX_FUZZER is ON - setting
+      TINT_BUILD_FUZZERS
+      TINT_BUILD_WGSL_READER
+      TINT_BUILD_WGSL_WRITER
+      TINT_BUILD_SPV_WRITER
+      TINT_BUILD_MSL_WRITER
+      TINT_BUILD_GLSL_WRITER
+      TINT_BUILD_HLSL_WRITER to ON")
+      set(TINT_BUILD_FUZZERS ON CACHE BOOL "Build tint fuzzers" FORCE)
+      set(TINT_BUILD_WGSL_READER ON CACHE BOOL "Build WGSL reader" FORCE)
+      set(TINT_BUILD_WGSL_WRITER ON CACHE BOOL "Build WGSL writer" FORCE)
+      set(TINT_BUILD_SPV_WRITER ON CACHE BOOL "Build SPIR-V writer" FORCE)
+      set(TINT_BUILD_MSL_WRITER ON CACHE BOOL "Build MSL writer" FORCE)
+      set(TINT_BUILD_GLSL_WRITER ON CACHE BOOL "Build GLSL writer" FORCE)
+      set(TINT_BUILD_HLSL_WRITER ON CACHE BOOL "Build HLSL writer" FORCE)
+endif()
+
+set(TINT_ROOT_SOURCE_DIR ${PROJECT_SOURCE_DIR})
+
+# CMake < 3.15 sets /W3 in CMAKE_CXX_FLAGS. Remove it if it's there.
+# See https://gitlab.kitware.com/cmake/cmake/-/issues/18317
+if (MSVC)
+  if (CMAKE_CXX_FLAGS MATCHES "/W3")
+    string(REPLACE "/W3" "" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
+  endif()
+endif()
+
+if (${TINT_CHECK_CHROMIUM_STYLE})
+   set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Xclang -add-plugin -Xclang find-bad-constructs")
+endif()
+
+if (${TINT_BUILD_SPV_READER})
+  include_directories("${DAWN_THIRD_PARTY_DIR}/vulkan-deps/spirv-tools/include")
+endif()
+
+if((CMAKE_CXX_COMPILER_ID STREQUAL "Clang") AND (CMAKE_CXX_SIMULATE_ID STREQUAL "MSVC"))
+  set(COMPILER_IS_CLANG_CL TRUE)
+endif()
+
+if((CMAKE_CXX_COMPILER_ID STREQUAL "GNU") OR
+    (CMAKE_CXX_COMPILER_ID STREQUAL "AppleClang") OR
+    ((CMAKE_CXX_COMPILER_ID STREQUAL "Clang") AND
+     (NOT COMPILER_IS_CLANG_CL)))
+  set(COMPILER_IS_LIKE_GNU TRUE)
+endif()
+
+# Enable msbuild multiprocessor builds
+if (MSVC AND NOT COMPILER_IS_CLANG_CL)
+  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /MP")
+endif()
+
+set(TINT_OS_CC_SUFFIX "other")
+if (NOT TINT_BUILD_AS_OTHER_OS)
+  if(UNIX OR APPLE)
+    set(TINT_OS_CC_SUFFIX "posix")
+    set(TINT_OS_CC_SUFFIX "posix")
+  elseif(WIN32)
+    set(TINT_OS_CC_SUFFIX "windows")
+    set(TINT_OS_CC_SUFFIX "windows")
+  endif()
+endif()
+
+if(${TINT_BUILD_DOCS})
+  find_package(Doxygen)
+  if(DOXYGEN_FOUND)
+    set(DOXYGEN_WARN_AS_ERROR NO)
+    if(TINT_DOCS_WARN_AS_ERROR)
+      set(DOXYGEN_WARN_AS_ERROR YES)
+    endif()
+
+    set(DOXYGEN_WARN_FORMAT "$file:$line: $text")
+    if (MSVC)
+      set(DOXYGEN_WARN_FORMAT "$file($line): $text")
+    endif()
+
+    add_custom_target(tint-docs ALL
+        COMMAND ${CMAKE_COMMAND}
+          -E env
+          "DOXYGEN_OUTPUT_DIRECTORY=${CMAKE_BINARY_DIR}/docs"
+          "DOXYGEN_WARN_AS_ERROR=${DOXYGEN_WARN_AS_ERROR}"
+          "DOXYGEN_WARN_FORMAT=${DOXYGEN_WARN_FORMAT}"
+          ${DOXYGEN_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/Doxyfile
+        WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
+        COMMENT "Generating API documentation"
+        VERBATIM)
+  else()
+    message("Doxygen not found. Skipping documentation")
+  endif(DOXYGEN_FOUND)
+endif()
+
+function(tint_core_compile_options TARGET)
+  target_include_directories(${TARGET} PUBLIC "${TINT_ROOT_SOURCE_DIR}")
+  target_include_directories(${TARGET} PUBLIC "${TINT_ROOT_SOURCE_DIR}/include")
+
+  if (${TINT_BUILD_SPV_READER} OR ${TINT_BUILD_SPV_WRITER})
+    target_include_directories(${TARGET} PUBLIC
+        "${DAWN_THIRD_PARTY_DIR}/spirv-headers/include")
+  endif()
+
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_SPV_READER=$<BOOL:${TINT_BUILD_SPV_READER}>)
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_WGSL_READER=$<BOOL:${TINT_BUILD_WGSL_READER}>)
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_GLSL_WRITER=$<BOOL:${TINT_BUILD_GLSL_WRITER}>)
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_HLSL_WRITER=$<BOOL:${TINT_BUILD_HLSL_WRITER}>)
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_MSL_WRITER=$<BOOL:${TINT_BUILD_MSL_WRITER}>)
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_SPV_WRITER=$<BOOL:${TINT_BUILD_SPV_WRITER}>)
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_WGSL_WRITER=$<BOOL:${TINT_BUILD_WGSL_WRITER}>)
+
+  if (COMPILER_IS_LIKE_GNU)
+    target_compile_options(${TARGET} PRIVATE
+      -std=c++17
+      -fno-exceptions
+      -fno-rtti
+    )
+
+    if (${TINT_ENABLE_MSAN})
+      target_compile_options(${TARGET} PRIVATE -fsanitize=memory)
+      target_link_options(${TARGET} PRIVATE -fsanitize=memory)
+    elseif (${TINT_ENABLE_ASAN})
+      target_compile_options(${TARGET} PRIVATE -fsanitize=address)
+      target_link_options(${TARGET} PRIVATE -fsanitize=address)
+    elseif (${TINT_ENABLE_UBSAN})
+      target_compile_options(${TARGET} PRIVATE -fsanitize=undefined)
+      target_link_options(${TARGET} PRIVATE -fsanitize=undefined)
+    endif()
+  endif(COMPILER_IS_LIKE_GNU)
+
+  if (TINT_EMIT_COVERAGE)
+    if(CMAKE_CXX_COMPILER_ID MATCHES "GNU")
+        target_compile_options(${TARGET} PRIVATE "--coverage")
+        target_link_options(${TARGET} PRIVATE "gcov")
+    elseif(CMAKE_CXX_COMPILER_ID MATCHES "Clang")
+        target_compile_options(${TARGET} PRIVATE "-fprofile-instr-generate" "-fcoverage-mapping")
+        target_link_options(${TARGET} PRIVATE "-fprofile-instr-generate" "-fcoverage-mapping")
+    else()
+        message(FATAL_ERROR "Coverage generation not supported for the ${CMAKE_CXX_COMPILER_ID} toolchain")
+    endif()
+  endif(TINT_EMIT_COVERAGE)
+endfunction()
+
+function(tint_default_compile_options TARGET)
+  tint_core_compile_options(${TARGET})
+
+  set(COMMON_GNU_OPTIONS
+    -Wall
+    -Werror
+    -Wextra
+    -Wno-documentation-unknown-command
+    -Wno-padded
+    -Wno-switch-enum
+    -Wno-unknown-pragmas
+  )
+
+  set(COMMON_CLANG_OPTIONS
+    -Wno-c++98-compat
+    -Wno-c++98-compat-pedantic
+    -Wno-format-pedantic
+    -Wno-return-std-move-in-c++11
+    -Wno-unknown-warning-option
+    -Wno-undefined-var-template
+    -Wno-used-but-marked-unused
+    -Weverything
+  )
+
+  if (COMPILER_IS_LIKE_GNU)
+    target_compile_options(${TARGET} PRIVATE
+      -pedantic-errors
+      ${COMMON_GNU_OPTIONS}
+    )
+
+    if (("${CMAKE_CXX_COMPILER_ID}" STREQUAL "Clang") OR
+        ("${CMAKE_CXX_COMPILER_ID}" STREQUAL "AppleClang"))
+      target_compile_options(${TARGET} PRIVATE
+        ${COMMON_CLANG_OPTIONS}
+      )
+    endif()
+  endif(COMPILER_IS_LIKE_GNU)
+
+  if (MSVC)
+    # Specify /EHs for exception handling.
+    target_compile_options(${TARGET} PRIVATE
+      /bigobj
+      /EHsc
+      /W4
+      /WX
+      /wd4068
+      /wd4127
+      /wd4244
+      /wd4267
+      /wd4324
+      /wd4458
+      /wd4514
+      /wd4571
+      /wd4625
+      /wd4626
+      /wd4710
+      /wd4774
+      /wd4820
+      /wd5026
+      /wd5027
+    )
+
+    # When building with clang-cl on Windows, try to match our clang build
+    # options as much as possible.
+    if (COMPILER_IS_CLANG_CL)
+      target_compile_options(${TARGET} PRIVATE
+        ${COMMON_GNU_OPTIONS}
+        ${COMMON_CLANG_OPTIONS}
+        # Disable warnings that are usually disabled in downstream deps for
+        # gcc/clang, but aren't for clang-cl.
+        -Wno-global-constructors
+        -Wno-zero-as-null-pointer-constant
+        -Wno-shorten-64-to-32
+        -Wno-shadow-field-in-constructor
+        -Wno-reserved-id-macro
+        -Wno-language-extension-token
+      )
+    endif()
+  endif()
+endfunction()
+
 ################################################################################
 # Run on all subdirectories
 ################################################################################
 
 add_subdirectory(third_party)
+add_subdirectory(src/tint)
 add_subdirectory(generator)
 add_subdirectory(src/dawn)
 
@@ -220,3 +569,42 @@ if (DAWN_BUILD_SAMPLES)
     #add_subdirectory(src/utils)
     add_subdirectory(samples/dawn)
 endif()
+
+if (TINT_BUILD_SAMPLES)
+  add_subdirectory(src/tint/cmd)
+endif()
+
+if (TINT_BUILD_FUZZERS)
+  add_subdirectory(src/tint/fuzzers)
+endif()
+
+add_custom_target(tint-lint
+  COMMAND ./tools/lint
+  WORKING_DIRECTORY ${TINT_ROOT_SOURCE_DIR}
+  COMMENT "Running linter"
+  VERBATIM)
+
+add_custom_target(tint-format
+  COMMAND ./tools/format
+  WORKING_DIRECTORY ${TINT_ROOT_SOURCE_DIR}
+  COMMENT "Running formatter"
+  VERBATIM)
+
+
+if (TINT_EMIT_COVERAGE AND CMAKE_CXX_COMPILER_ID MATCHES "Clang")
+  # Generates a lcov.info file at the project root.
+  # This can be used by tools such as VSCode's Coverage Gutters extension to
+  # visualize code coverage in the editor.
+  get_filename_component(CLANG_BIN_DIR ${CMAKE_C_COMPILER} DIRECTORY)
+  set(PATH_WITH_CLANG "${CLANG_BIN_DIR}:$ENV{PATH}")
+  add_custom_target(tint-generate-coverage
+    COMMAND ${CMAKE_COMMAND} -E env PATH=${PATH_WITH_CLANG} ./tools/tint-generate-coverage $<TARGET_FILE:tint_unittests>
+    DEPENDS tint_unittests
+    WORKING_DIRECTORY ${TINT_ROOT_SOURCE_DIR}
+    COMMENT "Generating tint coverage data"
+    VERBATIM)
+endif()
+
+if (TINT_BUILD_REMOTE_COMPILE)
+  add_subdirectory(tools/src/cmd/remote-compile)
+endif()

CMakeSettings.json

@@ -0,0 +1,100 @@
+{
+  "configurations": [
+    {
+      "name": "x64-Debug",
+      "generator": "Ninja",
+      "configurationType": "Debug",
+      "inheritEnvironments": [ "msvc_x64_x64" ],
+      "buildRoot": "${projectDir}\\out\\build\\${name}",
+      "installRoot": "${projectDir}\\out\\install\\${name}",
+      "cmakeCommandArgs": "",
+      "buildCommandArgs": "",
+      "ctestCommandArgs": "",
+      "variables": []
+    },
+    {
+      "name": "x64-Release",
+      "generator": "Ninja",
+      "configurationType": "RelWithDebInfo",
+      "buildRoot": "${projectDir}\\out\\build\\${name}",
+      "installRoot": "${projectDir}\\out\\install\\${name}",
+      "cmakeCommandArgs": "",
+      "buildCommandArgs": "",
+      "ctestCommandArgs": "",
+      "inheritEnvironments": [ "msvc_x64_x64" ],
+      "variables": []
+    },
+    {
+      "name": "x86-Debug",
+      "generator": "Ninja",
+      "configurationType": "Debug",
+      "buildRoot": "${projectDir}\\out\\build\\${name}",
+      "installRoot": "${projectDir}\\out\\install\\${name}",
+      "cmakeCommandArgs": "",
+      "buildCommandArgs": "",
+      "ctestCommandArgs": "",
+      "inheritEnvironments": [ "msvc_x86" ],
+      "variables": []
+    },
+    {
+      "name": "x86-Release",
+      "generator": "Ninja",
+      "configurationType": "RelWithDebInfo",
+      "buildRoot": "${projectDir}\\out\\build\\${name}",
+      "installRoot": "${projectDir}\\out\\install\\${name}",
+      "cmakeCommandArgs": "",
+      "buildCommandArgs": "",
+      "ctestCommandArgs": "",
+      "inheritEnvironments": [ "msvc_x86" ],
+      "variables": []
+    },
+    {
+      "name": "x64-Clang-Debug",
+      "generator": "Ninja",
+      "configurationType": "Debug",
+      "buildRoot": "${projectDir}\\out\\build\\${name}",
+      "installRoot": "${projectDir}\\out\\install\\${name}",
+      "cmakeCommandArgs": "",
+      "buildCommandArgs": "",
+      "ctestCommandArgs": "",
+      "inheritEnvironments": [ "clang_cl_x64_x64" ],
+      "variables": []
+    },
+    {
+      "name": "x64-Clang-Release",
+      "generator": "Ninja",
+      "configurationType": "RelWithDebInfo",
+      "buildRoot": "${projectDir}\\out\\build\\${name}",
+      "installRoot": "${projectDir}\\out\\install\\${name}",
+      "cmakeCommandArgs": "",
+      "buildCommandArgs": "",
+      "ctestCommandArgs": "",
+      "inheritEnvironments": [ "clang_cl_x64_x64" ],
+      "variables": []
+    },
+    {
+      "name": "x86-Clang-Debug",
+      "generator": "Ninja",
+      "configurationType": "Debug",
+      "buildRoot": "${projectDir}\\out\\build\\${name}",
+      "installRoot": "${projectDir}\\out\\install\\${name}",
+      "cmakeCommandArgs": "",
+      "buildCommandArgs": "",
+      "ctestCommandArgs": "",
+      "inheritEnvironments": [ "clang_cl_x86" ],
+      "variables": []
+    },
+    {
+      "name": "x86-Clang-Release",
+      "generator": "Ninja",
+      "configurationType": "RelWithDebInfo",
+      "buildRoot": "${projectDir}\\out\\build\\${name}",
+      "installRoot": "${projectDir}\\out\\install\\${name}",
+      "cmakeCommandArgs": "",
+      "buildCommandArgs": "",
+      "ctestCommandArgs": "",
+      "inheritEnvironments": [ "clang_cl_x86" ],
+      "variables": []
+    }
+  ]
+}

CODE_OF_CONDUCT.md

@@ -0,0 +1,93 @@
+# Code of Conduct
+
+## Our Pledge
+
+In the interest of fostering an open and welcoming environment, we as
+contributors and maintainers pledge to making participation in our project and
+our community a harassment-free experience for everyone, regardless of age, body
+size, disability, ethnicity, gender identity and expression, level of
+experience, education, socio-economic status, nationality, personal appearance,
+race, religion, or sexual identity and orientation.
+
+## Our Standards
+
+Examples of behavior that contributes to creating a positive environment
+include:
+
+*   Using welcoming and inclusive language
+*   Being respectful of differing viewpoints and experiences
+*   Gracefully accepting constructive criticism
+*   Focusing on what is best for the community
+*   Showing empathy towards other community members
+
+Examples of unacceptable behavior by participants include:
+
+*   The use of sexualized language or imagery and unwelcome sexual attention or
+    advances
+*   Trolling, insulting/derogatory comments, and personal or political attacks
+*   Public or private harassment
+*   Publishing others' private information, such as a physical or electronic
+    address, without explicit permission
+*   Other conduct which could reasonably be considered inappropriate in a
+    professional setting
+
+## Our Responsibilities
+
+Project maintainers are responsible for clarifying the standards of acceptable
+behavior and are expected to take appropriate and fair corrective action in
+response to any instances of unacceptable behavior.
+
+Project maintainers have the right and responsibility to remove, edit, or reject
+comments, commits, code, wiki edits, issues, and other contributions that are
+not aligned to this Code of Conduct, or to ban temporarily or permanently any
+contributor for other behaviors that they deem inappropriate, threatening,
+offensive, or harmful.
+
+## Scope
+
+This Code of Conduct applies both within project spaces and in public spaces
+when an individual is representing the project or its community. Examples of
+representing a project or community include using an official project e-mail
+address, posting via an official social media account, or acting as an appointed
+representative at an online or offline event. Representation of a project may be
+further defined and clarified by project maintainers.
+
+This Code of Conduct also applies outside the project spaces when the Project
+Steward has a reasonable belief that an individual's behavior may have a
+negative impact on the project or its community.
+
+## Conflict Resolution
+
+We do not believe that all conflict is bad; healthy debate and disagreement
+often yield positive results. However, it is never okay to be disrespectful or
+to engage in behavior that violates the project’s code of conduct.
+
+If you see someone violating the code of conduct, you are encouraged to address
+the behavior directly with those involved. Many issues can be resolved quickly
+and easily, and this gives people more control over the outcome of their
+dispute. If you are unable to resolve the matter for any reason, or if the
+behavior is threatening or harassing, report it. We are dedicated to providing
+an environment where participants feel welcome and safe.
+
+Reports should be directed to David Neto <dneto@google.com>, the
+Project Steward(s) for Tint. It is the Project Steward’s duty to
+receive and address reported violations of the code of conduct. They will then
+work with a committee consisting of representatives from the Open Source
+Programs Office and the Google Open Source Strategy team. If for any reason you
+are uncomfortable reaching out the Project Steward, please email
+opensource@google.com.
+
+We will investigate every complaint, but you may not receive a direct response.
+We will use our discretion in determining when and how to follow up on reported
+incidents, which may range from not taking action to permanent expulsion from
+the project and project-sponsored spaces. We will notify the accused of the
+report and provide them an opportunity to discuss it before any action is taken.
+The identity of the reporter will be omitted from the details of the report
+supplied to the accused. In potentially harmful situations, such as ongoing
+harassment or threats to anyone's safety, we may take action without notice.
+
+## Attribution
+
+This Code of Conduct is adapted from the Contributor Covenant, version 1.4,
+available at
+https://www.contributor-covenant.org/version/1/4/code-of-conduct.html

CONTRIBUTING.md

@@ -0,0 +1,45 @@
+# How to Contribute
+
+We'd love to accept your patches and contributions to this project. There are
+just a few small guidelines you need to follow.
+
+## Contributor License Agreement
+
+Contributions to this project must be accompanied by a Contributor License
+Agreement. You (or your employer) retain the copyright to your contribution;
+this simply gives us permission to use and redistribute your contributions as
+part of the project. Head over to <https://cla.developers.google.com/> to see
+your current agreements on file or to sign a new one.
+
+You generally only need to submit a CLA once, so if you've already submitted one
+(even if it was for a different project), you probably don't need to do it
+again.
+
+## Code reviews
+
+All submissions, including submissions by project members, require review. We
+use [Dawn's Gerrit](https://dawn-review.googlesource.com/) for this purpose.
+
+Submissions should follow the [Tint style guide](docs/tint/style_guide.md).
+
+## Pushing to Gerrit
+
+Each change requires a `Change-Id` field in the commit message, which is generated by the [Gerrit commit-msg hook](](https://gerrit-review.googlesource.com/Documentation/cmd-hook-commit-msg.html)). \
+In a bash terminal, with the current path set to your tint source tree, this can be obtained by running the following:
+
+```bash
+f=`git rev-parse --git-dir`/hooks/commit-msg ; mkdir -p $(dirname $f) ; curl -Lo $f https://gerrit-review.googlesource.com/tools/hooks/commit-msg ; chmod +x $f
+```
+
+If you've already locally committed a change without the `Change-Id`, running `git commit --amend` will add the missing `Change-Id`.
+
+To create a Gerrit change for review, type:
+
+```bash
+git push origin HEAD:refs/for/main
+```
+
+## Community Guidelines
+
+This project follows
+[Google's Open Source Community Guidelines](https://opensource.google.com/conduct/).

CPPLINT.cfg

@@ -0,0 +1 @@
+set noparent

DEPS

@@ -42,7 +42,6 @@ deps = {
     'url': '{chromium_git}/external/github.com/llvm/llvm-project/clang/tools/clang-format.git@99803d74e35962f63a775f29477882afd4d57d94',
     'condition': 'dawn_standalone',
   },
-
   'buildtools/linux64': {
     'packages': [{
       'package': 'gn/gn/linux-amd64',
@@ -116,11 +115,6 @@ deps = {
     'condition': 'dawn_standalone',
   },
 
-  # WGSL support
-  'third_party/tint': {
-    'url': '{dawn_git}/tint@a730eb738e9f00fb52e9ac38cebe978373602a1e',
-  },
-
   # GLFW for tests and samples
   'third_party/glfw': {
     'url': '{chromium_git}/external/github.com/glfw/glfw@94773111300fee0453844a4c9407af7e880b4df8',
@@ -176,6 +170,10 @@ deps = {
     'url': '{github_git}/gpuweb/gpuweb.git@881403b5fda2d9ac9ffc5daa24e34738205bf155',
     'condition': 'dawn_node',
   },
+  'third_party/gpuweb-cts': {
+    'url': '{chromium_git}/external/github.com/gpuweb/cts@b0291fd966b55a5efc496772555b94842bde1085',
+    'condition': 'dawn_standalone',
+  },
 
   'tools/golang': {
     'condition': 'dawn_node',
@@ -194,6 +192,16 @@ deps = {
     }],
     'dep_type': 'cipd',
   },
+
+  # Misc dependencies inherited from Tint
+  'third_party/benchmark': {
+    'url': '{chromium_git}/external/github.com/google/benchmark.git@e991355c02b93fe17713efe04cbc2e278e00fdbd',
+    'condition': 'dawn_standalone',
+  },
+  'third_party/protobuf': {
+    'url': '{chromium_git}/external/github.com/protocolbuffers/protobuf.git@fde7cf7358ec7cd69e8db9be4f1fa6a5c431386a',
+    'condition': 'dawn_standalone',
+  },
 }
 
 hooks = [
@@ -253,33 +261,22 @@ hooks = [
     'condition': 'dawn_standalone and host_os == "win"',
     'action': [ 'download_from_google_storage',
                 '--no_resume',
+                '--platform=win32',
                 '--no_auth',
                 '--bucket', 'chromium-clang-format',
                 '-s', 'buildtools/win/clang-format.exe.sha1',
     ],
   },
   {
-    'name': 'clang_format_mac_x64',
+    'name': 'clang_format_mac',
     'pattern': '.',
-    'condition': 'dawn_standalone and host_os == "mac" and host_cpu == "x64"',
+    'condition': 'dawn_standalone and host_os == "mac"',
     'action': [ 'download_from_google_storage',
                 '--no_resume',
+                '--platform=darwin',
                 '--no_auth',
                 '--bucket', 'chromium-clang-format',
-                '-s', 'buildtools/mac/clang-format.x64.sha1',
+                '-s', 'buildtools/mac/clang-format.sha1',
-                '-o', 'buildtools/mac/clang-format',
-    ],
-  },
-  {
-    'name': 'clang_format_mac_arm64',
-    'pattern': '.',
-    'condition': 'dawn_standalone and host_os == "mac" and host_cpu == "arm64"',
-    'action': [ 'download_from_google_storage',
-                '--no_resume',
-                '--no_auth',
-                '--bucket', 'chromium-clang-format',
-                '-s', 'buildtools/mac/clang-format.arm64.sha1',
-                '-o', 'buildtools/mac/clang-format',
     ],
   },
   {
@@ -288,11 +285,59 @@ hooks = [
     'condition': 'dawn_standalone and host_os == "linux"',
     'action': [ 'download_from_google_storage',
                 '--no_resume',
+                '--platform=linux*',
                 '--no_auth',
                 '--bucket', 'chromium-clang-format',
                 '-s', 'buildtools/linux64/clang-format.sha1',
     ],
   },
+  # Pull the compilers and system libraries for hermetic builds
+  {
+    'name': 'sysroot_x86',
+    'pattern': '.',
+    'condition': 'checkout_linux and ((checkout_x86 or checkout_x64))',
+    'action': ['python3', 'build/linux/sysroot_scripts/install-sysroot.py',
+               '--arch=x86'],
+  },
+  {
+    'name': 'sysroot_x64',
+    'pattern': '.',
+    'condition': 'checkout_linux and (checkout_x64)',
+    'action': ['python3', 'build/linux/sysroot_scripts/install-sysroot.py',
+               '--arch=x64'],
+  },
+  {
+    # Update the Mac toolchain if necessary.
+    'name': 'mac_toolchain',
+    'pattern': '.',
+    'condition': 'checkout_mac',
+    'action': ['python3', 'build/mac_toolchain.py'],
+  },
+  {
+    # Update the Windows toolchain if necessary. Must run before 'clang' below.
+    'name': 'win_toolchain',
+    'pattern': '.',
+    'condition': 'checkout_win',
+    'action': ['python3', 'build/vs_toolchain.py', 'update', '--force'],
+  },
+  {
+    # Note: On Win, this should run after win_toolchain, as it may use it.
+    'name': 'clang',
+    'pattern': '.',
+    'action': ['python3', 'tools/clang/scripts/update.py'],
+  },
+  {
+    # Pull rc binaries using checked-in hashes.
+    'name': 'rc_win',
+    'pattern': '.',
+    'condition': 'checkout_win and (host_os == "win")',
+    'action': [ 'download_from_google_storage',
+                '--no_resume',
+                '--no_auth',
+                '--bucket', 'chromium-browser-clang/rc',
+                '-s', 'build/toolchain/win/rc/win/rc.exe.sha1',
+    ],
+  },
   # Update build/util/LASTCHANGE.
   {
     'name': 'lastchange',

OWNERS

@@ -9,3 +9,12 @@ kainino@chromium.org
 per-file dawn.json=kainino@chromium.org
 per-file DEPS=*
 per-file README.md=file://docs/dawn/OWNERS
+
+# Tint specific OWNERS
+amaiorano@google.com
+bclayton@chromium.org
+bclayton@google.com
+cwallez@chromium.org
+dneto@google.com
+jrprice@google.com
+rharrison@chromium.org

OWNERS.dawn

@@ -0,0 +1,11 @@
+cwallez@chromium.org
+enga@chromium.org
+jiawei.shao@intel.com
+
+# Backup reviewers if needed.
+bclayton@google.com
+kainino@chromium.org
+
+per-file dawn.json=kainino@chromium.org
+per-file DEPS=*
+per-file README.md=file://docs/dawn/OWNERS

OWNERS.tint

@@ -0,0 +1,7 @@
+amaiorano@google.com
+bclayton@chromium.org
+bclayton@google.com
+cwallez@chromium.org
+dneto@google.com
+jrprice@google.com
+rharrison@chromium.org

PRESUBMIT.py

@@ -1,4 +1,4 @@
-# Copyright 2018 The Dawn Authors
+# Copyright 2022 The Dawn & Tint Authors
 #
 # Licensed under the Apache License, Version 2.0 (the "License");
 # you may not use this file except in compliance with the License.
@@ -12,12 +12,97 @@
 # See the License for the specific language governing permissions and
 # limitations under the License.
 
-import os
+import re
-import platform
-import subprocess
 
 USE_PYTHON3 = True
 
+NONINCLUSIVE_REGEXES = [
+    r"(?i)black[-_]?list",
+    r"(?i)white[-_]?list",
+    r"(?i)gr[ea]y[-_]?list",
+    r"(?i)(first class citizen)",
+    r"(?i)black[-_]?hat",
+    r"(?i)white[-_]?hat",
+    r"(?i)gr[ea]y[-_]?hat",
+    r"(?i)master",
+    r"(?i)slave",
+    r"(?i)\bhim\b",
+    r"(?i)\bhis\b",
+    r"(?i)\bshe\b",
+    r"(?i)\bher\b",
+    r"(?i)\bguys\b",
+    r"(?i)\bhers\b",
+    r"(?i)\bman\b",
+    r"(?i)\bwoman\b",
+    r"(?i)\she\s",
+    r"(?i)\she$",
+    r"(?i)^he\s",
+    r"(?i)^he$",
+    r"(?i)\she['|\u2019]d\s",
+    r"(?i)\she['|\u2019]d$",
+    r"(?i)^he['|\u2019]d\s",
+    r"(?i)^he['|\u2019]d$",
+    r"(?i)\she['|\u2019]s\s",
+    r"(?i)\she['|\u2019]s$",
+    r"(?i)^he['|\u2019]s\s",
+    r"(?i)^he['|\u2019]s$",
+    r"(?i)\she['|\u2019]ll\s",
+    r"(?i)\she['|\u2019]ll$",
+    r"(?i)^he['|\u2019]ll\s",
+    r"(?i)^he['|\u2019]ll$",
+    r"(?i)grandfather",
+    r"(?i)\bmitm\b",
+    r"(?i)\bcrazy\b",
+    r"(?i)\binsane\b",
+    r"(?i)\bblind\sto\b",
+    r"(?i)\bflying\sblind\b",
+    r"(?i)\bblind\seye\b",
+    r"(?i)\bcripple\b",
+    r"(?i)\bcrippled\b",
+    r"(?i)\bdumb\b",
+    r"(?i)\bdummy\b",
+    r"(?i)\bparanoid\b",
+    r"(?i)\bsane\b",
+    r"(?i)\bsanity\b",
+    r"(?i)red[-_]?line",
+]
+
+NONINCLUSIVE_REGEX_LIST = []
+for reg in NONINCLUSIVE_REGEXES:
+    NONINCLUSIVE_REGEX_LIST.append(re.compile(reg))
+
+
+def _CheckNonInclusiveLanguage(input_api, output_api, source_file_filter=None):
+    """Checks the files for non-inclusive language."""
+
+    matches = []
+    for f in input_api.AffectedFiles(include_deletes=False,
+                                     file_filter=source_file_filter):
+        for line_num, line in f.ChangedContents():
+            for reg in NONINCLUSIVE_REGEX_LIST:
+                match = reg.search(line)
+                if match:
+                    matches.append(
+                        "{} ({}): found non-inclusive language: {}".format(
+                            f.LocalPath(), line_num, match.group(0)))
+
+    if len(matches):
+        return [
+            output_api.PresubmitPromptWarning('Non-inclusive language found:',
+                                              items=matches)
+        ]
+
+    return []
+
+
+def _NonInclusiveFileFilter(file):
+    filter_list = [
+        "PRESUBMIT.py",  # Non-inclusive language check data
+        "docs/tint/spirv-input-output-variables.md",  # External URL
+        "test/tint/samples/compute_boids.wgsl ",  # External URL
+    ]
+    return file in filter_list
+
 
 def _DoCommonChecks(input_api, output_api):
     results = []
@@ -27,6 +112,30 @@ def _DoCommonChecks(input_api, output_api):
         input_api.canned_checks.CheckPatchFormatted(input_api,
                                                     output_api,
                                                     check_python=True))
+    results.extend(
+        input_api.canned_checks.CheckChangeHasDescription(
+            input_api, output_api))
+    results.extend(
+        input_api.canned_checks.CheckGNFormatted(input_api, output_api))
+    results.extend(
+        input_api.canned_checks.CheckChangeHasNoCrAndHasOnlyOneEol(
+            input_api, output_api))
+    results.extend(
+        input_api.canned_checks.CheckChangeHasNoTabs(input_api, output_api))
+    results.extend(
+        input_api.canned_checks.CheckChangeTodoHasOwner(input_api, output_api))
+    results.extend(
+        input_api.canned_checks.CheckChangeHasNoStrayWhitespace(
+            input_api, output_api))
+    results.extend(
+        input_api.canned_checks.CheckDoNotSubmit(input_api, output_api))
+    results.extend(
+        input_api.canned_checks.CheckChangeLintsClean(input_api,
+                                                      output_api,
+                                                      lint_filters=""))
+    results.extend(
+        _CheckNonInclusiveLanguage(input_api, output_api,
+                                   _NonInclusiveFileFilter))
     return results
 
 

PRESUBMIT.py.dawn

@@ -0,0 +1,38 @@
+# Copyright 2018 The Dawn Authors
+#
+# 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 os
+import platform
+import subprocess
+
+USE_PYTHON3 = True
+
+
+def _DoCommonChecks(input_api, output_api):
+    results = []
+    results.extend(
+        input_api.canned_checks.CheckChangedLUCIConfigs(input_api, output_api))
+    results.extend(
+        input_api.canned_checks.CheckPatchFormatted(input_api,
+                                                    output_api,
+                                                    check_python=True))
+    return results
+
+
+def CheckChangeOnUpload(input_api, output_api):
+    return _DoCommonChecks(input_api, output_api)
+
+
+def CheckChangeOnCommit(input_api, output_api):
+    return _DoCommonChecks(input_api, output_api)

PRESUBMIT.py.tint

@@ -0,0 +1,167 @@
+# Copyright 2020 The Tint Authors
+#
+# 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.
+"""Presubmit script for Tint.
+See http://dev.chromium.org/developers/how-tos/depottools/presubmit-scripts
+for more details about the presubmit API built into depot_tools.
+"""
+
+import re
+
+USE_PYTHON3 = True
+
+
+def _LicenseHeader(input_api):
+    """Returns the license header regexp."""
+    # Accept any year number from 2019 to the current year
+    current_year = int(input_api.time.strftime('%Y'))
+    allowed_years = (str(s) for s in reversed(xrange(2019, current_year + 1)))
+    years_re = '(' + '|'.join(allowed_years) + ')'
+    license_header = (
+        r'.*? Copyright( \(c\))? %(year)s The Tint [Aa]uthors\n '
+        r'.*?\n'
+        r'.*? Licensed under the Apache License, Version 2.0 (the "License");\n'
+        r'.*? you may not use this file except in compliance with the License.\n'
+        r'.*? You may obtain a copy of the License at\n'
+        r'.*?\n'
+        r'.*?     http://www.apache.org/licenses/LICENSE-2.0\n'
+        r'.*?\n'
+        r'.*? Unless required by applicable law or agreed to in writing, software\n'
+        r'.*? distributed under the License is distributed on an "AS IS" BASIS,\n'
+        r'.*? WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.\n'
+        r'.*? See the License for the specific language governing permissions and\n'
+        r'.*? limitations under the License.\n') % {
+            'year': years_re,
+        }
+    return license_header
+
+
+REGEXES = [
+    r"(?i)black[-_]?list",
+    r"(?i)white[-_]?list",
+    r"(?i)gr[ea]y[-_]?list",
+    r"(?i)(first class citizen)",
+    r"(?i)black[-_]?hat",
+    r"(?i)white[-_]?hat",
+    r"(?i)gr[ea]y[-_]?hat",
+    r"(?i)master",
+    r"(?i)slave",
+    r"(?i)\bhim\b",
+    r"(?i)\bhis\b",
+    r"(?i)\bshe\b",
+    r"(?i)\bher\b",
+    r"(?i)\bguys\b",
+    r"(?i)\bhers\b",
+    r"(?i)\bman\b",
+    r"(?i)\bwoman\b",
+    r"(?i)\she\s",
+    r"(?i)\she$",
+    r"(?i)^he\s",
+    r"(?i)^he$",
+    r"(?i)\she['|\u2019]d\s",
+    r"(?i)\she['|\u2019]d$",
+    r"(?i)^he['|\u2019]d\s",
+    r"(?i)^he['|\u2019]d$",
+    r"(?i)\she['|\u2019]s\s",
+    r"(?i)\she['|\u2019]s$",
+    r"(?i)^he['|\u2019]s\s",
+    r"(?i)^he['|\u2019]s$",
+    r"(?i)\she['|\u2019]ll\s",
+    r"(?i)\she['|\u2019]ll$",
+    r"(?i)^he['|\u2019]ll\s",
+    r"(?i)^he['|\u2019]ll$",
+    r"(?i)grandfather",
+    r"(?i)\bmitm\b",
+    r"(?i)\bcrazy\b",
+    r"(?i)\binsane\b",
+    r"(?i)\bblind\sto\b",
+    r"(?i)\bflying\sblind\b",
+    r"(?i)\bblind\seye\b",
+    r"(?i)\bcripple\b",
+    r"(?i)\bcrippled\b",
+    r"(?i)\bdumb\b",
+    r"(?i)\bdummy\b",
+    r"(?i)\bparanoid\b",
+    r"(?i)\bsane\b",
+    r"(?i)\bsanity\b",
+    r"(?i)red[-_]?line",
+]
+
+REGEX_LIST = []
+for reg in REGEXES:
+    REGEX_LIST.append(re.compile(reg))
+
+def CheckNonInclusiveLanguage(input_api, output_api, source_file_filter=None):
+    """Checks the files for non-inclusive language."""
+
+    matches = []
+    for f in input_api.AffectedFiles(include_deletes=False,
+                                     file_filter=source_file_filter):
+        for line_num, line in f.ChangedContents():
+            for reg in REGEX_LIST:
+                match = reg.search(line)
+                if match:
+                    matches.append(
+                        "{} ({}): found non-inclusive language: {}".format(
+                            f.LocalPath(), line_num, match.group(0)))
+
+    if len(matches):
+        return [
+            output_api.PresubmitPromptWarning('Non-inclusive language found:',
+                                              items=matches)
+        ]
+
+    return []
+
+
+def CheckChange(input_api, output_api):
+    results = []
+
+    results += input_api.canned_checks.CheckChangeHasDescription(
+        input_api, output_api)
+    results += input_api.canned_checks.CheckPatchFormatted(input_api,
+                                                           output_api,
+                                                           check_python=True)
+    results += input_api.canned_checks.CheckGNFormatted(input_api, output_api)
+    results += input_api.canned_checks.CheckChangeHasNoCrAndHasOnlyOneEol(
+        input_api, output_api)
+    results += input_api.canned_checks.CheckChangeHasNoTabs(
+        input_api, output_api)
+    results += input_api.canned_checks.CheckChangeTodoHasOwner(
+        input_api, output_api)
+    results += input_api.canned_checks.CheckChangeHasNoStrayWhitespace(
+        input_api, output_api)
+    results += input_api.canned_checks.CheckDoNotSubmit(input_api, output_api)
+    results += input_api.canned_checks.CheckChangeLintsClean(input_api,
+                                                             output_api,
+                                                             lint_filters="")
+
+    def NonInclusiveFileFilter(file):
+        filter_list = [
+            "docs/tint/spirv-input-output-variables.md",  # External URL
+            "test/tint/samples/compute_boids.wgsl ",  # External URL
+        ]
+        return file in filter_list
+
+    results += CheckNonInclusiveLanguage(input_api, output_api,
+                                         NonInclusiveFileFilter)
+
+    return results
+
+
+def CheckChangeOnUpload(input_api, output_api):
+    return CheckChange(input_api, output_api)
+
+
+def CheckChangeOnCommit(input_api, output_api):
+    return CheckChange(input_api, output_api)

README.md

@@ -50,3 +50,110 @@ Apache 2.0 Public License, please see [LICENSE](/LICENSE).
 ## Disclaimer
 
 This is not an officially supported Google product.
+
+# Tint
+
+Tint is a compiler for the WebGPU Shader Language (WGSL).
+
+This is not an officially supported Google product.
+
+## Requirements
+ * Git
+ * CMake (3.10.2 or later)
+ * Ninja (or other build tool)
+ * Python, for fetching dependencies
+ * [depot_tools] in your path
+
+## Build options
+ * `TINT_BUILD_SPV_READER` : enable the SPIR-V input reader (off by default)
+ * `TINT_BUILD_WGSL_READER` : enable the WGSL input reader (on by default)
+ * `TINT_BUILD_SPV_WRITER` : enable the SPIR-V output writer (on by default)
+ * `TINT_BUILD_WGSL_WRITER` : enable the WGSL output writer (on by default)
+ * `TINT_BUILD_FUZZERS` : enable building fuzzzers (off by default)
+
+## Building
+Tint uses Chromium dependency management so you need to install [depot_tools]
+and add it to your PATH.
+
+[depot_tools]: http://commondatastorage.googleapis.com/chrome-infra-docs/flat/depot_tools/docs/html/depot_tools_tutorial.html#_setting_up
+
+### Getting source & dependencies
+
+```sh
+# Clone the repo as "tint"
+git clone https://dawn.googlesource.com/tint tint
+cd tint
+
+# Bootstrap the gclient configuration
+cp standalone.gclient .gclient
+
+# Fetch external dependencies and toolchains with gclient
+gclient sync
+```
+
+### Compiling using CMake + Ninja
+```sh
+mkdir -p out/Debug
+cd out/Debug
+cmake -GNinja ../..
+ninja # or autoninja
+```
+
+### Compiling using CMake + make
+```sh
+mkdir -p out/Debug
+cd out/Debug
+cmake ../..
+make # -j N for N-way parallel build
+```
+
+### Compiling using gn + ninja
+```sh
+mkdir -p out/Debug
+gn gen out/Debug
+autoninja -C out/Debug
+```
+
+### Fuzzers on MacOS
+If you are attempting fuzz, using `TINT_BUILD_FUZZERS=ON`, the version of llvm
+in the XCode SDK does not have the needed libfuzzer functionality included.
+
+The build error that you will see from using the XCode SDK will look something
+like this:
+```
+ld: file not found:/Applications/Xcode.app/Contents/Developer/Toolchains/XcodeDefault.xctoolchain/usr/lib/clang/11.0.0/lib/darwin/libclang_rt.fuzzer_osx.a
+```
+
+The solution to this problem is to use a full version llvm, like what you would
+get via homebrew, `brew install llvm`, and use something like `CC=<path to full
+clang> cmake ..` to setup a build using that toolchain.
+
+### Checking [chromium-style] issues in CMake builds
+The gn based work flow uses the Chromium toolchain for building in anticipation
+of integration of Tint into Chromium based projects. This toolchain has
+additional plugins for checking for style issues, which are marked with
+[chromium-style] in log messages. This means that this toolchain is more strict
+then the default clang toolchain.
+
+In the future we will have a CQ that will build this work flow and flag issues
+automatically. Until that is in place, to avoid causing breakages you can run
+the [chromium-style] checks using the CMake based work flows. This requires
+setting `CC` to the version of clang checked out by `gclient sync` and setting
+the `TINT_CHECK_CHROMIUM_STYLE` to `ON`.
+
+```sh
+mkdir -p out/style
+cd out/style
+cmake ../..
+CC=../../third_party/llvm-build/Release+Asserts/bin/clang cmake -DTINT_CHECK_CHROMIUM_STYLE=ON ../../ # add -GNinja for ninja builds
+```
+
+## Issues
+Please file any issues or feature requests at
+https://bugs.chromium.org/p/tint/issues/entry
+
+## Contributing
+Please see the CONTRIBUTING and CODE_OF_CONDUCT files on how to contribute to
+Tint.
+
+Tint has a process for supporting [experimental extensions](docs/tint/experimental_extensions.md).

README.md.dawn

@@ -0,0 +1,52 @@
+![Dawn's logo: a sun rising behind a stylized mountain inspired by the WebGPU logo. The text "Dawn" is written below it.](docs/imgs/dawn_logo.png "Dawn's logo")
+
+# Dawn, a WebGPU implementation
+
+Dawn is an open-source and cross-platform implementation of the work-in-progress [WebGPU](https://webgpu.dev) standard.
+More precisely it implements [`webgpu.h`](https://github.com/webgpu-native/webgpu-headers/blob/master/webgpu.h) that is a one-to-one mapping with the WebGPU IDL.
+Dawn is meant to be integrated as part of a larger system and is the underlying implementation of WebGPU in Chromium.
+
+Dawn provides several WebGPU building blocks:
+ - **WebGPU C/C++ headers** that applications and other building blocks use.
+   - The `webgpu.h` version that Dawn implements.
+   - A C++ wrapper for the `webgpu.h`.
+ - **A "native" implementation of WebGPU** using platforms' GPU APIs:
+   - **D3D12** on Windows 10
+   - **Metal** on macOS and iOS
+   - **Vulkan** on Windows, Linux, ChromeOS, Android and Fuchsia
+   - OpenGL as best effort where available
+ - **A client-server implementation of WebGPU** for applications that are in a sandbox without access to native drivers
+
+Helpful links:
+
+ - [Dawn's bug tracker](https://bugs.chromium.org/p/dawn/issues/entry) if you find issues with Dawn.
+ - [Dawn's mailing list](https://groups.google.com/forum/#!members/dawn-graphics) for other discussions related to Dawn.
+ - [Dawn's source code](https://dawn.googlesource.com/dawn)
+ - [Dawn's Matrix chatroom](https://matrix.to/#/#webgpu-dawn:matrix.org) for live discussion around contributing or using Dawn.
+ - [WebGPU's Matrix chatroom](https://matrix.to/#/#WebGPU:matrix.org)
+
+## Documentation table of content
+
+Developer documentation:
+
+ - [Dawn overview](docs/dawn/overview.md)
+ - [Building Dawn](docs/dawn/building.md)
+ - [Contributing to Dawn](docs/dawn/contributing.md)
+ - [Testing Dawn](docs/dawn/testing.md)
+ - [Debugging Dawn](docs/dawn/debugging.md)
+ - [Dawn's infrastructure](docs/dawn/infra.md)
+ - [Dawn errors](docs/dawn/errors.md)
+
+User documentation: (TODO, figure out what overlaps with the webgpu.h docs)
+
+## Status
+
+(TODO)
+
+## License
+
+Apache 2.0 Public License, please see [LICENSE](/LICENSE).
+
+## Disclaimer
+
+This is not an officially supported Google product.

README.md.tint

@@ -0,0 +1,106 @@
+# Tint
+
+Tint is a compiler for the WebGPU Shader Language (WGSL).
+
+This is not an officially supported Google product.
+
+## Requirements
+ * Git
+ * CMake (3.10.2 or later)
+ * Ninja (or other build tool)
+ * Python, for fetching dependencies
+ * [depot_tools] in your path
+
+## Build options
+ * `TINT_BUILD_SPV_READER` : enable the SPIR-V input reader (off by default)
+ * `TINT_BUILD_WGSL_READER` : enable the WGSL input reader (on by default)
+ * `TINT_BUILD_SPV_WRITER` : enable the SPIR-V output writer (on by default)
+ * `TINT_BUILD_WGSL_WRITER` : enable the WGSL output writer (on by default)
+ * `TINT_BUILD_FUZZERS` : enable building fuzzzers (off by default)
+
+## Building
+Tint uses Chromium dependency management so you need to install [depot_tools]
+and add it to your PATH.
+
+[depot_tools]: http://commondatastorage.googleapis.com/chrome-infra-docs/flat/depot_tools/docs/html/depot_tools_tutorial.html#_setting_up
+
+### Getting source & dependencies
+
+```sh
+# Clone the repo as "tint"
+git clone https://dawn.googlesource.com/tint tint
+cd tint
+
+# Bootstrap the gclient configuration
+cp standalone.gclient .gclient
+
+# Fetch external dependencies and toolchains with gclient
+gclient sync
+```
+
+### Compiling using CMake + Ninja
+```sh
+mkdir -p out/Debug
+cd out/Debug
+cmake -GNinja ../..
+ninja # or autoninja
+```
+
+### Compiling using CMake + make
+```sh
+mkdir -p out/Debug
+cd out/Debug
+cmake ../..
+make # -j N for N-way parallel build
+```
+
+### Compiling using gn + ninja
+```sh
+mkdir -p out/Debug
+gn gen out/Debug
+autoninja -C out/Debug
+```
+
+### Fuzzers on MacOS
+If you are attempting fuzz, using `TINT_BUILD_FUZZERS=ON`, the version of llvm
+in the XCode SDK does not have the needed libfuzzer functionality included.
+
+The build error that you will see from using the XCode SDK will look something
+like this:
+```
+ld: file not found:/Applications/Xcode.app/Contents/Developer/Toolchains/XcodeDefault.xctoolchain/usr/lib/clang/11.0.0/lib/darwin/libclang_rt.fuzzer_osx.a
+```
+
+The solution to this problem is to use a full version llvm, like what you would
+get via homebrew, `brew install llvm`, and use something like `CC=<path to full
+clang> cmake ..` to setup a build using that toolchain.
+
+### Checking [chromium-style] issues in CMake builds
+The gn based work flow uses the Chromium toolchain for building in anticipation
+of integration of Tint into Chromium based projects. This toolchain has
+additional plugins for checking for style issues, which are marked with
+[chromium-style] in log messages. This means that this toolchain is more strict
+then the default clang toolchain.
+
+In the future we will have a CQ that will build this work flow and flag issues
+automatically. Until that is in place, to avoid causing breakages you can run
+the [chromium-style] checks using the CMake based work flows. This requires
+setting `CC` to the version of clang checked out by `gclient sync` and setting
+the `TINT_CHECK_CHROMIUM_STYLE` to `ON`.
+
+```sh
+mkdir -p out/style
+cd out/style
+cmake ../..
+CC=../../third_party/llvm-build/Release+Asserts/bin/clang cmake -DTINT_CHECK_CHROMIUM_STYLE=ON ../../ # add -GNinja for ninja builds
+```
+
+## Issues
+Please file any issues or feature requests at
+https://bugs.chromium.org/p/tint/issues/entry
+
+## Contributing
+Please see the CONTRIBUTING and CODE_OF_CONDUCT files on how to contribute to
+Tint.
+
+Tint has a process for supporting [experimental extensions](docs/tint/experimental_extensions.md).

build_overrides/build.gni

@@ -1,4 +1,4 @@
-# Copyright 2018 The Dawn Authors
+# Copyright 2022 The Dawn Authors
 #
 # Licensed under the Apache License, Version 2.0 (the "License");
 # you may not use this file except in compliance with the License.

build_overrides/dawn.gni

@@ -34,7 +34,6 @@ dawn_glfw_dir = "//third_party/glfw"
 dawn_googletest_dir = "//third_party/googletest"
 dawn_spirv_tools_dir = "//third_party/vulkan-deps/spirv-tools/src"
 dawn_swiftshader_dir = "//third_party/swiftshader"
-dawn_tint_dir = "//third_party/tint"
 dawn_vulkan_loader_dir = "//third_party/vulkan-deps/vulkan-loader/src"
 dawn_vulkan_validation_layers_dir =
     "//third_party/vulkan-deps/vulkan-validation-layers/src"

build_overrides/tint.gni

@@ -12,7 +12,6 @@
 # See the License for the specific language governing permissions and
 # limitations under the License.
 
-tint_root_dir = "//third_party/tint"
 tint_spirv_tools_dir = "//third_party/vulkan-deps/spirv-tools/src"
 tint_spirv_headers_dir = "//third_party/vulkan-deps/spirv-headers/src"
 

docs/tint/arch.md

@@ -0,0 +1,195 @@
+# Tint Architecture
+
+```
+                   ┏━━━━━━━━┓                   ┏━━━━━━┓
+                   ┃ SPIR━V ┃                   ┃ WGSL ┃
+                   ┗━━━━┃━━━┛                   ┗━━━┃━━┛
+                        ▼                           ▼
+              ┏━━━━━━━━━┃━━━━━━━━━━━━━━━━━━━━━━━━━━━┃━━━━━━━━┓
+              ┃         ┃          Reader           ┃        ┃
+              ┃         ┃                           ┃        ┃
+              ┃ ┏━━━━━━━┻━━━━━━┓             ┏━━━━━━┻━━━━━━┓ ┃
+              ┃ ┃ SPIRV-Reader ┃             ┃ WGSL-Reader ┃ ┃
+              ┃ ┗━━━━━━━━━━━━━━┛             ┗━━━━━━━━━━━━━┛ ┃
+              ┗━━━━━━━━━━━━━━━━━━━━━━━┳━━━━━━━━━━━━━━━━━━━━━━┛
+                                      ▼
+                    ┏━━━━━━━━━━━━━━━━━┻━━━━━━━━━━━━━━━━━┓
+                    ┃           ProgramBuilder          ┃
+                    ┃             (mutable)             ┃
+      ┏━━━━━━━━━━━━►┫   ┏━━━━━┓ ┏━━━━━━━┓ ┏━━━━━━━━━┓   ┃
+      ┃             ┃   ┃ AST ┃ ┃ Types ┃ ┃ Symbols ┃   ┃
+      ┃             ┃   ┗━━━━━┛ ┗━━━━━━━┛ ┗━━━━━━━━━┛   ┃
+      ┃             ┗━━━━━━━━━━━━━━━━━┳━━━━━━━━━━━━━━━━━┛
+      ┃                               ▼
+      ┃             ┌┄┄┄┄┄┄┄┄┄┄┄┄┄┄┄┄┄┃┄┄┄┄┄┄┄┄┄┄┄┄┄┄┄┄┐
+      ▲             ┆ Build           ▼                ┆
+  ┏━━━┻━━━┓         ┆        ┏━━━━━━━━┻━━━━━━━━┓       ┆
+  ┃ Clone ┃         ┆        ┃    Resolver     ┃       ┆
+  ┗━━━┳━━━┛         ┆        ┗━━━━━━━━━━━━━━━━━┛       ┆
+      ▲             └┄┄┄┄┄┄┄┄┄┄┄┄┄┄┄┄┄┃┄┄┄┄┄┄┄┄┄┄┄┄┄┄┄┄┘
+      ┃                               ▼
+      ┃       ┏━━━━━━━━━━━━━━━━━━━━━━━┻━━━━━━━━━━━━━━━━━━━━━━┓
+      ┃       ┃                    Program                   ┃
+      ┃       ┃                  (immutable)                 ┃
+      ┣━━━━━━◄┫  ┏━━━━━┓ ┏━━━━━━━┓ ┏━━━━━━━━━━┓ ┏━━━━━━━━━┓  ┃
+      ┃       ┃  ┃ AST ┃ ┃ Types ┃ ┃ Semantic ┃ ┃ Symbols ┃  ┃
+      ┃       ┃  ┗━━━━━┛ ┗━━━━━━━┛ ┗━━━━━━━━━━┛ ┗━━━━━━━━━┛  ┃
+      ┃       ┗━━━━━━━━━━━━━━━━━━━━━━━┳━━━━━━━━━━━━━━━━━━━━━━┛
+      ▲                               ▼
+┏━━━━━┻━━━━━┓                         ┃             ┏━━━━━━━━━━━┓
+┃ Transform ┃◄━━━━━━━━━━━━━━━━━━━━━━━━╋━━━━━━━━━━━━►┃ Inspector ┃
+┗━━━━━━━━━━━┛                         ┃             ┗━━━━━━━━━━━┛
+                                      ▼
+┏━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━┻━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━┓
+┃                                  Writers                                    ┃
+┃                                                                             ┃
+┃ ┏━━━━━━━━━━━━━━┓┏━━━━━━━━━━━━━┓┏━━━━━━━━━━━━━┓┏━━━━━━━━━━━━━┓┏━━━━━━━━━━━━┓ ┃
+┃ ┃ SPIRV-Writer ┃┃ WGSL-Writer ┃┃ HLSL-Writer ┃┃ GLSL-Writer ┃┃ MSL-Writer ┃ ┃
+┃ ┗━━━━━━━┳━━━━━━┛┗━━━━━━┳━━━━━━┛┗━━━━━━┳━━━━━━┛┗━━━━━━┳━━━━━━┛┗━━━━━━┳━━━━━┛ ┃
+┗━━━━━━━━━┃━━━━━━━━━━━━━━┃━━━━━━━━━━━━━━┃━━━━━━━━━━━━━━┃━━━━━━━━━━━━━━┃━━━━━━━┛
+          ▼              ▼              ▼              ▼              ▼
+     ┏━━━━┻━━━┓      ┏━━━┻━━┓       ┏━━━┻━━┓       ┏━━━┻━━┓        ┏━━┻━━┓
+     ┃ SPIR-V ┃      ┃ WGSL ┃       ┃ HLSL ┃       ┃ GLSL ┃        ┃ MSL ┃
+     ┗━━━━━━━━┛      ┗━━━━━━┛       ┗━━━━━━┛       ┗━━━━━━┛        ┗━━━━━┛
+```
+
+## Reader
+
+Readers are responsible for parsing a shader program and populating a
+`ProgramBuilder` with the parsed AST, type and symbol information.
+
+The WGSL reader is a recursive descent parser. It closely follows the WGSL
+grammar in the naming of the parse methods.
+
+## ProgramBuilder
+
+A `ProgramBuilder` is the primary interface to construct an immutable `Program`.
+There are a number of methods exposed which make creating of the `Program`
+simpler. A `ProgramBuilder` can only be used once, and must be discarded after
+the `Program` is constructed.
+
+A `Program` is built from the `ProgramBuilder` by `std::move()`ing the
+`ProgramBuilder` to a new `Program` object. When built, resolution is performed
+so the produced `Program` will contain all the needed semantic information.
+
+At any time before building the `Program`, `ProgramBuilder::IsValid()` may be
+called to ensure the AST is **structurally** correct. This checks that things
+like `if` statements have a condition and body attached.
+
+If further changes to the `Program` are needed (say via a `Transform`) then a
+new `ProgramBuilder` can be produced by cloning the `Program` into a new
+`ProgramBuilder`.
+
+Unlike `Program`s, `ProgramBuilder`s are not part of the public Tint API.
+
+## AST
+
+The Abstract Syntax Tree is a directed acyclic graph of `ast::Node`s which
+encode the syntactic structure of the WGSL program.
+
+The root of the AST is the `ast::Module` class which holds each of the declared
+functions, variables and user defined types (type aliases and structures).
+
+Each `ast::Node` represents a **single** part of the program's source, and so
+`ast::Node`s are not shared.
+
+The AST does not perform any verification of its content. For example, the
+`ast::StrideAttribute` node has numeric stride parameter, which is a count of
+the number of bytes from the start of one array element to the start of the
+next. The AST node itself does not constrain the set of stride values that you
+can set, aside from storing it as an unsigned integer.
+
+## Types
+
+Types are constructed during the Reader and resolution phases, and are
+held by the `Program` or `ProgramBuilder`. AST and semantic nodes can both
+reference types.
+
+Each `type::Type` node **uniquely** represents a particular spelling of a WGSL
+type within the program, so you can compare `type::Type*` pointers to check for
+equivalence of type expressions.
+For example, there is only one `type::Type` node for the `i32` type, no matter
+how many times it is mentioned in the source program.
+However, if `MyI32` is a type alias for `i32`, then they will have two different
+type nodes.
+
+## Semantic information
+
+Semantic information is held by `sem::Node`s which describe the program at
+a higher / more abstract level than the AST. This includes information such as
+the resolved type of each expression, the resolved overload of a builtin
+function call, and the module scoped variables used by each function.
+
+Semantic information is generated by the `Resolver` when the `Program`
+is built from a `ProgramBuilder`.
+
+The `sem::Info` class holds a map of `ast::Node`s to `sem::Node`s.
+This map is **many-to-one** - i.e. while a AST node might have a single
+corresponding semantic node, the reverse may not be true. For example:
+many `ast::IdentifierExpression` nodes may map to a single `sem::Variable`,
+and so the `sem::Variable` does not have a single corresponding
+`ast::Node`.
+
+Unlike `ast::Node`s, semantic nodes may not necessarily form a directed acyclic
+graph, and the semantic graph may contain diamonds.
+
+## Symbols
+
+Symbols represent a unique string identifier in the source program. These string
+identifiers are transformed into symbols within the `Reader`s.
+
+During the Writer phase, symbols may be emitted as strings using a `Namer`.
+A `Namer` may output the symbol in any form that preserves the uniqueness of
+that symbol.
+
+## Resolver
+
+The `Resolver` will automatically run when a `Program` is built.
+A `Resolver` creates the `Program`s semantic information by analyzing the
+`Program`s AST and type information.
+
+The `Resolver` will validate to make sure the generated `Program` is
+semantically valid.
+
+## Program
+
+A `Program` holds an immutable version of the information from the
+`ProgramBuilder` along with semantic information generated by the
+`Resolver`.
+
+Like `ProgramBuilder`, `Program::IsValid()` may be called to ensure the AST is
+structurally correct and semantically valid, and that the `Resolver` did not
+report any errors.
+
+Unlike the `ProgramBuilder`, a `Program` is fully immutable, and is part of the
+public Tint API. The immutable nature of `Program`s make these entirely safe
+to share between multiple threads without the use of synchronization primitives.
+
+## Inspector
+
+The inspectors job is to go through the `Program` and pull out various pieces of
+information. The information may be used to pass information into the downstream
+compilers (things like specialization constants) or may be used to pass into
+transforms to update the AST before generating the resulting code.
+
+The input `Program` to the inspector must be valid (pass validation).
+
+## Transforms
+
+There maybe various transforms we want to run over the `Program`.
+This is for things like Vertex Pulling or Robust Buffer Access.
+
+A transform operates by cloning the input `Program` into a new `ProgramBuilder`,
+applying the required changes, and then finally building and returning a new
+output `Program`. As the resolver is always run when a `Program` is built,
+Transforms will always emit a `Program` with semantic information.
+
+The input `Program` to a transform must be valid (pass validation).
+If the input `Program` of a transform is valid then the transform must guarantee
+that the output program is also valid.
+
+## Writers
+
+A writer is responsible for writing the `Program` in the target shader language.
+
+The input `Program` to a writer must be valid (pass validation).

docs/tint/compound_statements.md

@@ -0,0 +1,119 @@
+# Compound Statements
+
+Compound statements are statements that can hold other statements.
+
+This document maps the WGSL compound statements to their semantic tree representations.
+
+## if statement
+
+WGSL:
+```
+if (condition_a) {
+    statement_a;
+} else if (condition_b) {
+    statement_b;
+} else {
+    statement_c;
+}
+```
+
+Semantic tree:
+```
+sem::IfStatement {
+    condition_a
+    sem::BlockStatement {
+        statement_a
+    }
+    sem::ElseStatement {
+        condition_b
+        sem::BlockStatement {
+            statement_b
+        }
+    }
+    sem::ElseStatement {
+        sem::BlockStatement {
+            statement_c
+        }
+    }
+}
+```
+
+## for loop
+
+WGSL:
+```
+for (initializer; condition; continuing) {
+    statement;
+}
+```
+
+Semantic tree:
+```
+sem::ForLoopStatement {
+    sem::Statement  initializer
+    sem::Expression condition
+    sem::Statement  continuing
+
+    sem::LoopBlockStatement {
+        sem::Statement statement
+    }
+}
+```
+
+## loop
+
+WGSL:
+```
+loop (condition) {
+    statement_a;
+    continuing {
+        statement_b;
+    }
+}
+```
+
+Semantic tree:
+```
+sem::LoopStatement {
+    sem::Expression condition
+
+    sem::LoopBlockStatement {
+        sem::Statement statement_a
+        sem::LoopContinuingBlockStatement {
+            sem::Statement statement_b
+        }
+    }
+}
+```
+
+
+## switch statement
+
+WGSL:
+```
+switch (condition) {
+    case literal_a, literal_b: {
+        statement_a;
+    }
+    default {
+        statement_b;
+    }
+}
+```
+
+Semantic tree:
+```
+sem::SwitchStatement {
+    sem::Expression condition
+    sem::CaseStatement {
+        sem::BlockStatement {
+            sem::Statement statement_a
+        }
+    }
+    sem::CaseStatement {
+        sem::BlockStatement {
+            sem::Statement statement_b
+        }
+    }
+}
+```

docs/tint/coverage-info.md

@@ -0,0 +1,24 @@
+# Generating and viewing Tint code-coverage
+
+Requirements:
+
+* Host running Linux or macOS
+* Clang toolchain on the `PATH` environment variable
+
+## Building Tint with coverage generation enabled
+
+Follow the steps [to build Tint with CMake](../README.md), but include the additional `-DTINT_EMIT_COVERAGE=1` CMake flag.
+
+## Generate coverage information
+
+Use the `<tint>/tools/tint-generate-coverage` script to run the tint executable or unit tests and generate the coverage information.
+
+The script takes the executable to invoke as the first command line argument, followed by additional arguments to pass to the executable.
+
+For example, to see the code coverage for all unit tests, run:
+`<tint>/tools/tint-generate-coverage <build>/tint_unittests --gtest_brief`
+
+The script will emit two files at the root of the tint directory:
+
+* `coverage.summary` - A text file giving a coverage summary for all Tint source files.
+* `lcov.info` - A binary coverage file that can be consumed with the [VSCode Coverage Gutters](https://marketplace.visualstudio.com/items?itemName=ryanluker.vscode-coverage-gutters) extension.

docs/tint/diagnostics_guide.md

@@ -0,0 +1,126 @@
+# Tint diagnostic style guide
+
+This guide provides a set of best practices when writing code that emits
+diagnostic messages in Tint. These diagnostics are messages presented to the
+user in case of error or warning.
+
+The goal of this document is to have our diagnostic messages be clear and
+understandable to our users, so that problems are easy to fix, and to try and
+keep a consistent style.
+
+## Message style
+
+* Start diagnostic messages with a lower-case letter
+* Try to keep the message to a single sentence, if possible
+* Do not end the message with punctuation (full stop, exclamation mark, etc)
+
+**Don't:**
+
+```
+shader.wgsl:7:1 error: Cannot take the address of expression.
+```
+
+**Do:**
+
+```
+shader.wgsl:7:1 error: cannot take the address of expression
+```
+
+**Justification:**
+
+Succinct messages are more important than grammatical correctness. \
+This style matches the style found in most other compilers.
+
+## Prefer to use a `Source` location instead of quoting the code in the message
+
+**Don't:**
+
+```
+shader.wgsl:5:7 error: cannot multiply 'expr_a * expr_b' with types i32 and f32
+
+var res : f32 = expr_a * expr_b
+                ^^^^^^^^^^^^^^^
+```
+
+**Do:**
+
+```
+shader.wgsl:5:7 error: cannot multiply types i32 and f32
+
+var res : f32 = expr_a * expr_b
+                ^^^^^^^^^^^^^^^
+```
+
+**Justification:**
+
+The highlighted line provides even more contextual information than the quoted
+source, and duplicating this information doesn't provide any more help to the
+developer. \
+Quoting single word identifiers or keywords from the source is not discouraged.
+
+## Use `note` diagnostics for providing additional links to relevant code
+
+**Don't:**
+
+```
+shader.wgsl:5:11 error: type cannot be used in storage class 'storage' as it is non-host-shareable
+
+    cond : bool;
+           ^^^^
+```
+
+**Do:**
+
+```
+shader.wgsl:5:11 error: type cannot be used in storage class 'storage' as it is non-host-shareable
+
+    cond : bool;
+           ^^^^
+
+shader.wgsl:8:4 note: while instantiating variable 'StorageBuffer'
+
+var<storage> sb : StorageBuffer;
+             ^^
+```
+
+**Justification:**
+
+To properly understand some diagnostics requires looking at more than a single
+line. \
+Multi-source links can greatly reduce the time it takes to properly
+understand a diagnostic message. \
+This is especially important for diagnostics raised from complex whole-program
+analysis, but can also greatly aid simple diagnostics like symbol collision errors.
+
+## Use simple terminology
+
+**Don't:**
+
+```
+shader.wgsl:7:1 error: the originating variable of the left-hand side of an assignment expression must not be declared with read access control.
+```
+
+**Do:**
+
+```
+shader.wgsl:7:1 error: cannot assign to variable with read access control
+
+x.y = 1;
+^^^^^^^
+
+shader.wgsl:2:8 note: read access control declared here
+
+var<storage, read> x : i32;
+             ^^^^
+```
+
+**Justification:**
+
+Diagnostics will be read by Web developers who may not be native English
+speakers and are unlikely to be familiar with WGSL specification terminology.
+Too much technical jargon can be intimidating and confusing. \
+Diagnostics should give enough information to explain what's wrong, and most
+importantly, give enough information so that a fix actionable.
+
+**Caution:** Be careful to not over simplify. Use the specification terminology
+if there's potential ambiguity by not including it.

docs/tint/end-to-end-tests.md

@@ -0,0 +1,35 @@
+# Tint end-to-end tests
+
+This repo contains a large number of end-to-end tests at `<tint>/test`.
+
+## Test files
+
+Test input files have either the `.wgsl`, `.spv` or `.spvasm` file extension.
+
+Each test input file is tested against each of the Tint backends. There are `<number-of-input-files>` &times; `<number-of-tint-backends>` tests that are performed on an unfiltered end-to-end test run.
+
+Each backend test can have an **expectation file**. This expectation file sits next to the input file, with a `<input-file>.expected.<format>` extension. For example the test `test/foo.wgsl` would have the HLSL expectation file `test/foo.wgsl.expected.hlsl`.
+
+An expectation file contains the expected output of Tint, when passed the input file for the given backend.
+
+If the first line of the expectation file starts `SKIP`, then the test will be skipped instead of failing the end-to-end test run. It is good practice to include after the `SKIP` a reason for why the test is being skipped, along with any additional details, such as compiler error messages.
+
+## Running
+
+To run the end-to-end tests use the `<tint>/test/test-all.sh` script, passing the path to the tint executable as the first command line argument.
+
+You can pass `--help` to see the full list of command line flags.\
+The most commonly used flags are:
+
+| flag                 | description |
+|----------------------|-------------|
+|`--filter`            | Filters the testing to subset of the tests. The filter argument is a glob pattern that can include `*` for any substring of a file or directory, and `**` for any number of directories.<br>Example: `--filter 'expressions/**/i32.wgsl'` will test all the `i32.wgsl` expression tests.
+|`--format`            | Filters the tests to the particular backend.<br>Example: `--format hlsl` will just test the HLSL backend.
+|`--generate-expected` | Generate expectation files for the tests that previously had no expectation file, or were marked as `SKIP` but now pass.
+|`--generate-skip`     | Generate `SKIP` expectation files for tests that are not currently passing.
+
+## Authoring guidelines
+
+Each test should be as small as possible, and focused on the particular feature being tested.
+
+Use sub-directories whenever possible to group similar tests, and try to keep the pattern of directories as consistent as possible between different tests. This helps filter tests using the `--filter` glob patterns.

docs/tint/experimental_extensions.md

@@ -0,0 +1,44 @@
+# Experimental extensions
+
+Sometimes a language feature proposed for WGSL requires experiementation
+to prove its worth.  Tint needs to support these, in general to enable
+that experimentation.
+
+The steps for doing so are:
+
+1. Choose a name for the feature, to be used in an `enable` directive.
+   An experimental extension should use prefix of `google_experimental_`
+   Example:
+
+      enable google_experimental_f16;
+
+2. Write down what the feature is supposed to mean.
+   This informs the Tint implementation, and tells shader authors what
+   has changed.
+   Ideally, this will take the form of one of the following:
+
+   - A PR against the WGSL spec.
+
+   - A description of what the contents of that PR would be, committed
+     as a document in this Tint repository.
+
+3. File a tracking bug for adding the feature.
+   Note: Should the Tint repo have a label for experimental features?
+
+4. File a tracking bug for removing the feature or converting it to
+   non-experimental.
+
+5. Write a plan for removal of the experiment.
+   - Ideally, this plan is committed to this repository, especially the
+     description of public activities and commitments. However, we recognize
+     that some internal goals or metrics may be sensitive, and can be hidden.
+   - The plan is about process, not technical details.  It should include:
+       - Who is the point of contact for this feature? The point of contact
+         is responsible when the feature causes an issue or gets in the way.
+       - What is your target date for declaring the experiment a success or
+         failure. In Chrome an experiment must be shipped or removed, in
+         finite time.
+       - What experience are you hoping to gain?  Do you have target metrics?
+       - What approvals, if any, do you need from W3C? What is your plan to
+         present your case to W3C?
+       - The bug tracking removal of the experiment.

docs/tint/origin-trial-changes.md

@@ -0,0 +1,133 @@
+# Tint changes during Origin Trial
+
+## Changes for M102
+
+### New Features
+
+* Parentheses are no longer required around expressions for if and switch statements [tint:1424](crbug.com/tint/1424)
+* Compound assignment statements are now supported. [tint:1325](https://crbug.com/tint/1325)
+* The colon in case statements is now optional. [tint:1485](crbug.com/tint/1485)
+
+### Breaking changes
+
+* Struct members are now separated by commas. [tint:1475](crbug.com/tint/1475)
+* The `@block` attribute has been removed. [tint:1324](crbug.com/tint/1324)
+* The `@stride` attribute has been removed. [tint:1381](crbug.com/tint/1381)
+* Attributes using `[[attribute]]` syntax are no longer supported. [tint:1382](crbug.com/tint/1382)
+* The `elseif` keyword is no longer supported. [tint:1289](crbug.com/tint/1289)
+
+### Deprecated Features
+
+* The `smoothStep()` builtin has been renamed to `smoothstep()`. [tint:1483](crbug.com/tint/1483)
+
+## Changes for M101
+
+### New Features
+
+* Tint now supports unicode identifiers. [tint:1437](crbug.com/tint/1437)
+
+### Breaking changes
+
+* The `isNan()`, `isInf()`, `isFinite()`, and `isNormal()` builtins have been removed. [tint:1312](https://crbug.com/tint/1312)
+
+## Changes for M100
+
+### Breaking changes
+
+* The `@interpolate(flat)` attribute must now be specified on integral user-defined IO. [tint:1224](crbug.com/tint/1224)
+* The `ignore()` intrinsic has been removed. Use phoney-assignment instead: `ignore(expr);` -> `_ = expr;`.
+* `break` statements in `continuing` blocks are now correctly validated.
+
+### New Features
+
+* Module-scope declarations can now be declared in any order. [tint:1266](crbug.com/tint/1266)
+* The `override` keyword and `@id()` attribute for pipeline-overridable constants are now supported, replacing the `@override` attribute. [tint:1403](crbug.com/tint/1403)
+
+## Changes for M99
+
+### Breaking changes
+
+Obviously infinite loops (no condition, no break) are now a validation error.
+
+### Deprecated Features
+
+The following features have been deprecated and will be removed in M102:
+
+* The `[[block]]` attribute has been deprecated. [tint:1324](https://crbug.com/tint/1324)
+* Attributes now use the `@decoration` syntax instead of the `[[decoration]]` syntax. [tint:1382](https://crbug.com/tint/1382)
+* `elseif` has been replaced with `else if`. [tint:1289](https://crbug.com/tint/1289)
+* The `[[stride]]` attribute has been deprecated. [tint:1381](https://crbug.com/tint/1381)
+
+### New Features
+
+* Vector and matrix element type can now be inferred from constructor argument types. [tint:1334](https://crbug.com/tint/1334)
+* Added builtins `degrees()` and `radians()` for converting between degrees and radians. [tint:1329](https://crbug.com/tint/1329)
+* `let` arrays and matrices can now be dynamically indexed. [tint:1352](https://crbug.com/tint/1352)
+* Storage and Uniform buffer types no longer have to be structures. [tint:1372](crbug.com/tint/1372)
+* A struct declaration does not have to be followed by a semicolon. [tint:1380](crbug.com/tint/1380)
+
+### Fixes
+
+* Fixed an issue where for-loops that contain array or structure constructors in the loop initializer statements, condition expressions or continuing statements could fail to compile. [tint:1364](https://crbug.com/tint/1364)
+
+## Changes for M98
+
+### Breaking Changes
+
+* Taking the address of a vector component is no longer allowed.
+* Module-scope declarations can no longer alias a builtin name. [tint:1318](https://crbug.com/tint/1318)
+* It is now an error to call a function either directly or transitively, from a loop continuing block, that uses `discard`. [tint:1302](https://crbug.com/tint/1302)
+
+### Deprecated Features
+
+* The `isNan()`, `isInf()`, `isFinite()` and `isNormal()` builtins has been deprecated and will be removed in M101. [tint:1312](https://crbug.com/tint/1312)
+
+### New Features
+
+* New texture gather builtins: `textureGather()` and `textureGatherCompare()`. [tint:1330](https://crbug.com/tint/1330)
+* Shadowing is now fully supported. [tint:819](https://crbug.com/tint/819)
+* The `dot()` builtin now supports integer vector types.
+* Identifiers can now start with a single leading underscore.  [tint:1292](https://crbug.com/tint/1292)
+* Control flow analysis has been improved, and functions no longer need to `return` if the statement is unreachable. [tint:1302](https://crbug.com/tint/1302)
+* Unreachable statements now produce a warning instead of an error, to allow WGSL code to be updated to the new analysis behavior. These warnings may become errors in the future [gpuweb#2378](https://github.com/gpuweb/gpuweb/issues/2378)
+
+### Fixes
+
+* Fixed an issue where using a module-scoped `let` in a `workgroup_size` may result in a compilation error. [tint:1320](https://crbug.com/tint/1320)
+
+## Changes for M97
+
+### Breaking Changes
+
+* Deprecated `modf()` and `frexp()` builtin overloads that take a pointer second parameter have been removed.
+* Deprecated texture builtin functions that accepted a `read` access controlled storage texture have been removed.
+* Storage textures must now only use the `write` access control.
+
+### Deprecated Features
+
+* The `ignore()` builtin has been replaced with phony-assignment. [gpuweb#2127](https://github.com/gpuweb/gpuweb/pull/2127)
+
+### New Features
+
+* `any()` and `all()` now support a `bool` parameter. These simply return the passed argument. [tint:1253](https://crbug.com/tint/1253)
+* Call statements may now include functions that return a value (`ignore()` is no longer needed).
+* The `interpolate(flat)` attribute can now be specified on integral user-defined IO. It will eventually become an error to define integral user-defined IO without this attribute.
+* Matrix construction from scalar element values is now supported.
+
+### Fixes
+
+* Swizzling of `vec3` types in `storage` and `uniform` buffers has been fixed for Metal 1.x. [tint:1249](https://crbug.com/tint/1249)
+* Calling a function that returns an unused value no longer produces an FXC compilation error. [tint:1259](https://crbug.com/tint/1259)
+* `abs()` fixed for unsigned integers on SPIR-V backend
+
+## Changes for M95
+
+### New Features
+
+* The size of an array can now be defined using a non-overridable module-scope constant
+* The `num_workgroups` builtin is now supported.
+
+### Fixes
+
+* Hex floats: now correctly errors when the magnitude is non-zero, and the exponent would cause overflow. [tint:1150](https://crbug.com/tint/1150), [tint:1166](https://crbug.com/tint/1166)
+* Identifiers beginning with an underscore are now correctly rejected.  [tint:1179](https://crbug.com/tint/1179)

docs/tint/spirv-input-output-variables.md

@@ -0,0 +1,267 @@
+# SPIR-V translation of shader input and output variables
+
+WGSL [MR 1315](https://github.com/gpuweb/gpuweb/issues/1315) changed WGSL so
+that pipeline inputs and outputs are handled similar to HLSL:
+
+- Shader pipeline inputs are the WGSL entry point function arguments.
+- Shader pipeline outputs are the WGSL entry point return value.
+
+Note: In both cases, a struct may be used to pack multiple values together.
+In that case, I/O specific attributes appear on struct members at the struct declaration.
+
+Resource variables, e.g. buffers, samplers, and textures, are still declared
+as variables at module scope.
+
+## Vulkan SPIR-V today
+
+SPIR-V for Vulkan models inputs and outputs as module-scope variables in
+the Input and Output storage classes, respectively.
+
+The `OpEntryPoint` instruction has a list of module-scope variables that must
+be a superset of all the input and output variables that are statically
+accessed in the shader call tree.
+From SPIR-V 1.4 onward, all interface variables that might be statically accessed
+must appear on that list.
+So that includes all resource variables that might be statically accessed
+by the shader call tree.
+
+## Translation scheme for SPIR-V to WGSL
+
+A translation scheme from SPIR-V to WGSL is as follows:
+
+Each SPIR-V entry point maps to a set of Private variables proxying the
+inputs and outputs, and two functions:
+
+- An inner function with no arguments or return values, and whose body
+  is the same as the original SPIR-V entry point.
+- Original input variables are mapped to pseudo-in Private variables
+  with the same store types, but no other attributes or properties copied.
+  In Vulkan, Input variables don't have initalizers.
+- Original output variables are mapped to pseudo-out Private variables
+  with the same store types and optional initializer, but no other attributes
+  or properties are copied.
+- A wrapper entry point function whose arguments correspond in type, location
+  and builtin attributes the original input variables, and whose return type is
+  a structure containing members correspond in type, location, and builtin
+  attributes to the original output variables.
+  The body of the wrapper function the following phases:
+  - Copy formal parameter values into pseudo-in variables.
+    - Insert a bitcast if the WGSL builtin variable has different signedness
+      from the SPIR-V declared type.
+  - Execute the inner function.
+  - Copy pseudo-out variables into the return structure.
+    - Insert a bitcast if the WGSL builtin variable has different signedness
+      from the SPIR-V declared type.
+  - Return the return structure.
+
+- Replace uses of the the original input/output variables to the pseudo-in and
+  pseudo-out variables, respectively.
+- Remap pointer-to-Input with pointer-to-Private
+- Remap pointer-to-Output with pointer-to-Private
+
+We are not concerned with the cost of extra copying input/output values.
+First, the pipeline inputs/outputs tend to be small.
+Second, we expect the backend compiler in the driver will be able to see
+through the copying and optimize the result.
+
+### Example
+
+
+```glsl
+    #version 450
+
+    layout(location = 0) out vec4 frag_colour;
+    layout(location = 0) in vec4 the_colour;
+
+    void bar() {
+      frag_colour = the_colour;
+    }
+
+    void main() {
+        bar();
+    }
+```
+
+Current translation, through SPIR-V, SPIR-V reader, WGSL writer:
+
+```groovy
+    @location(0) var<out> frag_colour : vec4<f32>;
+    @location(0) var<in> the_colour : vec4<f32>;
+
+    fn bar_() -> void {
+      const x_14 : vec4<f32> = the_colour;
+      frag_colour = x_14;
+      return;
+    }
+
+    @stage(fragment)
+    fn main() -> void {
+      bar_();
+      return;
+    }
+```
+
+Proposed translation, through SPIR-V, SPIR-V reader, WGSL writer:
+
+```groovy
+    // 'in' variables are now 'private'.
+    var<private> frag_colour : vec4<f32>;
+    var<private> the_colour : vec4<f32>;
+
+    fn bar_() -> void {
+      // Accesses to the module-scope variables do not change.
+      // This is a big simplifying advantage.
+      const x_14 : vec4<f32> = the_colour;
+      frag_colour = x_14;
+      return;
+    }
+
+    fn main_inner() -> void {
+      bar_();
+      return;
+    }
+
+    // Declare a structure type to collect the return values.
+    struct main_result_type {
+      @location(0) frag_color : vec4<f32>;
+    };
+
+    @stage(fragment)
+    fn main(
+
+      // 'in' variables are entry point parameters
+      @location(0) the_color_arg : vec4<f32>
+
+    ) -> main_result_type {
+
+      // Save 'in' arguments to 'private' variables.
+      the_color = the_color_arg;
+
+      // Initialize 'out' variables.
+      // Use the zero value, since no initializer was specified.
+      frag_color = vec4<f32>();
+
+      // Invoke the original entry point.
+      main_inner();
+
+      // Collect outputs into a structure and return it.
+      var result : main_outer_result_type;
+      result.frag_color = frag_color;
+      return result;
+    }
+```
+
+Alternately, we could emit the body of the original entry point at
+the point of invocation.
+However that is more complex because the original entry point function
+may return from multiple locations, and we would like to have only
+a single exit path to construct and return the result value.
+
+### Handling fragment discard
+
+In SPIR-V `OpKill` causes immediate termination of the shader.
+Is the shader obligated to write its outputs when `OpKill` is executed?
+
+The Vulkan fragment operations are as follows:
+(see [6. Fragment operations](https://www.khronos.org/registry/vulkan/specs/1.2/html/vkspec.html#fragops)).
+
+* Scissor test
+* Sample mask test
+* Fragment shading
+* Multisample coverage
+* Depth bounds test
+* Stencil test
+* Depth test
+* Sample counting
+* Coverage reduction
+
+After that, the fragment results are used to update output attachments, including
+colour, depth, and stencil attachments.
+
+Vulkan says:
+
+> If a fragment operation results in all bits of the coverage mask being 0,
+> the fragment is discarded, and no further operations are performed.
+> Fragments can also be programmatically discarded in a fragment shader by executing one of
+>
+>     OpKill.
+
+I interpret this to mean that the outputs of a discarded fragment are ignored.
+
+Therefore, `OpKill` does not require us to modify the basic scheme from the previous
+section.
+
+The `OpDemoteToHelperInvocationEXT`
+instruction is an alternative way to throw away a fragment, but which
+does not immediately terminate execution of the invocation.
+It is introduced in the [`SPV_EXT_demote_to_helper_invocation](http://htmlpreview.github.io/?https://github.com/KhronosGroup/SPIRV-Registry/blob/master/extensions/EXT/SPV_EXT_demote_to_helper_invocation.html)
+extension.  WGSL does not have this feature, but we expect it will be introduced by a
+future WGSL extension.  The same analysis applies to demote-to-helper.  When introduced,
+it will not affect translation of pipeline outputs.
+
+### Handling depth-replacing mode
+
+A Vulkan fragment shader must write to the fragment depth builtin if and only if it
+has a `DepthReplacing` execution mode. Otherwise behaviour is undefined.
+
+We will ignore the case where the SPIR-V shader writes to the `FragDepth` builtin
+and then discards the fragment.
+This is justified because "no further operations" are performed by the pipeline
+after the fragment is discarded, and that includes writing to depth output attachments.
+
+Assuming the shader is valid, no special translation is required.
+
+### Handling output sample mask
+
+By the same reasoning as for depth-replacing, it is ok to incidentally not write
+to the sample-mask builtin variable when the fragment is discarded.
+
+### Handling clip distance and cull distance
+
+Most builtin variables are scalars or vectors.
+However, the `ClipDistance` and `CullDistance` builtin variables are arrays of 32-bit float values.
+Each entry defines a clip half-plane (respectively cull half-plane)
+A Vulkan implementation must support array sizes of up to 8 elements.
+
+How prevalent are shaders that use these features?
+These variables are supported when Vulkan features `shaderClipDistance` and `shaderCullDistance`
+are supported.
+According to gpuinfo.org as of this writing, those
+Vulkan features appear to be nearly universally supported on Windows devices (>99%),
+but by only 70% on Android.
+It appears that Qualcomm devices support them, but Mali devices do not (e.g. Mali-G77).
+
+The proposed translation scheme forces a copy of each array from private
+variables into the return value of a vertex shader, or into a private
+variable of a fragment shader.
+In addition to the register pressure, there may be a performance degradation
+due to the bulk copying of data.
+
+We think this is an acceptable tradeoff for the gain in usability and
+consistency with other pipeline inputs and outputs.
+
+## Translation scheme for WGSL AST to SPIR-V
+
+To translate from the WGSL AST to SPIR-V, do the following:
+
+- Each entry point formal parameter is mapped to a SPIR-V `Input` variable.
+  - Struct and array inputs may have to be broken down into individual variables.
+- The return of the entry point is broken down into fields, with one
+  `Output` variable per field.
+- In the above, builtins must be separated from user attributes.
+  - Builtin attributes are moved to the corresponding variable.
+  - Location and interpolation attributes are moved to the corresponding
+    variables.
+- This translation relies on the fact that pipeline inputs and pipeline
+  outputs are IO-shareable types. IO-shareable types are always storable,
+  and can be the store type of input/output variables.
+- Input function parameters will be automatically initialized by the system
+  as part of setting up the pipeline inputs to the entry point.
+- Replace each return statement in the entry point with a code sequence
+  which writes the return value components to the synthesized output variables,
+  and then executes an `OpReturn` (without value).
+
+This translation is sufficient even for fragment shaders with discard.
+In that case, outputs will be ignored because downstream pipeline
+operations will not be performed.
+This is the same rationale as for translation from SPIR-V to WGSL AST.

docs/tint/spirv-ptr-ref.md

@@ -0,0 +1,115 @@
+# SPIR-V translation of WGSL pointers and references
+
+WGSL was updated to have two kinds of memory views: pointers and references.
+See https://github.com/gpuweb/gpuweb/pull/1569
+
+In summary:
+
+* Reference types are never explicitly mentioned in WGSL source.
+* A use of a variable is a value of reference type corresponding
+  to the reference memory view of the storage allocated for the
+  variable.
+* Let-declared constants can be of pointer type, but not reference
+  type.
+* Function parameter can be of pointer type, but not reference type.
+* A variable's store type is never a pointer type, and never a
+  reference type.
+* The "Load Rule" allows a reference to decay to the underlying
+  store type, by issuing a load of the value in the underlying memory.
+* For an assignment:
+  * The right-hand side evaluates to a non-reference type (atomic-free
+    plain type).
+  * The left-hand side evaluates to a reference type, whose store
+    type is the same as the result of evaluating the right hand side.
+* The address-of (unary `&`) operator converts a reference to a
+  pointer.
+* The dereference (unary `*`) operator converts a pointer to a
+  reference.
+
+TODO: Passing textures and samplers to helper functions might be
+done by "handler value", or by pointer-to-handle.
+
+## Writing SPIR-V from WGSL
+
+The distinction in WGSL between reference and pointer disappears
+at the SPIR-V level.  Both types map into pointer types in SPIR-V.
+
+To translate a valid WGSL program to SPIR-V:
+
+* The dereference operator (unary `*`) is the identity operation.
+* The address-of operator (unary `&`) is the identity operation.
+* Assignment maps to OpStore.
+* The Load Rule translates to OpLoad.
+
+## Reading SPIR-V to create WGSL
+
+The main changes to the SPIR-V reader are:
+
+* When translating a SPIR-V pointer expression, track whether the
+  corresponding WGSL expression is of corresponding WGSL pointer
+  type or correspoinding WGSL type.
+* Insert dereference (unary-`*`) or address-of (unary-`&`) operators
+  as needed to generate valid WGSL expressions.
+
+The choices can be made deterministic, as described below.
+
+The SPIR-V reader only supports baseline functionality in Vulkan.
+Therefore we assume no VariablePointers or VariablePointersStorageBuffer
+capabilities.  All pointers are
+[SPIR-V logical pointers](https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html#LogicalPointerType).
+The [SPIR-V Universal Validation Rules](https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html#_universal_validation_rules)
+specify where logical pointers can appear as results of instructions
+or operands of instructions.
+
+Each SPIR-V pointer result expression is a logical pointer, and
+therefore is one of:
+
+* OpVariable: map to the reference type.
+* OpFunctionParameter: map to the pointer type.
+* OpCopyObject:
+   * When these only have one use, then these often fold away.
+     Otherwise, they map to a a let-declared constant.
+   * Map to the pointer type.
+* OpAccessChain, OpInBoundsAccessChain:
+   * This could map to either pointer or reference, and adjustments
+     in other areas could make it work.  However, we recommend mapping
+     this to the reference type.
+* OpImageTexelPointer is not supported in WGSL.
+   It is used to get a pointer into a storage texture, for use with
+   atomic instructions.  But image atomics is not supported in
+   WebGPU/WGSL.
+
+Each SPIR-V pointer operand is also a logical pointer, and is an
+operand to one of:
+* OpLoad Pointer operand:
+   * Map to reference, inserting a dereference operator if needed.
+* OpStore Pointer operand:
+   * Map to reference, inserting a dereference operator if needed.
+* OpStore Pointer operand:
+* OpAccessChain, OpInBoundsAccessChain Base operand:
+   * WGSL array-access and subfield access only works on references.
+      * [Gpuweb issue 1530](https://github.com/gpuweb/gpuweb/issues/1530)
+        is filed to allow those operations to work on pointers.
+   * Map to reference, inserting a dereference operator if needed.
+* OpFunctionCall function argument pointer operands
+   * Function operands can't be references.
+   * Map to pointer, inserting an address-of operator if needed.
+* OpAtomic instruction Pointer operand
+   * These map to WGSL atomic builtins.
+   * Map to pointer, inserting an address-of operator if needed.
+   * Note: As of this writing, the atomic instructions are not supported
+     by the SPIR-V reader.
+* OpCopyObject source operand
+   * This could have been mapped either way, but it's easiest to
+     map to pointer, to match the choice for OpCopyObject result type.
+   * Map to pointer, inserting an address-of operator if needed.
+* OpCopyMemory, source and destination operands
+   * This acts as an assignment.
+   * Map both source and destination to reference, inserting dereference
+     operators if needed.
+   * Note: As of this writing, OpCopyMemory is not supported by the
+     SPIR-V reader.
+* Extended instruction set instructions Modf and Frexp
+   * These map to builtins.
+   * Map the pointer operand to pointer, inserting an address-of
+     operator if needed.

docs/tint/style_guide.md

@@ -0,0 +1,47 @@
+# Tint style guide
+
+* Generally, follow the [Chromium style guide for C++](https://chromium.googlesource.com/chromium/src/+/HEAD/styleguide/c++/c++.md)
+  which itself is built on the [Google C++ style guide](https://google.github.io/styleguide/cppguide.html).
+
+* Overall try to use the same style and convention as code around your change.
+
+* Code must be formatted. Use `clang-format` with the provided [.clang-format](../.clang-format)
+  file.  The `tools/format` script runs the formatter.
+
+* Code should not have linting errors.
+    The `tools/lint` script runs the linter. So does `git cl upload`.
+
+* Do not use C++ exceptions
+
+* Do not use C++ RTTI.
+   Instead, use `tint::Castable::As<T>()` from
+   [src/castable.h](../src/castable.h)
+
+* Generally, avoid `assert`.  Instead, issue a [diagnostic](../src/diagnostic.h)
+  and fail gracefully, possibly by returning an error sentinel value.
+  Code that should not be reachable should call `TINT_UNREACHABLE` macro
+  and other internal error conditions should call the `TINT_ICE` macro.
+  See [src/debug.h](../src/debug.h)
+
+* Use `type` as part of a name only when the name refers to a type
+  in WGSL or another shader language processed by Tint.  If the concept you are
+  trying to name is about distinguishing between alternatives, use `kind` instead.
+
+## Compiler support
+
+Tint requires C++17.
+
+Tint uses the Chromium build system and will stay synchronized with that system.
+Compiler configurations beyond that baseline is on a best-effort basis.
+We strive to support recent GCC and MSVC compilers.
+
+## Test code
+
+We might relax the above rules rules for test code, since test code
+shouldn't ship to users.
+
+However, test code should still be readable and maintainable.
+
+For test code, the tradeoff between readability and maintainability
+and other factors is weighted even more strongly toward readability
+and maintainability.

docs/tint/translations.md

@@ -0,0 +1,187 @@
+# Translations
+
+This document attempts to document how WGSL translates into the various backends
+for the cases where the translation is not a direct mapping.
+
+# Access Control
+
+## HLSL
+ * ReadOnly -> `ByteAddressBuffer`
+ * ReadWrite -> `RWByteAddressBuffer`
+
+## MSL
+ * ReadOnly -> `const`
+
+## SPIR-V
+There are two ways this can be achieved in SPIR-V. Either the variable can be
+decorated with `NonWritable` or each member of the struct can be decorated with
+`NonWritable`. We chose to go the struct member route.
+ * The read-only becomes part of the type in this case. Otherwise, you are
+   treating the readonly type information as part of the variable which is
+   confusing.
+ * Treating the readonly as part of the variable means we should be
+   deduplicating the types behind the access control, which causes confusing
+   with the type_names and various tracking systems within Tint.
+
+
+# Builtin Decorations
+| Name | SPIR-V | MSL | HLSL |
+|------|--------|-----|------|
+| position | SpvBuiltInPosition |position | SV_Position |
+| vertex_index | SpvBuiltInVertexIndex |vertex_id | SV_VertexID |
+| instance_index | SpvBuiltInInstanceIndex | instance_id| SV_InstanceID |
+| front_facing | SpvBuiltInFrontFacing | front_facing | SV_IsFrontFacing |
+| frag_coord | SpvBuiltInFragCoord | position | SV_Position |
+| frag_depth | SpvBuiltInFragDepth | depth(any) | SV_Depth |
+| local_invocation_id | SpvBuiltInLocalInvocationId | thread_position_in_threadgroup | SV_GroupThreadID |
+| local_invocation_index | SpvBuiltInLocalInvocationIndex | thread_index_in_threadgroup | SV_GroupIndex |
+| global_invocation_id | SpvBuiltInGlobalInvocationId | thread_position_in_grid | SV_DispatchThreadID |
+
+
+# Builtins Methods
+| Name | SPIR-V | MSL | HLSL |
+| ------|--------|-----|------ |
+| abs | GLSLstd450FAbs or GLSLstd450SAbs| fabs or abs | abs |
+| acos | GLSLstd450Acos | acos | acos |
+| all | SpvOpAll | all | all |
+| any | SpvOpAny | any | any |
+| arrayLength | SpvOpArrayLength | | |
+| asin | GLSLstd450Asin | asin | asin |
+| atan | GLSLstd450Atan | atan | atan |
+| atan2 | GLSLstd450Atan2| atan2 | atan2 |
+| ceil | GLSLstd450Ceil| ceil | ceil |
+| clamp | GLSLstd450NClamp or GLSLstd450UClamp or GLSLstd450SClamp| clamp | clamp |
+| cos | GLSLstd450Cos | cos | cos |
+| cosh | GLSLstd450Cosh | cosh | cosh |
+| countOneBits | SpvOpBitCount | popcount | countbits |
+| cross | GLSLstd450Cross | cross | cross |
+| determinant | GLSLstd450Determinant | determinant | determinant |
+| distance | GLSLstd450Distance | distance | distance |
+| dot | SpOpDot | dot | dot |
+| dpdx | SpvOpDPdx | dpdx | ddx |
+| dpdxCoarse | SpvOpDPdxCoarse | dpdx | ddx_coarse |
+| dpdxFine | SpvOpDPdxFine | dpdx | ddx_fine |
+| dpdy | SpvOpDPdy | dpdy | ddy |
+| dpdyCoarse | SpvOpDPdyCoarse | dpdy | ddy_coarse |
+| dpdyFine | SpvOpDPdyFine | dpdy | ddy_fine |
+| exp | GLSLstd450Exp | exp |  exp |
+| exp2 | GLSLstd450Exp2 | exp2 | exp2 |
+| faceForward | GLSLstd450FaceForward | faceforward | faceforward |
+| floor | GLSLstd450Floor | floor | floor |
+| fma | GLSLstd450Fma | fma | fma |
+| fract | GLSLstd450Fract | fract | frac |
+| frexp | GLSLstd450Frexp | | |
+| fwidth | SpvOpFwidth | fwidth | fwidth |
+| fwidthCoarse | SpvOpFwidthCoarse | fwidth | fwidth |
+| fwidthFine | SpvOpFwidthFine | fwidth | fwidth |
+| inverseSqrt | GLSLstd450InverseSqrt | rsqrt | rsqrt |
+| ldexp | GLSLstd450Ldexp | | |
+| length | GLSLstd450Length | length | length |
+| log | GLSLstd450Log | log | log |
+| log2 | GLSLstd450Log2 | log2 | log2 |
+| max | GLSLstd450NMax or GLSLstd450SMax or GLSLstd450UMax | fmax or max | max |
+| min | GLSLstd450NMin or GLSLstd450SMin or GLSLstd450UMin | fmin or min | min |
+| mix | GLSLstd450FMix | mix | mix |
+| modf | GLSLstd450Modf | | |
+| normalize | GLSLstd450Normalize | normalize | normalize |
+| pow | GLSLstd450Pow | pow | pow |
+| reflect | GLSLstd450Reflect | reflect | reflect |
+| reverseBits | SpvOpBitReverse | reverse_bits | reversebits |
+| round | GLSLstd450Round | round | round |
+| select | SpvOpSelect | select | |
+| sign | GLSLstd450FSign | sign | sign |
+| sin | GLSLstd450Sin | sin | sin |
+| sinh | GLSLstd450Sinh | sinh | sinh |
+| smoothStep | GLSLstd450SmoothStep | smoothstep | smoothstep |
+| sqrt | GLSLstd450Sqrt | sqrt | sqrt |
+| step | GLSLstd450Step | step | step |
+| tan | GLSLstd450Tan | tan | tan |
+| tanh | GLSLstd450Tanh | tanh | tanh |
+| trunc | GLSLstd450Trunc | trunc | trunc |
+
+# Types
+## Sampler Types
+| WGSL | SPIR-V | MSL | HLSL |
+|------|--------|-----|------|
+| sampler | OpTypeSampler | sampler | SamplerState |
+| sampler_comparison | OpTypeSampler | sampler | SamplerComparisonState |
+
+## Texture Types
+| WGSL | SPIR-V | MSL | HLSL |
+|------|--------|-----|------|
+| texture_1d<type> | OpTypeImage 1D Sampled=1 | texture1d<type, access::sample> | Texture1D |
+| texture_2d<type> | OpTypeImage 2D Sampled=1 | texture2d<type, access::sample> | Texture2D |
+| texture_2d_array<type> | OpTypeImage 2D Arrayed=1 Sampled=1 | texture2d_array<type, access::sample> | Texture2DArray |
+| texture_3d<type> | OpTypeImage 3D Sampled=1 | texture3d<type, access::sample> | Texture3D |
+| texture_cube<type> | OpTypeImage Cube Sampled=1 | texturecube<type, access::sample> | TextureCube |
+| texture_cube_array<type> | OpTypeImage Cube Arrayed=1 Sampled=1 | texturecube_array<type, access::sample> | TextureCubeArray |
+| | | |
+| texture_multisampled_2d<type> | OpTypeImage 2D MS=1 Sampled=1 | texture2d_ms<type, access::sample> | Texture2D |
+| | | |
+| texture_depth_2d | OpTypeImage 2D Depth=1 Sampled=1 | depth2d<float, access::sample>| Texture2D |
+| texture_depth_2d_array | OpTypeImage 2D Depth=1 Arrayed=1 Sampled=1 | depth2d_array<float, access::sample> | Texture2DArray |
+| texture_depth_cube | OpTypeImage Cube Depth=1 Sampled=1 | depthcube<float, access::sample> | TextureCube |
+| texture_depth_cube_array | OpTypeImage Cube Depth=1 Arrayed=1 Sampled=1 | depthcube_array<float, access::sample> | TextureCubeArray |
+| texture_depth_multisampled_2d | OpTypeImage 2D Depth=1 MS=1 Sampled=1 | depth2d<float, access::sample>| Texture2DMSArray |
+| | | |
+| texture_storage_1d<image_storage_type> | OpTypeImage 1D Sampled=2| texture1d<type, access::read> | RWTexture1D |
+| texture_storage_2d<image_storage_type> | OpTypeImage 2D Sampled=2 | texture2d<type, access::read> | RWTexture2D |
+| texture_storage_2d_array<image_storage_type> | OpTypeImage 2D Arrayed=1 Sampled=2 | texture2d_array<type, access::read> | RWTexture2DArray |
+| texture_storage_3d<image_storage_type> | OpTypeImage 3D Sampled=2 | texture3d<type, access::read> | RWTexture3D |
+| | | |
+| texture_storage_1d<image_storage_type> | OpTypeImage 1D Sampled=2 | texture1d<type, access::write> | RWTexture1D |
+| texture_storage_2d<image_storage_type> | OpTypeImage 2D Sampled=1 | texture2d<type, access::write> | RWTexture2D |
+| texture_storage_2d_array<image_storage_type> | OpTypeImage 2D Arrayed=1 Sampled=2 | texture2d_array<type, access::write> | RWTexture2DArray |
+| texture_storage_3d<image_storage_type> | OpTypeImage 3D Sampled=2 | texture3d<type, access::write> | RWTexture3D|
+
+# Short-circuting
+## HLSL
+TODO(dsinclair): Nested if's
+
+## SPIR-V
+TODO(dsinclair): Nested if's
+
+# Storage classes
+TODO(dsinclair): do ...
+
+# Storage buffers
+## HLSL
+TODO(dsinclair): Rewriting of accessors to loads
+
+# Loop blocks
+## HLSL
+TODO(dsinclair): Rewrite with bools
+
+## MSL
+TODO(dsinclair): Rewrite with bools
+
+# Input / Output storage class
+## HLSL
+TODO(dsinclair): Structs and params
+
+## MSL
+TODO(dsinclair): Structs and params
+
+# Discard
+## HLSL
+ * `discard`
+
+## MSL
+ * `discard_fragment()`
+
+
+# Specialization constants
+## HLSL
+```
+#ifndef WGSL_SPEC_CONSTANT_<id>
+-- if default provided
+#define WGSL_SPEC_CONSTANT_<id> default value
+-- else
+#error spec constant required for constant id
+--
+#endif
+static const <type> <name> = WGSL_SPEC_CONSTANT_<id>
+```
+
+## MSL
+`@function_constant(<id>)`

include/tint/tint.h

@@ -0,0 +1,68 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#ifndef INCLUDE_TINT_TINT_H_
+#define INCLUDE_TINT_TINT_H_
+
+// TODO(tint:88): When implementing support for an install target, all of these
+//                headers will need to be moved to include/tint/.
+
+#include "src/tint/ast/pipeline_stage.h"
+#include "src/tint/demangler.h"
+#include "src/tint/diagnostic/printer.h"
+#include "src/tint/inspector/inspector.h"
+#include "src/tint/reader/reader.h"
+#include "src/tint/sem/type_manager.h"
+#include "src/tint/transform/binding_remapper.h"
+#include "src/tint/transform/first_index_offset.h"
+#include "src/tint/transform/fold_trivial_single_use_lets.h"
+#include "src/tint/transform/manager.h"
+#include "src/tint/transform/multiplanar_external_texture.h"
+#include "src/tint/transform/renamer.h"
+#include "src/tint/transform/robustness.h"
+#include "src/tint/transform/single_entry_point.h"
+#include "src/tint/transform/vertex_pulling.h"
+#include "src/tint/writer/writer.h"
+
+#if TINT_BUILD_SPV_READER
+#include "src/tint/reader/spirv/parser.h"
+#endif  // TINT_BUILD_SPV_READER
+
+#if TINT_BUILD_WGSL_READER
+#include "src/tint/reader/wgsl/parser.h"
+#endif  // TINT_BUILD_WGSL_READER
+
+#if TINT_BUILD_SPV_WRITER
+#include "spirv-tools/libspirv.hpp"
+#include "src/tint/writer/spirv/generator.h"
+#endif  // TINT_BUILD_SPV_WRITER
+
+#if TINT_BUILD_WGSL_WRITER
+#include "src/tint/writer/wgsl/generator.h"
+#endif  // TINT_BUILD_WGSL_WRITER
+
+#if TINT_BUILD_MSL_WRITER
+#include "src/tint/writer/msl/generator.h"
+#endif  // TINT_BUILD_MSL_WRITER
+
+#if TINT_BUILD_HLSL_WRITER
+#include "src/tint/writer/hlsl/generator.h"
+#endif  // TINT_BUILD_HLSL_WRITER
+
+#if TINT_BUILD_GLSL_WRITER
+#include "src/tint/transform/glsl.h"
+#include "src/tint/writer/glsl/generator.h"
+#endif  // TINT_BUILD_GLSL_WRITER
+
+#endif  // INCLUDE_TINT_TINT_H_

kokoro/linux/build.sh

@@ -0,0 +1,48 @@
+#!/bin/bash
+
+# Copyright 2021 The Tint Authors.
+#
+# 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.
+
+set -e # Fail on any error.
+
+SCRIPT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}")" >/dev/null 2>&1 && pwd )"
+ROOT_DIR="$( cd "${SCRIPT_DIR}/../.." >/dev/null 2>&1 && pwd )"
+
+# Inside the docker VM, we clone the project to a new directory.
+# We do this so that the docker script can be tested in a local development
+# checkout, without having the build litter the local checkout with artifacts.
+# This directory is mapped to the host temporary directory.
+# Kokoro uses a '/tmpfs' root, where as most linux enviroments just have '/tmp'
+if [ -d "/tmpfs" ]; then
+    TMP_DIR=/tmpfs
+else
+    TMP_DIR=/tmp
+fi
+
+
+# --privileged is required for some sanitizer builds, as they seem to require PTRACE privileges
+docker run --rm -i \
+  --privileged \
+  --volume "${ROOT_DIR}:${ROOT_DIR}" \
+  --volume "${TMP_DIR}/kokoro/tint:/tint" \
+  --volume "${KOKORO_ARTIFACTS_DIR}:/mnt/artifacts" \
+  --workdir "${ROOT_DIR}" \
+  --env SRC_DIR="/tint/src" \
+  --env BUILD_DIR="/tint/build" \
+  --env BUILD_TYPE=$BUILD_TYPE \
+  --env BUILD_SYSTEM=$BUILD_SYSTEM \
+  --env BUILD_SANITIZER=$BUILD_SANITIZER \
+  --env BUILD_TOOLCHAIN=$BUILD_TOOLCHAIN \
+  --entrypoint "${SCRIPT_DIR}/docker.sh" \
+  "gcr.io/shaderc-build/radial-build:latest"

kokoro/linux/cmake-clang-debug-asan/build.sh

@@ -0,0 +1,26 @@
+#!/bin/bash
+
+# Copyright 2021 The Tint Authors.
+#
+# 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.
+
+set -e # Fail on any error.
+
+SCRIPT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}")" >/dev/null 2>&1 && pwd )"
+
+export BUILD_SYSTEM=cmake
+export BUILD_TOOLCHAIN=clang
+export BUILD_TYPE=Debug
+export BUILD_SANITIZER=asan
+
+${SCRIPT_DIR}/../build.sh

kokoro/linux/cmake-clang-debug-asan/presubmit.cfg

@@ -0,0 +1,3 @@
+# Format: //devtools/kokoro/config/proto/build.proto
+
+build_file: "tint/kokoro/linux/cmake-clang-debug-asan/build.sh"

kokoro/linux/cmake-clang-debug-ubsan/build.sh

@@ -0,0 +1,26 @@
+#!/bin/bash
+
+# Copyright 2021 The Tint Authors.
+#
+# 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.
+
+set -e # Fail on any error.
+
+SCRIPT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}")" >/dev/null 2>&1 && pwd )"
+
+export BUILD_SYSTEM=cmake
+export BUILD_TOOLCHAIN=clang
+export BUILD_TYPE=Debug
+export BUILD_SANITIZER=ubsan
+
+${SCRIPT_DIR}/../build.sh

kokoro/linux/cmake-clang-debug-ubsan/presubmit.cfg

@@ -0,0 +1,3 @@
+# Format: //devtools/kokoro/config/proto/build.proto
+
+build_file: "tint/kokoro/linux/cmake-clang-debug-ubsan/build.sh"

kokoro/linux/cmake-clang-debug/build.sh

@@ -0,0 +1,25 @@
+#!/bin/bash
+
+# Copyright 2021 The Tint Authors.
+#
+# 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.
+
+set -e # Fail on any error.
+
+SCRIPT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}")" >/dev/null 2>&1 && pwd )"
+
+export BUILD_SYSTEM=cmake
+export BUILD_TOOLCHAIN=clang
+export BUILD_TYPE=Debug
+
+${SCRIPT_DIR}/../build.sh

kokoro/linux/cmake-clang-debug/presubmit.cfg

@@ -0,0 +1,3 @@
+# Format: //devtools/kokoro/config/proto/build.proto
+
+build_file: "tint/kokoro/linux/cmake-clang-debug/build.sh"

kokoro/linux/cmake-clang-release-asan/build.sh

@@ -0,0 +1,26 @@
+#!/bin/bash
+
+# Copyright 2021 The Tint Authors.
+#
+# 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.
+
+set -e # Fail on any error.
+
+SCRIPT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}")" >/dev/null 2>&1 && pwd )"
+
+export BUILD_SYSTEM=cmake
+export BUILD_TOOLCHAIN=clang
+export BUILD_TYPE=Release
+export BUILD_SANITIZER=asan
+
+${SCRIPT_DIR}/../build.sh

kokoro/linux/cmake-clang-release-asan/presubmit.cfg

@@ -0,0 +1,3 @@
+# Format: //devtools/kokoro/config/proto/build.proto
+
+build_file: "tint/kokoro/linux/cmake-clang-release-asan/build.sh"

kokoro/linux/cmake-clang-release-ubsan/build.sh

@@ -0,0 +1,26 @@
+#!/bin/bash
+
+# Copyright 2021 The Tint Authors.
+#
+# 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.
+
+set -e # Fail on any error.
+
+SCRIPT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}")" >/dev/null 2>&1 && pwd )"
+
+export BUILD_SYSTEM=cmake
+export BUILD_TOOLCHAIN=clang
+export BUILD_TYPE=Release
+export BUILD_SANITIZER=ubsan
+
+${SCRIPT_DIR}/../build.sh

kokoro/linux/cmake-clang-release-ubsan/presubmit.cfg

@@ -0,0 +1,3 @@
+# Format: //devtools/kokoro/config/proto/build.proto
+
+build_file: "tint/kokoro/linux/cmake-clang-release-ubsan/build.sh"

kokoro/linux/cmake-clang-release/build.sh

@@ -0,0 +1,25 @@
+#!/bin/bash
+
+# Copyright 2021 The Tint Authors.
+#
+# 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.
+
+set -e # Fail on any error.
+
+SCRIPT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}")" >/dev/null 2>&1 && pwd )"
+
+export BUILD_SYSTEM=cmake
+export BUILD_TOOLCHAIN=clang
+export BUILD_TYPE=Release
+
+${SCRIPT_DIR}/../build.sh

kokoro/linux/cmake-clang-release/presubmit.cfg

@@ -0,0 +1,3 @@
+# Format: //devtools/kokoro/config/proto/build.proto
+
+build_file: "tint/kokoro/linux/cmake-clang-release/build.sh"

kokoro/linux/cmake-gcc-debug/build.sh

@@ -0,0 +1,25 @@
+#!/bin/bash
+
+# Copyright 2021 The Tint Authors.
+#
+# 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.
+
+set -e # Fail on any error.
+
+SCRIPT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}")" >/dev/null 2>&1 && pwd )"
+
+export BUILD_SYSTEM=cmake
+export BUILD_TOOLCHAIN=gcc
+export BUILD_TYPE=Debug
+
+${SCRIPT_DIR}/../build.sh

kokoro/linux/cmake-gcc-debug/presubmit.cfg

@@ -0,0 +1,3 @@
+# Format: //devtools/kokoro/config/proto/build.proto
+
+build_file: "tint/kokoro/linux/cmake-gcc-debug/build.sh"

kokoro/linux/cmake-gcc-release/build.sh

@@ -0,0 +1,25 @@
+#!/bin/bash
+
+# Copyright 2021 The Tint Authors.
+#
+# 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.
+
+set -e # Fail on any error.
+
+SCRIPT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}")" >/dev/null 2>&1 && pwd )"
+
+export BUILD_SYSTEM=cmake
+export BUILD_TOOLCHAIN=gcc
+export BUILD_TYPE=Release
+
+${SCRIPT_DIR}/../build.sh

kokoro/linux/cmake-gcc-release/presubmit.cfg

@@ -0,0 +1,3 @@
+# Format: //devtools/kokoro/config/proto/build.proto
+
+build_file: "tint/kokoro/linux/cmake-gcc-release/build.sh"

kokoro/linux/docker.sh

@@ -0,0 +1,183 @@
+#!/bin/bash
+
+# Copyright 2021 The Tint Authors.
+#
+# 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
+#
+#     https://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.
+
+# This is the bash script invoked inside a docker container.
+# The script expects that the CWD points to a clean checkout of Tint.
+# As `gclient sync` will litter the tint checkout with fetched tools and
+# projects, this script will first clone the pristine tint checkout to
+# ${SRC_DIR}. This allows developers to locally run this script without having
+# to worry about their local tint copy being touched.
+#
+# This script expects the following environment variables to be set on entry:
+#
+# SRC_DIR         - Path to where the local Tint copy will be made. See above.
+# BUILD_DIR       - Path to where Tint will be built.
+# BUILD_TYPE      - Either: 'Debug' or 'Release'
+# BUILD_SYSTEM    - Must be 'cmake'
+# BUILD_SANITIZER - Either: '', 'asan', or 'ubstan'
+# BUILD_TOOLCHAIN - Either: 'clang' or 'gcc'
+
+set -e # Fail on any error.
+
+function show_cmds { set -x; }
+function hide_cmds { { set +x; } 2>/dev/null; }
+function task_begin {
+    TASK_NAME="$@"
+    SECONDS=0
+}
+function print_last_task_duration {
+    if [ ! -z "${TASK_NAME}" ]; then
+        echo "${TASK_NAME} completed in $(($SECONDS / 3600))h$((($SECONDS / 60) % 60))m$(($SECONDS % 60))s"
+    fi
+}
+function status {
+    echo ""
+    echo ""
+    print_last_task_duration
+    echo ""
+    echo "*****************************************************************"
+    echo "* $@"
+    echo "*****************************************************************"
+    echo ""
+    task_begin $@
+}
+function with_retry {
+  local MAX_ATTEMPTS=5
+  local RETRY_DELAY_SECS=5
+  local ATTEMPT=1
+  while true; do
+    "$@" && break
+    if [[ $ATTEMPT -ge $MAX_ATTEMPTS ]]; then
+        echo "The command has failed after $ATTEMPT attempts."
+        exit $?
+    fi
+    ((ATTEMPT++))
+    echo "'$@' failed. Attempt ($ATTEMPT/$MAX_ATTEMPTS). Retrying..."
+    sleep $RETRY_DELAY_SECS;
+  done
+}
+
+CLONE_SRC_DIR="$(pwd)"
+
+. /bin/using.sh # Declare the bash `using` function for configuring toolchains.
+
+using depot_tools
+using go-1.14.4      # Speeds up ./tools/lint
+using doxygen-1.8.18
+
+status "Creating source directory '${SRC_DIR}' and build directory '${BUILD_DIR}'"
+mkdir -p ${SRC_DIR}
+mkdir -p ${BUILD_DIR}
+
+status "Cloning to source directory '${SRC_DIR}'"
+cd ${SRC_DIR}
+git clone ${CLONE_SRC_DIR} .
+
+status "Fetching dependencies"
+cp standalone.gclient .gclient
+with_retry gclient sync
+
+status "Linting"
+./tools/lint
+
+status "Configuring build system"
+if [ "$BUILD_SYSTEM" == "cmake" ]; then
+    using cmake-3.17.2
+
+    COMMON_CMAKE_FLAGS=""
+    COMMON_CMAKE_FLAGS+=" -DCMAKE_BUILD_TYPE=${BUILD_TYPE}"
+    COMMON_CMAKE_FLAGS+=" -DTINT_DOCS_WARN_AS_ERROR=1"
+    COMMON_CMAKE_FLAGS+=" -DTINT_BUILD_BENCHMARKS=1"
+
+    if [ "$BUILD_TOOLCHAIN" == "clang" ]; then
+        using clang-10.0.0
+        COMMON_CMAKE_FLAGS+=" -DTINT_BUILD_FUZZERS=1"
+        COMMON_CMAKE_FLAGS+=" -DTINT_BUILD_SPIRV_TOOLS_FUZZER=1"
+        COMMON_CMAKE_FLAGS+=" -DTINT_BUILD_AST_FUZZER=1"
+        COMMON_CMAKE_FLAGS+=" -DTINT_BUILD_REGEX_FUZZER=1"
+    elif [ "$BUILD_TOOLCHAIN" == "gcc" ]; then
+        using gcc-9
+    fi
+
+    if [ "$BUILD_SANITIZER" == "asan" ]; then
+        COMMON_CMAKE_FLAGS+=" -DTINT_ENABLE_ASAN=1"
+    elif [ "$BUILD_SANITIZER" == "ubsan" ]; then
+        COMMON_CMAKE_FLAGS+=" -DTINT_ENABLE_UBSAN=1"
+        export UBSAN_OPTIONS=print_stacktrace=1:halt_on_error=1
+    fi
+
+    cd ${BUILD_DIR}
+
+    status "Running Doxygen"
+    echo "NOTE: This will fail on first warning. Run with -DTINT_DOCS_WARN_AS_ERROR=OFF to see all warnings".
+    echo ""
+    show_cmds
+        # NOTE: If we upgrade Doxygen to a more recent version, we can set DOXYGEN_WARN_AS_ERROR to
+        # "FAIL_ON_WARNINGS" instead of "YES" in our CMakeLists.txt so see all warnings, and then
+        # fail. See https://www.doxygen.nl/manual/config.html#cfg_warn_as_error
+        cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS}
+        cmake --build . --target tint-docs
+    hide_cmds
+
+    status "Building tint in '${BUILD_DIR}'"
+    show_cmds
+        cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS}
+        cmake --build . -- --jobs=$(nproc)
+    hide_cmds
+
+    status "Running tint_unittests"
+    show_cmds
+        ./tint_unittests
+    hide_cmds
+
+    if [ -f ./tint_ast_fuzzer_unittests ]; then
+        status "Running tint_ast_fuzzer_unittests"
+        show_cmds
+            ./tint_ast_fuzzer_unittests
+        hide_cmds
+    fi
+
+    if [ -f ./tint_regex_fuzzer_unittests ]; then
+        status "Running tint_regex_fuzzer_unittests"
+        show_cmds
+            ./tint_regex_fuzzer_unittests
+        hide_cmds
+    fi
+
+    status "Testing test/tint/test-all.sh"
+    show_cmds
+        ${SRC_DIR}/test/tint/test-all.sh "${BUILD_DIR}/tint" --verbose
+    hide_cmds
+
+    status "Checking _other.cc files also build"
+    show_cmds
+        cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS} -DTINT_BUILD_AS_OTHER_OS=ON
+        cmake --build . -- --jobs=$(nproc)
+        cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS} -DTINT_BUILD_AS_OTHER_OS=OFF
+    hide_cmds
+
+    status "Checking disabling all readers and writers also builds"
+    show_cmds
+        cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS} -DTINT_BUILD_SPV_READER=OFF -DTINT_BUILD_SPV_WRITER=OFF -DTINT_BUILD_WGSL_READER=OFF -DTINT_BUILD_WGSL_WRITER=OFF -DTINT_BUILD_MSL_WRITER=OFF -DTINT_BUILD_HLSL_WRITER=OFF -DTINT_BUILD_BENCHMARKS=OFF
+        cmake --build . -- --jobs=$(nproc)
+        cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS} -DTINT_BUILD_SPV_READER=ON -DTINT_BUILD_SPV_WRITER=ON -DTINT_BUILD_WGSL_READER=ON -DTINT_BUILD_WGSL_WRITER=ON -DTINT_BUILD_MSL_WRITER=ON -DTINT_BUILD_HLSL_WRITER=ON -DTINT_BUILD_BENCHMARKS=ON
+    hide_cmds
+else
+    status "Unsupported build system: $BUILD_SYSTEM"
+    exit 1
+fi
+
+status "Done"

kokoro/windows/build.bat

@@ -0,0 +1,167 @@
+@rem Copyright 2021 The Tint Authors.
+@rem
+@rem Licensed under the Apache License, Version 2.0 (the "License");
+@rem you may not use this file except in compliance with the License.
+@rem You may obtain a copy of the License at
+@rem
+@rem     http://www.apache.org/licenses/LICENSE-2.0
+@rem
+@rem Unless required by applicable law or agreed to in writing, software
+@rem distributed under the License is distributed on an "AS IS" BASIS,
+@rem WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+@rem See the License for the specific language governing permissions and
+@rem limitations under the License.
+
+@echo off
+SETLOCAL ENABLEDELAYEDEXPANSION
+
+goto :main
+
+:task_begin
+set TASK_NAME=%~1
+echo %TASK_NAME% starting at %Time%
+exit /b 0
+
+:print_last_task_duration
+if not "%TASK_NAME%" == "" (
+    echo %TASK_NAME% completed at %Time%
+)
+exit /b 0
+
+:status
+echo.
+echo.
+call :print_last_task_duration
+echo.
+echo *****************************************************************
+echo %~1
+echo *****************************************************************
+echo.
+call :task_begin "%~1"
+exit /b 0
+
+:main
+
+set ORIGINAL_SRC_DIR= %~dp0\..\..
+set TEMP_DIR=%TEMP%\tint-temp
+set SRC_DIR="%TEMP_DIR%\tint-src"
+set BUILD_DIR="%TEMP_DIR%\tint-build"
+
+cd /d %ORIGINAL_SRC_DIR%
+if not exist ".git\" (
+    echo "ORIGINAL_SRC_DIR should point to project root: %ORIGINAL_SRC_DIR%"
+    goto :error
+)
+
+if exist %TEMP_DIR% (
+    call :status "Deleting %TEMP_DIR%"
+    del /q/f/s %TEMP_DIR% > NUL || goto :error
+    rmdir /q/s %TEMP_DIR% > NUL || goto :error
+)
+mkdir %TEMP_DIR% || goto :error
+
+call :status "Fetching and installing DXC"
+@echo on
+set DXC_RELEASE="https://github.com/microsoft/DirectXShaderCompiler/releases/download/v1.6.2112/dxc_2021_12_08.zip"
+curl -k -L %DXC_RELEASE% --output "%TEMP_DIR%\dxc_release.zip" || goto :error
+powershell.exe -Command "Expand-Archive -LiteralPath '%TEMP_DIR%\dxc_release.zip' -DestinationPath '%TEMP_DIR%\dxc'" || goto :error
+set DXC_PATH=%TEMP_DIR%\dxc\bin\x64
+
+rem Patch with artifact build that contains fixes not present in the release build
+set DXC_ARTIFACT="https://ci.appveyor.com/api/projects/dnovillo/directxshadercompiler/artifacts/build%%2FRelease%%2Fdxc-artifacts.zip?branch=master&pr=false&job=image%%3A%%20Visual%%20Studio%%202019"
+curl -k -L %DXC_ARTIFACT% --output "%TEMP_DIR%\dxc_artifact.zip" || goto :error
+powershell.exe -Command "Expand-Archive -Force -LiteralPath '%TEMP_DIR%\dxc_artifact.zip' -DestinationPath '%TEMP_DIR%\dxc_artifact'" || goto :error
+move /Y %TEMP_DIR%\dxc_artifact\bin\* %DXC_PATH%
+@echo off
+
+call :status "Fetching and installing Windows SDK for d3dcompiler DLL"
+@echo on
+set WINSDK_DLL_INSTALLER=https://go.microsoft.com/fwlink/?linkid=2164145
+set WINSDK_VERSION=10.0.20348.0
+curl -k -L %WINSDK_DLL_INSTALLER% --output "%TEMP_DIR%\winsdksetup.exe" || goto :error
+start "download" /wait "%TEMP_DIR%\winsdksetup.exe" /quiet /norestart /ceip off /features OptionId.DesktopCPPx64 /layout "%TEMP_DIR%\winsdkinstall" || goto :error
+start "install" /wait "%TEMP_DIR%\winsdkinstall\Installers\Windows SDK for Windows Store Apps Tools-x86_en-us.msi" || goto :error
+set D3DCOMPILER_PATH=C:\Program Files (x86)\Windows Kits\10\bin\%WINSDK_VERSION%\x64
+@echo off
+
+call :status "Installing depot_tools"
+@echo on
+pushd %TEMP_DIR%
+rem For Windows, we must download and extract a bundle.
+rem See https://chromium.googlesource.com/chromium/src/+/HEAD/docs/windows_build_instructions.md#install
+powershell -Command "(New-Object Net.WebClient).DownloadFile('https://storage.googleapis.com/chrome-infra/depot_tools.zip', 'depot_tools.zip')" || goto :error
+powershell -Command "Expand-Archive -Force 'depot_tools.zip' 'depot_tools'" || goto :error
+rem Run gclient once to install deps
+set PATH=%TEMP_DIR%\depot_tools;%PATH%
+set DEPOT_TOOLS_UPDATE=1
+set DEPOT_TOOLS_WIN_TOOLCHAIN=0
+call gclient || goto :error
+@echo off
+popd
+
+call :status "Cloning to clean source directory"
+@echo on
+mkdir %SRC_DIR% || goto :error
+cd /d %SRC_DIR% || goto :error
+call git clone %ORIGINAL_SRC_DIR% . || goto :error
+@echo off
+
+call :status "Fetching dependencies"
+@echo on
+copy standalone.gclient .gclient || goto :error
+call gclient sync || goto :error
+@echo off
+
+call :status "Configuring build system"
+@echo on
+mkdir %BUILD_DIR%
+cd /d %BUILD_DIR%
+set COMMON_CMAKE_FLAGS=-DTINT_BUILD_DOCS=O -DTINT_BUILD_BENCHMARKS=1 -DCMAKE_BUILD_TYPE=%BUILD_TYPE%
+@echo off
+
+call :status "Building tint"
+@echo on
+rem Disable msbuild "Intermediate or Output directory cannot reside in Temporary directory"
+set IgnoreWarnIntDirInTempDetected=true
+rem Add Python3 to path as this Kokoro image only has Python2 in it
+set PATH=C:\Python37;%PATH%
+rem To use ninja with CMake requires VC env vars
+call "C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Auxiliary\Build\vcvars64.bat"
+@echo on
+rem Note that we need to specify the C and C++ compiler only because Cygwin is in PATH and CMake finds GCC and picks that over MSVC
+cmake %SRC_DIR% -G "Ninja" -DCMAKE_C_COMPILER="cl.exe" -DCMAKE_CXX_COMPILER="cl.exe" %COMMON_CMAKE_FLAGS% || goto :error
+cmake --build . || goto :error
+call "C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Auxiliary\Build\vcvars64.bat" /clean_env
+@echo off
+
+call :status "Running tint_unittests"
+@echo on
+tint_unittests.exe || goto :error
+@echo off
+
+call :status "Testing test/tint/test-all.sh"
+@echo on
+cd /d %SRC_DIR% || goto :error
+rem Run tests with DXC and Metal validation
+set OLD_PATH=%PATH%
+set PATH=C:\Program Files\Metal Developer Tools\macos\bin;%PATH%
+where metal.exe
+set PATH=%DXC_PATH%;%OLD_PATH%
+where dxc.exe dxil.dll
+call git bash -- ./test/tint/test-all.sh ../tint-build/tint.exe --verbose || goto :error
+@echo on
+set PATH=%OLD_PATH%
+rem Run again to test with FXC validation
+set PATH=%D3DCOMPILER_PATH%;%OLD_PATH%
+where d3dcompiler_47.dll
+call git bash -- ./test/tint/test-all.sh ../tint-build/tint.exe --verbose --format hlsl --fxc || goto :error
+@echo on
+set PATH=%OLD_PATH%
+@echo off
+
+call :status "Done"
+exit /b 0
+
+:error
+echo BUILD FAILED! errorlevel: %errorlevel%
+exit /b %errorlevel%

kokoro/windows/cmake-msvc2019-debug/build.bat

@@ -0,0 +1,18 @@
+@rem Copyright 2021 The Tint Authors.
+@rem
+@rem Licensed under the Apache License, Version 2.0 (the "License");
+@rem you may not use this file except in compliance with the License.
+@rem You may obtain a copy of the License at
+@rem
+@rem     http://www.apache.org/licenses/LICENSE-2.0
+@rem
+@rem Unless required by applicable law or agreed to in writing, software
+@rem distributed under the License is distributed on an "AS IS" BASIS,
+@rem WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+@rem See the License for the specific language governing permissions and
+@rem limitations under the License.
+
+@echo on
+set BUILD_TYPE=Debug
+call %~dp0\..\build.bat
+exit /b %errorlevel%

kokoro/windows/cmake-msvc2019-debug/presubmit.cfg

@@ -0,0 +1,3 @@
+# Format: //devtools/kokoro/config/proto/build.proto
+
+build_file: "tint/kokoro/windows/cmake-msvc2019-debug/build.bat"

kokoro/windows/cmake-msvc2019-release/build.bat

@@ -0,0 +1,18 @@
+@rem Copyright 2021 The Tint Authors.
+@rem
+@rem Licensed under the Apache License, Version 2.0 (the "License");
+@rem you may not use this file except in compliance with the License.
+@rem You may obtain a copy of the License at
+@rem
+@rem     http://www.apache.org/licenses/LICENSE-2.0
+@rem
+@rem Unless required by applicable law or agreed to in writing, software
+@rem distributed under the License is distributed on an "AS IS" BASIS,
+@rem WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+@rem See the License for the specific language governing permissions and
+@rem limitations under the License.
+
+@echo on
+set BUILD_TYPE=Release
+call %~dp0\..\build.bat
+exit /b %errorlevel%

kokoro/windows/cmake-msvc2019-release/presubmit.cfg

@@ -0,0 +1,3 @@
+# Format: //devtools/kokoro/config/proto/build.proto
+
+build_file: "tint/kokoro/windows/cmake-msvc2019-release/build.bat"

scripts/dawn_overrides_with_defaults.gni

@@ -75,11 +75,6 @@ if (!defined(dawn_vulkan_validation_layers_dir)) {
   dawn_vulkan_validation_layers_dir = ""
 }
 
-if (!defined(dawn_tint_dir)) {
-  # Default to Tint being Dawn's DEPS
-  dawn_tint_dir = "${dawn_root}/third_party/tint"
-}
-
 if (!defined(dawn_abseil_dir)) {
   dawn_abseil_dir = "//third_party/abseil-cpp"
 }

src/dawn/native/BUILD.gn

@@ -163,7 +163,7 @@ source_set("sources") {
     "${dawn_root}/src/dawn/common",
     "${dawn_spirv_tools_dir}:spvtools_opt",
     "${dawn_spirv_tools_dir}:spvtools_val",
-    "${dawn_tint_dir}/src/tint:libtint",
+    "${dawn_root}/src/tint:libtint",
   ]
   defines = []
   libs = []

src/tint/BUILD.gn

@@ -0,0 +1,799 @@
+# Copyright 2021 The Tint Authors
+#
+# 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("//build_overrides/build.gni")
+import("../../tint_overrides_with_defaults.gni")
+
+###############################################################################
+# Common - Configs, etc. shared across targets
+###############################################################################
+
+config("tint_common_config") {
+  include_dirs = [
+    "${target_gen_dir}",
+    "${tint_root_dir}/",
+    "${tint_spirv_headers_dir}/include",
+    "${tint_spirv_tools_dir}/",
+    "${tint_spirv_tools_dir}/include",
+  ]
+}
+
+config("tint_public_config") {
+  defines = []
+  if (tint_build_spv_reader) {
+    defines += [ "TINT_BUILD_SPV_READER=1" ]
+  } else {
+    defines += [ "TINT_BUILD_SPV_READER=0" ]
+  }
+
+  if (tint_build_spv_writer) {
+    defines += [ "TINT_BUILD_SPV_WRITER=1" ]
+  } else {
+    defines += [ "TINT_BUILD_SPV_WRITER=0" ]
+  }
+
+  if (tint_build_wgsl_reader) {
+    defines += [ "TINT_BUILD_WGSL_READER=1" ]
+  } else {
+    defines += [ "TINT_BUILD_WGSL_READER=0" ]
+  }
+
+  if (tint_build_wgsl_writer) {
+    defines += [ "TINT_BUILD_WGSL_WRITER=1" ]
+  } else {
+    defines += [ "TINT_BUILD_WGSL_WRITER=0" ]
+  }
+
+  if (tint_build_msl_writer) {
+    defines += [ "TINT_BUILD_MSL_WRITER=1" ]
+  } else {
+    defines += [ "TINT_BUILD_MSL_WRITER=0" ]
+  }
+
+  if (tint_build_hlsl_writer) {
+    defines += [ "TINT_BUILD_HLSL_WRITER=1" ]
+  } else {
+    defines += [ "TINT_BUILD_HLSL_WRITER=0" ]
+  }
+
+  if (tint_build_glsl_writer) {
+    defines += [ "TINT_BUILD_GLSL_WRITER=1" ]
+  } else {
+    defines += [ "TINT_BUILD_GLSL_WRITER=0" ]
+  }
+
+  include_dirs = [
+    "${tint_root_dir}/",
+    "${tint_root_dir}/include/",
+    "${tint_spirv_headers_dir}/include",
+  ]
+}
+
+config("tint_config") {
+  include_dirs = []
+  if (tint_build_spv_reader || tint_build_spv_writer) {
+    include_dirs += [ "${tint_spirv_tools_dir}/include/" ]
+  }
+}
+
+###############################################################################
+# Helper library for IO operations
+# Only to be used by tests and sample executable
+###############################################################################
+source_set("tint_utils_io") {
+  sources = [
+    "utils/io/command.h",
+    "utils/io/tmpfile.h",
+  ]
+
+  if (is_linux || is_mac) {
+    sources += [ "utils/io/command_posix.cc" ]
+    sources += [ "utils/io/tmpfile_posix.cc" ]
+  } else if (is_win) {
+    sources += [ "utils/io/command_windows.cc" ]
+    sources += [ "utils/io/tmpfile_windows.cc" ]
+  } else {
+    sources += [ "utils/io/command_other.cc" ]
+    sources += [ "utils/io/tmpfile_other.cc" ]
+  }
+
+  public_deps = [ ":libtint_core_all_src" ]
+}
+
+###############################################################################
+# Helper library for validating generated shaders
+# As this depends on tint_utils_io, this is only to be used by tests and sample
+# executable
+###############################################################################
+source_set("tint_val") {
+  sources = [
+    "val/hlsl.cc",
+    "val/msl.cc",
+    "val/val.h",
+  ]
+  public_deps = [ ":tint_utils_io" ]
+}
+
+###############################################################################
+# Library - Tint core and optional modules of libtint
+###############################################################################
+# libtint source sets are divided into a non-optional core in :libtint_core_src
+# and optional :libtint_*_src subsets, because ninja does not like having
+# multiple source files with the same name, like function.cc, in the same
+# source set
+# target.
+#
+# Targets that want to use tint as a library should depend on ":libtint" and
+# use the build flags to control what is included, instead of trying to specify
+# the subsets that they want.
+
+template("libtint_source_set") {
+  source_set(target_name) {
+    forward_variables_from(invoker, "*", [ "configs" ])
+
+    if (!defined(invoker.deps)) {
+      deps = []
+    }
+    deps += [
+      "${tint_spirv_headers_dir}:spv_headers",
+      "${tint_spirv_tools_dir}:spvtools_core_enums_unified1",
+      "${tint_spirv_tools_dir}:spvtools_core_tables_unified1",
+      "${tint_spirv_tools_dir}:spvtools_headers",
+      "${tint_spirv_tools_dir}:spvtools_language_header_cldebuginfo100",
+      "${tint_spirv_tools_dir}:spvtools_language_header_debuginfo",
+      "${tint_spirv_tools_dir}:spvtools_language_header_vkdebuginfo100",
+    ]
+
+    if (defined(invoker.configs)) {
+      configs += invoker.configs
+    }
+    configs += [ ":tint_common_config" ]
+    if (build_with_chromium) {
+      configs -= [ "//build/config/compiler:chromium_code" ]
+      configs += [ "//build/config/compiler:no_chromium_code" ]
+    }
+
+    if (!defined(invoker.public_configs)) {
+      public_configs = []
+    }
+    public_configs += [ ":tint_public_config" ]
+  }
+}
+
+libtint_source_set("libtint_core_all_src") {
+  sources = [
+    "ast/access.cc",
+    "ast/access.h",
+    "ast/alias.cc",
+    "ast/alias.h",
+    "ast/array.cc",
+    "ast/array.h",
+    "ast/assignment_statement.cc",
+    "ast/assignment_statement.h",
+    "ast/ast_type.cc",  # TODO(bclayton) - rename to type.cc
+    "ast/atomic.cc",
+    "ast/atomic.h",
+    "ast/attribute.cc",
+    "ast/attribute.h",
+    "ast/binary_expression.cc",
+    "ast/binary_expression.h",
+    "ast/binding_attribute.cc",
+    "ast/binding_attribute.h",
+    "ast/bitcast_expression.cc",
+    "ast/bitcast_expression.h",
+    "ast/block_statement.cc",
+    "ast/block_statement.h",
+    "ast/bool.cc",
+    "ast/bool.h",
+    "ast/bool_literal_expression.cc",
+    "ast/bool_literal_expression.h",
+    "ast/break_statement.cc",
+    "ast/break_statement.h",
+    "ast/builtin.cc",
+    "ast/builtin.h",
+    "ast/builtin_attribute.cc",
+    "ast/builtin_attribute.h",
+    "ast/call_expression.cc",
+    "ast/call_expression.h",
+    "ast/call_statement.cc",
+    "ast/call_statement.h",
+    "ast/case_statement.cc",
+    "ast/case_statement.h",
+    "ast/compound_assignment_statement.cc",
+    "ast/compound_assignment_statement.h",
+    "ast/continue_statement.cc",
+    "ast/continue_statement.h",
+    "ast/depth_multisampled_texture.cc",
+    "ast/depth_multisampled_texture.h",
+    "ast/depth_texture.cc",
+    "ast/depth_texture.h",
+    "ast/disable_validation_attribute.cc",
+    "ast/disable_validation_attribute.h",
+    "ast/discard_statement.cc",
+    "ast/discard_statement.h",
+    "ast/else_statement.cc",
+    "ast/else_statement.h",
+    "ast/expression.cc",
+    "ast/expression.h",
+    "ast/external_texture.cc",
+    "ast/external_texture.h",
+    "ast/f32.cc",
+    "ast/f32.h",
+    "ast/fallthrough_statement.cc",
+    "ast/fallthrough_statement.h",
+    "ast/float_literal_expression.cc",
+    "ast/float_literal_expression.h",
+    "ast/for_loop_statement.cc",
+    "ast/for_loop_statement.h",
+    "ast/function.cc",
+    "ast/function.h",
+    "ast/group_attribute.cc",
+    "ast/group_attribute.h",
+    "ast/i32.cc",
+    "ast/i32.h",
+    "ast/id_attribute.cc",
+    "ast/id_attribute.h",
+    "ast/identifier_expression.cc",
+    "ast/identifier_expression.h",
+    "ast/if_statement.cc",
+    "ast/if_statement.h",
+    "ast/index_accessor_expression.cc",
+    "ast/index_accessor_expression.h",
+    "ast/int_literal_expression.cc",
+    "ast/int_literal_expression.h",
+    "ast/internal_attribute.cc",
+    "ast/internal_attribute.h",
+    "ast/interpolate_attribute.cc",
+    "ast/interpolate_attribute.h",
+    "ast/invariant_attribute.cc",
+    "ast/invariant_attribute.h",
+    "ast/literal_expression.cc",
+    "ast/literal_expression.h",
+    "ast/location_attribute.cc",
+    "ast/location_attribute.h",
+    "ast/loop_statement.cc",
+    "ast/loop_statement.h",
+    "ast/matrix.cc",
+    "ast/matrix.h",
+    "ast/member_accessor_expression.cc",
+    "ast/member_accessor_expression.h",
+    "ast/module.cc",
+    "ast/module.h",
+    "ast/multisampled_texture.cc",
+    "ast/multisampled_texture.h",
+    "ast/node.cc",
+    "ast/node.h",
+    "ast/phony_expression.cc",
+    "ast/phony_expression.h",
+    "ast/pipeline_stage.cc",
+    "ast/pipeline_stage.h",
+    "ast/pointer.cc",
+    "ast/pointer.h",
+    "ast/return_statement.cc",
+    "ast/return_statement.h",
+    "ast/sampled_texture.cc",
+    "ast/sampled_texture.h",
+    "ast/sampler.cc",
+    "ast/sampler.h",
+    "ast/sint_literal_expression.cc",
+    "ast/sint_literal_expression.h",
+    "ast/stage_attribute.cc",
+    "ast/stage_attribute.h",
+    "ast/statement.cc",
+    "ast/statement.h",
+    "ast/storage_class.cc",
+    "ast/storage_class.h",
+    "ast/storage_texture.cc",
+    "ast/storage_texture.h",
+    "ast/stride_attribute.cc",
+    "ast/stride_attribute.h",
+    "ast/struct.cc",
+    "ast/struct.h",
+    "ast/struct_member.cc",
+    "ast/struct_member.h",
+    "ast/struct_member_align_attribute.cc",
+    "ast/struct_member_align_attribute.h",
+    "ast/struct_member_offset_attribute.cc",
+    "ast/struct_member_offset_attribute.h",
+    "ast/struct_member_size_attribute.cc",
+    "ast/struct_member_size_attribute.h",
+    "ast/switch_statement.cc",
+    "ast/switch_statement.h",
+    "ast/texture.cc",
+    "ast/texture.h",
+    "ast/traverse_expressions.h",
+    "ast/type.h",
+    "ast/type_decl.cc",
+    "ast/type_decl.h",
+    "ast/type_name.cc",
+    "ast/type_name.h",
+    "ast/u32.cc",
+    "ast/u32.h",
+    "ast/uint_literal_expression.cc",
+    "ast/uint_literal_expression.h",
+    "ast/unary_op.cc",
+    "ast/unary_op.h",
+    "ast/unary_op_expression.cc",
+    "ast/unary_op_expression.h",
+    "ast/variable.cc",
+    "ast/variable.h",
+    "ast/variable_decl_statement.cc",
+    "ast/variable_decl_statement.h",
+    "ast/vector.cc",
+    "ast/vector.h",
+    "ast/void.cc",
+    "ast/void.h",
+    "ast/workgroup_attribute.cc",
+    "ast/workgroup_attribute.h",
+    "builtin_table.cc",
+    "builtin_table.h",
+    "builtin_table.inl",
+    "castable.cc",
+    "castable.h",
+    "clone_context.cc",
+    "clone_context.h",
+    "debug.cc",
+    "debug.h",
+    "demangler.cc",
+    "demangler.h",
+    "diagnostic/diagnostic.cc",
+    "diagnostic/diagnostic.h",
+    "diagnostic/formatter.cc",
+    "diagnostic/formatter.h",
+    "diagnostic/printer.cc",
+    "diagnostic/printer.h",
+    "inspector/entry_point.cc",
+    "inspector/entry_point.h",
+    "inspector/inspector.cc",
+    "inspector/inspector.h",
+    "inspector/resource_binding.cc",
+    "inspector/resource_binding.h",
+    "inspector/scalar.cc",
+    "inspector/scalar.h",
+    "program.cc",
+    "program.h",
+    "program_builder.cc",
+    "program_builder.h",
+    "program_id.cc",
+    "program_id.h",
+    "reader/reader.cc",
+    "reader/reader.h",
+    "resolver/dependency_graph.cc",
+    "resolver/dependency_graph.h",
+    "resolver/resolver.cc",
+    "resolver/resolver.h",
+    "resolver/resolver_constants.cc",
+    "resolver/resolver_validation.cc",
+    "scope_stack.h",
+    "sem/array.h",
+    "sem/atomic_type.h",
+    "sem/behavior.h",
+    "sem/binding_point.h",
+    "sem/bool_type.h",
+    "sem/builtin.h",
+    "sem/builtin_type.h",
+    "sem/call.h",
+    "sem/call_target.h",
+    "sem/constant.h",
+    "sem/depth_multisampled_texture_type.h",
+    "sem/depth_texture_type.h",
+    "sem/expression.h",
+    "sem/external_texture_type.h",
+    "sem/f32_type.h",
+    "sem/for_loop_statement.h",
+    "sem/i32_type.h",
+    "sem/if_statement.h",
+    "sem/info.h",
+    "sem/loop_statement.h",
+    "sem/matrix_type.h",
+    "sem/module.h",
+    "sem/multisampled_texture_type.h",
+    "sem/node.h",
+    "sem/parameter_usage.h",
+    "sem/pipeline_stage_set.h",
+    "sem/pointer_type.h",
+    "sem/reference_type.h",
+    "sem/sampled_texture_type.h",
+    "sem/sampler_texture_pair.h",
+    "sem/sampler_type.h",
+    "sem/storage_texture_type.h",
+    "sem/switch_statement.h",
+    "sem/texture_type.h",
+    "sem/type.h",
+    "sem/type_constructor.h",
+    "sem/type_conversion.h",
+    "sem/type_manager.h",
+    "sem/type_mappings.h",
+    "sem/u32_type.h",
+    "sem/vector_type.h",
+    "sem/void_type.h",
+    "source.cc",
+    "source.h",
+    "symbol.cc",
+    "symbol.h",
+    "symbol_table.cc",
+    "symbol_table.h",
+    "text/unicode.cc",
+    "text/unicode.h",
+    "traits.h",
+    "transform/add_empty_entry_point.cc",
+    "transform/add_empty_entry_point.h",
+    "transform/add_spirv_block_attribute.cc",
+    "transform/add_spirv_block_attribute.h",
+    "transform/array_length_from_uniform.cc",
+    "transform/array_length_from_uniform.h",
+    "transform/binding_remapper.cc",
+    "transform/binding_remapper.h",
+    "transform/builtin_polyfill.cc",
+    "transform/builtin_polyfill.h",
+    "transform/calculate_array_length.cc",
+    "transform/calculate_array_length.h",
+    "transform/canonicalize_entry_point_io.cc",
+    "transform/canonicalize_entry_point_io.h",
+    "transform/combine_samplers.cc",
+    "transform/combine_samplers.h",
+    "transform/decompose_memory_access.cc",
+    "transform/decompose_memory_access.h",
+    "transform/decompose_strided_array.cc",
+    "transform/decompose_strided_array.h",
+    "transform/decompose_strided_matrix.cc",
+    "transform/decompose_strided_matrix.h",
+    "transform/first_index_offset.cc",
+    "transform/first_index_offset.h",
+    "transform/fold_constants.cc",
+    "transform/fold_constants.h",
+    "transform/fold_trivial_single_use_lets.cc",
+    "transform/fold_trivial_single_use_lets.h",
+    "transform/for_loop_to_loop.cc",
+    "transform/for_loop_to_loop.h",
+    "transform/expand_compound_assignment.cc",
+    "transform/expand_compound_assignment.h",
+    "transform/localize_struct_array_assignment.cc",
+    "transform/localize_struct_array_assignment.h",
+    "transform/loop_to_for_loop.cc",
+    "transform/loop_to_for_loop.h",
+    "transform/manager.cc",
+    "transform/manager.h",
+    "transform/module_scope_var_to_entry_point_param.cc",
+    "transform/module_scope_var_to_entry_point_param.h",
+    "transform/multiplanar_external_texture.cc",
+    "transform/multiplanar_external_texture.h",
+    "transform/num_workgroups_from_uniform.cc",
+    "transform/num_workgroups_from_uniform.h",
+    "transform/promote_initializers_to_const_var.cc",
+    "transform/promote_initializers_to_const_var.h",
+    "transform/promote_side_effects_to_decl.cc",
+    "transform/promote_side_effects_to_decl.h",
+    "transform/remove_continue_in_switch.cc",
+    "transform/remove_continue_in_switch.h",
+    "transform/remove_phonies.cc",
+    "transform/remove_phonies.h",
+    "transform/remove_unreachable_statements.cc",
+    "transform/remove_unreachable_statements.h",
+    "transform/renamer.cc",
+    "transform/renamer.h",
+    "transform/robustness.cc",
+    "transform/robustness.h",
+    "transform/simplify_pointers.cc",
+    "transform/simplify_pointers.h",
+    "transform/single_entry_point.cc",
+    "transform/single_entry_point.h",
+    "transform/transform.cc",
+    "transform/transform.h",
+    "transform/unshadow.cc",
+    "transform/unshadow.h",
+    "transform/unwind_discard_functions.cc",
+    "transform/unwind_discard_functions.h",
+    "transform/utils/get_insertion_point.cc",
+    "transform/utils/get_insertion_point.h",
+    "transform/utils/hoist_to_decl_before.cc",
+    "transform/utils/hoist_to_decl_before.h",
+    "transform/var_for_dynamic_index.cc",
+    "transform/var_for_dynamic_index.h",
+    "transform/vectorize_scalar_matrix_constructors.cc",
+    "transform/vectorize_scalar_matrix_constructors.h",
+    "transform/vertex_pulling.cc",
+    "transform/vertex_pulling.h",
+    "transform/wrap_arrays_in_structs.cc",
+    "transform/wrap_arrays_in_structs.h",
+    "transform/zero_init_workgroup_memory.cc",
+    "transform/zero_init_workgroup_memory.h",
+    "utils/block_allocator.h",
+    "utils/crc32.h",
+    "utils/debugger.cc",
+    "utils/debugger.h",
+    "utils/enum_set.h",
+    "utils/hash.h",
+    "utils/map.h",
+    "utils/math.h",
+    "utils/scoped_assignment.h",
+    "utils/string.h",
+    "utils/unique_allocator.h",
+    "utils/unique_vector.h",
+    "writer/append_vector.cc",
+    "writer/append_vector.h",
+    "writer/array_length_from_uniform_options.cc",
+    "writer/array_length_from_uniform_options.h",
+    "writer/float_to_string.cc",
+    "writer/float_to_string.h",
+    "writer/generate_external_texture_bindings.cc",
+    "writer/generate_external_texture_bindings.h",
+    "writer/text.cc",
+    "writer/text.h",
+    "writer/text_generator.cc",
+    "writer/text_generator.h",
+    "writer/writer.cc",
+    "writer/writer.h",
+  ]
+
+  if (is_linux) {
+    sources += [ "diagnostic/printer_linux.cc" ]
+  } else if (is_win) {
+    sources += [ "diagnostic/printer_windows.cc" ]
+  } else {
+    sources += [ "diagnostic/printer_other.cc" ]
+  }
+}
+
+libtint_source_set("libtint_sem_src") {
+  sources = [
+    "sem/array.cc",
+    "sem/array.h",
+    "sem/atomic_type.cc",
+    "sem/atomic_type.h",
+    "sem/behavior.cc",
+    "sem/behavior.h",
+    "sem/binding_point.h",
+    "sem/block_statement.cc",
+    "sem/bool_type.cc",
+    "sem/bool_type.h",
+    "sem/builtin.cc",
+    "sem/builtin.h",
+    "sem/builtin_type.cc",
+    "sem/builtin_type.h",
+    "sem/call.cc",
+    "sem/call.h",
+    "sem/call_target.cc",
+    "sem/call_target.h",
+    "sem/constant.cc",
+    "sem/constant.h",
+    "sem/depth_multisampled_texture_type.cc",
+    "sem/depth_multisampled_texture_type.h",
+    "sem/depth_texture_type.cc",
+    "sem/depth_texture_type.h",
+    "sem/expression.cc",
+    "sem/expression.h",
+    "sem/external_texture_type.cc",
+    "sem/external_texture_type.h",
+    "sem/f32_type.cc",
+    "sem/f32_type.h",
+    "sem/for_loop_statement.cc",
+    "sem/for_loop_statement.h",
+    "sem/function.cc",
+    "sem/i32_type.cc",
+    "sem/i32_type.h",
+    "sem/if_statement.cc",
+    "sem/if_statement.h",
+    "sem/info.cc",
+    "sem/info.h",
+    "sem/loop_statement.cc",
+    "sem/loop_statement.h",
+    "sem/matrix_type.cc",
+    "sem/matrix_type.h",
+    "sem/member_accessor_expression.cc",
+    "sem/module.cc",
+    "sem/module.h",
+    "sem/multisampled_texture_type.cc",
+    "sem/multisampled_texture_type.h",
+    "sem/node.cc",
+    "sem/node.h",
+    "sem/parameter_usage.cc",
+    "sem/parameter_usage.h",
+    "sem/pipeline_stage_set.h",
+    "sem/pointer_type.cc",
+    "sem/pointer_type.h",
+    "sem/reference_type.cc",
+    "sem/reference_type.h",
+    "sem/sampled_texture_type.cc",
+    "sem/sampled_texture_type.h",
+    "sem/sampler_type.cc",
+    "sem/sampler_type.h",
+    "sem/statement.cc",
+    "sem/storage_texture_type.cc",
+    "sem/storage_texture_type.h",
+    "sem/struct.cc",
+    "sem/switch_statement.cc",
+    "sem/switch_statement.h",
+    "sem/texture_type.cc",
+    "sem/texture_type.h",
+    "sem/type.cc",
+    "sem/type.h",
+    "sem/type_constructor.cc",
+    "sem/type_constructor.h",
+    "sem/type_conversion.cc",
+    "sem/type_conversion.h",
+    "sem/type_manager.cc",
+    "sem/type_manager.h",
+    "sem/type_mappings.h",
+    "sem/u32_type.cc",
+    "sem/u32_type.h",
+    "sem/variable.cc",
+    "sem/vector_type.cc",
+    "sem/vector_type.h",
+    "sem/void_type.cc",
+    "sem/void_type.h",
+  ]
+
+  public_deps = [ ":libtint_core_all_src" ]
+}
+
+libtint_source_set("libtint_core_src") {
+  public_deps = [
+    ":libtint_core_all_src",
+    ":libtint_sem_src",
+  ]
+}
+
+libtint_source_set("libtint_spv_reader_src") {
+  sources = [
+    "reader/spirv/construct.cc",
+    "reader/spirv/construct.h",
+    "reader/spirv/entry_point_info.cc",
+    "reader/spirv/entry_point_info.h",
+    "reader/spirv/enum_converter.cc",
+    "reader/spirv/enum_converter.h",
+    "reader/spirv/fail_stream.h",
+    "reader/spirv/function.cc",
+    "reader/spirv/function.h",
+    "reader/spirv/namer.cc",
+    "reader/spirv/namer.h",
+    "reader/spirv/parser.cc",
+    "reader/spirv/parser.h",
+    "reader/spirv/parser_impl.cc",
+    "reader/spirv/parser_impl.h",
+    "reader/spirv/parser_type.cc",
+    "reader/spirv/parser_type.h",
+    "reader/spirv/usage.cc",
+    "reader/spirv/usage.h",
+  ]
+
+  public_deps = [
+    ":libtint_core_src",
+    "${tint_spirv_tools_dir}/:spvtools_opt",
+  ]
+
+  public_configs = [ "${tint_spirv_tools_dir}/:spvtools_internal_config" ]
+}
+
+libtint_source_set("libtint_spv_writer_src") {
+  sources = [
+    "writer/spirv/binary_writer.cc",
+    "writer/spirv/binary_writer.h",
+    "writer/spirv/builder.cc",
+    "writer/spirv/builder.h",
+    "writer/spirv/function.cc",
+    "writer/spirv/function.h",
+    "writer/spirv/generator.cc",
+    "writer/spirv/generator.h",
+    "writer/spirv/instruction.cc",
+    "writer/spirv/instruction.h",
+    "writer/spirv/operand.cc",
+    "writer/spirv/operand.h",
+    "writer/spirv/scalar_constant.h",
+  ]
+
+  public_deps = [ ":libtint_core_src" ]
+}
+
+libtint_source_set("libtint_wgsl_reader_src") {
+  sources = [
+    "reader/wgsl/lexer.cc",
+    "reader/wgsl/lexer.h",
+    "reader/wgsl/parser.cc",
+    "reader/wgsl/parser.h",
+    "reader/wgsl/parser_impl.cc",
+    "reader/wgsl/parser_impl.h",
+    "reader/wgsl/parser_impl_detail.h",
+    "reader/wgsl/token.cc",
+    "reader/wgsl/token.h",
+  ]
+
+  public_deps = [ ":libtint_core_src" ]
+}
+
+libtint_source_set("libtint_wgsl_writer_src") {
+  sources = [
+    "writer/wgsl/generator.cc",
+    "writer/wgsl/generator.h",
+    "writer/wgsl/generator_impl.cc",
+    "writer/wgsl/generator_impl.h",
+  ]
+
+  public_deps = [ ":libtint_core_src" ]
+}
+
+libtint_source_set("libtint_msl_writer_src") {
+  sources = [
+    "writer/msl/generator.cc",
+    "writer/msl/generator.h",
+    "writer/msl/generator_impl.cc",
+    "writer/msl/generator_impl.h",
+  ]
+
+  public_deps = [ ":libtint_core_src" ]
+}
+
+libtint_source_set("libtint_hlsl_writer_src") {
+  sources = [
+    "writer/hlsl/generator.cc",
+    "writer/hlsl/generator.h",
+    "writer/hlsl/generator_impl.cc",
+    "writer/hlsl/generator_impl.h",
+  ]
+
+  public_deps = [ ":libtint_core_src" ]
+}
+
+libtint_source_set("libtint_glsl_writer_src") {
+  sources = [
+    "transform/glsl.cc",
+    "transform/glsl.h",
+    "writer/glsl/generator.cc",
+    "writer/glsl/generator.h",
+    "writer/glsl/generator_impl.cc",
+    "writer/glsl/generator_impl.h",
+  ]
+
+  public_deps = [ ":libtint_core_src" ]
+}
+
+source_set("libtint") {
+  public_deps = [ ":libtint_core_src" ]
+
+  if (tint_build_spv_reader) {
+    public_deps += [ ":libtint_spv_reader_src" ]
+  }
+
+  if (tint_build_spv_writer) {
+    public_deps += [ ":libtint_spv_writer_src" ]
+  }
+
+  if (tint_build_wgsl_reader) {
+    public_deps += [ ":libtint_wgsl_reader_src" ]
+  }
+
+  if (tint_build_wgsl_writer) {
+    public_deps += [ ":libtint_wgsl_writer_src" ]
+  }
+
+  if (tint_build_msl_writer) {
+    public_deps += [ ":libtint_msl_writer_src" ]
+  }
+
+  if (tint_build_hlsl_writer) {
+    public_deps += [ ":libtint_hlsl_writer_src" ]
+  }
+
+  if (tint_build_glsl_writer) {
+    public_deps += [ ":libtint_glsl_writer_src" ]
+  }
+
+  configs += [ ":tint_common_config" ]
+  public_configs = [ ":tint_public_config" ]
+
+  if (build_with_chromium) {
+    configs -= [ "//build/config/compiler:chromium_code" ]
+    configs += [ "//build/config/compiler:no_chromium_code" ]
+  }
+}

src/tint/ast/access.cc

@@ -0,0 +1,43 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/access.h"
+
+namespace tint {
+namespace ast {
+
+std::ostream& operator<<(std::ostream& out, Access access) {
+  switch (access) {
+    case ast::Access::kUndefined: {
+      out << "undefined";
+      break;
+    }
+    case ast::Access::kRead: {
+      out << "read";
+      break;
+    }
+    case ast::Access::kReadWrite: {
+      out << "read_write";
+      break;
+    }
+    case ast::Access::kWrite: {
+      out << "write";
+      break;
+    }
+  }
+  return out;
+}
+
+}  // namespace ast
+}  // namespace tint

src/tint/ast/access.h

@@ -0,0 +1,46 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#ifndef SRC_TINT_AST_ACCESS_H_
+#define SRC_TINT_AST_ACCESS_H_
+
+#include <ostream>
+#include <string>
+
+namespace tint {
+namespace ast {
+
+/// The access control settings
+enum Access {
+  /// Not declared in the source
+  kUndefined = 0,
+  /// Read only
+  kRead,
+  /// Write only
+  kWrite,
+  /// Read write
+  kReadWrite,
+  // Last valid access mode
+  kLastValid = kReadWrite,
+};
+
+/// @param out the std::ostream to write to
+/// @param access the Access
+/// @return the std::ostream so calls can be chained
+std::ostream& operator<<(std::ostream& out, Access access);
+
+}  // namespace ast
+}  // namespace tint
+
+#endif  // SRC_TINT_AST_ACCESS_H_

src/tint/ast/alias.cc

@@ -0,0 +1,45 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/alias.h"
+
+#include "src/tint/program_builder.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::Alias);
+
+namespace tint {
+namespace ast {
+
+Alias::Alias(ProgramID pid,
+             const Source& src,
+             const Symbol& n,
+             const Type* subtype)
+    : Base(pid, src, n), type(subtype) {
+  TINT_ASSERT(AST, type);
+}
+
+Alias::Alias(Alias&&) = default;
+
+Alias::~Alias() = default;
+
+const Alias* Alias::Clone(CloneContext* ctx) const {
+  // Clone arguments outside of create() call to have deterministic ordering
+  auto src = ctx->Clone(source);
+  auto sym = ctx->Clone(name);
+  auto* ty = ctx->Clone(type);
+  return ctx->dst->create<Alias>(src, sym, ty);
+}
+
+}  // namespace ast
+}  // namespace tint

src/tint/ast/alias.h

@@ -0,0 +1,54 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#ifndef SRC_TINT_AST_ALIAS_H_
+#define SRC_TINT_AST_ALIAS_H_
+
+#include <string>
+
+#include "src/tint/ast/type_decl.h"
+
+namespace tint {
+namespace ast {
+
+/// A type alias type. Holds a name and pointer to another type.
+class Alias final : public Castable<Alias, TypeDecl> {
+ public:
+  /// Constructor
+  /// @param pid the identifier of the program that owns this node
+  /// @param src the source of this node
+  /// @param name the symbol for the alias
+  /// @param subtype the alias'd type
+  Alias(ProgramID pid,
+        const Source& src,
+        const Symbol& name,
+        const Type* subtype);
+  /// Move constructor
+  Alias(Alias&&);
+  /// Destructor
+  ~Alias() override;
+
+  /// Clones this type and all transitive types using the `CloneContext` `ctx`.
+  /// @param ctx the clone context
+  /// @return the newly cloned type
+  const Alias* Clone(CloneContext* ctx) const override;
+
+  /// the alias type
+  const Type* const type;
+};
+
+}  // namespace ast
+}  // namespace tint
+
+#endif  // SRC_TINT_AST_ALIAS_H_

src/tint/ast/alias_test.cc

@@ -0,0 +1,45 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/alias.h"
+#include "src/tint/ast/access.h"
+#include "src/tint/ast/array.h"
+#include "src/tint/ast/bool.h"
+#include "src/tint/ast/f32.h"
+#include "src/tint/ast/i32.h"
+#include "src/tint/ast/matrix.h"
+#include "src/tint/ast/pointer.h"
+#include "src/tint/ast/sampler.h"
+#include "src/tint/ast/struct.h"
+#include "src/tint/ast/test_helper.h"
+#include "src/tint/ast/texture.h"
+#include "src/tint/ast/u32.h"
+#include "src/tint/ast/vector.h"
+
+namespace tint {
+namespace ast {
+namespace {
+
+using AstAliasTest = TestHelper;
+
+TEST_F(AstAliasTest, Create) {
+  auto* u32 = create<U32>();
+  auto* a = Alias("a_type", u32);
+  EXPECT_EQ(a->name, Symbol(1, ID()));
+  EXPECT_EQ(a->type, u32);
+}
+
+}  // namespace
+}  // namespace ast
+}  // namespace tint

src/tint/ast/array.cc

@@ -0,0 +1,78 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/array.h"
+
+#include <cmath>
+
+#include "src/tint/program_builder.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::Array);
+
+namespace tint {
+namespace ast {
+
+namespace {
+// Returns the string representation of an array size expression.
+std::string SizeExprToString(const Expression* size,
+                             const SymbolTable& symbols) {
+  if (auto* ident = size->As<IdentifierExpression>()) {
+    return symbols.NameFor(ident->symbol);
+  }
+  if (auto* literal = size->As<IntLiteralExpression>()) {
+    return std::to_string(literal->ValueAsU32());
+  }
+  // This will never be exposed to the user as the Resolver will reject this
+  // expression for array size.
+  return "<invalid>";
+}
+}  // namespace
+
+Array::Array(ProgramID pid,
+             const Source& src,
+             const Type* subtype,
+             const Expression* cnt,
+             AttributeList attrs)
+    : Base(pid, src), type(subtype), count(cnt), attributes(attrs) {}
+
+Array::Array(Array&&) = default;
+
+Array::~Array() = default;
+
+std::string Array::FriendlyName(const SymbolTable& symbols) const {
+  std::ostringstream out;
+  for (auto* attr : attributes) {
+    if (auto* stride = attr->As<ast::StrideAttribute>()) {
+      out << "@stride(" << stride->stride << ") ";
+    }
+  }
+  out << "array<" << type->FriendlyName(symbols);
+  if (!IsRuntimeArray()) {
+    out << ", " << SizeExprToString(count, symbols);
+  }
+  out << ">";
+  return out.str();
+}
+
+const Array* Array::Clone(CloneContext* ctx) const {
+  // Clone arguments outside of create() call to have deterministic ordering
+  auto src = ctx->Clone(source);
+  auto* ty = ctx->Clone(type);
+  auto* cnt = ctx->Clone(count);
+  auto attrs = ctx->Clone(attributes);
+  return ctx->dst->create<Array>(src, ty, cnt, attrs);
+}
+
+}  // namespace ast
+}  // namespace tint

src/tint/ast/array.h

@@ -0,0 +1,75 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#ifndef SRC_TINT_AST_ARRAY_H_
+#define SRC_TINT_AST_ARRAY_H_
+
+#include <string>
+
+#include "src/tint/ast/attribute.h"
+#include "src/tint/ast/type.h"
+
+namespace tint {
+namespace ast {
+
+// Forward declarations.
+class Expression;
+
+/// An array type. If size is zero then it is a runtime array.
+class Array final : public Castable<Array, Type> {
+ public:
+  /// Constructor
+  /// @param pid the identifier of the program that owns this node
+  /// @param src the source of this node
+  /// @param subtype the type of the array elements
+  /// @param count the number of elements in the array. nullptr represents a
+  /// runtime-sized array.
+  /// @param attributes the array attributes
+  Array(ProgramID pid,
+        const Source& src,
+        const Type* subtype,
+        const Expression* count,
+        AttributeList attributes);
+  /// Move constructor
+  Array(Array&&);
+  ~Array() override;
+
+  /// @returns true if this is a runtime array.
+  /// i.e. the size is determined at runtime
+  bool IsRuntimeArray() const { return count == nullptr; }
+
+  /// @param symbols the program's symbol table
+  /// @returns the name for this type that closely resembles how it would be
+  /// declared in WGSL.
+  std::string FriendlyName(const SymbolTable& symbols) const override;
+
+  /// Clones this type and all transitive types using the `CloneContext` `ctx`.
+  /// @param ctx the clone context
+  /// @return the newly cloned type
+  const Array* Clone(CloneContext* ctx) const override;
+
+  /// the array element type
+  const Type* const type;
+
+  /// the array size in elements, or nullptr for a runtime array
+  const Expression* const count;
+
+  /// the array attributes
+  const AttributeList attributes;
+};
+
+}  // namespace ast
+}  // namespace tint
+
+#endif  // SRC_TINT_AST_ARRAY_H_

src/tint/ast/array_test.cc

@@ -0,0 +1,71 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/array.h"
+
+#include "src/tint/ast/test_helper.h"
+
+namespace tint {
+namespace ast {
+namespace {
+
+using AstArrayTest = TestHelper;
+
+TEST_F(AstArrayTest, CreateSizedArray) {
+  auto* u32 = create<U32>();
+  auto* count = Expr(3);
+  auto* arr = create<Array>(u32, count, AttributeList{});
+  EXPECT_EQ(arr->type, u32);
+  EXPECT_EQ(arr->count, count);
+  EXPECT_TRUE(arr->Is<Array>());
+  EXPECT_FALSE(arr->IsRuntimeArray());
+}
+
+TEST_F(AstArrayTest, CreateRuntimeArray) {
+  auto* u32 = create<U32>();
+  auto* arr = create<Array>(u32, nullptr, AttributeList{});
+  EXPECT_EQ(arr->type, u32);
+  EXPECT_EQ(arr->count, nullptr);
+  EXPECT_TRUE(arr->Is<Array>());
+  EXPECT_TRUE(arr->IsRuntimeArray());
+}
+
+TEST_F(AstArrayTest, FriendlyName_RuntimeSized) {
+  auto* i32 = create<I32>();
+  auto* arr = create<Array>(i32, nullptr, AttributeList{});
+  EXPECT_EQ(arr->FriendlyName(Symbols()), "array<i32>");
+}
+
+TEST_F(AstArrayTest, FriendlyName_LiteralSized) {
+  auto* i32 = create<I32>();
+  auto* arr = create<Array>(i32, Expr(5), AttributeList{});
+  EXPECT_EQ(arr->FriendlyName(Symbols()), "array<i32, 5>");
+}
+
+TEST_F(AstArrayTest, FriendlyName_ConstantSized) {
+  auto* i32 = create<I32>();
+  auto* arr = create<Array>(i32, Expr("size"), AttributeList{});
+  EXPECT_EQ(arr->FriendlyName(Symbols()), "array<i32, size>");
+}
+
+TEST_F(AstArrayTest, FriendlyName_WithStride) {
+  auto* i32 = create<I32>();
+  auto* arr =
+      create<Array>(i32, Expr(5), AttributeList{create<StrideAttribute>(32)});
+  EXPECT_EQ(arr->FriendlyName(Symbols()), "@stride(32) array<i32, 5>");
+}
+
+}  // namespace
+}  // namespace ast
+}  // namespace tint

src/tint/ast/assignment_statement.cc

@@ -0,0 +1,48 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/assignment_statement.h"
+
+#include "src/tint/program_builder.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::AssignmentStatement);
+
+namespace tint {
+namespace ast {
+
+AssignmentStatement::AssignmentStatement(ProgramID pid,
+                                         const Source& src,
+                                         const Expression* l,
+                                         const Expression* r)
+    : Base(pid, src), lhs(l), rhs(r) {
+  TINT_ASSERT(AST, lhs);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, lhs, program_id);
+  TINT_ASSERT(AST, rhs);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, rhs, program_id);
+}
+
+AssignmentStatement::AssignmentStatement(AssignmentStatement&&) = default;
+
+AssignmentStatement::~AssignmentStatement() = default;
+
+const AssignmentStatement* AssignmentStatement::Clone(CloneContext* ctx) const {
+  // Clone arguments outside of create() call to have deterministic ordering
+  auto src = ctx->Clone(source);
+  auto* l = ctx->Clone(lhs);
+  auto* r = ctx->Clone(rhs);
+  return ctx->dst->create<AssignmentStatement>(src, l, r);
+}
+
+}  // namespace ast
+}  // namespace tint

src/tint/ast/assignment_statement.h

@@ -0,0 +1,57 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#ifndef SRC_TINT_AST_ASSIGNMENT_STATEMENT_H_
+#define SRC_TINT_AST_ASSIGNMENT_STATEMENT_H_
+
+#include "src/tint/ast/expression.h"
+#include "src/tint/ast/statement.h"
+
+namespace tint {
+namespace ast {
+
+/// An assignment statement
+class AssignmentStatement final
+    : public Castable<AssignmentStatement, Statement> {
+ public:
+  /// Constructor
+  /// @param program_id the identifier of the program that owns this node
+  /// @param source the assignment statement source
+  /// @param lhs the left side of the expression
+  /// @param rhs the right side of the expression
+  AssignmentStatement(ProgramID program_id,
+                      const Source& source,
+                      const Expression* lhs,
+                      const Expression* rhs);
+  /// Move constructor
+  AssignmentStatement(AssignmentStatement&&);
+  ~AssignmentStatement() override;
+
+  /// Clones this node and all transitive child nodes using the `CloneContext`
+  /// `ctx`.
+  /// @param ctx the clone context
+  /// @return the newly cloned node
+  const AssignmentStatement* Clone(CloneContext* ctx) const override;
+
+  /// left side expression
+  const Expression* const lhs;
+
+  /// right side expression
+  const Expression* const rhs;
+};
+
+}  // namespace ast
+}  // namespace tint
+
+#endif  // SRC_TINT_AST_ASSIGNMENT_STATEMENT_H_

src/tint/ast/assignment_statement_test.cc

@@ -0,0 +1,94 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/assignment_statement.h"
+
+#include "gtest/gtest-spi.h"
+#include "src/tint/ast/test_helper.h"
+
+namespace tint {
+namespace ast {
+namespace {
+
+using AssignmentStatementTest = TestHelper;
+
+TEST_F(AssignmentStatementTest, Creation) {
+  auto* lhs = Expr("lhs");
+  auto* rhs = Expr("rhs");
+
+  auto* stmt = create<AssignmentStatement>(lhs, rhs);
+  EXPECT_EQ(stmt->lhs, lhs);
+  EXPECT_EQ(stmt->rhs, rhs);
+}
+
+TEST_F(AssignmentStatementTest, CreationWithSource) {
+  auto* lhs = Expr("lhs");
+  auto* rhs = Expr("rhs");
+
+  auto* stmt =
+      create<AssignmentStatement>(Source{Source::Location{20, 2}}, lhs, rhs);
+  auto src = stmt->source;
+  EXPECT_EQ(src.range.begin.line, 20u);
+  EXPECT_EQ(src.range.begin.column, 2u);
+}
+
+TEST_F(AssignmentStatementTest, IsAssign) {
+  auto* lhs = Expr("lhs");
+  auto* rhs = Expr("rhs");
+
+  auto* stmt = create<AssignmentStatement>(lhs, rhs);
+  EXPECT_TRUE(stmt->Is<AssignmentStatement>());
+}
+
+TEST_F(AssignmentStatementTest, Assert_Null_LHS) {
+  EXPECT_FATAL_FAILURE(
+      {
+        ProgramBuilder b;
+        b.create<AssignmentStatement>(nullptr, b.Expr(1));
+      },
+      "internal compiler error");
+}
+
+TEST_F(AssignmentStatementTest, Assert_Null_RHS) {
+  EXPECT_FATAL_FAILURE(
+      {
+        ProgramBuilder b;
+        b.create<AssignmentStatement>(b.Expr(1), nullptr);
+      },
+      "internal compiler error");
+}
+
+TEST_F(AssignmentStatementTest, Assert_DifferentProgramID_LHS) {
+  EXPECT_FATAL_FAILURE(
+      {
+        ProgramBuilder b1;
+        ProgramBuilder b2;
+        b1.create<AssignmentStatement>(b2.Expr("lhs"), b1.Expr("rhs"));
+      },
+      "internal compiler error");
+}
+
+TEST_F(AssignmentStatementTest, Assert_DifferentProgramID_RHS) {
+  EXPECT_FATAL_FAILURE(
+      {
+        ProgramBuilder b1;
+        ProgramBuilder b2;
+        b1.create<AssignmentStatement>(b1.Expr("lhs"), b2.Expr("rhs"));
+      },
+      "internal compiler error");
+}
+
+}  // namespace
+}  // namespace ast
+}  // namespace tint

src/tint/ast/ast_type.cc

@@ -0,0 +1,41 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/type.h"
+
+#include "src/tint/ast/alias.h"
+#include "src/tint/ast/bool.h"
+#include "src/tint/ast/f32.h"
+#include "src/tint/ast/i32.h"
+#include "src/tint/ast/matrix.h"
+#include "src/tint/ast/pointer.h"
+#include "src/tint/ast/sampler.h"
+#include "src/tint/ast/texture.h"
+#include "src/tint/ast/u32.h"
+#include "src/tint/ast/vector.h"
+#include "src/tint/symbol_table.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::Type);
+
+namespace tint {
+namespace ast {
+
+Type::Type(ProgramID pid, const Source& src) : Base(pid, src) {}
+
+Type::Type(Type&&) = default;
+
+Type::~Type() = default;
+
+}  // namespace ast
+}  // namespace tint

src/tint/ast/atomic.cc

@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/atomic.h"
+
+#include "src/tint/program_builder.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::Atomic);
+
+namespace tint {
+namespace ast {
+
+Atomic::Atomic(ProgramID pid, const Source& src, const Type* const subtype)
+    : Base(pid, src), type(subtype) {}
+
+std::string Atomic::FriendlyName(const SymbolTable& symbols) const {
+  std::ostringstream out;
+  out << "atomic<" << type->FriendlyName(symbols) << ">";
+  return out.str();
+}
+
+Atomic::Atomic(Atomic&&) = default;
+
+Atomic::~Atomic() = default;
+
+const Atomic* Atomic::Clone(CloneContext* ctx) const {
+  // Clone arguments outside of create() call to have deterministic ordering
+  auto src = ctx->Clone(source);
+  auto* ty = ctx->Clone(type);
+  return ctx->dst->create<Atomic>(src, ty);
+}
+
+}  // namespace ast
+}  // namespace tint

src/tint/ast/atomic.h

@@ -0,0 +1,54 @@
+// Copyright 2021 The Tint Authors.
+//
+// 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.
+
+#ifndef SRC_TINT_AST_ATOMIC_H_
+#define SRC_TINT_AST_ATOMIC_H_
+
+#include <string>
+
+#include "src/tint/ast/type.h"
+
+namespace tint {
+namespace ast {
+
+/// An atomic type.
+class Atomic final : public Castable<Atomic, Type> {
+ public:
+  /// Constructor
+  /// @param pid the identifier of the program that owns this node
+  /// @param src the source of this node
+  /// @param subtype the pointee type
+  Atomic(ProgramID pid, const Source& src, const Type* const subtype);
+  /// Move constructor
+  Atomic(Atomic&&);
+  ~Atomic() override;
+
+  /// @param symbols the program's symbol table
+  /// @returns the name for this type that closely resembles how it would be
+  /// declared in WGSL.
+  std::string FriendlyName(const SymbolTable& symbols) const override;
+
+  /// Clones this type and all transitive types using the `CloneContext` `ctx`.
+  /// @param ctx the clone context
+  /// @return the newly cloned type
+  const Atomic* Clone(CloneContext* ctx) const override;
+
+  /// the pointee type
+  const Type* const type;
+};
+
+}  // namespace ast
+}  // namespace tint
+
+#endif  // SRC_TINT_AST_ATOMIC_H_

src/tint/ast/atomic_test.cc

@@ -0,0 +1,40 @@
+// Copyright 2021 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/atomic.h"
+
+#include "src/tint/ast/i32.h"
+#include "src/tint/ast/test_helper.h"
+
+namespace tint {
+namespace ast {
+namespace {
+
+using AstAtomicTest = TestHelper;
+
+TEST_F(AstAtomicTest, Creation) {
+  auto* i32 = create<I32>();
+  auto* p = create<Atomic>(i32);
+  EXPECT_EQ(p->type, i32);
+}
+
+TEST_F(AstAtomicTest, FriendlyName) {
+  auto* i32 = create<I32>();
+  auto* p = create<Atomic>(i32);
+  EXPECT_EQ(p->FriendlyName(Symbols()), "atomic<i32>");
+}
+
+}  // namespace
+}  // namespace ast
+}  // namespace tint

src/tint/ast/attribute.cc

@@ -0,0 +1,25 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/attribute.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::Attribute);
+
+namespace tint {
+namespace ast {
+
+Attribute::~Attribute() = default;
+
+}  // namespace ast
+}  // namespace tint

src/tint/ast/attribute.h

@@ -0,0 +1,71 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#ifndef SRC_TINT_AST_ATTRIBUTE_H_
+#define SRC_TINT_AST_ATTRIBUTE_H_
+
+#include <string>
+#include <vector>
+
+#include "src/tint/ast/node.h"
+
+namespace tint {
+namespace ast {
+
+/// The base class for all attributes
+class Attribute : public Castable<Attribute, Node> {
+ public:
+  ~Attribute() override;
+
+  /// @returns the WGSL name for the attribute
+  virtual std::string Name() const = 0;
+
+ protected:
+  /// Constructor
+  /// @param pid the identifier of the program that owns this node
+  /// @param src the source of this node
+  Attribute(ProgramID pid, const Source& src) : Base(pid, src) {}
+};
+
+/// A list of attributes
+using AttributeList = std::vector<const Attribute*>;
+
+/// @param attributes the list of attributes to search
+/// @returns true if `attributes` includes a attribute of type `T`
+template <typename T>
+bool HasAttribute(const AttributeList& attributes) {
+  for (auto* attr : attributes) {
+    if (attr->Is<T>()) {
+      return true;
+    }
+  }
+  return false;
+}
+
+/// @param attributes the list of attributes to search
+/// @returns a pointer to `T` from `attributes` if found, otherwise nullptr.
+template <typename T>
+const T* GetAttribute(const AttributeList& attributes) {
+  for (auto* attr : attributes) {
+    if (attr->Is<T>()) {
+      return attr->As<T>();
+    }
+  }
+  return nullptr;
+}
+
+}  // namespace ast
+}  // namespace tint
+
+#endif  // SRC_TINT_AST_ATTRIBUTE_H_

src/tint/ast/binary_expression.cc

@@ -0,0 +1,50 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/binary_expression.h"
+
+#include "src/tint/program_builder.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::BinaryExpression);
+
+namespace tint {
+namespace ast {
+
+BinaryExpression::BinaryExpression(ProgramID pid,
+                                   const Source& src,
+                                   BinaryOp o,
+                                   const Expression* l,
+                                   const Expression* r)
+    : Base(pid, src), op(o), lhs(l), rhs(r) {
+  TINT_ASSERT(AST, lhs);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, lhs, program_id);
+  TINT_ASSERT(AST, rhs);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, rhs, program_id);
+  TINT_ASSERT(AST, op != BinaryOp::kNone);
+}
+
+BinaryExpression::BinaryExpression(BinaryExpression&&) = default;
+
+BinaryExpression::~BinaryExpression() = default;
+
+const BinaryExpression* BinaryExpression::Clone(CloneContext* ctx) const {
+  // Clone arguments outside of create() call to have deterministic ordering
+  auto src = ctx->Clone(source);
+  auto* l = ctx->Clone(lhs);
+  auto* r = ctx->Clone(rhs);
+  return ctx->dst->create<BinaryExpression>(src, op, l, r);
+}
+
+}  // namespace ast
+}  // namespace tint

src/tint/ast/binary_expression.h

@@ -0,0 +1,264 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#ifndef SRC_TINT_AST_BINARY_EXPRESSION_H_
+#define SRC_TINT_AST_BINARY_EXPRESSION_H_
+
+#include "src/tint/ast/expression.h"
+
+namespace tint {
+namespace ast {
+
+/// The operator type
+enum class BinaryOp {
+  kNone = 0,
+  kAnd,  // &
+  kOr,   // |
+  kXor,
+  kLogicalAnd,  // &&
+  kLogicalOr,   // ||
+  kEqual,
+  kNotEqual,
+  kLessThan,
+  kGreaterThan,
+  kLessThanEqual,
+  kGreaterThanEqual,
+  kShiftLeft,
+  kShiftRight,
+  kAdd,
+  kSubtract,
+  kMultiply,
+  kDivide,
+  kModulo,
+};
+
+/// An binary expression
+class BinaryExpression final : public Castable<BinaryExpression, Expression> {
+ public:
+  /// Constructor
+  /// @param program_id the identifier of the program that owns this node
+  /// @param source the binary expression source
+  /// @param op the operation type
+  /// @param lhs the left side of the expression
+  /// @param rhs the right side of the expression
+  BinaryExpression(ProgramID program_id,
+                   const Source& source,
+                   BinaryOp op,
+                   const Expression* lhs,
+                   const Expression* rhs);
+  /// Move constructor
+  BinaryExpression(BinaryExpression&&);
+  ~BinaryExpression() override;
+
+  /// @returns true if the op is and
+  bool IsAnd() const { return op == BinaryOp::kAnd; }
+  /// @returns true if the op is or
+  bool IsOr() const { return op == BinaryOp::kOr; }
+  /// @returns true if the op is xor
+  bool IsXor() const { return op == BinaryOp::kXor; }
+  /// @returns true if the op is logical and
+  bool IsLogicalAnd() const { return op == BinaryOp::kLogicalAnd; }
+  /// @returns true if the op is logical or
+  bool IsLogicalOr() const { return op == BinaryOp::kLogicalOr; }
+  /// @returns true if the op is equal
+  bool IsEqual() const { return op == BinaryOp::kEqual; }
+  /// @returns true if the op is not equal
+  bool IsNotEqual() const { return op == BinaryOp::kNotEqual; }
+  /// @returns true if the op is less than
+  bool IsLessThan() const { return op == BinaryOp::kLessThan; }
+  /// @returns true if the op is greater than
+  bool IsGreaterThan() const { return op == BinaryOp::kGreaterThan; }
+  /// @returns true if the op is less than equal
+  bool IsLessThanEqual() const { return op == BinaryOp::kLessThanEqual; }
+  /// @returns true if the op is greater than equal
+  bool IsGreaterThanEqual() const { return op == BinaryOp::kGreaterThanEqual; }
+  /// @returns true if the op is shift left
+  bool IsShiftLeft() const { return op == BinaryOp::kShiftLeft; }
+  /// @returns true if the op is shift right
+  bool IsShiftRight() const { return op == BinaryOp::kShiftRight; }
+  /// @returns true if the op is add
+  bool IsAdd() const { return op == BinaryOp::kAdd; }
+  /// @returns true if the op is subtract
+  bool IsSubtract() const { return op == BinaryOp::kSubtract; }
+  /// @returns true if the op is multiply
+  bool IsMultiply() const { return op == BinaryOp::kMultiply; }
+  /// @returns true if the op is divide
+  bool IsDivide() const { return op == BinaryOp::kDivide; }
+  /// @returns true if the op is modulo
+  bool IsModulo() const { return op == BinaryOp::kModulo; }
+  /// @returns true if the op is an arithmetic operation
+  bool IsArithmetic() const;
+  /// @returns true if the op is a comparison operation
+  bool IsComparison() const;
+  /// @returns true if the op is a bitwise operation
+  bool IsBitwise() const;
+  /// @returns true if the op is a bit shift operation
+  bool IsBitshift() const;
+  /// @returns true if the op is a logical expression
+  bool IsLogical() const;
+
+  /// Clones this node and all transitive child nodes using the `CloneContext`
+  /// `ctx`.
+  /// @param ctx the clone context
+  /// @return the newly cloned node
+  const BinaryExpression* Clone(CloneContext* ctx) const override;
+
+  /// the binary op type
+  const BinaryOp op;
+  /// the left side expression
+  const Expression* const lhs;
+  /// the right side expression
+  const Expression* const rhs;
+};
+
+/// @param op the operator
+/// @returns true if the op is an arithmetic operation
+inline bool IsArithmetic(BinaryOp op) {
+  switch (op) {
+    case ast::BinaryOp::kAdd:
+    case ast::BinaryOp::kSubtract:
+    case ast::BinaryOp::kMultiply:
+    case ast::BinaryOp::kDivide:
+    case ast::BinaryOp::kModulo:
+      return true;
+    default:
+      return false;
+  }
+}
+
+/// @param op the operator
+/// @returns true if the op is a comparison operation
+inline bool IsComparison(BinaryOp op) {
+  switch (op) {
+    case ast::BinaryOp::kEqual:
+    case ast::BinaryOp::kNotEqual:
+    case ast::BinaryOp::kLessThan:
+    case ast::BinaryOp::kLessThanEqual:
+    case ast::BinaryOp::kGreaterThan:
+    case ast::BinaryOp::kGreaterThanEqual:
+      return true;
+    default:
+      return false;
+  }
+}
+
+/// @param op the operator
+/// @returns true if the op is a bitwise operation
+inline bool IsBitwise(BinaryOp op) {
+  switch (op) {
+    case ast::BinaryOp::kAnd:
+    case ast::BinaryOp::kOr:
+    case ast::BinaryOp::kXor:
+      return true;
+    default:
+      return false;
+  }
+}
+
+/// @param op the operator
+/// @returns true if the op is a bit shift operation
+inline bool IsBitshift(BinaryOp op) {
+  switch (op) {
+    case ast::BinaryOp::kShiftLeft:
+    case ast::BinaryOp::kShiftRight:
+      return true;
+    default:
+      return false;
+  }
+}
+
+inline bool BinaryExpression::IsLogical() const {
+  switch (op) {
+    case ast::BinaryOp::kLogicalAnd:
+    case ast::BinaryOp::kLogicalOr:
+      return true;
+    default:
+      return false;
+  }
+}
+
+inline bool BinaryExpression::IsArithmetic() const {
+  return ast::IsArithmetic(op);
+}
+
+inline bool BinaryExpression::IsComparison() const {
+  return ast::IsComparison(op);
+}
+
+inline bool BinaryExpression::IsBitwise() const {
+  return ast::IsBitwise(op);
+}
+
+inline bool BinaryExpression::IsBitshift() const {
+  return ast::IsBitshift(op);
+}
+
+/// @returns the human readable name of the given BinaryOp
+/// @param op the BinaryOp
+constexpr const char* FriendlyName(BinaryOp op) {
+  switch (op) {
+    case BinaryOp::kNone:
+      return "none";
+    case BinaryOp::kAnd:
+      return "and";
+    case BinaryOp::kOr:
+      return "or";
+    case BinaryOp::kXor:
+      return "xor";
+    case BinaryOp::kLogicalAnd:
+      return "logical_and";
+    case BinaryOp::kLogicalOr:
+      return "logical_or";
+    case BinaryOp::kEqual:
+      return "equal";
+    case BinaryOp::kNotEqual:
+      return "not_equal";
+    case BinaryOp::kLessThan:
+      return "less_than";
+    case BinaryOp::kGreaterThan:
+      return "greater_than";
+    case BinaryOp::kLessThanEqual:
+      return "less_than_equal";
+    case BinaryOp::kGreaterThanEqual:
+      return "greater_than_equal";
+    case BinaryOp::kShiftLeft:
+      return "shift_left";
+    case BinaryOp::kShiftRight:
+      return "shift_right";
+    case BinaryOp::kAdd:
+      return "add";
+    case BinaryOp::kSubtract:
+      return "subtract";
+    case BinaryOp::kMultiply:
+      return "multiply";
+    case BinaryOp::kDivide:
+      return "divide";
+    case BinaryOp::kModulo:
+      return "modulo";
+  }
+  return "INVALID";
+}
+
+/// @param out the std::ostream to write to
+/// @param op the BinaryOp
+/// @return the std::ostream so calls can be chained
+inline std::ostream& operator<<(std::ostream& out, BinaryOp op) {
+  out << FriendlyName(op);
+  return out;
+}
+
+}  // namespace ast
+}  // namespace tint
+
+#endif  // SRC_TINT_AST_BINARY_EXPRESSION_H_

src/tint/ast/binary_expression_test.cc

@@ -0,0 +1,95 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "gtest/gtest-spi.h"
+#include "src/tint/ast/test_helper.h"
+
+namespace tint {
+namespace ast {
+namespace {
+
+using BinaryExpressionTest = TestHelper;
+
+TEST_F(BinaryExpressionTest, Creation) {
+  auto* lhs = Expr("lhs");
+  auto* rhs = Expr("rhs");
+
+  auto* r = create<BinaryExpression>(BinaryOp::kEqual, lhs, rhs);
+  EXPECT_EQ(r->lhs, lhs);
+  EXPECT_EQ(r->rhs, rhs);
+  EXPECT_EQ(r->op, BinaryOp::kEqual);
+}
+
+TEST_F(BinaryExpressionTest, Creation_WithSource) {
+  auto* lhs = Expr("lhs");
+  auto* rhs = Expr("rhs");
+
+  auto* r = create<BinaryExpression>(Source{Source::Location{20, 2}},
+                                     BinaryOp::kEqual, lhs, rhs);
+  auto src = r->source;
+  EXPECT_EQ(src.range.begin.line, 20u);
+  EXPECT_EQ(src.range.begin.column, 2u);
+}
+
+TEST_F(BinaryExpressionTest, IsBinary) {
+  auto* lhs = Expr("lhs");
+  auto* rhs = Expr("rhs");
+
+  auto* r = create<BinaryExpression>(BinaryOp::kEqual, lhs, rhs);
+  EXPECT_TRUE(r->Is<BinaryExpression>());
+}
+
+TEST_F(BinaryExpressionTest, Assert_Null_LHS) {
+  EXPECT_FATAL_FAILURE(
+      {
+        ProgramBuilder b;
+        b.create<BinaryExpression>(BinaryOp::kEqual, nullptr, b.Expr("rhs"));
+      },
+      "internal compiler error");
+}
+
+TEST_F(BinaryExpressionTest, Assert_Null_RHS) {
+  EXPECT_FATAL_FAILURE(
+      {
+        ProgramBuilder b;
+        b.create<BinaryExpression>(BinaryOp::kEqual, b.Expr("lhs"), nullptr);
+      },
+      "internal compiler error");
+}
+
+TEST_F(BinaryExpressionTest, Assert_DifferentProgramID_LHS) {
+  EXPECT_FATAL_FAILURE(
+      {
+        ProgramBuilder b1;
+        ProgramBuilder b2;
+        b1.create<BinaryExpression>(BinaryOp::kEqual, b2.Expr("lhs"),
+                                    b1.Expr("rhs"));
+      },
+      "internal compiler error");
+}
+
+TEST_F(BinaryExpressionTest, Assert_DifferentProgramID_RHS) {
+  EXPECT_FATAL_FAILURE(
+      {
+        ProgramBuilder b1;
+        ProgramBuilder b2;
+        b1.create<BinaryExpression>(BinaryOp::kEqual, b1.Expr("lhs"),
+                                    b2.Expr("rhs"));
+      },
+      "internal compiler error");
+}
+
+}  // namespace
+}  // namespace ast
+}  // namespace tint

src/tint/ast/binding_attribute.cc

@@ -0,0 +1,44 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/binding_attribute.h"
+
+#include <string>
+
+#include "src/tint/program_builder.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::BindingAttribute);
+
+namespace tint {
+namespace ast {
+
+BindingAttribute::BindingAttribute(ProgramID pid,
+                                   const Source& src,
+                                   uint32_t val)
+    : Base(pid, src), value(val) {}
+
+BindingAttribute::~BindingAttribute() = default;
+
+std::string BindingAttribute::Name() const {
+  return "binding";
+}
+
+const BindingAttribute* BindingAttribute::Clone(CloneContext* ctx) const {
+  // Clone arguments outside of create() call to have deterministic ordering
+  auto src = ctx->Clone(source);
+  return ctx->dst->create<BindingAttribute>(src, value);
+}
+
+}  // namespace ast
+}  // namespace tint

src/tint/ast/binding_attribute.h

@@ -0,0 +1,51 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#ifndef SRC_TINT_AST_BINDING_ATTRIBUTE_H_
+#define SRC_TINT_AST_BINDING_ATTRIBUTE_H_
+
+#include <string>
+
+#include "src/tint/ast/attribute.h"
+
+namespace tint {
+namespace ast {
+
+/// A binding attribute
+class BindingAttribute final : public Castable<BindingAttribute, Attribute> {
+ public:
+  /// Constructor
+  /// @param pid the identifier of the program that owns this node
+  /// @param src the source of this node
+  /// @param value the binding value
+  BindingAttribute(ProgramID pid, const Source& src, uint32_t value);
+  ~BindingAttribute() override;
+
+  /// @returns the WGSL name for the attribute
+  std::string Name() const override;
+
+  /// Clones this node and all transitive child nodes using the `CloneContext`
+  /// `ctx`.
+  /// @param ctx the clone context
+  /// @return the newly cloned node
+  const BindingAttribute* Clone(CloneContext* ctx) const override;
+
+  /// the binding value
+  const uint32_t value;
+};
+
+}  // namespace ast
+}  // namespace tint
+
+#endif  // SRC_TINT_AST_BINDING_ATTRIBUTE_H_

src/tint/ast/binding_attribute_test.cc

@@ -0,0 +1,30 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/test_helper.h"
+
+namespace tint {
+namespace ast {
+namespace {
+
+using BindingAttributeTest = TestHelper;
+
+TEST_F(BindingAttributeTest, Creation) {
+  auto* d = create<BindingAttribute>(2);
+  EXPECT_EQ(2u, d->value);
+}
+
+}  // namespace
+}  // namespace ast
+}  // namespace tint

src/tint/ast/bitcast_expression.cc

@@ -0,0 +1,46 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/bitcast_expression.h"
+
+#include "src/tint/program_builder.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::BitcastExpression);
+
+namespace tint {
+namespace ast {
+
+BitcastExpression::BitcastExpression(ProgramID pid,
+                                     const Source& src,
+                                     const Type* t,
+                                     const Expression* e)
+    : Base(pid, src), type(t), expr(e) {
+  TINT_ASSERT(AST, type);
+  TINT_ASSERT(AST, expr);
+  TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, expr, program_id);
+}
+
+BitcastExpression::BitcastExpression(BitcastExpression&&) = default;
+BitcastExpression::~BitcastExpression() = default;
+
+const BitcastExpression* BitcastExpression::Clone(CloneContext* ctx) const {
+  // Clone arguments outside of create() call to have deterministic ordering
+  auto src = ctx->Clone(source);
+  auto* t = ctx->Clone(type);
+  auto* e = ctx->Clone(expr);
+  return ctx->dst->create<BitcastExpression>(src, t, e);
+}
+
+}  // namespace ast
+}  // namespace tint

src/tint/ast/bitcast_expression.h

@@ -0,0 +1,57 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#ifndef SRC_TINT_AST_BITCAST_EXPRESSION_H_
+#define SRC_TINT_AST_BITCAST_EXPRESSION_H_
+
+#include "src/tint/ast/expression.h"
+
+namespace tint {
+namespace ast {
+
+// Forward declaration
+class Type;
+
+/// A bitcast expression
+class BitcastExpression final : public Castable<BitcastExpression, Expression> {
+ public:
+  /// Constructor
+  /// @param program_id the identifier of the program that owns this node
+  /// @param source the bitcast expression source
+  /// @param type the type
+  /// @param expr the expr
+  BitcastExpression(ProgramID program_id,
+                    const Source& source,
+                    const Type* type,
+                    const Expression* expr);
+  /// Move constructor
+  BitcastExpression(BitcastExpression&&);
+  ~BitcastExpression() override;
+
+  /// Clones this node and all transitive child nodes using the `CloneContext`
+  /// `ctx`.
+  /// @param ctx the clone context
+  /// @return the newly cloned node
+  const BitcastExpression* Clone(CloneContext* ctx) const override;
+
+  /// the target cast type
+  const Type* const type;
+  /// the expression
+  const Expression* const expr;
+};
+
+}  // namespace ast
+}  // namespace tint
+
+#endif  // SRC_TINT_AST_BITCAST_EXPRESSION_H_

src/tint/ast/bitcast_expression_test.cc

@@ -0,0 +1,81 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/bitcast_expression.h"
+
+#include "gtest/gtest-spi.h"
+#include "src/tint/ast/test_helper.h"
+
+namespace tint {
+namespace ast {
+namespace {
+
+using BitcastExpressionTest = TestHelper;
+
+TEST_F(BitcastExpressionTest, Create) {
+  auto* expr = Expr("expr");
+
+  auto* exp = create<BitcastExpression>(ty.f32(), expr);
+  EXPECT_TRUE(exp->type->Is<ast::F32>());
+  ASSERT_EQ(exp->expr, expr);
+}
+
+TEST_F(BitcastExpressionTest, CreateWithSource) {
+  auto* expr = Expr("expr");
+
+  auto* exp = create<BitcastExpression>(Source{Source::Location{20, 2}},
+                                        ty.f32(), expr);
+  auto src = exp->source;
+  EXPECT_EQ(src.range.begin.line, 20u);
+  EXPECT_EQ(src.range.begin.column, 2u);
+}
+
+TEST_F(BitcastExpressionTest, IsBitcast) {
+  auto* expr = Expr("expr");
+
+  auto* exp = create<BitcastExpression>(ty.f32(), expr);
+  EXPECT_TRUE(exp->Is<BitcastExpression>());
+}
+
+TEST_F(BitcastExpressionTest, Assert_Null_Type) {
+  EXPECT_FATAL_FAILURE(
+      {
+        ProgramBuilder b;
+        b.create<BitcastExpression>(nullptr, b.Expr("idx"));
+      },
+      "internal compiler error");
+}
+
+TEST_F(BitcastExpressionTest, Assert_Null_Expr) {
+  EXPECT_FATAL_FAILURE(
+      {
+        ProgramBuilder b;
+        b.create<BitcastExpression>(b.ty.f32(), nullptr);
+      },
+      "internal compiler error");
+}
+
+TEST_F(BitcastExpressionTest, Assert_DifferentProgramID_Expr) {
+  EXPECT_FATAL_FAILURE(
+      {
+        ProgramBuilder b1;
+        ProgramBuilder b2;
+        b1.create<BitcastExpression>(b1.ty.f32(), b2.Expr("idx"));
+      },
+      "internal compiler error");
+}
+
+}  // namespace
+}  // namespace ast
+}  // namespace tint

src/tint/ast/block_statement.cc

@@ -0,0 +1,46 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/block_statement.h"
+
+#include "src/tint/program_builder.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::BlockStatement);
+
+namespace tint {
+namespace ast {
+
+BlockStatement::BlockStatement(ProgramID pid,
+                               const Source& src,
+                               const StatementList& stmts)
+    : Base(pid, src), statements(std::move(stmts)) {
+  for (auto* stmt : statements) {
+    TINT_ASSERT(AST, stmt);
+    TINT_ASSERT_PROGRAM_IDS_EQUAL_IF_VALID(AST, stmt, program_id);
+  }
+}
+
+BlockStatement::BlockStatement(BlockStatement&&) = default;
+
+BlockStatement::~BlockStatement() = default;
+
+const BlockStatement* BlockStatement::Clone(CloneContext* ctx) const {
+  // Clone arguments outside of create() call to have deterministic ordering
+  auto src = ctx->Clone(source);
+  auto stmts = ctx->Clone(statements);
+  return ctx->dst->create<BlockStatement>(src, stmts);
+}
+
+}  // namespace ast
+}  // namespace tint

src/tint/ast/block_statement.h

@@ -0,0 +1,60 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#ifndef SRC_TINT_AST_BLOCK_STATEMENT_H_
+#define SRC_TINT_AST_BLOCK_STATEMENT_H_
+
+#include <utility>
+
+#include "src/tint/ast/statement.h"
+
+namespace tint {
+namespace ast {
+
+/// A block statement
+class BlockStatement final : public Castable<BlockStatement, Statement> {
+ public:
+  /// Constructor
+  /// @param program_id the identifier of the program that owns this node
+  /// @param source the block statement source
+  /// @param statements the statements
+  BlockStatement(ProgramID program_id,
+                 const Source& source,
+                 const StatementList& statements);
+  /// Move constructor
+  BlockStatement(BlockStatement&&);
+  ~BlockStatement() override;
+
+  /// @returns true if the block has no statements
+  bool Empty() const { return statements.empty(); }
+
+  /// @returns the last statement in the block or nullptr if block empty
+  const Statement* Last() const {
+    return statements.empty() ? nullptr : statements.back();
+  }
+
+  /// Clones this node and all transitive child nodes using the `CloneContext`
+  /// `ctx`.
+  /// @param ctx the clone context
+  /// @return the newly cloned node
+  const BlockStatement* Clone(CloneContext* ctx) const override;
+
+  /// the statement list
+  const StatementList statements;
+};
+
+}  // namespace ast
+}  // namespace tint
+
+#endif  // SRC_TINT_AST_BLOCK_STATEMENT_H_

src/tint/ast/block_statement_test.cc

@@ -0,0 +1,71 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "gtest/gtest-spi.h"
+#include "src/tint/ast/discard_statement.h"
+#include "src/tint/ast/if_statement.h"
+#include "src/tint/ast/test_helper.h"
+
+namespace tint {
+namespace ast {
+namespace {
+
+using BlockStatementTest = TestHelper;
+
+TEST_F(BlockStatementTest, Creation) {
+  auto* d = create<DiscardStatement>();
+  auto* ptr = d;
+
+  auto* b = create<BlockStatement>(StatementList{d});
+
+  ASSERT_EQ(b->statements.size(), 1u);
+  EXPECT_EQ(b->statements[0], ptr);
+}
+
+TEST_F(BlockStatementTest, Creation_WithSource) {
+  auto* b = create<BlockStatement>(Source{Source::Location{20, 2}},
+                                   ast::StatementList{});
+  auto src = b->source;
+  EXPECT_EQ(src.range.begin.line, 20u);
+  EXPECT_EQ(src.range.begin.column, 2u);
+}
+
+TEST_F(BlockStatementTest, IsBlock) {
+  auto* b = create<BlockStatement>(ast::StatementList{});
+  EXPECT_TRUE(b->Is<BlockStatement>());
+}
+
+TEST_F(BlockStatementTest, Assert_Null_Statement) {
+  EXPECT_FATAL_FAILURE(
+      {
+        ProgramBuilder b;
+        b.create<BlockStatement>(ast::StatementList{nullptr});
+      },
+      "internal compiler error");
+}
+
+TEST_F(BlockStatementTest, Assert_DifferentProgramID_Statement) {
+  EXPECT_FATAL_FAILURE(
+      {
+        ProgramBuilder b1;
+        ProgramBuilder b2;
+        b1.create<BlockStatement>(
+            ast::StatementList{b2.create<DiscardStatement>()});
+      },
+      "internal compiler error");
+}
+
+}  // namespace
+}  // namespace ast
+}  // namespace tint

src/tint/ast/bool.cc

@@ -0,0 +1,40 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/bool.h"
+
+#include "src/tint/program_builder.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::Bool);
+
+namespace tint {
+namespace ast {
+
+Bool::Bool(ProgramID pid, const Source& src) : Base(pid, src) {}
+
+Bool::Bool(Bool&&) = default;
+
+Bool::~Bool() = default;
+
+std::string Bool::FriendlyName(const SymbolTable&) const {
+  return "bool";
+}
+
+const Bool* Bool::Clone(CloneContext* ctx) const {
+  auto src = ctx->Clone(source);
+  return ctx->dst->create<Bool>(src);
+}
+
+}  // namespace ast
+}  // namespace tint

src/tint/ast/bool.h

@@ -0,0 +1,56 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#ifndef SRC_TINT_AST_BOOL_H_
+#define SRC_TINT_AST_BOOL_H_
+
+#include <string>
+
+#include "src/tint/ast/type.h"
+
+// X11 likes to #define Bool leading to confusing error messages.
+// If its defined, undefine it.
+#ifdef Bool
+#undef Bool
+#endif
+
+namespace tint {
+namespace ast {
+
+/// A boolean type
+class Bool final : public Castable<Bool, Type> {
+ public:
+  /// Constructor
+  /// @param pid the identifier of the program that owns this node
+  /// @param src the source of this node
+  Bool(ProgramID pid, const Source& src);
+  /// Move constructor
+  Bool(Bool&&);
+  ~Bool() override;
+
+  /// @param symbols the program's symbol table
+  /// @returns the name for this type that closely resembles how it would be
+  /// declared in WGSL.
+  std::string FriendlyName(const SymbolTable& symbols) const override;
+
+  /// Clones this type and all transitive types using the `CloneContext` `ctx`.
+  /// @param ctx the clone context
+  /// @return the newly cloned type
+  const Bool* Clone(CloneContext* ctx) const override;
+};
+
+}  // namespace ast
+}  // namespace tint
+
+#endif  // SRC_TINT_AST_BOOL_H_

src/tint/ast/bool_literal_expression.cc

@@ -0,0 +1,39 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/bool_literal_expression.h"
+
+#include "src/tint/program_builder.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::BoolLiteralExpression);
+
+namespace tint {
+namespace ast {
+
+BoolLiteralExpression::BoolLiteralExpression(ProgramID pid,
+                                             const Source& src,
+                                             bool val)
+    : Base(pid, src), value(val) {}
+
+BoolLiteralExpression::~BoolLiteralExpression() = default;
+
+const BoolLiteralExpression* BoolLiteralExpression::Clone(
+    CloneContext* ctx) const {
+  // Clone arguments outside of create() call to have deterministic ordering
+  auto src = ctx->Clone(source);
+  return ctx->dst->create<BoolLiteralExpression>(src, value);
+}
+
+}  // namespace ast
+}  // namespace tint

src/tint/ast/bool_literal_expression.h

@@ -0,0 +1,49 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#ifndef SRC_TINT_AST_BOOL_LITERAL_EXPRESSION_H_
+#define SRC_TINT_AST_BOOL_LITERAL_EXPRESSION_H_
+
+#include <string>
+
+#include "src/tint/ast/literal_expression.h"
+
+namespace tint {
+namespace ast {
+
+/// A boolean literal
+class BoolLiteralExpression final
+    : public Castable<BoolLiteralExpression, LiteralExpression> {
+ public:
+  /// Constructor
+  /// @param pid the identifier of the program that owns this node
+  /// @param src the source of this node
+  /// @param value the bool literals value
+  BoolLiteralExpression(ProgramID pid, const Source& src, bool value);
+  ~BoolLiteralExpression() override;
+
+  /// Clones this node and all transitive child nodes using the `CloneContext`
+  /// `ctx`.
+  /// @param ctx the clone context
+  /// @return the newly cloned node
+  const BoolLiteralExpression* Clone(CloneContext* ctx) const override;
+
+  /// The boolean literal value
+  const bool value;
+};
+
+}  // namespace ast
+}  // namespace tint
+
+#endif  // SRC_TINT_AST_BOOL_LITERAL_EXPRESSION_H_

src/tint/ast/bool_literal_expression_test.cc

@@ -0,0 +1,37 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/test_helper.h"
+
+namespace tint {
+namespace ast {
+namespace {
+
+using BoolLiteralExpressionTest = TestHelper;
+
+TEST_F(BoolLiteralExpressionTest, True) {
+  auto* b = create<BoolLiteralExpression>(true);
+  ASSERT_TRUE(b->Is<BoolLiteralExpression>());
+  ASSERT_TRUE(b->value);
+}
+
+TEST_F(BoolLiteralExpressionTest, False) {
+  auto* b = create<BoolLiteralExpression>(false);
+  ASSERT_TRUE(b->Is<BoolLiteralExpression>());
+  ASSERT_FALSE(b->value);
+}
+
+}  // namespace
+}  // namespace ast
+}  // namespace tint

src/tint/ast/bool_test.cc

@@ -0,0 +1,32 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/bool.h"
+
+#include "src/tint/ast/test_helper.h"
+
+namespace tint {
+namespace ast {
+namespace {
+
+using AstBoolTest = TestHelper;
+
+TEST_F(AstBoolTest, FriendlyName) {
+  auto* b = create<Bool>();
+  EXPECT_EQ(b->FriendlyName(Symbols()), "bool");
+}
+
+}  // namespace
+}  // namespace ast
+}  // namespace tint

src/tint/ast/break_statement.cc

@@ -0,0 +1,38 @@
+// Copyright 2020 The Tint Authors.
+//
+// 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.
+
+#include "src/tint/ast/break_statement.h"
+
+#include "src/tint/program_builder.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::BreakStatement);
+
+namespace tint {
+namespace ast {
+
+BreakStatement::BreakStatement(ProgramID pid, const Source& src)
+    : Base(pid, src) {}
+
+BreakStatement::BreakStatement(BreakStatement&&) = default;
+
+BreakStatement::~BreakStatement() = default;
+
+const BreakStatement* BreakStatement::Clone(CloneContext* ctx) const {
+  // Clone arguments outside of create() call to have deterministic ordering
+  auto src = ctx->Clone(source);
+  return ctx->dst->create<BreakStatement>(src);
+}
+
+}  // namespace ast
+}  // namespace tint