Implement textureGather, textureGatherCompare

All writers implemented, along with resolving and validation.

TODO: SPIR-V Reader.

Bug: tint:1330
Change-Id: I8ba2f6023749474f80efb8a5422ac187e6c73a69
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/71820
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>

docs/origin-trial-changes.md

@@ -14,6 +14,7 @@
 
 ### 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)

src/ast/intrinsic_texture_helper_test.cc

@@ -461,6 +461,349 @@ std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
           "textureDimensions",
           [](ProgramBuilder* b) { return b->ExprList("texture"); },
       },
+
+      {
+          ValidTextureOverload::kGather2dF32,
+          "textureGather(component : i32,\n"
+          "              t         : texture_2d<T>,\n"
+          "              s         : sampler,\n"
+          "              coords    : vec2<f32>) -> vec4<T>",
+          TextureKind::kRegular,
+          ast::SamplerKind::kSampler,
+          ast::TextureDimension::k2d,
+          TextureDataType::kF32,
+          "textureGather",
+          [](ProgramBuilder* b) {
+            return b->ExprList(0,                        // component
+                               "texture",                // t
+                               "sampler",                // s
+                               b->vec2<f32>(1.f, 2.f));  // coords
+          },
+      },
+      {
+          ValidTextureOverload::kGather2dOffsetF32,
+          "textureGather(component : i32,\n"
+          "              t         : texture_2d<T>,\n"
+          "              s         : sampler,\n"
+          "              coords    : vec2<f32>,\n"
+          "              offset    : vec2<i32>) -> vec4<T>",
+          TextureKind::kRegular,
+          ast::SamplerKind::kSampler,
+          ast::TextureDimension::k2d,
+          TextureDataType::kF32,
+          "textureGather",
+          [](ProgramBuilder* b) {
+            return b->ExprList(0,                       // component
+                               "texture",               // t
+                               "sampler",               // s
+                               b->vec2<f32>(1.f, 2.f),  // coords
+                               b->vec2<i32>(3, 4));     // offset
+          },
+      },
+      {
+          ValidTextureOverload::kGather2dArrayF32,
+          "textureGather(component   : i32,\n"
+          "              t           : texture_2d_array<T>,\n"
+          "              s           : sampler,\n"
+          "              coords      : vec2<f32>,\n"
+          "              array_index : i32) -> vec4<T>",
+          TextureKind::kRegular,
+          ast::SamplerKind::kSampler,
+          ast::TextureDimension::k2dArray,
+          TextureDataType::kF32,
+          "textureGather",
+          [](ProgramBuilder* b) {
+            return b->ExprList(0,                       // component
+                               "texture",               // t
+                               "sampler",               // s
+                               b->vec2<f32>(1.f, 2.f),  // coords
+                               3);                      // array index
+          },
+      },
+      {
+          ValidTextureOverload::kGather2dArrayOffsetF32,
+          "textureGather(component   : i32,\n"
+          "              t           : texture_2d_array<T>,\n"
+          "              s           : sampler,\n"
+          "              coords      : vec2<f32>,\n"
+          "              array_index : i32,\n"
+          "              offset      : vec2<i32>) -> vec4<T>",
+          TextureKind::kRegular,
+          ast::SamplerKind::kSampler,
+          ast::TextureDimension::k2dArray,
+          TextureDataType::kF32,
+          "textureGather",
+          [](ProgramBuilder* b) {
+            return b->ExprList(0,                       // component
+                               "texture",               // t
+                               "sampler",               // s
+                               b->vec2<f32>(1.f, 2.f),  // coords
+                               3,                       // array_index
+                               b->vec2<i32>(4, 5));     // offset
+          },
+      },
+      {
+          ValidTextureOverload::kGatherCubeF32,
+          "textureGather(component : i32,\n"
+          "              t         : texture_cube<T>,\n"
+          "              s         : sampler,\n"
+          "              coords    : vec3<f32>) -> vec4<T>",
+          TextureKind::kRegular,
+          ast::SamplerKind::kSampler,
+          ast::TextureDimension::kCube,
+          TextureDataType::kF32,
+          "textureGather",
+          [](ProgramBuilder* b) {
+            return b->ExprList(0,                             // component
+                               "texture",                     // t
+                               "sampler",                     // s
+                               b->vec3<f32>(1.f, 2.f, 3.f));  // coords
+          },
+      },
+      {
+          ValidTextureOverload::kGatherCubeArrayF32,
+          "textureGather(component   : i32,\n"
+          "              t           : texture_cube_array<T>,\n"
+          "              s           : sampler,\n"
+          "              coords      : vec3<f32>,\n"
+          "              array_index : i32) -> vec4<T>",
+          TextureKind::kRegular,
+          ast::SamplerKind::kSampler,
+          ast::TextureDimension::kCubeArray,
+          TextureDataType::kF32,
+          "textureGather",
+          [](ProgramBuilder* b) {
+            return b->ExprList(0,                            // component
+                               "texture",                    // t
+                               "sampler",                    // s
+                               b->vec3<f32>(1.f, 2.f, 3.f),  // coords
+                               4);                           // array_index
+          },
+      },
+      {
+          ValidTextureOverload::kGatherDepth2dF32,
+          "textureGather(t      : texture_depth_2d,\n"
+          "              s      : sampler,\n"
+          "              coords : vec2<f32>) -> vec4<f32>",
+          TextureKind::kDepth,
+          ast::SamplerKind::kSampler,
+          ast::TextureDimension::k2d,
+          TextureDataType::kF32,
+          "textureGather",
+          [](ProgramBuilder* b) {
+            return b->ExprList("texture",                // t
+                               "sampler",                // s
+                               b->vec2<f32>(1.f, 2.f));  // coords
+          },
+      },
+      {
+          ValidTextureOverload::kGatherDepth2dOffsetF32,
+          "textureGather(t      : texture_depth_2d,\n"
+          "              s      : sampler,\n"
+          "              coords : vec2<f32>,\n"
+          "              offset : vec2<i32>) -> vec4<f32>",
+          TextureKind::kDepth,
+          ast::SamplerKind::kSampler,
+          ast::TextureDimension::k2d,
+          TextureDataType::kF32,
+          "textureGather",
+          [](ProgramBuilder* b) {
+            return b->ExprList("texture",               // t
+                               "sampler",               // s
+                               b->vec2<f32>(1.f, 2.f),  // coords
+                               b->vec2<i32>(3, 4));     // offset
+          },
+      },
+      {
+          ValidTextureOverload::kGatherDepth2dArrayF32,
+          "textureGather(t           : texture_depth_2d_array,\n"
+          "              s           : sampler,\n"
+          "              coords      : vec2<f32>,\n"
+          "              array_index : i32) -> vec4<f32>",
+          TextureKind::kDepth,
+          ast::SamplerKind::kSampler,
+          ast::TextureDimension::k2dArray,
+          TextureDataType::kF32,
+          "textureGather",
+          [](ProgramBuilder* b) {
+            return b->ExprList("texture",               // t
+                               "sampler",               // s
+                               b->vec2<f32>(1.f, 2.f),  // coords
+                               3);                      // array_index
+          },
+      },
+      {
+          ValidTextureOverload::kGatherDepth2dArrayOffsetF32,
+          "textureGather(t           : texture_depth_2d_array,\n"
+          "              s           : sampler,\n"
+          "              coords      : vec2<f32>,\n"
+          "              array_index : i32,\n"
+          "              offset      : vec2<i32>) -> vec4<f32>",
+          TextureKind::kDepth,
+          ast::SamplerKind::kSampler,
+          ast::TextureDimension::k2dArray,
+          TextureDataType::kF32,
+          "textureGather",
+          [](ProgramBuilder* b) {
+            return b->ExprList("texture",               // t
+                               "sampler",               // s
+                               b->vec2<f32>(1.f, 2.f),  // coords
+                               3,                       // array_index
+                               b->vec2<i32>(4, 5));     // offset
+          },
+      },
+      {
+          ValidTextureOverload::kGatherDepthCubeF32,
+          "textureGather(t      : texture_depth_cube,\n"
+          "              s      : sampler,\n"
+          "              coords : vec3<f32>) -> vec4<f32>",
+          TextureKind::kDepth,
+          ast::SamplerKind::kSampler,
+          ast::TextureDimension::kCube,
+          TextureDataType::kF32,
+          "textureGather",
+          [](ProgramBuilder* b) {
+            return b->ExprList("texture",                     // t
+                               "sampler",                     // s
+                               b->vec3<f32>(1.f, 2.f, 3.f));  // coords
+          },
+      },
+      {
+          ValidTextureOverload::kGatherDepthCubeArrayF32,
+          "textureGather(t           : texture_depth_cube_array,\n"
+          "              s           : sampler,\n"
+          "              coords      : vec3<f32>,\n"
+          "              array_index : i32) -> vec4<f32>",
+          TextureKind::kDepth,
+          ast::SamplerKind::kSampler,
+          ast::TextureDimension::kCubeArray,
+          TextureDataType::kF32,
+          "textureGather",
+          [](ProgramBuilder* b) {
+            return b->ExprList("texture",                    // t
+                               "sampler",                    // s
+                               b->vec3<f32>(1.f, 2.f, 3.f),  // coords
+                               4);                           // array_index
+          },
+      },
+      {
+          ValidTextureOverload::kGatherCompareDepth2dF32,
+          "textureGatherCompare(t         : texture_depth_2d,\n"
+          "                     s         : sampler_comparison,\n"
+          "                     coords    : vec2<f32>,\n"
+          "                     depth_ref : f32) -> vec4<f32>",
+          TextureKind::kDepth,
+          ast::SamplerKind::kComparisonSampler,
+          ast::TextureDimension::k2d,
+          TextureDataType::kF32,
+          "textureGatherCompare",
+          [](ProgramBuilder* b) {
+            return b->ExprList("texture",               // t
+                               "sampler",               // s
+                               b->vec2<f32>(1.f, 2.f),  // coords
+                               3.f);                    // depth_ref
+          },
+      },
+      {
+          ValidTextureOverload::kGatherCompareDepth2dOffsetF32,
+          "textureGatherCompare(t         : texture_depth_2d,\n"
+          "                     s         : sampler_comparison,\n"
+          "                     coords    : vec2<f32>,\n"
+          "                     depth_ref : f32,\n"
+          "                     offset    : vec2<i32>) -> vec4<f32>",
+          TextureKind::kDepth,
+          ast::SamplerKind::kComparisonSampler,
+          ast::TextureDimension::k2d,
+          TextureDataType::kF32,
+          "textureGatherCompare",
+          [](ProgramBuilder* b) {
+            return b->ExprList("texture",               // t
+                               "sampler",               // s
+                               b->vec2<f32>(1.f, 2.f),  // coords
+                               3.f,                     // depth_ref
+                               b->vec2<i32>(4, 5));     // offset
+          },
+      },
+      {
+          ValidTextureOverload::kGatherCompareDepth2dArrayF32,
+          "textureGatherCompare(t           : texture_depth_2d_array,\n"
+          "                     s           : sampler_comparison,\n"
+          "                     coords      : vec2<f32>,\n"
+          "                     array_index : i32,\n"
+          "                     depth_ref   : f32) -> vec4<f32>",
+          TextureKind::kDepth,
+          ast::SamplerKind::kComparisonSampler,
+          ast::TextureDimension::k2dArray,
+          TextureDataType::kF32,
+          "textureGatherCompare",
+          [](ProgramBuilder* b) {
+            return b->ExprList("texture",               // t
+                               "sampler",               // s
+                               b->vec2<f32>(1.f, 2.f),  // coords
+                               3,                       // array_index
+                               4.f);                    // depth_ref
+          },
+      },
+      {
+          ValidTextureOverload::kGatherCompareDepth2dArrayOffsetF32,
+          "textureGatherCompare(t           : texture_depth_2d_array,\n"
+          "                     s           : sampler_comparison,\n"
+          "                     coords      : vec2<f32>,\n"
+          "                     array_index : i32,\n"
+          "                     depth_ref   : f32,\n"
+          "                     offset      : vec2<i32>) -> vec4<f32>",
+          TextureKind::kDepth,
+          ast::SamplerKind::kComparisonSampler,
+          ast::TextureDimension::k2dArray,
+          TextureDataType::kF32,
+          "textureGatherCompare",
+          [](ProgramBuilder* b) {
+            return b->ExprList("texture",               // t
+                               "sampler",               // s
+                               b->vec2<f32>(1.f, 2.f),  // coords
+                               3,                       // array_index
+                               4.f,                     // depth_ref
+                               b->vec2<i32>(5, 6));     // offset
+          },
+      },
+      {
+          ValidTextureOverload::kGatherCompareDepthCubeF32,
+          "textureGatherCompare(t         : texture_depth_cube,\n"
+          "                     s         : sampler_comparison,\n"
+          "                     coords    : vec3<f32>,\n"
+          "                     depth_ref : f32) -> vec4<f32>",
+          TextureKind::kDepth,
+          ast::SamplerKind::kComparisonSampler,
+          ast::TextureDimension::kCube,
+          TextureDataType::kF32,
+          "textureGatherCompare",
+          [](ProgramBuilder* b) {
+            return b->ExprList("texture",                    // t
+                               "sampler",                    // s
+                               b->vec3<f32>(1.f, 2.f, 3.f),  // coords
+                               4.f);                         // depth_ref
+          },
+      },
+      {
+          ValidTextureOverload::kGatherCompareDepthCubeArrayF32,
+          "textureGatherCompare(t           : texture_depth_cube_array,\n"
+          "                     s           : sampler_comparison,\n"
+          "                     coords      : vec3<f32>,\n"
+          "                     array_index : i32,\n"
+          "                     depth_ref   : f32) -> vec4<f32>",
+          TextureKind::kDepth,
+          ast::SamplerKind::kComparisonSampler,
+          ast::TextureDimension::kCubeArray,
+          TextureDataType::kF32,
+          "textureGatherCompare",
+          [](ProgramBuilder* b) {
+            return b->ExprList("texture",                    // t
+                               "sampler",                    // s
+                               b->vec3<f32>(1.f, 2.f, 3.f),  // coords
+                               4,                            // array_index
+                               5.f);                         // depth_ref
+          },
+      },
       {
           ValidTextureOverload::kNumLayers2dArray,
           "textureNumLayers(t : texture_2d_array<f32>) -> i32",

src/ast/intrinsic_texture_helper_test.h

@@ -65,6 +65,24 @@ enum class ValidTextureOverload {
   kDimensionsStorageWO2d,
   kDimensionsStorageWO2dArray,
   kDimensionsStorageWO3d,
+  kGather2dF32,
+  kGather2dOffsetF32,
+  kGather2dArrayF32,
+  kGather2dArrayOffsetF32,
+  kGatherCubeF32,
+  kGatherCubeArrayF32,
+  kGatherDepth2dF32,
+  kGatherDepth2dOffsetF32,
+  kGatherDepth2dArrayF32,
+  kGatherDepth2dArrayOffsetF32,
+  kGatherDepthCubeF32,
+  kGatherDepthCubeArrayF32,
+  kGatherCompareDepth2dF32,
+  kGatherCompareDepth2dOffsetF32,
+  kGatherCompareDepth2dArrayF32,
+  kGatherCompareDepth2dArrayOffsetF32,
+  kGatherCompareDepthCubeF32,
+  kGatherCompareDepthCubeArrayF32,
   kNumLayers2dArray,
   kNumLayersCubeArray,
   kNumLayersDepth2dArray,

src/intrinsics.def

@@ -423,6 +423,24 @@ fn textureDimensions<F: texel_format, A: write_only>(texture: texture_storage_2d
 fn textureDimensions<F: texel_format, A: write_only>(texture: texture_storage_2d_array<F, A>) -> vec2<i32>
 fn textureDimensions<F: texel_format, A: write_only>(texture: texture_storage_3d<F, A>) -> vec3<i32>
 fn textureDimensions(texture: texture_external) -> vec2<i32>
+fn textureGather<T: fiu32>(component: i32, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>) -> vec4<T>
+fn textureGather<T: fiu32>(component: i32, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<T>
+fn textureGather<T: fiu32>(component: i32, texture: texture_2d_array<T>, sampler: sampler, coords: vec2<f32>, array_index: i32) -> vec4<T>
+fn textureGather<T: fiu32>(component: i32, texture: texture_2d_array<T>, sampler: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<T>
+fn textureGather<T: fiu32>(component: i32, texture: texture_cube<T>, sampler: sampler, coords: vec3<f32>) -> vec4<T>
+fn textureGather<T: fiu32>(component: i32, texture: texture_cube_array<T>, sampler: sampler, coords: vec3<f32>, array_index: i32) -> vec4<T>
+fn textureGather(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
+fn textureGather(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<f32>
+fn textureGather(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32) -> vec4<f32>
+fn textureGather(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<f32>
+fn textureGather(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
+fn textureGather(texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: i32) -> vec4<f32>
+fn textureGatherCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> vec4<f32>
+fn textureGatherCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, offset: vec2<i32>) -> vec4<f32>
+fn textureGatherCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32) -> vec4<f32>
+fn textureGatherCompare(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: i32, depth_ref: f32, offset: vec2<i32>) -> vec4<f32>
+fn textureGatherCompare(texture: texture_depth_cube, sampler: sampler_comparison, coords: vec3<f32>, depth_ref: f32) -> vec4<f32>
+fn textureGatherCompare(texture: texture_depth_cube_array, sampler: sampler_comparison, coords: vec3<f32>, array_index: i32, depth_ref: f32) -> vec4<f32>
 fn textureNumLayers<T: fiu32>(texture: texture_2d_array<T>) -> i32
 fn textureNumLayers<T: fiu32>(texture: texture_cube_array<T>) -> i32
 fn textureNumLayers(texture: texture_depth_2d_array) -> i32

src/resolver/intrinsic_test.cc

@@ -1744,6 +1744,42 @@ const char* expected_texture_overload(
     case ValidTextureOverload::kDimensionsStorageWO2dArray:
     case ValidTextureOverload::kDimensionsStorageWO3d:
       return R"(textureDimensions(texture))";
+    case ValidTextureOverload::kGather2dF32:
+      return R"(textureGather(component, texture, sampler, coords))";
+    case ValidTextureOverload::kGather2dOffsetF32:
+      return R"(textureGather(component, texture, sampler, coords, offset))";
+    case ValidTextureOverload::kGather2dArrayF32:
+      return R"(textureGather(component, texture, sampler, coords, array_index))";
+    case ValidTextureOverload::kGather2dArrayOffsetF32:
+      return R"(textureGather(component, texture, sampler, coords, array_index, offset))";
+    case ValidTextureOverload::kGatherCubeF32:
+      return R"(textureGather(component, texture, sampler, coords))";
+    case ValidTextureOverload::kGatherCubeArrayF32:
+      return R"(textureGather(component, texture, sampler, coords, array_index))";
+    case ValidTextureOverload::kGatherDepth2dF32:
+      return R"(textureGather(texture, sampler, coords))";
+    case ValidTextureOverload::kGatherDepth2dOffsetF32:
+      return R"(textureGather(texture, sampler, coords, offset))";
+    case ValidTextureOverload::kGatherDepth2dArrayF32:
+      return R"(textureGather(texture, sampler, coords, array_index))";
+    case ValidTextureOverload::kGatherDepth2dArrayOffsetF32:
+      return R"(textureGather(texture, sampler, coords, array_index, offset))";
+    case ValidTextureOverload::kGatherDepthCubeF32:
+      return R"(textureGather(texture, sampler, coords))";
+    case ValidTextureOverload::kGatherDepthCubeArrayF32:
+      return R"(textureGather(texture, sampler, coords, array_index))";
+    case ValidTextureOverload::kGatherCompareDepth2dF32:
+      return R"(textureGatherCompare(texture, sampler, coords, depth_ref))";
+    case ValidTextureOverload::kGatherCompareDepth2dOffsetF32:
+      return R"(textureGatherCompare(texture, sampler, coords, depth_ref, offset))";
+    case ValidTextureOverload::kGatherCompareDepth2dArrayF32:
+      return R"(textureGatherCompare(texture, sampler, coords, array_index, depth_ref))";
+    case ValidTextureOverload::kGatherCompareDepth2dArrayOffsetF32:
+      return R"(textureGatherCompare(texture, sampler, coords, array_index, depth_ref, offset))";
+    case ValidTextureOverload::kGatherCompareDepthCubeF32:
+      return R"(textureGatherCompare(texture, sampler, coords, depth_ref))";
+    case ValidTextureOverload::kGatherCompareDepthCubeArrayF32:
+      return R"(textureGatherCompare(texture, sampler, coords, array_index, depth_ref))";
     case ValidTextureOverload::kNumLayers2dArray:
     case ValidTextureOverload::kNumLayersCubeArray:
     case ValidTextureOverload::kNumLayersDepth2dArray:
@@ -1965,6 +2001,26 @@ TEST_P(ResolverIntrinsicTest_Texture, Call) {
     EXPECT_TRUE(TypeOf(call)->Is<sem::I32>());
   } else if (std::string(param.function) == "textureStore") {
     EXPECT_TRUE(TypeOf(call)->Is<sem::Void>());
+  } else if (std::string(param.function) == "textureGather") {
+    auto* vec = As<sem::Vector>(TypeOf(call));
+    ASSERT_NE(vec, nullptr);
+    EXPECT_EQ(vec->Width(), 4u);
+    switch (param.texture_data_type) {
+      case ast::intrinsic::test::TextureDataType::kF32:
+        EXPECT_TRUE(vec->type()->Is<sem::F32>());
+        break;
+      case ast::intrinsic::test::TextureDataType::kU32:
+        EXPECT_TRUE(vec->type()->Is<sem::U32>());
+        break;
+      case ast::intrinsic::test::TextureDataType::kI32:
+        EXPECT_TRUE(vec->type()->Is<sem::I32>());
+        break;
+    }
+  } else if (std::string(param.function) == "textureGatherCompare") {
+    auto* vec = As<sem::Vector>(TypeOf(call));
+    ASSERT_NE(vec, nullptr);
+    EXPECT_EQ(vec->Width(), 4u);
+    EXPECT_TRUE(vec->type()->Is<sem::F32>());
   } else {
     switch (param.texture_kind) {
       case ast::intrinsic::test::TextureKind::kRegular:

src/resolver/intrinsic_validation_test.cc

@@ -121,7 +121,7 @@ TEST_F(ResolverIntrinsicValidationTest, IntrinsicRedeclaredAsStruct) {
       R"(12:34 error: 'mix' is a builtin and cannot be redeclared as a struct)");
 }
 
-namespace TextureSamplerOffset {
+namespace texture_constexpr_args {
 
 using TextureOverloadCase = ast::intrinsic::test::TextureOverloadCase;
 using ValidTextureOverload = ast::intrinsic::test::ValidTextureOverload;
@@ -131,194 +131,272 @@ using u32 = ProgramBuilder::u32;
 using i32 = ProgramBuilder::i32;
 using f32 = ProgramBuilder::f32;
 
-static std::vector<TextureOverloadCase> ValidCases() {
+static std::vector<TextureOverloadCase> TextureCases(
+    std::unordered_set<ValidTextureOverload> overloads) {
   std::vector<TextureOverloadCase> cases;
   for (auto c : TextureOverloadCase::ValidCases()) {
-    if (std::string(c.function).find("textureSample") == 0) {
+    if (overloads.count(c.overload)) {
-      if (std::string(c.description).find("offset ") != std::string::npos) {
       cases.push_back(c);
     }
   }
-  }
   return cases;
 }
 
-struct OffsetCase {
+enum class Position {
-  bool is_valid;
+  kFirst,
-  int32_t x;
+  kLast,
-  int32_t y;
-  int32_t z;
-  int32_t illegal_value = 0;
 };
 
-static std::vector<OffsetCase> OffsetCases() {
+struct Parameter {
-  return {
+  const char* const name;
-      {true, 0, 1, 2},          //
+  const Position position;
-      {true, 7, -8, 7},         //
+  int min;
-      {false, 10, 10, 20, 10},  //
+  int max;
-      {false, -9, 0, 0, -9},    //
-      {false, 0, 8, 0, 8},      //
 };
+
+class Constexpr {
+ public:
+  enum class Kind {
+    kScalar,
+    kVec2,
+    kVec3,
+    kVec3_Scalar_Vec2,
+    kVec3_Vec2_Scalar,
+    kEmptyVec2,
+    kEmptyVec3,
+  };
+
+  Constexpr(int32_t invalid_idx,
+            Kind k,
+            int32_t x = 0,
+            int32_t y = 0,
+            int32_t z = 0)
+      : invalid_index(invalid_idx), kind(k), values{x, y, z} {}
+
+  const ast::Expression* operator()(Source src, ProgramBuilder& b) {
+    switch (kind) {
+      case Kind::kScalar:
+        return b.Expr(src, values[0]);
+      case Kind::kVec2:
+        return b.Construct(src, b.ty.vec2<i32>(), values[0], values[1]);
+      case Kind::kVec3:
+        return b.Construct(src, b.ty.vec3<i32>(), values[0], values[1],
+                           values[2]);
+      case Kind::kVec3_Scalar_Vec2:
+        return b.Construct(src, b.ty.vec3<i32>(), values[0],
+                           b.vec2<i32>(values[1], values[2]));
+      case Kind::kVec3_Vec2_Scalar:
+        return b.Construct(src, b.ty.vec3<i32>(),
+                           b.vec2<i32>(values[0], values[1]), values[2]);
+      case Kind::kEmptyVec2:
+        return b.Construct(src, b.ty.vec2<i32>());
+      case Kind::kEmptyVec3:
+        return b.Construct(src, b.ty.vec3<i32>());
+    }
+    return nullptr;
   }
 
-using IntrinsicTextureSamplerValidationTest =
+  static const constexpr int32_t kValid = -1;
-    ResolverTestWithParam<std::tuple<TextureOverloadCase,  // texture info
+  const int32_t invalid_index;  // Expected error value, or kValid
-                                     OffsetCase            // offset info
+  const Kind kind;
-                                     >>;
+  const std::array<int32_t, 3> values;
-TEST_P(IntrinsicTextureSamplerValidationTest, ConstExpr) {
+};
+
+static std::ostream& operator<<(std::ostream& out, Parameter param) {
+  return out << param.name;
+}
+
+static std::ostream& operator<<(std::ostream& out, Constexpr expr) {
+  switch (expr.kind) {
+    case Constexpr::Kind::kScalar:
+      return out << expr.values[0];
+    case Constexpr::Kind::kVec2:
+      return out << "vec2(" << expr.values[0] << ", " << expr.values[1] << ")";
+    case Constexpr::Kind::kVec3:
+      return out << "vec3(" << expr.values[0] << ", " << expr.values[1] << ", "
+                 << expr.values[2] << ")";
+    case Constexpr::Kind::kVec3_Scalar_Vec2:
+      return out << "vec3(" << expr.values[0] << ", vec2(" << expr.values[1]
+                 << ", " << expr.values[2] << "))";
+    case Constexpr::Kind::kVec3_Vec2_Scalar:
+      return out << "vec3(vec2(" << expr.values[0] << ", " << expr.values[1]
+                 << "), " << expr.values[2] << ")";
+    case Constexpr::Kind::kEmptyVec2:
+      return out << "vec2()";
+    case Constexpr::Kind::kEmptyVec3:
+      return out << "vec3()";
+  }
+  return out;
+}
+
+using IntrinsicTextureConstExprArgValidationTest = ResolverTestWithParam<
+    std::tuple<TextureOverloadCase, Parameter, Constexpr>>;
+
+TEST_P(IntrinsicTextureConstExprArgValidationTest, Immediate) {
   auto& p = GetParam();
-  auto param = std::get<0>(p);
+  auto overload = std::get<0>(p);
-  auto offset = std::get<1>(p);
+  auto param = std::get<1>(p);
-  param.BuildTextureVariable(this);
+  auto expr = std::get<2>(p);
-  param.BuildSamplerVariable(this);
 
-  auto args = param.args(this);
+  overload.BuildTextureVariable(this);
-  // Make Resolver visit the Node about to be removed
+  overload.BuildSamplerVariable(this);
-  WrapInFunction(args.back());
-  args.pop_back();
-  if (NumCoordinateAxes(param.texture_dimension) == 2) {
-    args.push_back(
-        Construct(Source{{12, 34}}, ty.vec2<i32>(), offset.x, offset.y));
-  } else if (NumCoordinateAxes(param.texture_dimension) == 3) {
-    args.push_back(Construct(Source{{12, 34}}, ty.vec3<i32>(), offset.x,
-                             offset.y, offset.z));
-  }
 
-  auto* call = Call(param.function, args);
+  auto args = overload.args(this);
-  Func("func", {}, ty.void_(), {CallStmt(call)},
+  auto*& arg_to_replace =
-       {create<ast::StageDecoration>(ast::PipelineStage::kFragment)});
+      (param.position == Position::kFirst) ? args.front() : args.back();
 
-  if (offset.is_valid) {
+  // BuildTextureVariable() uses a Literal for scalars, and a CallExpression for
+  // a vector constructor.
+  bool is_vector = arg_to_replace->Is<ast::CallExpression>();
+
+  // Make the expression to be replaced, reachable. This keeps the resolver
+  // happy.
+  WrapInFunction(arg_to_replace);
+
+  arg_to_replace = expr(Source{{12, 34}}, *this);
+
+  // Call the intrinsic with the constexpr argument replaced
+  Func("func", {}, ty.void_(), {CallStmt(Call(overload.function, args))},
+       {Stage(ast::PipelineStage::kFragment)});
+
+  if (expr.invalid_index == Constexpr::kValid) {
     EXPECT_TRUE(r()->Resolve()) << r()->error();
   } else {
     EXPECT_FALSE(r()->Resolve());
     std::stringstream err;
-    err << "12:34 error: each offset component of '" << param.function
+    if (is_vector) {
-        << "' must be at least -8 and at most 7. found: '"
+      err << "12:34 error: each component of the " << param.name
-        << std::to_string(offset.illegal_value) << "'";
+          << " argument must be at least " << param.min << " and at most "
-    EXPECT_EQ(r()->error(), err.str());
+          << param.max << ". " << param.name << " component "
-  }
+          << expr.invalid_index << " is "
-}
+          << std::to_string(expr.values[expr.invalid_index]);
-
-TEST_P(IntrinsicTextureSamplerValidationTest, ConstExprOfConstExpr) {
-  auto& p = GetParam();
-  auto param = std::get<0>(p);
-  auto offset = std::get<1>(p);
-  param.BuildTextureVariable(this);
-  param.BuildSamplerVariable(this);
-
-  auto args = param.args(this);
-  // Make Resolver visit the Node about to be removed
-  WrapInFunction(args.back());
-  args.pop_back();
-  if (NumCoordinateAxes(param.texture_dimension) == 2) {
-    args.push_back(Construct(Source{{12, 34}}, ty.vec2<i32>(),
-                             Construct(ty.i32(), offset.x), offset.y));
-  } else if (NumCoordinateAxes(param.texture_dimension) == 3) {
-    args.push_back(Construct(Source{{12, 34}}, ty.vec3<i32>(), offset.x,
-                             Construct(ty.vec2<i32>(), offset.y, offset.z)));
-  }
-  auto* call = Call(param.function, args);
-  Func("func", {}, ty.void_(), {CallStmt(call)},
-       {create<ast::StageDecoration>(ast::PipelineStage::kFragment)});
-  if (offset.is_valid) {
-    EXPECT_TRUE(r()->Resolve()) << r()->error();
     } else {
-    EXPECT_FALSE(r()->Resolve());
+      err << "12:34 error: the " << param.name << " argument must be at least "
-    std::stringstream err;
+          << param.min << " and at most " << param.max << ". " << param.name
-    err << "12:34 error: each offset component of '" << param.function
+          << " is " << std::to_string(expr.values[expr.invalid_index]);
-        << "' must be at least -8 and at most 7. found: '"
+    }
-        << std::to_string(offset.illegal_value) << "'";
     EXPECT_EQ(r()->error(), err.str());
   }
 }
 
-TEST_P(IntrinsicTextureSamplerValidationTest, EmptyVectorConstructor) {
+TEST_P(IntrinsicTextureConstExprArgValidationTest, GlobalConst) {
   auto& p = GetParam();
-  auto param = std::get<0>(p);
+  auto overload = std::get<0>(p);
-  param.BuildTextureVariable(this);
+  auto param = std::get<1>(p);
-  param.BuildSamplerVariable(this);
+  auto expr = std::get<2>(p);
 
-  auto args = param.args(this);
+  // Build the global texture and sampler variables
-  // Make Resolver visit the Node about to be removed
+  overload.BuildTextureVariable(this);
-  WrapInFunction(args.back());
+  overload.BuildSamplerVariable(this);
-  args.pop_back();
-  if (NumCoordinateAxes(param.texture_dimension) == 2) {
-    args.push_back(Construct(Source{{12, 34}}, ty.vec2<i32>()));
-  } else if (NumCoordinateAxes(param.texture_dimension) == 3) {
-    args.push_back(Construct(Source{{12, 34}}, ty.vec3<i32>()));
-  }
 
-  auto* call = Call(param.function, args);
+  // Build the module-scope let 'G' with the offset value
-  Func("func", {}, ty.void_(), {CallStmt(call)},
+  GlobalConst("G", nullptr, expr({}, *this));
-       {create<ast::StageDecoration>(ast::PipelineStage::kFragment)});
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
-}
 
-TEST_P(IntrinsicTextureSamplerValidationTest, GlobalConst) {
+  auto args = overload.args(this);
-  auto& p = GetParam();
+  auto*& arg_to_replace =
-  auto param = std::get<0>(p);
+      (param.position == Position::kFirst) ? args.front() : args.back();
-  auto offset = std::get<1>(p);
-  param.BuildTextureVariable(this);
-  param.BuildSamplerVariable(this);
 
-  auto args = param.args(this);
+  // Make the expression to be replaced, reachable. This keeps the resolver
-  // Make Resolver visit the Node about to be removed
+  // happy.
-  WrapInFunction(args.back());
+  WrapInFunction(arg_to_replace);
-  args.pop_back();
+
-  GlobalConst("offset_2d", ty.vec2<i32>(), vec2<i32>(offset.x, offset.y));
+  arg_to_replace = Expr(Source{{12, 34}}, "G");
-  GlobalConst("offset_3d", ty.vec3<i32>(),
+
-              vec3<i32>(offset.x, offset.y, offset.z));
+  // Call the intrinsic with the constexpr argument replaced
-  if (NumCoordinateAxes(param.texture_dimension) == 2) {
+  Func("func", {}, ty.void_(), {CallStmt(Call(overload.function, args))},
-    args.push_back(Expr(Source{{12, 34}}, "offset_2d"));
+       {Stage(ast::PipelineStage::kFragment)});
-  } else if (NumCoordinateAxes(param.texture_dimension) == 3) {
-    args.push_back(Expr(Source{{12, 34}}, "offset_3d"));
-  }
 
-  auto* call = Call(param.function, args);
-  Func("func", {}, ty.void_(), {CallStmt(call)},
-       {create<ast::StageDecoration>(ast::PipelineStage::kFragment)});
   EXPECT_FALSE(r()->Resolve());
   std::stringstream err;
-  err << "12:34 error: '" << param.function
+  err << "12:34 error: the " << param.name
-      << "' offset parameter must be a const_expression";
+      << " argument must be a const_expression";
   EXPECT_EQ(r()->error(), err.str());
 }
 
-TEST_P(IntrinsicTextureSamplerValidationTest, ScalarConst) {
+INSTANTIATE_TEST_SUITE_P(
-  auto& p = GetParam();
+    Offset2D,
-  auto param = std::get<0>(p);
+    IntrinsicTextureConstExprArgValidationTest,
-  auto offset = std::get<1>(p);
+    testing::Combine(
-  param.BuildTextureVariable(this);
+        testing::ValuesIn(TextureCases({
-  param.BuildSamplerVariable(this);
+            ValidTextureOverload::kSample2dOffsetF32,
-  auto* x = Const("x", ty.i32(), Construct(ty.i32(), offset.x));
+            ValidTextureOverload::kSample2dArrayOffsetF32,
+            ValidTextureOverload::kSampleDepth2dOffsetF32,
+            ValidTextureOverload::kSampleDepth2dArrayOffsetF32,
+            ValidTextureOverload::kSampleBias2dOffsetF32,
+            ValidTextureOverload::kSampleBias2dArrayOffsetF32,
+            ValidTextureOverload::kSampleLevel2dOffsetF32,
+            ValidTextureOverload::kSampleLevel2dArrayOffsetF32,
+            ValidTextureOverload::kSampleLevelDepth2dOffsetF32,
+            ValidTextureOverload::kSampleLevelDepth2dArrayOffsetF32,
+            ValidTextureOverload::kSampleGrad2dOffsetF32,
+            ValidTextureOverload::kSampleGrad2dArrayOffsetF32,
+            ValidTextureOverload::kSampleCompareDepth2dOffsetF32,
+            ValidTextureOverload::kSampleCompareDepth2dArrayOffsetF32,
+            ValidTextureOverload::kSampleCompareLevelDepth2dOffsetF32,
+            ValidTextureOverload::kSampleCompareLevelDepth2dArrayOffsetF32,
+        })),
+        testing::Values(Parameter{"offset", Position::kLast, -8, 7}),
+        testing::Values(
+            Constexpr{Constexpr::kValid, Constexpr::Kind::kEmptyVec2},
+            Constexpr{Constexpr::kValid, Constexpr::Kind::kVec2, -1, 1},
+            Constexpr{Constexpr::kValid, Constexpr::Kind::kVec2, 7, -8},
+            Constexpr{0, Constexpr::Kind::kVec2, 8, 0},
+            Constexpr{1, Constexpr::Kind::kVec2, 0, 8},
+            Constexpr{0, Constexpr::Kind::kVec2, -9, 0},
+            Constexpr{1, Constexpr::Kind::kVec2, 0, -9},
+            Constexpr{0, Constexpr::Kind::kVec2, 8, 8},
+            Constexpr{0, Constexpr::Kind::kVec2, -9, -9})));
 
-  auto args = param.args(this);
+INSTANTIATE_TEST_SUITE_P(
-  // Make Resolver visit the Node about to be removed
+    Offset3D,
-  WrapInFunction(args.back());
+    IntrinsicTextureConstExprArgValidationTest,
-  args.pop_back();
+    testing::Combine(
-  if (NumCoordinateAxes(param.texture_dimension) == 2) {
+        testing::ValuesIn(TextureCases({
-    args.push_back(Construct(Source{{12, 34}}, ty.vec2<i32>(), x, offset.y));
+            ValidTextureOverload::kSample3dOffsetF32,
-  } else if (NumCoordinateAxes(param.texture_dimension) == 3) {
+            ValidTextureOverload::kSampleBias3dOffsetF32,
-    args.push_back(
+            ValidTextureOverload::kSampleLevel3dOffsetF32,
-        Construct(Source{{12, 34}}, ty.vec3<i32>(), x, offset.y, offset.z));
+            ValidTextureOverload::kSampleGrad3dOffsetF32,
-  }
+        })),
+        testing::Values(Parameter{"offset", Position::kLast, -8, 7}),
+        testing::Values(
+            Constexpr{Constexpr::kValid, Constexpr::Kind::kEmptyVec3},
+            Constexpr{Constexpr::kValid, Constexpr::Kind::kVec3, 0, 0, 0},
+            Constexpr{Constexpr::kValid, Constexpr::Kind::kVec3, 7, -8, 7},
+            Constexpr{0, Constexpr::Kind::kVec3, 10, 0, 0},
+            Constexpr{1, Constexpr::Kind::kVec3, 0, 10, 0},
+            Constexpr{2, Constexpr::Kind::kVec3, 0, 0, 10},
+            Constexpr{0, Constexpr::Kind::kVec3, 10, 11, 12},
+            Constexpr{0, Constexpr::Kind::kVec3_Scalar_Vec2, 10, 0, 0},
+            Constexpr{1, Constexpr::Kind::kVec3_Scalar_Vec2, 0, 10, 0},
+            Constexpr{2, Constexpr::Kind::kVec3_Scalar_Vec2, 0, 0, 10},
+            Constexpr{0, Constexpr::Kind::kVec3_Scalar_Vec2, 10, 11, 12},
+            Constexpr{0, Constexpr::Kind::kVec3_Vec2_Scalar, 10, 0, 0},
+            Constexpr{1, Constexpr::Kind::kVec3_Vec2_Scalar, 0, 10, 0},
+            Constexpr{2, Constexpr::Kind::kVec3_Vec2_Scalar, 0, 0, 10},
+            Constexpr{0, Constexpr::Kind::kVec3_Vec2_Scalar, 10, 11, 12})));
 
-  auto* call = Call(param.function, args);
+INSTANTIATE_TEST_SUITE_P(
-  Func("func", {}, ty.void_(), {Decl(x), CallStmt(call)},
+    Component,
-       {create<ast::StageDecoration>(ast::PipelineStage::kFragment)});
+    IntrinsicTextureConstExprArgValidationTest,
-  EXPECT_FALSE(r()->Resolve());
+    testing::Combine(
-  std::stringstream err;
+        testing::ValuesIn(
-  err << "12:34 error: '" << param.function
+            TextureCases({ValidTextureOverload::kGather2dF32,
-      << "' offset parameter must be a const_expression";
+                          ValidTextureOverload::kGather2dOffsetF32,
-  EXPECT_EQ(r()->error(), err.str());
+                          ValidTextureOverload::kGather2dArrayF32,
-}
+                          ValidTextureOverload::kGather2dArrayOffsetF32,
+                          ValidTextureOverload::kGatherCubeF32,
+                          ValidTextureOverload::kGatherCubeArrayF32})),
+        testing::Values(Parameter{"component", Position::kFirst, 0, 3}),
+        testing::Values(
+            Constexpr{Constexpr::kValid, Constexpr::Kind::kScalar, 0},
+            Constexpr{Constexpr::kValid, Constexpr::Kind::kScalar, 1},
+            Constexpr{Constexpr::kValid, Constexpr::Kind::kScalar, 2},
+            Constexpr{Constexpr::kValid, Constexpr::Kind::kScalar, 3},
+            Constexpr{0, Constexpr::Kind::kScalar, 4},
+            Constexpr{0, Constexpr::Kind::kScalar, 123},
+            Constexpr{0, Constexpr::Kind::kScalar, -1})));
 
-INSTANTIATE_TEST_SUITE_P(IntrinsicTextureSamplerValidationTest,
+}  // namespace texture_constexpr_args
-                         IntrinsicTextureSamplerValidationTest,
-                         testing::Combine(testing::ValuesIn(ValidCases()),
-                                          testing::ValuesIn(OffsetCases())));
-}  // namespace TextureSamplerOffset
 
 }  // namespace
 }  // namespace resolver

src/resolver/resolver.h

@@ -45,7 +45,6 @@ class BitcastExpression;
 class CallExpression;
 class CallStatement;
 class CaseStatement;
-class ConstructorExpression;
 class ForLoopStatement;
 class Function;
 class IdentifierExpression;

src/resolver/resolver_validation.cc

@@ -1505,22 +1505,30 @@ bool Resolver::ValidateTextureIntrinsicFunction(const sem::Call* call) {
   if (!intrinsic) {
     return false;
   }
+
   std::string func_name = intrinsic->str();
   auto& signature = intrinsic->Signature();
-  auto index = signature.IndexOf(sem::ParameterUsage::kOffset);
+
-  if (index > -1) {
+  auto check_arg_is_constexpr = [&](sem::ParameterUsage usage, int min,
+                                    int max) {
+    auto index = signature.IndexOf(usage);
+    if (index < 0) {
+      return true;
+    }
+    std::string name = sem::str(usage);
     auto* arg = call->Arguments()[index];
     if (auto values = arg->ConstantValue()) {
       // Assert that the constant values are of the expected type.
-      if (!values.Type()->Is<sem::Vector>() ||
+      if (!values.Type()->IsAnyOf<sem::I32, sem::Vector>() ||
           !values.ElementType()->Is<sem::I32>()) {
         TINT_ICE(Resolver, diagnostics_)
-            << "failed to resolve '" + func_name + "' offset parameter type";
+            << "failed to resolve '" + func_name + "' " << name
+            << " parameter type";
         return false;
       }
 
       // Currently const_expr is restricted to literals and type constructors.
-      // Check that that's all we have for the offset parameter.
+      // Check that that's all we have for the parameter.
       bool is_const_expr = true;
       ast::TraverseExpressions(
           arg->Declaration(), diagnostics_, [&](const ast::Expression* e) {
@@ -1531,25 +1539,37 @@ bool Resolver::ValidateTextureIntrinsicFunction(const sem::Call* call) {
             return ast::TraverseAction::Stop;
           });
       if (is_const_expr) {
-        for (auto offset : values.Elements()) {
+        auto vector = intrinsic->Parameters()[index]->Type()->Is<sem::Vector>();
-          auto offset_value = offset.i32;
+        for (size_t i = 0; i < values.Elements().size(); i++) {
-          if (offset_value < -8 || offset_value > 7) {
+          auto value = values.Elements()[i].i32;
-            AddError("each offset component of '" + func_name +
+          if (value < min || value > max) {
-                         "' must be at least -8 and at most 7. "
+            if (vector) {
-                         "found: '" +
+              AddError("each component of the " + name +
-                         std::to_string(offset_value) + "'",
+                           " argument must be at least " + std::to_string(min) +
+                           " and at most " + std::to_string(max) + ". " + name +
+                           " component " + std::to_string(i) + " is " +
+                           std::to_string(value),
                        arg->Declaration()->source);
+            } else {
+              AddError("the " + name + " argument must be at least " +
+                           std::to_string(min) + " and at most " +
+                           std::to_string(max) + ". " + name + " is " +
+                           std::to_string(value),
+                       arg->Declaration()->source);
+            }
             return false;
           }
         }
         return true;
       }
     }
-    AddError("'" + func_name + "' offset parameter must be a const_expression",
+    AddError("the " + name + " argument must be a const_expression",
              arg->Declaration()->source);
     return false;
-  }
+  };
-  return true;
+
+  return check_arg_is_constexpr(sem::ParameterUsage::kOffset, -8, 7) &&
+         check_arg_is_constexpr(sem::ParameterUsage::kComponent, 0, 3);
 }
 
 bool Resolver::ValidateFunctionCall(const sem::Call* call) {

src/sem/intrinsic.cc

@@ -50,6 +50,8 @@ bool IsFloatClassificationIntrinsic(IntrinsicType i) {
 
 bool IsTextureIntrinsic(IntrinsicType i) {
   return IsImageQueryIntrinsic(i) || i == IntrinsicType::kTextureLoad ||
+         i == IntrinsicType::kTextureGather ||
+         i == IntrinsicType::kTextureGatherCompare ||
          i == IntrinsicType::kTextureSample ||
          i == IntrinsicType::kTextureSampleLevel ||
          i == IntrinsicType::kTextureSampleBias ||

src/sem/intrinsic_type.cc

@@ -261,6 +261,12 @@ IntrinsicType ParseIntrinsicType(const std::string& name) {
   if (name == "textureDimensions") {
     return IntrinsicType::kTextureDimensions;
   }
+  if (name == "textureGather") {
+    return IntrinsicType::kTextureGather;
+  }
+  if (name == "textureGatherCompare") {
+    return IntrinsicType::kTextureGatherCompare;
+  }
   if (name == "textureNumLayers") {
     return IntrinsicType::kTextureNumLayers;
   }
@@ -488,6 +494,10 @@ const char* str(IntrinsicType i) {
       return "workgroupBarrier";
     case IntrinsicType::kTextureDimensions:
       return "textureDimensions";
+    case IntrinsicType::kTextureGather:
+      return "textureGather";
+    case IntrinsicType::kTextureGatherCompare:
+      return "textureGatherCompare";
     case IntrinsicType::kTextureNumLayers:
       return "textureNumLayers";
     case IntrinsicType::kTextureNumLevels:

src/sem/intrinsic_type.h

@@ -111,6 +111,8 @@ enum class IntrinsicType {
   kUnpack4x8unorm,
   kWorkgroupBarrier,
   kTextureDimensions,
+  kTextureGather,
+  kTextureGatherCompare,
   kTextureNumLayers,
   kTextureNumLevels,
   kTextureNumSamples,

src/sem/parameter_usage.cc

@@ -35,6 +35,8 @@ const char* str(ParameterUsage usage) {
       return "array_index";
     case ParameterUsage::kBias:
       return "bias";
+    case ParameterUsage::kComponent:
+      return "component";
     case ParameterUsage::kCoords:
       return "coords";
     case ParameterUsage::kDdx:

src/sem/parameter_usage.h

@@ -34,6 +34,7 @@ enum class ParameterUsage {
   kNone = -1,
   kArrayIndex,
   kBias,
+  kComponent,
   kCoords,
   kDdx,
   kDdy,

src/writer/glsl/generator_impl.cc

@@ -480,7 +480,7 @@ bool GeneratorImpl::EmitIntrinsicCall(std::ostream& out,
                                       const sem::Intrinsic* intrinsic) {
   auto* expr = call->Declaration();
   if (intrinsic->IsTexture()) {
-    return EmitTextureCall(out, expr, intrinsic);
+    return EmitTextureCall(out, call, intrinsic);
   }
   if (intrinsic->Type() == sem::IntrinsicType::kSelect) {
     return EmitSelectCall(out, expr);
@@ -1098,11 +1098,12 @@ bool GeneratorImpl::EmitBarrierCall(std::ostream& out,
 }
 
 bool GeneratorImpl::EmitTextureCall(std::ostream& out,
-                                    const ast::CallExpression* expr,
+                                    const sem::Call* call,
                                     const sem::Intrinsic* intrinsic) {
   using Usage = sem::ParameterUsage;
 
   auto& signature = intrinsic->Signature();
+  auto* expr = call->Declaration();
   auto arguments = expr->args;
 
   // Returns the argument with the given usage
@@ -1184,6 +1185,12 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
     case sem::IntrinsicType::kTextureSampleLevel:
       out << "textureLod(";
       break;
+    case sem::IntrinsicType::kTextureGather:
+    case sem::IntrinsicType::kTextureGatherCompare:
+      out << (intrinsic->Signature().IndexOf(sem::ParameterUsage::kOffset) < 0
+                  ? "textureGather("
+                  : "textureGatherOffset(");
+      break;
     case sem::IntrinsicType::kTextureSampleGrad:
       out << "textureGrad(";
       break;
@@ -1232,9 +1239,9 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
     }
   }
 
-  for (auto usage :
+  for (auto usage : {Usage::kDepthRef, Usage::kBias, Usage::kLevel, Usage::kDdx,
-       {Usage::kDepthRef, Usage::kBias, Usage::kLevel, Usage::kDdx, Usage::kDdy,
+                     Usage::kDdy, Usage::kSampleIndex, Usage::kOffset,
-        Usage::kSampleIndex, Usage::kOffset, Usage::kValue}) {
+                     Usage::kComponent, Usage::kValue}) {
     if (auto* e = arg(usage)) {
       out << ", ";
       if (!EmitExpression(out, e)) {

src/writer/glsl/generator_impl.h

@@ -165,11 +165,11 @@ class GeneratorImpl : public TextGenerator {
   /// Handles generating a call to a texture function (`textureSample`,
   /// `textureSampleGrad`, etc)
   /// @param out the output of the expression stream
-  /// @param expr the call expression
+  /// @param call the call expression
   /// @param intrinsic the semantic information for the texture intrinsic
   /// @returns true if the call expression is emitted
   bool EmitTextureCall(std::ostream& out,
-                       const ast::CallExpression* expr,
+                       const sem::Call* call,
                        const sem::Intrinsic* intrinsic);
   /// Handles generating a call to the `select()` intrinsic
   /// @param out the output of the expression stream

src/writer/glsl/generator_impl_intrinsic_texture_test.cc

@@ -63,6 +63,42 @@ ExpectedResult expected_texture_overload(
     case ValidTextureOverload::kDimensionsStorageWO2dArray:
     case ValidTextureOverload::kDimensionsStorageWO3d:
       return {"imageSize"};
+    case ValidTextureOverload::kGather2dF32:
+      return R"(textureGather(tint_symbol, vec2(1.0f, 2.0f), 0))";
+    case ValidTextureOverload::kGather2dOffsetF32:
+      return R"(textureGatherOffset(tint_symbol, vec2(1.0f, 2.0f), ivec2(3, 4), 0))";
+    case ValidTextureOverload::kGather2dArrayF32:
+      return R"(textureGather(tint_symbol, vec3(1.0f, 2.0f, float(3)), 0))";
+    case ValidTextureOverload::kGather2dArrayOffsetF32:
+      return R"(textureGatherOffset(tint_symbol, vec3(1.0f, 2.0f, float(3)), ivec2(4, 5), 0))";
+    case ValidTextureOverload::kGatherCubeF32:
+      return R"(textureGather(tint_symbol, vec3(1.0f, 2.0f, 3.0f), 0))";
+    case ValidTextureOverload::kGatherCubeArrayF32:
+      return R"(textureGather(tint_symbol, vec4(1.0f, 2.0f, 3.0f, float(4)), 0))";
+    case ValidTextureOverload::kGatherDepth2dF32:
+      return R"(textureGather(tint_symbol, vec2(1.0f, 2.0f)))";
+    case ValidTextureOverload::kGatherDepth2dOffsetF32:
+      return R"(textureGatherOffset(tint_symbol, vec2(1.0f, 2.0f), ivec2(3, 4)))";
+    case ValidTextureOverload::kGatherDepth2dArrayF32:
+      return R"(textureGather(tint_symbol, vec3(1.0f, 2.0f, float(3))))";
+    case ValidTextureOverload::kGatherDepth2dArrayOffsetF32:
+      return R"(textureGatherOffset(tint_symbol, vec3(1.0f, 2.0f, float(3)), ivec2(4, 5)))";
+    case ValidTextureOverload::kGatherDepthCubeF32:
+      return R"(textureGather(tint_symbol, vec3(1.0f, 2.0f, 3.0f)))";
+    case ValidTextureOverload::kGatherDepthCubeArrayF32:
+      return R"(textureGather(tint_symbol, vec4(1.0f, 2.0f, 3.0f, float(4))))";
+    case ValidTextureOverload::kGatherCompareDepth2dF32:
+      return R"(textureGather(tint_symbol, vec2(1.0f, 2.0f), 3.0f))";
+    case ValidTextureOverload::kGatherCompareDepth2dOffsetF32:
+      return R"(textureGatherOffset(tint_symbol, vec2(1.0f, 2.0f), 3.0f, ivec2(4, 5)))";
+    case ValidTextureOverload::kGatherCompareDepth2dArrayF32:
+      return R"(textureGather(tint_symbol, vec3(1.0f, 2.0f, float(3)), 4.0f))";
+    case ValidTextureOverload::kGatherCompareDepth2dArrayOffsetF32:
+      return R"(textureGatherOffset(tint_symbol, vec3(1.0f, 2.0f, float(3)), 4.0f, ivec2(5, 6)))";
+    case ValidTextureOverload::kGatherCompareDepthCubeF32:
+      return R"(textureGather(tint_symbol, vec3(1.0f, 2.0f, 3.0f), 4.0f))";
+    case ValidTextureOverload::kGatherCompareDepthCubeArrayF32:
+      return R"(textureGather(tint_symbol, vec4(1.0f, 2.0f, 3.0f, float(4)), 5.0f))";
     case ValidTextureOverload::kNumLayers2dArray:
     case ValidTextureOverload::kNumLayersDepth2dArray:
     case ValidTextureOverload::kNumLayersCubeArray:

src/writer/hlsl/generator_impl.cc

@@ -671,7 +671,7 @@ bool GeneratorImpl::EmitIntrinsicCall(std::ostream& out,
                                       const sem::Intrinsic* intrinsic) {
   auto* expr = call->Declaration();
   if (intrinsic->IsTexture()) {
-    return EmitTextureCall(out, expr, intrinsic);
+    return EmitTextureCall(out, call, intrinsic);
   }
   if (intrinsic->Type() == sem::IntrinsicType::kSelect) {
     return EmitSelectCall(out, expr);
@@ -1749,11 +1749,12 @@ bool GeneratorImpl::EmitBarrierCall(std::ostream& out,
 }
 
 bool GeneratorImpl::EmitTextureCall(std::ostream& out,
-                                    const ast::CallExpression* expr,
+                                    const sem::Call* call,
                                     const sem::Intrinsic* intrinsic) {
   using Usage = sem::ParameterUsage;
 
   auto& signature = intrinsic->Signature();
+  auto* expr = call->Declaration();
   auto arguments = expr->args;
 
   // Returns the argument with the given usage
@@ -1987,6 +1988,30 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
         pack_level_in_coords = true;
       }
       break;
+    case sem::IntrinsicType::kTextureGather:
+      out << ".Gather";
+      if (intrinsic->Parameters()[0]->Usage() ==
+          sem::ParameterUsage::kComponent) {
+        switch (call->Arguments()[0]->ConstantValue().Elements()[0].i32) {
+          case 0:
+            out << "Red";
+            break;
+          case 1:
+            out << "Green";
+            break;
+          case 2:
+            out << "Blue";
+            break;
+          case 3:
+            out << "Alpha";
+            break;
+        }
+      }
+      out << "(";
+      break;
+    case sem::IntrinsicType::kTextureGatherCompare:
+      out << ".GatherCmp(";
+      break;
     case sem::IntrinsicType::kTextureStore:
       out << "[";
       break;

src/writer/hlsl/generator_impl.h

@@ -208,11 +208,11 @@ class GeneratorImpl : public TextGenerator {
   /// Handles generating a call to a texture function (`textureSample`,
   /// `textureSampleGrad`, etc)
   /// @param out the output of the expression stream
-  /// @param expr the call expression
+  /// @param call the call expression
   /// @param intrinsic the semantic information for the texture intrinsic
   /// @returns true if the call expression is emitted
   bool EmitTextureCall(std::ostream& out,
-                       const ast::CallExpression* expr,
+                       const sem::Call* call,
                        const sem::Intrinsic* intrinsic);
   /// Handles generating a call to the `select()` intrinsic
   /// @param out the output of the expression stream

src/writer/hlsl/generator_impl_intrinsic_texture_test.cc

@@ -135,6 +135,42 @@ ExpectedResult expected_texture_overload(
 )",
           "tint_tmp.xy;",
       };
+    case ValidTextureOverload::kGather2dF32:
+      return R"(tint_symbol.GatherRed(tint_symbol_1, float2(1.0f, 2.0f)))";
+    case ValidTextureOverload::kGather2dOffsetF32:
+      return R"(tint_symbol.GatherRed(tint_symbol_1, float2(1.0f, 2.0f), int2(3, 4)))";
+    case ValidTextureOverload::kGather2dArrayF32:
+      return R"(tint_symbol.GatherRed(tint_symbol_1, float3(1.0f, 2.0f, float(3))))";
+    case ValidTextureOverload::kGather2dArrayOffsetF32:
+      return R"(tint_symbol.GatherRed(tint_symbol_1, float3(1.0f, 2.0f, float(3)), int2(4, 5)))";
+    case ValidTextureOverload::kGatherCubeF32:
+      return R"(tint_symbol.GatherRed(tint_symbol_1, float3(1.0f, 2.0f, 3.0f)))";
+    case ValidTextureOverload::kGatherCubeArrayF32:
+      return R"(tint_symbol.GatherRed(tint_symbol_1, float4(1.0f, 2.0f, 3.0f, float(4))))";
+    case ValidTextureOverload::kGatherDepth2dF32:
+      return R"(tint_symbol.Gather(tint_symbol_1, float2(1.0f, 2.0f)))";
+    case ValidTextureOverload::kGatherDepth2dOffsetF32:
+      return R"(tint_symbol.Gather(tint_symbol_1, float2(1.0f, 2.0f), int2(3, 4)))";
+    case ValidTextureOverload::kGatherDepth2dArrayF32:
+      return R"(tint_symbol.Gather(tint_symbol_1, float3(1.0f, 2.0f, float(3))))";
+    case ValidTextureOverload::kGatherDepth2dArrayOffsetF32:
+      return R"(tint_symbol.Gather(tint_symbol_1, float3(1.0f, 2.0f, float(3)), int2(4, 5)))";
+    case ValidTextureOverload::kGatherDepthCubeF32:
+      return R"(tint_symbol.Gather(tint_symbol_1, float3(1.0f, 2.0f, 3.0f)))";
+    case ValidTextureOverload::kGatherDepthCubeArrayF32:
+      return R"(tint_symbol.Gather(tint_symbol_1, float4(1.0f, 2.0f, 3.0f, float(4))))";
+    case ValidTextureOverload::kGatherCompareDepth2dF32:
+      return R"(tint_symbol.GatherCmp(tint_symbol_1, float2(1.0f, 2.0f), 3.0f))";
+    case ValidTextureOverload::kGatherCompareDepth2dOffsetF32:
+      return R"(tint_symbol.GatherCmp(tint_symbol_1, float2(1.0f, 2.0f), 3.0f, int2(4, 5)))";
+    case ValidTextureOverload::kGatherCompareDepth2dArrayF32:
+      return R"(tint_symbol.GatherCmp(tint_symbol_1, float3(1.0f, 2.0f, float(3)), 4.0f))";
+    case ValidTextureOverload::kGatherCompareDepth2dArrayOffsetF32:
+      return R"(tint_symbol.GatherCmp(tint_symbol_1, float3(1.0f, 2.0f, float(3)), 4.0f, int2(5, 6)))";
+    case ValidTextureOverload::kGatherCompareDepthCubeF32:
+      return R"(tint_symbol.GatherCmp(tint_symbol_1, float3(1.0f, 2.0f, 3.0f), 4.0f))";
+    case ValidTextureOverload::kGatherCompareDepthCubeArrayF32:
+      return R"(tint_symbol.GatherCmp(tint_symbol_1, float4(1.0f, 2.0f, 3.0f, float(4)), 5.0f))";
     case ValidTextureOverload::kNumLayers2dArray:
     case ValidTextureOverload::kNumLayersDepth2dArray:
     case ValidTextureOverload::kNumLayersCubeArray:

src/writer/msl/generator_impl.cc

@@ -582,7 +582,7 @@ bool GeneratorImpl::EmitIntrinsicCall(std::ostream& out,
     return EmitAtomicCall(out, expr, intrinsic);
   }
   if (intrinsic->IsTexture()) {
-    return EmitTextureCall(out, expr, intrinsic);
+    return EmitTextureCall(out, call, intrinsic);
   }
 
   auto name = generate_builtin_name(intrinsic);
@@ -839,12 +839,13 @@ bool GeneratorImpl::EmitAtomicCall(std::ostream& out,
 }
 
 bool GeneratorImpl::EmitTextureCall(std::ostream& out,
-                                    const ast::CallExpression* expr,
+                                    const sem::Call* call,
                                     const sem::Intrinsic* intrinsic) {
   using Usage = sem::ParameterUsage;
 
   auto& signature = intrinsic->Signature();
-  auto arguments = expr->args;
+  auto* expr = call->Declaration();
+  auto& arguments = call->Arguments();
 
   // Returns the argument with the given usage
   auto arg = [&](Usage usage) {
@@ -852,7 +853,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
     return (idx >= 0) ? arguments[idx] : nullptr;
   };
 
-  auto* texture = arg(Usage::kTexture);
+  auto* texture = arg(Usage::kTexture)->Declaration();
   if (!texture) {
     TINT_ICE(Writer, diagnostics_) << "missing texture arg";
     return false;
@@ -908,7 +909,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
         }
         out << ".get_" << name << "(";
         if (auto* level = arg(Usage::kLevel)) {
-          if (!EmitExpression(out, level)) {
+          if (!EmitExpression(out, level->Declaration())) {
             return false;
           }
         }
@@ -978,6 +979,12 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
     case sem::IntrinsicType::kTextureSampleCompareLevel:
       out << ".sample_compare(";
       break;
+    case sem::IntrinsicType::kTextureGather:
+      out << ".gather(";
+      break;
+    case sem::IntrinsicType::kTextureGatherCompare:
+      out << ".gather_compare(";
+      break;
     case sem::IntrinsicType::kTextureLoad:
       out << ".read(";
       lod_param_is_named = false;
@@ -1003,14 +1010,12 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
        {Usage::kValue, Usage::kSampler, Usage::kCoords, Usage::kArrayIndex,
         Usage::kDepthRef, Usage::kSampleIndex}) {
     if (auto* e = arg(usage)) {
-      auto* sem_e = program_->Sem().Get(e);
-
       maybe_write_comma();
 
       // Cast the coordinates to unsigned integers if necessary.
       bool casted = false;
       if (usage == Usage::kCoords &&
-          sem_e->Type()->UnwrapRef()->is_integer_scalar_or_vector()) {
+          e->Type()->UnwrapRef()->is_integer_scalar_or_vector()) {
         casted = true;
         switch (texture_type->dim()) {
           case ast::TextureDimension::k1d:
@@ -1030,7 +1035,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
         }
       }
 
-      if (!EmitExpression(out, e))
+      if (!EmitExpression(out, e->Declaration()))
         return false;
 
       if (casted) {
@@ -1042,7 +1047,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
   if (auto* bias = arg(Usage::kBias)) {
     maybe_write_comma();
     out << "bias(";
-    if (!EmitExpression(out, bias)) {
+    if (!EmitExpression(out, bias->Declaration())) {
       return false;
     }
     out << ")";
@@ -1052,7 +1057,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
     if (lod_param_is_named) {
       out << "level(";
     }
-    if (!EmitExpression(out, level)) {
+    if (!EmitExpression(out, level->Declaration())) {
       return false;
     }
     if (lod_param_is_named) {
@@ -1087,23 +1092,59 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
         return false;
       }
     }
-    if (!EmitExpression(out, ddx)) {
+    if (!EmitExpression(out, ddx->Declaration())) {
       return false;
     }
     out << ", ";
-    if (!EmitExpression(out, arg(Usage::kDdy))) {
+    if (!EmitExpression(out, arg(Usage::kDdy)->Declaration())) {
       return false;
     }
     out << ")";
   }
 
+  bool has_offset = false;
   if (auto* offset = arg(Usage::kOffset)) {
+    has_offset = true;
     maybe_write_comma();
-    if (!EmitExpression(out, offset)) {
+    if (!EmitExpression(out, offset->Declaration())) {
       return false;
     }
   }
 
+  if (auto* component = arg(Usage::kComponent)) {
+    maybe_write_comma();
+    if (!has_offset) {
+      // offset argument may need to be provided if we have a component.
+      switch (texture_type->dim()) {
+        case ast::TextureDimension::k2d:
+        case ast::TextureDimension::k2dArray:
+          out << "int2(0), ";
+          break;
+        default:
+          break;  // Other texture dimensions don't have an offset
+      }
+    }
+    auto c = component->ConstantValue().Elements()[0].i32;
+    switch (c) {
+      case 0:
+        out << "component::x";
+        break;
+      case 1:
+        out << "component::y";
+        break;
+      case 2:
+        out << "component::z";
+        break;
+      case 3:
+        out << "component::w";
+        break;
+      default:
+        TINT_ICE(Writer, diagnostics_)
+            << "invalid textureGather component: " << c;
+        break;
+    }
+  }
+
   out << ")";
 
   return true;

src/writer/msl/generator_impl.h

@@ -187,11 +187,11 @@ class GeneratorImpl : public TextGenerator {
   /// Handles generating a call to a texture function (`textureSample`,
   /// `textureSampleGrad`, etc)
   /// @param out the output of the expression stream
-  /// @param expr the call expression
+  /// @param call the call expression
   /// @param intrinsic the semantic information for the texture intrinsic
   /// @returns true if the call expression is emitted
   bool EmitTextureCall(std::ostream& out,
-                       const ast::CallExpression* expr,
+                       const sem::Call* call,
                        const sem::Intrinsic* intrinsic);
   /// Handles generating a call to the `dot()` intrinsic
   /// @param out the output of the expression stream

src/writer/msl/generator_impl_intrinsic_texture_test.cc

@@ -55,6 +55,42 @@ std::string expected_texture_overload(
       return R"(int2(texture.get_width(1), texture.get_height(1)))";
     case ValidTextureOverload::kDimensions3dLevel:
       return R"(int3(texture.get_width(1), texture.get_height(1), texture.get_depth(1)))";
+    case ValidTextureOverload::kGather2dF32:
+      return R"(texture.gather(sampler, float2(1.0f, 2.0f), int2(0), component::x))";
+    case ValidTextureOverload::kGather2dOffsetF32:
+      return R"(texture.gather(sampler, float2(1.0f, 2.0f), int2(3, 4), component::x))";
+    case ValidTextureOverload::kGather2dArrayF32:
+      return R"(texture.gather(sampler, float2(1.0f, 2.0f), 3, int2(0), component::x))";
+    case ValidTextureOverload::kGather2dArrayOffsetF32:
+      return R"(texture.gather(sampler, float2(1.0f, 2.0f), 3, int2(4, 5), component::x))";
+    case ValidTextureOverload::kGatherCubeF32:
+      return R"(texture.gather(sampler, float3(1.0f, 2.0f, 3.0f), component::x))";
+    case ValidTextureOverload::kGatherCubeArrayF32:
+      return R"(texture.gather(sampler, float3(1.0f, 2.0f, 3.0f), 4, component::x))";
+    case ValidTextureOverload::kGatherDepth2dF32:
+      return R"(texture.gather(sampler, float2(1.0f, 2.0f)))";
+    case ValidTextureOverload::kGatherDepth2dOffsetF32:
+      return R"(texture.gather(sampler, float2(1.0f, 2.0f), int2(3, 4)))";
+    case ValidTextureOverload::kGatherDepth2dArrayF32:
+      return R"(texture.gather(sampler, float2(1.0f, 2.0f), 3))";
+    case ValidTextureOverload::kGatherDepth2dArrayOffsetF32:
+      return R"(texture.gather(sampler, float2(1.0f, 2.0f), 3, int2(4, 5)))";
+    case ValidTextureOverload::kGatherDepthCubeF32:
+      return R"(texture.gather(sampler, float3(1.0f, 2.0f, 3.0f)))";
+    case ValidTextureOverload::kGatherDepthCubeArrayF32:
+      return R"(texture.gather(sampler, float3(1.0f, 2.0f, 3.0f), 4))";
+    case ValidTextureOverload::kGatherCompareDepth2dF32:
+      return R"(texture.gather_compare(sampler, float2(1.0f, 2.0f), 3.0f))";
+    case ValidTextureOverload::kGatherCompareDepth2dOffsetF32:
+      return R"(texture.gather_compare(sampler, float2(1.0f, 2.0f), 3.0f, int2(4, 5)))";
+    case ValidTextureOverload::kGatherCompareDepth2dArrayF32:
+      return R"(texture.gather_compare(sampler, float2(1.0f, 2.0f), 3, 4.0f))";
+    case ValidTextureOverload::kGatherCompareDepth2dArrayOffsetF32:
+      return R"(texture.gather_compare(sampler, float2(1.0f, 2.0f), 3, 4.0f, int2(5, 6)))";
+    case ValidTextureOverload::kGatherCompareDepthCubeF32:
+      return R"(texture.gather_compare(sampler, float3(1.0f, 2.0f, 3.0f), 4.0f))";
+    case ValidTextureOverload::kGatherCompareDepthCubeArrayF32:
+      return R"(texture.gather_compare(sampler, float3(1.0f, 2.0f, 3.0f), 4, 5.0f))";
     case ValidTextureOverload::kNumLayers2dArray:
     case ValidTextureOverload::kNumLayersCubeArray:
     case ValidTextureOverload::kNumLayersDepth2dArray:

src/writer/spirv/builder.cc

@@ -2987,6 +2987,29 @@ bool Builder::GenerateTextureIntrinsic(const sem::Call* call,
       spirv_params.emplace_back(gen_arg(Usage::kValue));
       break;
     }
+    case IntrinsicType::kTextureGather: {
+      op = spv::Op::OpImageGather;
+      append_result_type_and_id_to_spirv_params();
+      if (!append_image_and_coords_to_spirv_params()) {
+        return false;
+      }
+      if (signature.IndexOf(Usage::kComponent) < 0) {
+        spirv_params.emplace_back(
+            Operand::Int(GenerateConstantIfNeeded(ScalarConstant::I32(0))));
+      } else {
+        spirv_params.emplace_back(gen_arg(Usage::kComponent));
+      }
+      break;
+    }
+    case IntrinsicType::kTextureGatherCompare: {
+      op = spv::Op::OpImageDrefGather;
+      append_result_type_and_id_to_spirv_params();
+      if (!append_image_and_coords_to_spirv_params()) {
+        return false;
+      }
+      spirv_params.emplace_back(gen_arg(Usage::kDepthRef));
+      break;
+    }
     case IntrinsicType::kTextureSample: {
       op = spv::Op::OpImageSampleImplicitLod;
       append_result_type_and_id_to_spirv_params_for_read();

src/writer/spirv/builder_intrinsic_texture_test.cc

@@ -574,6 +574,520 @@ OpCapability ImageQuery
 )",
           R"(
 OpCapability ImageQuery
+)"};
+    case ValidTextureOverload::kGather2dF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 2D 0 0 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 2
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%17 = OpConstantComposite %14 %15 %16
+%18 = OpTypeInt 32 1
+%19 = OpConstant %18 0
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%8 = OpImageGather %9 %13 %17 %19
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGather2dOffsetF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 2D 0 0 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 2
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%17 = OpConstantComposite %14 %15 %16
+%18 = OpTypeInt 32 1
+%19 = OpConstant %18 0
+%20 = OpTypeVector %18 2
+%21 = OpConstant %18 3
+%22 = OpConstant %18 4
+%23 = OpConstantComposite %20 %21 %22
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%8 = OpImageGather %9 %13 %17 %19 ConstOffset %23
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGather2dArrayF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 2D 0 1 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 3
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%18 = OpTypeInt 32 1
+%19 = OpConstant %18 3
+%21 = OpConstant %18 0
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%17 = OpConvertSToF %4 %19
+%20 = OpCompositeConstruct %14 %15 %16 %17
+%8 = OpImageGather %9 %13 %20 %21
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGather2dArrayOffsetF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 2D 0 1 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 3
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%18 = OpTypeInt 32 1
+%19 = OpConstant %18 3
+%21 = OpConstant %18 0
+%22 = OpTypeVector %18 2
+%23 = OpConstant %18 4
+%24 = OpConstant %18 5
+%25 = OpConstantComposite %22 %23 %24
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%17 = OpConvertSToF %4 %19
+%20 = OpCompositeConstruct %14 %15 %16 %17
+%8 = OpImageGather %9 %13 %20 %21 ConstOffset %25
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGatherCubeF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 Cube 0 0 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 3
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%17 = OpConstant %4 3
+%18 = OpConstantComposite %14 %15 %16 %17
+%19 = OpTypeInt 32 1
+%20 = OpConstant %19 0
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%8 = OpImageGather %9 %13 %18 %20
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGatherCubeArrayF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 Cube 0 1 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpConstant %4 1
+%15 = OpConstant %4 2
+%16 = OpConstant %4 3
+%18 = OpTypeInt 32 1
+%19 = OpConstant %18 4
+%21 = OpConstant %18 0
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%17 = OpConvertSToF %4 %19
+%20 = OpCompositeConstruct %9 %14 %15 %16 %17
+%8 = OpImageGather %9 %13 %20 %21
+)",
+              R"(
+OpCapability SampledCubeArray
+)"};
+    case ValidTextureOverload::kGatherDepth2dF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 2D 1 0 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 2
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%17 = OpConstantComposite %14 %15 %16
+%18 = OpTypeInt 32 1
+%19 = OpConstant %18 0
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%8 = OpImageGather %9 %13 %17 %19
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGatherDepth2dOffsetF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 2D 1 0 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 2
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%17 = OpConstantComposite %14 %15 %16
+%18 = OpTypeInt 32 1
+%19 = OpConstant %18 0
+%20 = OpTypeVector %18 2
+%21 = OpConstant %18 3
+%22 = OpConstant %18 4
+%23 = OpConstantComposite %20 %21 %22
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%8 = OpImageGather %9 %13 %17 %19 ConstOffset %23
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGatherDepth2dArrayF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 2D 1 1 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 3
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%18 = OpTypeInt 32 1
+%19 = OpConstant %18 3
+%21 = OpConstant %18 0
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%17 = OpConvertSToF %4 %19
+%20 = OpCompositeConstruct %14 %15 %16 %17
+%8 = OpImageGather %9 %13 %20 %21
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGatherDepth2dArrayOffsetF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 2D 1 1 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 3
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%18 = OpTypeInt 32 1
+%19 = OpConstant %18 3
+%21 = OpConstant %18 0
+%22 = OpTypeVector %18 2
+%23 = OpConstant %18 4
+%24 = OpConstant %18 5
+%25 = OpConstantComposite %22 %23 %24
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%17 = OpConvertSToF %4 %19
+%20 = OpCompositeConstruct %14 %15 %16 %17
+%8 = OpImageGather %9 %13 %20 %21 ConstOffset %25
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGatherDepthCubeF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 Cube 1 0 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 3
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%17 = OpConstant %4 3
+%18 = OpConstantComposite %14 %15 %16 %17
+%19 = OpTypeInt 32 1
+%20 = OpConstant %19 0
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%8 = OpImageGather %9 %13 %18 %20
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGatherDepthCubeArrayF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 Cube 1 1 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpConstant %4 1
+%15 = OpConstant %4 2
+%16 = OpConstant %4 3
+%18 = OpTypeInt 32 1
+%19 = OpConstant %18 4
+%21 = OpConstant %18 0
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%17 = OpConvertSToF %4 %19
+%20 = OpCompositeConstruct %9 %14 %15 %16 %17
+%8 = OpImageGather %9 %13 %20 %21
+)",
+              R"(
+OpCapability SampledCubeArray
+)"};
+    case ValidTextureOverload::kGatherCompareDepth2dF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 2D 1 0 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 2
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%17 = OpConstantComposite %14 %15 %16
+%18 = OpConstant %4 3
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%8 = OpImageDrefGather %9 %13 %17 %18
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGatherCompareDepth2dOffsetF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 2D 1 0 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 2
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%17 = OpConstantComposite %14 %15 %16
+%18 = OpConstant %4 3
+%20 = OpTypeInt 32 1
+%19 = OpTypeVector %20 2
+%21 = OpConstant %20 4
+%22 = OpConstant %20 5
+%23 = OpConstantComposite %19 %21 %22
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%8 = OpImageDrefGather %9 %13 %17 %18 ConstOffset %23
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGatherCompareDepth2dArrayF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 2D 1 1 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 3
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%18 = OpTypeInt 32 1
+%19 = OpConstant %18 3
+%21 = OpConstant %4 4
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%17 = OpConvertSToF %4 %19
+%20 = OpCompositeConstruct %14 %15 %16 %17
+%8 = OpImageDrefGather %9 %13 %20 %21
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGatherCompareDepth2dArrayOffsetF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 2D 1 1 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 3
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%18 = OpTypeInt 32 1
+%19 = OpConstant %18 3
+%21 = OpConstant %4 4
+%22 = OpTypeVector %18 2
+%23 = OpConstant %18 5
+%24 = OpConstant %18 6
+%25 = OpConstantComposite %22 %23 %24
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%17 = OpConvertSToF %4 %19
+%20 = OpCompositeConstruct %14 %15 %16 %17
+%8 = OpImageDrefGather %9 %13 %20 %21 ConstOffset %25
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGatherCompareDepthCubeF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 Cube 1 0 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpTypeVector %4 3
+%15 = OpConstant %4 1
+%16 = OpConstant %4 2
+%17 = OpConstant %4 3
+%18 = OpConstantComposite %14 %15 %16 %17
+%19 = OpConstant %4 4
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%8 = OpImageDrefGather %9 %13 %18 %19
+)",
+              R"(
+)"};
+    case ValidTextureOverload::kGatherCompareDepthCubeArrayF32:
+      return {R"(
+%4 = OpTypeFloat 32
+%3 = OpTypeImage %4 Cube 1 1 0 1 Unknown
+%2 = OpTypePointer UniformConstant %3
+%1 = OpVariable %2 UniformConstant
+%7 = OpTypeSampler
+%6 = OpTypePointer UniformConstant %7
+%5 = OpVariable %6 UniformConstant
+%9 = OpTypeVector %4 4
+%12 = OpTypeSampledImage %3
+%14 = OpConstant %4 1
+%15 = OpConstant %4 2
+%16 = OpConstant %4 3
+%18 = OpTypeInt 32 1
+%19 = OpConstant %18 4
+%21 = OpConstant %4 5
+)",
+              R"(
+%10 = OpLoad %7 %5
+%11 = OpLoad %3 %1
+%13 = OpSampledImage %12 %11 %10
+%17 = OpConvertSToF %4 %19
+%20 = OpCompositeConstruct %9 %14 %15 %16 %17
+%8 = OpImageDrefGather %9 %13 %20 %21
+)",
+              R"(
+OpCapability SampledCubeArray
 )"};
     case ValidTextureOverload::kNumLayers2dArray:
       return {R"(

test/intrinsics/gen/textureGather/01305f.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(1)]] var arg_1: texture_2d_array<u32>;
+[[group(1), binding(2)]] var arg_2: sampler;
+
+// fn textureGather(component: i32, texture: texture_2d_array<u32>, sampler: sampler, coords: vec2<f32>, array_index: i32) -> vec4<u32>
+fn textureGather_01305f() {
+  var res: vec4<u32> = textureGather(1, arg_1, arg_2, vec2<f32>(), 1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_01305f();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_01305f();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_01305f();
+}

test/intrinsics/gen/textureGather/01305f.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+Texture2DArray<uint4> arg_1 : register(t1, space1);
+SamplerState arg_2 : register(s2, space1);
+
+void textureGather_01305f() {
+  uint4 res = arg_1.GatherGreen(arg_2, float3(0.0f, 0.0f, float(1)));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_01305f();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_01305f();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_01305f();
+  return;
+}

test/intrinsics/gen/textureGather/01305f.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_01305f(texture2d_array<uint, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  uint4 res = tint_symbol_1.gather(tint_symbol_2, float2(), 1, int2(0), component::y);
+}
+
+float4 vertex_main_inner(texture2d_array<uint, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_01305f(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(texture2d_array<uint, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texture2d_array<uint, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_01305f(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(texture2d_array<uint, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_01305f(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/01305f.wgsl.expected.spvasm

@@ -0,0 +1,89 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 49
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_1 "arg_1"
+               OpName %arg_2 "arg_2"
+               OpName %textureGather_01305f "textureGather_01305f"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+               OpDecorate %arg_2 DescriptorSet 1
+               OpDecorate %arg_2 Binding 2
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %uint = OpTypeInt 32 0
+         %11 = OpTypeImage %uint 2D 0 1 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %15 = OpTypeSampler
+%_ptr_UniformConstant_15 = OpTypePointer UniformConstant %15
+      %arg_2 = OpVariable %_ptr_UniformConstant_15 UniformConstant
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+     %v4uint = OpTypeVector %uint 4
+         %24 = OpTypeSampledImage %11
+    %v3float = OpTypeVector %float 3
+    %float_0 = OpConstant %float 0
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+         %34 = OpConstantNull %v4uint
+         %35 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_01305f = OpFunction %void None %16
+         %19 = OpLabel
+        %res = OpVariable %_ptr_Function_v4uint Function %34
+         %22 = OpLoad %15 %arg_2
+         %23 = OpLoad %11 %arg_1
+         %25 = OpSampledImage %24 %23 %22
+         %28 = OpConvertSToF %float %int_1
+         %31 = OpCompositeConstruct %v3float %float_0 %float_0 %28
+         %20 = OpImageGather %v4uint %25 %31 %int_1
+               OpStore %res %20
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %35
+         %37 = OpLabel
+         %38 = OpFunctionCall %void %textureGather_01305f
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %16
+         %40 = OpLabel
+         %41 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %41
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %16
+         %44 = OpLabel
+         %45 = OpFunctionCall %void %textureGather_01305f
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %16
+         %47 = OpLabel
+         %48 = OpFunctionCall %void %textureGather_01305f
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/01305f.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(1)]] var arg_1 : texture_2d_array<u32>;
+
+[[group(1), binding(2)]] var arg_2 : sampler;
+
+fn textureGather_01305f() {
+  var res : vec4<u32> = textureGather(1, arg_1, arg_2, vec2<f32>(), 1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_01305f();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_01305f();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_01305f();
+}

test/intrinsics/gen/textureGather/06030a.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(1)]] var arg_1: texture_2d_array<f32>;
+[[group(1), binding(2)]] var arg_2: sampler;
+
+// fn textureGather(component: i32, texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<f32>
+fn textureGather_06030a() {
+  var res: vec4<f32> = textureGather(1, arg_1, arg_2, vec2<f32>(), 1, vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_06030a();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_06030a();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_06030a();
+}

test/intrinsics/gen/textureGather/06030a.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+Texture2DArray<float4> arg_1 : register(t1, space1);
+SamplerState arg_2 : register(s2, space1);
+
+void textureGather_06030a() {
+  float4 res = arg_1.GatherGreen(arg_2, float3(0.0f, 0.0f, float(1)), int2(0, 0));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_06030a();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_06030a();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_06030a();
+  return;
+}

test/intrinsics/gen/textureGather/06030a.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_06030a(texture2d_array<float, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  float4 res = tint_symbol_1.gather(tint_symbol_2, float2(), 1, int2(), component::y);
+}
+
+float4 vertex_main_inner(texture2d_array<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_06030a(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(texture2d_array<float, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texture2d_array<float, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_06030a(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(texture2d_array<float, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_06030a(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/06030a.wgsl.expected.spvasm

@@ -0,0 +1,88 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_1 "arg_1"
+               OpName %arg_2 "arg_2"
+               OpName %textureGather_06030a "textureGather_06030a"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+               OpDecorate %arg_2 DescriptorSet 1
+               OpDecorate %arg_2 Binding 2
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+         %11 = OpTypeImage %float 2D 0 1 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %14 = OpTypeSampler
+%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
+      %arg_2 = OpVariable %_ptr_UniformConstant_14 UniformConstant
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+         %22 = OpTypeSampledImage %11
+    %v3float = OpTypeVector %float 3
+    %float_0 = OpConstant %float 0
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+      %v2int = OpTypeVector %int 2
+         %31 = OpConstantNull %v2int
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %34 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_06030a = OpFunction %void None %15
+         %18 = OpLabel
+        %res = OpVariable %_ptr_Function_v4float Function %5
+         %20 = OpLoad %14 %arg_2
+         %21 = OpLoad %11 %arg_1
+         %23 = OpSampledImage %22 %21 %20
+         %26 = OpConvertSToF %float %int_1
+         %29 = OpCompositeConstruct %v3float %float_0 %float_0 %26
+         %19 = OpImageGather %v4float %23 %29 %int_1 ConstOffset %31
+               OpStore %res %19
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %textureGather_06030a
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %15
+         %39 = OpLabel
+         %40 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %40
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %15
+         %43 = OpLabel
+         %44 = OpFunctionCall %void %textureGather_06030a
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %15
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %textureGather_06030a
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/06030a.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(1)]] var arg_1 : texture_2d_array<f32>;
+
+[[group(1), binding(2)]] var arg_2 : sampler;
+
+fn textureGather_06030a() {
+  var res : vec4<f32> = textureGather(1, arg_1, arg_2, vec2<f32>(), 1, vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_06030a();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_06030a();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_06030a();
+}

test/intrinsics/gen/textureGather/10c554.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(0)]] var arg_0: texture_depth_cube;
+[[group(1), binding(1)]] var arg_1: sampler;
+
+// fn textureGather(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
+fn textureGather_10c554() {
+  var res: vec4<f32> = textureGather(arg_0, arg_1, vec3<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_10c554();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_10c554();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_10c554();
+}

test/intrinsics/gen/textureGather/10c554.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+TextureCube arg_0 : register(t0, space1);
+SamplerState arg_1 : register(s1, space1);
+
+void textureGather_10c554() {
+  float4 res = arg_0.Gather(arg_1, float3(0.0f, 0.0f, 0.0f));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_10c554();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_10c554();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_10c554();
+  return;
+}

test/intrinsics/gen/textureGather/10c554.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_10c554(depthcube<float, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  float4 res = tint_symbol_1.gather(tint_symbol_2, float3());
+}
+
+float4 vertex_main_inner(depthcube<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_10c554(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(depthcube<float, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(depthcube<float, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_10c554(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(depthcube<float, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_10c554(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/10c554.wgsl.expected.spvasm

@@ -0,0 +1,84 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 44
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_0 "arg_0"
+               OpName %arg_1 "arg_1"
+               OpName %textureGather_10c554 "textureGather_10c554"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+         %11 = OpTypeImage %float Cube 1 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %14 = OpTypeSampler
+%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
+      %arg_1 = OpVariable %_ptr_UniformConstant_14 UniformConstant
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+         %22 = OpTypeSampledImage %11
+    %v3float = OpTypeVector %float 3
+         %25 = OpConstantNull %v3float
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %30 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_10c554 = OpFunction %void None %15
+         %18 = OpLabel
+        %res = OpVariable %_ptr_Function_v4float Function %5
+         %20 = OpLoad %14 %arg_1
+         %21 = OpLoad %11 %arg_0
+         %23 = OpSampledImage %22 %21 %20
+         %19 = OpImageGather %v4float %23 %25 %int_0
+               OpStore %res %19
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %30
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %textureGather_10c554
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %15
+         %35 = OpLabel
+         %36 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %36
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %15
+         %39 = OpLabel
+         %40 = OpFunctionCall %void %textureGather_10c554
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %15
+         %42 = OpLabel
+         %43 = OpFunctionCall %void %textureGather_10c554
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/10c554.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(0)]] var arg_0 : texture_depth_cube;
+
+[[group(1), binding(1)]] var arg_1 : sampler;
+
+fn textureGather_10c554() {
+  var res : vec4<f32> = textureGather(arg_0, arg_1, vec3<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_10c554();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_10c554();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_10c554();
+}

test/intrinsics/gen/textureGather/15d79c.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(1)]] var arg_1: texture_2d<f32>;
+[[group(1), binding(2)]] var arg_2: sampler;
+
+// fn textureGather(component: i32, texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<f32>
+fn textureGather_15d79c() {
+  var res: vec4<f32> = textureGather(1, arg_1, arg_2, vec2<f32>(), vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_15d79c();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_15d79c();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_15d79c();
+}

test/intrinsics/gen/textureGather/15d79c.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+Texture2D<float4> arg_1 : register(t1, space1);
+SamplerState arg_2 : register(s2, space1);
+
+void textureGather_15d79c() {
+  float4 res = arg_1.GatherGreen(arg_2, float2(0.0f, 0.0f), int2(0, 0));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_15d79c();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_15d79c();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_15d79c();
+  return;
+}

test/intrinsics/gen/textureGather/15d79c.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_15d79c(texture2d<float, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  float4 res = tint_symbol_1.gather(tint_symbol_2, float2(), int2(), component::y);
+}
+
+float4 vertex_main_inner(texture2d<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_15d79c(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(texture2d<float, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texture2d<float, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_15d79c(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(texture2d<float, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_15d79c(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/15d79c.wgsl.expected.spvasm

@@ -0,0 +1,86 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 46
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_1 "arg_1"
+               OpName %arg_2 "arg_2"
+               OpName %textureGather_15d79c "textureGather_15d79c"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+               OpDecorate %arg_2 DescriptorSet 1
+               OpDecorate %arg_2 Binding 2
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+         %11 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %14 = OpTypeSampler
+%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
+      %arg_2 = OpVariable %_ptr_UniformConstant_14 UniformConstant
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+         %22 = OpTypeSampledImage %11
+    %v2float = OpTypeVector %float 2
+         %25 = OpConstantNull %v2float
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+      %v2int = OpTypeVector %int 2
+         %29 = OpConstantNull %v2int
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %32 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_15d79c = OpFunction %void None %15
+         %18 = OpLabel
+        %res = OpVariable %_ptr_Function_v4float Function %5
+         %20 = OpLoad %14 %arg_2
+         %21 = OpLoad %11 %arg_1
+         %23 = OpSampledImage %22 %21 %20
+         %19 = OpImageGather %v4float %23 %25 %int_1 ConstOffset %29
+               OpStore %res %19
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %32
+         %34 = OpLabel
+         %35 = OpFunctionCall %void %textureGather_15d79c
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %15
+         %37 = OpLabel
+         %38 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %38
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %15
+         %41 = OpLabel
+         %42 = OpFunctionCall %void %textureGather_15d79c
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %15
+         %44 = OpLabel
+         %45 = OpFunctionCall %void %textureGather_15d79c
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/15d79c.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(1)]] var arg_1 : texture_2d<f32>;
+
+[[group(1), binding(2)]] var arg_2 : sampler;
+
+fn textureGather_15d79c() {
+  var res : vec4<f32> = textureGather(1, arg_1, arg_2, vec2<f32>(), vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_15d79c();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_15d79c();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_15d79c();
+}

test/intrinsics/gen/textureGather/2e0ed5.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(0)]] var arg_0: texture_depth_2d;
+[[group(1), binding(1)]] var arg_1: sampler;
+
+// fn textureGather(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
+fn textureGather_2e0ed5() {
+  var res: vec4<f32> = textureGather(arg_0, arg_1, vec2<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_2e0ed5();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_2e0ed5();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_2e0ed5();
+}

test/intrinsics/gen/textureGather/2e0ed5.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+Texture2D arg_0 : register(t0, space1);
+SamplerState arg_1 : register(s1, space1);
+
+void textureGather_2e0ed5() {
+  float4 res = arg_0.Gather(arg_1, float2(0.0f, 0.0f));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_2e0ed5();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_2e0ed5();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_2e0ed5();
+  return;
+}

test/intrinsics/gen/textureGather/2e0ed5.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_2e0ed5(depth2d<float, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  float4 res = tint_symbol_1.gather(tint_symbol_2, float2());
+}
+
+float4 vertex_main_inner(depth2d<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_2e0ed5(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(depth2d<float, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(depth2d<float, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_2e0ed5(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(depth2d<float, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_2e0ed5(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/2e0ed5.wgsl.expected.spvasm

@@ -0,0 +1,84 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 44
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_0 "arg_0"
+               OpName %arg_1 "arg_1"
+               OpName %textureGather_2e0ed5 "textureGather_2e0ed5"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+         %11 = OpTypeImage %float 2D 1 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %14 = OpTypeSampler
+%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
+      %arg_1 = OpVariable %_ptr_UniformConstant_14 UniformConstant
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+         %22 = OpTypeSampledImage %11
+    %v2float = OpTypeVector %float 2
+         %25 = OpConstantNull %v2float
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %30 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_2e0ed5 = OpFunction %void None %15
+         %18 = OpLabel
+        %res = OpVariable %_ptr_Function_v4float Function %5
+         %20 = OpLoad %14 %arg_1
+         %21 = OpLoad %11 %arg_0
+         %23 = OpSampledImage %22 %21 %20
+         %19 = OpImageGather %v4float %23 %25 %int_0
+               OpStore %res %19
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %30
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %textureGather_2e0ed5
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %15
+         %35 = OpLabel
+         %36 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %36
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %15
+         %39 = OpLabel
+         %40 = OpFunctionCall %void %textureGather_2e0ed5
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %15
+         %42 = OpLabel
+         %43 = OpFunctionCall %void %textureGather_2e0ed5
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/2e0ed5.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(0)]] var arg_0 : texture_depth_2d;
+
+[[group(1), binding(1)]] var arg_1 : sampler;
+
+fn textureGather_2e0ed5() {
+  var res : vec4<f32> = textureGather(arg_0, arg_1, vec2<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_2e0ed5();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_2e0ed5();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_2e0ed5();
+}

test/intrinsics/gen/textureGather/3112e8.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(1)]] var arg_1: texture_cube_array<f32>;
+[[group(1), binding(2)]] var arg_2: sampler;
+
+// fn textureGather(component: i32, texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: i32) -> vec4<f32>
+fn textureGather_3112e8() {
+  var res: vec4<f32> = textureGather(1, arg_1, arg_2, vec3<f32>(), 1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_3112e8();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_3112e8();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_3112e8();
+}

test/intrinsics/gen/textureGather/3112e8.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+TextureCubeArray<float4> arg_1 : register(t1, space1);
+SamplerState arg_2 : register(s2, space1);
+
+void textureGather_3112e8() {
+  float4 res = arg_1.GatherGreen(arg_2, float4(0.0f, 0.0f, 0.0f, float(1)));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_3112e8();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_3112e8();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_3112e8();
+  return;
+}

test/intrinsics/gen/textureGather/3112e8.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_3112e8(texturecube_array<float, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  float4 res = tint_symbol_1.gather(tint_symbol_2, float3(), 1, component::y);
+}
+
+float4 vertex_main_inner(texturecube_array<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_3112e8(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(texturecube_array<float, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texturecube_array<float, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_3112e8(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(texturecube_array<float, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_3112e8(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/3112e8.wgsl.expected.spvasm

@@ -0,0 +1,86 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 45
+; Schema: 0
+               OpCapability Shader
+               OpCapability SampledCubeArray
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_1 "arg_1"
+               OpName %arg_2 "arg_2"
+               OpName %textureGather_3112e8 "textureGather_3112e8"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+               OpDecorate %arg_2 DescriptorSet 1
+               OpDecorate %arg_2 Binding 2
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+         %11 = OpTypeImage %float Cube 0 1 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %14 = OpTypeSampler
+%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
+      %arg_2 = OpVariable %_ptr_UniformConstant_14 UniformConstant
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+         %22 = OpTypeSampledImage %11
+    %float_0 = OpConstant %float 0
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %31 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_3112e8 = OpFunction %void None %15
+         %18 = OpLabel
+        %res = OpVariable %_ptr_Function_v4float Function %5
+         %20 = OpLoad %14 %arg_2
+         %21 = OpLoad %11 %arg_1
+         %23 = OpSampledImage %22 %21 %20
+         %25 = OpConvertSToF %float %int_1
+         %28 = OpCompositeConstruct %v4float %float_0 %float_0 %float_0 %25
+         %19 = OpImageGather %v4float %23 %28 %int_1
+               OpStore %res %19
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %31
+         %33 = OpLabel
+         %34 = OpFunctionCall %void %textureGather_3112e8
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %15
+         %36 = OpLabel
+         %37 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %37
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %15
+         %40 = OpLabel
+         %41 = OpFunctionCall %void %textureGather_3112e8
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %15
+         %43 = OpLabel
+         %44 = OpFunctionCall %void %textureGather_3112e8
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/3112e8.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(1)]] var arg_1 : texture_cube_array<f32>;
+
+[[group(1), binding(2)]] var arg_2 : sampler;
+
+fn textureGather_3112e8() {
+  var res : vec4<f32> = textureGather(1, arg_1, arg_2, vec3<f32>(), 1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_3112e8();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_3112e8();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_3112e8();
+}

test/intrinsics/gen/textureGather/3c527e.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(1)]] var arg_1: texture_cube_array<u32>;
+[[group(1), binding(2)]] var arg_2: sampler;
+
+// fn textureGather(component: i32, texture: texture_cube_array<u32>, sampler: sampler, coords: vec3<f32>, array_index: i32) -> vec4<u32>
+fn textureGather_3c527e() {
+  var res: vec4<u32> = textureGather(1, arg_1, arg_2, vec3<f32>(), 1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_3c527e();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_3c527e();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_3c527e();
+}

test/intrinsics/gen/textureGather/3c527e.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+TextureCubeArray<uint4> arg_1 : register(t1, space1);
+SamplerState arg_2 : register(s2, space1);
+
+void textureGather_3c527e() {
+  uint4 res = arg_1.GatherGreen(arg_2, float4(0.0f, 0.0f, 0.0f, float(1)));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_3c527e();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_3c527e();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_3c527e();
+  return;
+}

test/intrinsics/gen/textureGather/3c527e.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_3c527e(texturecube_array<uint, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  uint4 res = tint_symbol_1.gather(tint_symbol_2, float3(), 1, component::y);
+}
+
+float4 vertex_main_inner(texturecube_array<uint, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_3c527e(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(texturecube_array<uint, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texturecube_array<uint, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_3c527e(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(texturecube_array<uint, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_3c527e(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/3c527e.wgsl.expected.spvasm

@@ -0,0 +1,89 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+               OpCapability Shader
+               OpCapability SampledCubeArray
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_1 "arg_1"
+               OpName %arg_2 "arg_2"
+               OpName %textureGather_3c527e "textureGather_3c527e"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+               OpDecorate %arg_2 DescriptorSet 1
+               OpDecorate %arg_2 Binding 2
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %uint = OpTypeInt 32 0
+         %11 = OpTypeImage %uint Cube 0 1 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %15 = OpTypeSampler
+%_ptr_UniformConstant_15 = OpTypePointer UniformConstant %15
+      %arg_2 = OpVariable %_ptr_UniformConstant_15 UniformConstant
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+     %v4uint = OpTypeVector %uint 4
+         %24 = OpTypeSampledImage %11
+    %float_0 = OpConstant %float 0
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+         %33 = OpConstantNull %v4uint
+         %34 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_3c527e = OpFunction %void None %16
+         %19 = OpLabel
+        %res = OpVariable %_ptr_Function_v4uint Function %33
+         %22 = OpLoad %15 %arg_2
+         %23 = OpLoad %11 %arg_1
+         %25 = OpSampledImage %24 %23 %22
+         %27 = OpConvertSToF %float %int_1
+         %30 = OpCompositeConstruct %v4float %float_0 %float_0 %float_0 %27
+         %20 = OpImageGather %v4uint %25 %30 %int_1
+               OpStore %res %20
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %textureGather_3c527e
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %16
+         %39 = OpLabel
+         %40 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %40
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %16
+         %43 = OpLabel
+         %44 = OpFunctionCall %void %textureGather_3c527e
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %16
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %textureGather_3c527e
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/3c527e.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(1)]] var arg_1 : texture_cube_array<u32>;
+
+[[group(1), binding(2)]] var arg_2 : sampler;
+
+fn textureGather_3c527e() {
+  var res : vec4<u32> = textureGather(1, arg_1, arg_2, vec3<f32>(), 1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_3c527e();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_3c527e();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_3c527e();
+}

test/intrinsics/gen/textureGather/43025d.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(0)]] var arg_0: texture_depth_cube_array;
+[[group(1), binding(1)]] var arg_1: sampler;
+
+// fn textureGather(texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: i32) -> vec4<f32>
+fn textureGather_43025d() {
+  var res: vec4<f32> = textureGather(arg_0, arg_1, vec3<f32>(), 1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_43025d();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_43025d();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_43025d();
+}

test/intrinsics/gen/textureGather/43025d.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+TextureCubeArray arg_0 : register(t0, space1);
+SamplerState arg_1 : register(s1, space1);
+
+void textureGather_43025d() {
+  float4 res = arg_0.Gather(arg_1, float4(0.0f, 0.0f, 0.0f, float(1)));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_43025d();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_43025d();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_43025d();
+  return;
+}

test/intrinsics/gen/textureGather/43025d.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_43025d(depthcube_array<float, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  float4 res = tint_symbol_1.gather(tint_symbol_2, float3(), 1);
+}
+
+float4 vertex_main_inner(depthcube_array<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_43025d(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(depthcube_array<float, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(depthcube_array<float, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_43025d(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(depthcube_array<float, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_43025d(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/43025d.wgsl.expected.spvasm

@@ -0,0 +1,87 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 46
+; Schema: 0
+               OpCapability Shader
+               OpCapability SampledCubeArray
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_0 "arg_0"
+               OpName %arg_1 "arg_1"
+               OpName %textureGather_43025d "textureGather_43025d"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+         %11 = OpTypeImage %float Cube 1 1 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %14 = OpTypeSampler
+%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
+      %arg_1 = OpVariable %_ptr_UniformConstant_14 UniformConstant
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+         %22 = OpTypeSampledImage %11
+    %float_0 = OpConstant %float 0
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+      %int_0 = OpConstant %int 0
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %32 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_43025d = OpFunction %void None %15
+         %18 = OpLabel
+        %res = OpVariable %_ptr_Function_v4float Function %5
+         %20 = OpLoad %14 %arg_1
+         %21 = OpLoad %11 %arg_0
+         %23 = OpSampledImage %22 %21 %20
+         %25 = OpConvertSToF %float %int_1
+         %28 = OpCompositeConstruct %v4float %float_0 %float_0 %float_0 %25
+         %19 = OpImageGather %v4float %23 %28 %int_0
+               OpStore %res %19
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %32
+         %34 = OpLabel
+         %35 = OpFunctionCall %void %textureGather_43025d
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %15
+         %37 = OpLabel
+         %38 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %38
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %15
+         %41 = OpLabel
+         %42 = OpFunctionCall %void %textureGather_43025d
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %15
+         %44 = OpLabel
+         %45 = OpFunctionCall %void %textureGather_43025d
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/43025d.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(0)]] var arg_0 : texture_depth_cube_array;
+
+[[group(1), binding(1)]] var arg_1 : sampler;
+
+fn textureGather_43025d() {
+  var res : vec4<f32> = textureGather(arg_0, arg_1, vec3<f32>(), 1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_43025d();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_43025d();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_43025d();
+}

test/intrinsics/gen/textureGather/4f2350.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(1)]] var arg_1: texture_2d_array<i32>;
+[[group(1), binding(2)]] var arg_2: sampler;
+
+// fn textureGather(component: i32, texture: texture_2d_array<i32>, sampler: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<i32>
+fn textureGather_4f2350() {
+  var res: vec4<i32> = textureGather(1, arg_1, arg_2, vec2<f32>(), 1, vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_4f2350();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_4f2350();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_4f2350();
+}

test/intrinsics/gen/textureGather/4f2350.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+Texture2DArray<int4> arg_1 : register(t1, space1);
+SamplerState arg_2 : register(s2, space1);
+
+void textureGather_4f2350() {
+  int4 res = arg_1.GatherGreen(arg_2, float3(0.0f, 0.0f, float(1)), int2(0, 0));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_4f2350();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_4f2350();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_4f2350();
+  return;
+}

test/intrinsics/gen/textureGather/4f2350.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_4f2350(texture2d_array<int, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  int4 res = tint_symbol_1.gather(tint_symbol_2, float2(), 1, int2(), component::y);
+}
+
+float4 vertex_main_inner(texture2d_array<int, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_4f2350(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(texture2d_array<int, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texture2d_array<int, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_4f2350(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(texture2d_array<int, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_4f2350(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/4f2350.wgsl.expected.spvasm

@@ -0,0 +1,90 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 50
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_1 "arg_1"
+               OpName %arg_2 "arg_2"
+               OpName %textureGather_4f2350 "textureGather_4f2350"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+               OpDecorate %arg_2 DescriptorSet 1
+               OpDecorate %arg_2 Binding 2
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+        %int = OpTypeInt 32 1
+         %11 = OpTypeImage %int 2D 0 1 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %15 = OpTypeSampler
+%_ptr_UniformConstant_15 = OpTypePointer UniformConstant %15
+      %arg_2 = OpVariable %_ptr_UniformConstant_15 UniformConstant
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+      %v4int = OpTypeVector %int 4
+         %24 = OpTypeSampledImage %11
+    %v3float = OpTypeVector %float 3
+    %float_0 = OpConstant %float 0
+      %int_1 = OpConstant %int 1
+      %v2int = OpTypeVector %int 2
+         %32 = OpConstantNull %v2int
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+         %35 = OpConstantNull %v4int
+         %36 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_4f2350 = OpFunction %void None %16
+         %19 = OpLabel
+        %res = OpVariable %_ptr_Function_v4int Function %35
+         %22 = OpLoad %15 %arg_2
+         %23 = OpLoad %11 %arg_1
+         %25 = OpSampledImage %24 %23 %22
+         %28 = OpConvertSToF %float %int_1
+         %30 = OpCompositeConstruct %v3float %float_0 %float_0 %28
+         %20 = OpImageGather %v4int %25 %30 %int_1 ConstOffset %32
+               OpStore %res %20
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %36
+         %38 = OpLabel
+         %39 = OpFunctionCall %void %textureGather_4f2350
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %16
+         %41 = OpLabel
+         %42 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %42
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %16
+         %45 = OpLabel
+         %46 = OpFunctionCall %void %textureGather_4f2350
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %16
+         %48 = OpLabel
+         %49 = OpFunctionCall %void %textureGather_4f2350
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/4f2350.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(1)]] var arg_1 : texture_2d_array<i32>;
+
+[[group(1), binding(2)]] var arg_2 : sampler;
+
+fn textureGather_4f2350() {
+  var res : vec4<i32> = textureGather(1, arg_1, arg_2, vec2<f32>(), 1, vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_4f2350();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_4f2350();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_4f2350();
+}

test/intrinsics/gen/textureGather/51cf0b.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(1)]] var arg_1: texture_2d_array<i32>;
+[[group(1), binding(2)]] var arg_2: sampler;
+
+// fn textureGather(component: i32, texture: texture_2d_array<i32>, sampler: sampler, coords: vec2<f32>, array_index: i32) -> vec4<i32>
+fn textureGather_51cf0b() {
+  var res: vec4<i32> = textureGather(1, arg_1, arg_2, vec2<f32>(), 1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_51cf0b();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_51cf0b();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_51cf0b();
+}

test/intrinsics/gen/textureGather/51cf0b.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+Texture2DArray<int4> arg_1 : register(t1, space1);
+SamplerState arg_2 : register(s2, space1);
+
+void textureGather_51cf0b() {
+  int4 res = arg_1.GatherGreen(arg_2, float3(0.0f, 0.0f, float(1)));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_51cf0b();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_51cf0b();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_51cf0b();
+  return;
+}

test/intrinsics/gen/textureGather/51cf0b.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_51cf0b(texture2d_array<int, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  int4 res = tint_symbol_1.gather(tint_symbol_2, float2(), 1, int2(0), component::y);
+}
+
+float4 vertex_main_inner(texture2d_array<int, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_51cf0b(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(texture2d_array<int, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texture2d_array<int, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_51cf0b(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(texture2d_array<int, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_51cf0b(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/51cf0b.wgsl.expected.spvasm

@@ -0,0 +1,88 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_1 "arg_1"
+               OpName %arg_2 "arg_2"
+               OpName %textureGather_51cf0b "textureGather_51cf0b"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+               OpDecorate %arg_2 DescriptorSet 1
+               OpDecorate %arg_2 Binding 2
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+        %int = OpTypeInt 32 1
+         %11 = OpTypeImage %int 2D 0 1 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %15 = OpTypeSampler
+%_ptr_UniformConstant_15 = OpTypePointer UniformConstant %15
+      %arg_2 = OpVariable %_ptr_UniformConstant_15 UniformConstant
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+      %v4int = OpTypeVector %int 4
+         %24 = OpTypeSampledImage %11
+    %v3float = OpTypeVector %float 3
+    %float_0 = OpConstant %float 0
+      %int_1 = OpConstant %int 1
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+         %33 = OpConstantNull %v4int
+         %34 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_51cf0b = OpFunction %void None %16
+         %19 = OpLabel
+        %res = OpVariable %_ptr_Function_v4int Function %33
+         %22 = OpLoad %15 %arg_2
+         %23 = OpLoad %11 %arg_1
+         %25 = OpSampledImage %24 %23 %22
+         %28 = OpConvertSToF %float %int_1
+         %30 = OpCompositeConstruct %v3float %float_0 %float_0 %28
+         %20 = OpImageGather %v4int %25 %30 %int_1
+               OpStore %res %20
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %textureGather_51cf0b
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %16
+         %39 = OpLabel
+         %40 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %40
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %16
+         %43 = OpLabel
+         %44 = OpFunctionCall %void %textureGather_51cf0b
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %16
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %textureGather_51cf0b
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/51cf0b.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(1)]] var arg_1 : texture_2d_array<i32>;
+
+[[group(1), binding(2)]] var arg_2 : sampler;
+
+fn textureGather_51cf0b() {
+  var res : vec4<i32> = textureGather(1, arg_1, arg_2, vec2<f32>(), 1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_51cf0b();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_51cf0b();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_51cf0b();
+}

test/intrinsics/gen/textureGather/53ece6.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(0)]] var arg_0: texture_depth_2d_array;
+[[group(1), binding(1)]] var arg_1: sampler;
+
+// fn textureGather(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, offset: vec2<i32>) -> vec4<f32>
+fn textureGather_53ece6() {
+  var res: vec4<f32> = textureGather(arg_0, arg_1, vec2<f32>(), 1, vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_53ece6();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_53ece6();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_53ece6();
+}

test/intrinsics/gen/textureGather/53ece6.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+Texture2DArray arg_0 : register(t0, space1);
+SamplerState arg_1 : register(s1, space1);
+
+void textureGather_53ece6() {
+  float4 res = arg_0.Gather(arg_1, float3(0.0f, 0.0f, float(1)), int2(0, 0));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_53ece6();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_53ece6();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_53ece6();
+  return;
+}

test/intrinsics/gen/textureGather/53ece6.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_53ece6(depth2d_array<float, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  float4 res = tint_symbol_1.gather(tint_symbol_2, float2(), 1, int2());
+}
+
+float4 vertex_main_inner(depth2d_array<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_53ece6(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(depth2d_array<float, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(depth2d_array<float, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_53ece6(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(depth2d_array<float, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_53ece6(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/53ece6.wgsl.expected.spvasm

@@ -0,0 +1,89 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 49
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_0 "arg_0"
+               OpName %arg_1 "arg_1"
+               OpName %textureGather_53ece6 "textureGather_53ece6"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+         %11 = OpTypeImage %float 2D 1 1 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %14 = OpTypeSampler
+%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
+      %arg_1 = OpVariable %_ptr_UniformConstant_14 UniformConstant
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+         %22 = OpTypeSampledImage %11
+    %v3float = OpTypeVector %float 3
+    %float_0 = OpConstant %float 0
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+      %int_0 = OpConstant %int 0
+      %v2int = OpTypeVector %int 2
+         %32 = OpConstantNull %v2int
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %35 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_53ece6 = OpFunction %void None %15
+         %18 = OpLabel
+        %res = OpVariable %_ptr_Function_v4float Function %5
+         %20 = OpLoad %14 %arg_1
+         %21 = OpLoad %11 %arg_0
+         %23 = OpSampledImage %22 %21 %20
+         %26 = OpConvertSToF %float %int_1
+         %29 = OpCompositeConstruct %v3float %float_0 %float_0 %26
+         %19 = OpImageGather %v4float %23 %29 %int_0 ConstOffset %32
+               OpStore %res %19
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %35
+         %37 = OpLabel
+         %38 = OpFunctionCall %void %textureGather_53ece6
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %15
+         %40 = OpLabel
+         %41 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %41
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %15
+         %44 = OpLabel
+         %45 = OpFunctionCall %void %textureGather_53ece6
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %15
+         %47 = OpLabel
+         %48 = OpFunctionCall %void %textureGather_53ece6
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/53ece6.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(0)]] var arg_0 : texture_depth_2d_array;
+
+[[group(1), binding(1)]] var arg_1 : sampler;
+
+fn textureGather_53ece6() {
+  var res : vec4<f32> = textureGather(arg_0, arg_1, vec2<f32>(), 1, vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_53ece6();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_53ece6();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_53ece6();
+}

test/intrinsics/gen/textureGather/57bfc6.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(1)]] var arg_1: texture_cube<f32>;
+[[group(1), binding(2)]] var arg_2: sampler;
+
+// fn textureGather(component: i32, texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
+fn textureGather_57bfc6() {
+  var res: vec4<f32> = textureGather(1, arg_1, arg_2, vec3<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_57bfc6();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_57bfc6();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_57bfc6();
+}

test/intrinsics/gen/textureGather/57bfc6.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+TextureCube<float4> arg_1 : register(t1, space1);
+SamplerState arg_2 : register(s2, space1);
+
+void textureGather_57bfc6() {
+  float4 res = arg_1.GatherGreen(arg_2, float3(0.0f, 0.0f, 0.0f));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_57bfc6();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_57bfc6();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_57bfc6();
+  return;
+}

test/intrinsics/gen/textureGather/57bfc6.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_57bfc6(texturecube<float, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  float4 res = tint_symbol_1.gather(tint_symbol_2, float3(), component::y);
+}
+
+float4 vertex_main_inner(texturecube<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_57bfc6(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(texturecube<float, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texturecube<float, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_57bfc6(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(texturecube<float, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_57bfc6(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/57bfc6.wgsl.expected.spvasm

@@ -0,0 +1,84 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 44
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_1 "arg_1"
+               OpName %arg_2 "arg_2"
+               OpName %textureGather_57bfc6 "textureGather_57bfc6"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+               OpDecorate %arg_2 DescriptorSet 1
+               OpDecorate %arg_2 Binding 2
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+         %11 = OpTypeImage %float Cube 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %14 = OpTypeSampler
+%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
+      %arg_2 = OpVariable %_ptr_UniformConstant_14 UniformConstant
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+         %22 = OpTypeSampledImage %11
+    %v3float = OpTypeVector %float 3
+         %25 = OpConstantNull %v3float
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %30 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_57bfc6 = OpFunction %void None %15
+         %18 = OpLabel
+        %res = OpVariable %_ptr_Function_v4float Function %5
+         %20 = OpLoad %14 %arg_2
+         %21 = OpLoad %11 %arg_1
+         %23 = OpSampledImage %22 %21 %20
+         %19 = OpImageGather %v4float %23 %25 %int_1
+               OpStore %res %19
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %30
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %textureGather_57bfc6
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %15
+         %35 = OpLabel
+         %36 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %36
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %15
+         %39 = OpLabel
+         %40 = OpFunctionCall %void %textureGather_57bfc6
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %15
+         %42 = OpLabel
+         %43 = OpFunctionCall %void %textureGather_57bfc6
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/57bfc6.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(1)]] var arg_1 : texture_cube<f32>;
+
+[[group(1), binding(2)]] var arg_2 : sampler;
+
+fn textureGather_57bfc6() {
+  var res : vec4<f32> = textureGather(1, arg_1, arg_2, vec3<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_57bfc6();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_57bfc6();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_57bfc6();
+}

test/intrinsics/gen/textureGather/587ba3.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(1)]] var arg_1: texture_2d<i32>;
+[[group(1), binding(2)]] var arg_2: sampler;
+
+// fn textureGather(component: i32, texture: texture_2d<i32>, sampler: sampler, coords: vec2<f32>) -> vec4<i32>
+fn textureGather_587ba3() {
+  var res: vec4<i32> = textureGather(1, arg_1, arg_2, vec2<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_587ba3();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_587ba3();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_587ba3();
+}

test/intrinsics/gen/textureGather/587ba3.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+Texture2D<int4> arg_1 : register(t1, space1);
+SamplerState arg_2 : register(s2, space1);
+
+void textureGather_587ba3() {
+  int4 res = arg_1.GatherGreen(arg_2, float2(0.0f, 0.0f));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_587ba3();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_587ba3();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_587ba3();
+  return;
+}

test/intrinsics/gen/textureGather/587ba3.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_587ba3(texture2d<int, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  int4 res = tint_symbol_1.gather(tint_symbol_2, float2(), int2(0), component::y);
+}
+
+float4 vertex_main_inner(texture2d<int, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_587ba3(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(texture2d<int, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texture2d<int, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_587ba3(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(texture2d<int, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_587ba3(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/587ba3.wgsl.expected.spvasm

@@ -0,0 +1,86 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 46
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_1 "arg_1"
+               OpName %arg_2 "arg_2"
+               OpName %textureGather_587ba3 "textureGather_587ba3"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+               OpDecorate %arg_2 DescriptorSet 1
+               OpDecorate %arg_2 Binding 2
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+        %int = OpTypeInt 32 1
+         %11 = OpTypeImage %int 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %15 = OpTypeSampler
+%_ptr_UniformConstant_15 = OpTypePointer UniformConstant %15
+      %arg_2 = OpVariable %_ptr_UniformConstant_15 UniformConstant
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+      %v4int = OpTypeVector %int 4
+         %24 = OpTypeSampledImage %11
+    %v2float = OpTypeVector %float 2
+         %27 = OpConstantNull %v2float
+      %int_1 = OpConstant %int 1
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+         %31 = OpConstantNull %v4int
+         %32 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_587ba3 = OpFunction %void None %16
+         %19 = OpLabel
+        %res = OpVariable %_ptr_Function_v4int Function %31
+         %22 = OpLoad %15 %arg_2
+         %23 = OpLoad %11 %arg_1
+         %25 = OpSampledImage %24 %23 %22
+         %20 = OpImageGather %v4int %25 %27 %int_1
+               OpStore %res %20
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %32
+         %34 = OpLabel
+         %35 = OpFunctionCall %void %textureGather_587ba3
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %16
+         %37 = OpLabel
+         %38 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %38
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %16
+         %41 = OpLabel
+         %42 = OpFunctionCall %void %textureGather_587ba3
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %16
+         %44 = OpLabel
+         %45 = OpFunctionCall %void %textureGather_587ba3
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/587ba3.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(1)]] var arg_1 : texture_2d<i32>;
+
+[[group(1), binding(2)]] var arg_2 : sampler;
+
+fn textureGather_587ba3() {
+  var res : vec4<i32> = textureGather(1, arg_1, arg_2, vec2<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_587ba3();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_587ba3();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_587ba3();
+}

test/intrinsics/gen/textureGather/69e0fb.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(1)]] var arg_1: texture_2d<i32>;
+[[group(1), binding(2)]] var arg_2: sampler;
+
+// fn textureGather(component: i32, texture: texture_2d<i32>, sampler: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<i32>
+fn textureGather_69e0fb() {
+  var res: vec4<i32> = textureGather(1, arg_1, arg_2, vec2<f32>(), vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_69e0fb();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_69e0fb();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_69e0fb();
+}

test/intrinsics/gen/textureGather/69e0fb.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+Texture2D<int4> arg_1 : register(t1, space1);
+SamplerState arg_2 : register(s2, space1);
+
+void textureGather_69e0fb() {
+  int4 res = arg_1.GatherGreen(arg_2, float2(0.0f, 0.0f), int2(0, 0));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_69e0fb();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_69e0fb();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_69e0fb();
+  return;
+}

test/intrinsics/gen/textureGather/69e0fb.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_69e0fb(texture2d<int, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  int4 res = tint_symbol_1.gather(tint_symbol_2, float2(), int2(), component::y);
+}
+
+float4 vertex_main_inner(texture2d<int, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_69e0fb(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(texture2d<int, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texture2d<int, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_69e0fb(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(texture2d<int, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_69e0fb(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/69e0fb.wgsl.expected.spvasm

@@ -0,0 +1,88 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_1 "arg_1"
+               OpName %arg_2 "arg_2"
+               OpName %textureGather_69e0fb "textureGather_69e0fb"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+               OpDecorate %arg_2 DescriptorSet 1
+               OpDecorate %arg_2 Binding 2
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+        %int = OpTypeInt 32 1
+         %11 = OpTypeImage %int 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %15 = OpTypeSampler
+%_ptr_UniformConstant_15 = OpTypePointer UniformConstant %15
+      %arg_2 = OpVariable %_ptr_UniformConstant_15 UniformConstant
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+      %v4int = OpTypeVector %int 4
+         %24 = OpTypeSampledImage %11
+    %v2float = OpTypeVector %float 2
+         %27 = OpConstantNull %v2float
+      %int_1 = OpConstant %int 1
+      %v2int = OpTypeVector %int 2
+         %30 = OpConstantNull %v2int
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+         %33 = OpConstantNull %v4int
+         %34 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_69e0fb = OpFunction %void None %16
+         %19 = OpLabel
+        %res = OpVariable %_ptr_Function_v4int Function %33
+         %22 = OpLoad %15 %arg_2
+         %23 = OpLoad %11 %arg_1
+         %25 = OpSampledImage %24 %23 %22
+         %20 = OpImageGather %v4int %25 %27 %int_1 ConstOffset %30
+               OpStore %res %20
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %textureGather_69e0fb
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %16
+         %39 = OpLabel
+         %40 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %40
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %16
+         %43 = OpLabel
+         %44 = OpFunctionCall %void %textureGather_69e0fb
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %16
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %textureGather_69e0fb
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/69e0fb.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(1)]] var arg_1 : texture_2d<i32>;
+
+[[group(1), binding(2)]] var arg_2 : sampler;
+
+fn textureGather_69e0fb() {
+  var res : vec4<i32> = textureGather(1, arg_1, arg_2, vec2<f32>(), vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_69e0fb();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_69e0fb();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_69e0fb();
+}

test/intrinsics/gen/textureGather/93003d.wgsl

@@ -0,0 +1,47 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(1)]] var arg_1: texture_2d<u32>;
+[[group(1), binding(2)]] var arg_2: sampler;
+
+// fn textureGather(component: i32, texture: texture_2d<u32>, sampler: sampler, coords: vec2<f32>, offset: vec2<i32>) -> vec4<u32>
+fn textureGather_93003d() {
+  var res: vec4<u32> = textureGather(1, arg_1, arg_2, vec2<f32>(), vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_93003d();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_93003d();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_93003d();
+}

test/intrinsics/gen/textureGather/93003d.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+Texture2D<uint4> arg_1 : register(t1, space1);
+SamplerState arg_2 : register(s2, space1);
+
+void textureGather_93003d() {
+  uint4 res = arg_1.GatherGreen(arg_2, float2(0.0f, 0.0f), int2(0, 0));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureGather_93003d();
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureGather_93003d();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureGather_93003d();
+  return;
+}

test/intrinsics/gen/textureGather/93003d.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureGather_93003d(texture2d<uint, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  uint4 res = tint_symbol_1.gather(tint_symbol_2, float2(), int2(), component::y);
+}
+
+float4 vertex_main_inner(texture2d<uint, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureGather_93003d(tint_symbol_3, tint_symbol_4);
+  return float4();
+}
+
+vertex tint_symbol vertex_main(texture2d<uint, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texture2d<uint, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureGather_93003d(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(texture2d<uint, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureGather_93003d(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/intrinsics/gen/textureGather/93003d.wgsl.expected.spvasm

@@ -0,0 +1,89 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 49
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_1 "arg_1"
+               OpName %arg_2 "arg_2"
+               OpName %textureGather_93003d "textureGather_93003d"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+               OpDecorate %arg_2 DescriptorSet 1
+               OpDecorate %arg_2 Binding 2
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %uint = OpTypeInt 32 0
+         %11 = OpTypeImage %uint 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %15 = OpTypeSampler
+%_ptr_UniformConstant_15 = OpTypePointer UniformConstant %15
+      %arg_2 = OpVariable %_ptr_UniformConstant_15 UniformConstant
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+     %v4uint = OpTypeVector %uint 4
+         %24 = OpTypeSampledImage %11
+    %v2float = OpTypeVector %float 2
+         %27 = OpConstantNull %v2float
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+      %v2int = OpTypeVector %int 2
+         %31 = OpConstantNull %v2int
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+         %34 = OpConstantNull %v4uint
+         %35 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%textureGather_93003d = OpFunction %void None %16
+         %19 = OpLabel
+        %res = OpVariable %_ptr_Function_v4uint Function %34
+         %22 = OpLoad %15 %arg_2
+         %23 = OpLoad %11 %arg_1
+         %25 = OpSampledImage %24 %23 %22
+         %20 = OpImageGather %v4uint %25 %27 %int_1 ConstOffset %31
+               OpStore %res %20
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %35
+         %37 = OpLabel
+         %38 = OpFunctionCall %void %textureGather_93003d
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %16
+         %40 = OpLabel
+         %41 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %41
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %16
+         %44 = OpLabel
+         %45 = OpFunctionCall %void %textureGather_93003d
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %16
+         %47 = OpLabel
+         %48 = OpFunctionCall %void %textureGather_93003d
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureGather/93003d.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[group(1), binding(1)]] var arg_1 : texture_2d<u32>;
+
+[[group(1), binding(2)]] var arg_2 : sampler;
+
+fn textureGather_93003d() {
+  var res : vec4<u32> = textureGather(1, arg_1, arg_2, vec2<f32>(), vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureGather_93003d();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureGather_93003d();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureGather_93003d();
+}