tint/SingleEntryPoint: Handle diagnostic directive

Add some E2E tests to make sure that we actually produce valid code
for each backend when diagnostic filtering is present.

Bug: tint:1809
Change-Id: I5e903ac0d2ca385967211bb889f86cb85de8f418
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/117590
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

src/tint/transform/single_entry_point.cc

@@ -114,6 +114,7 @@ Transform::ApplyResult SingleEntryPoint::Apply(const Program* src,
                 }
             },
             [&](const ast::Enable* ext) { b.AST().AddEnable(ctx.Clone(ext)); },
+            [&](const ast::DiagnosticControl* dc) { b.AST().AddDiagnosticControl(ctx.Clone(dc)); },
             [&](Default) {
                 TINT_UNREACHABLE(Transform, b.Diagnostics())
                     << "unhandled global declaration: " << decl->TypeInfo().name;

src/tint/transform/single_entry_point_test.cc

@@ -611,5 +611,25 @@ fn main() {
     EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(SingleEntryPointTest, Directives) {
+    // Make sure that directives are preserved.
+    auto* src = R"(
+enable f16;
+diagnostic(off, derivative_uniformity);
+
+@compute @workgroup_size(1)
+fn main() {
+}
+)";
+
+    SingleEntryPoint::Config cfg("main");
+
+    DataMap data;
+    data.Add<SingleEntryPoint::Config>(cfg);
+    auto got = Run<SingleEntryPoint>(src, data);
+
+    EXPECT_EQ(src, str(got));
+}
+
 }  // namespace
 }  // namespace tint::transform

test/tint/diagnostic_filtering/directive.wgsl

@@ -0,0 +1,11 @@
+diagnostic(warning, derivative_uniformity);
+
+@group(0) @binding(1) var t : texture_2d<f32>;
+@group(0) @binding(2) var s : sampler;
+
+@fragment
+fn main(@location(0) x : f32) {
+  if (x > 0) {
+    _ = textureSample(t, s, vec2(0, 0));
+  }
+}

test/tint/diagnostic_filtering/directive.wgsl.expected.dxc.hlsl

@@ -0,0 +1,28 @@
+diagnostic_filtering/directive.wgsl:9:9 warning: 'textureSample' must only be called from uniform control flow
+    _ = textureSample(t, s, vec2(0, 0));
+        ^^^^^^^^^^^^^
+
+diagnostic_filtering/directive.wgsl:8:3 note: control flow depends on possibly non-uniform value
+  if (x > 0) {
+  ^^
+
+diagnostic_filtering/directive.wgsl:8:7 note: user-defined input 'x' of 'main' may be non-uniform
+  if (x > 0) {
+      ^
+
+Texture2D<float4> t : register(t1, space0);
+SamplerState s : register(s2, space0);
+
+struct tint_symbol_1 {
+  float x : TEXCOORD0;
+};
+
+void main_inner(float x) {
+  if ((x > 0.0f)) {
+  }
+}
+
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.x);
+  return;
+}

test/tint/diagnostic_filtering/directive.wgsl.expected.fxc.hlsl

@@ -0,0 +1,28 @@
+diagnostic_filtering/directive.wgsl:9:9 warning: 'textureSample' must only be called from uniform control flow
+    _ = textureSample(t, s, vec2(0, 0));
+        ^^^^^^^^^^^^^
+
+diagnostic_filtering/directive.wgsl:8:3 note: control flow depends on possibly non-uniform value
+  if (x > 0) {
+  ^^
+
+diagnostic_filtering/directive.wgsl:8:7 note: user-defined input 'x' of 'main' may be non-uniform
+  if (x > 0) {
+      ^
+
+Texture2D<float4> t : register(t1, space0);
+SamplerState s : register(s2, space0);
+
+struct tint_symbol_1 {
+  float x : TEXCOORD0;
+};
+
+void main_inner(float x) {
+  if ((x > 0.0f)) {
+  }
+}
+
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.x);
+  return;
+}

test/tint/diagnostic_filtering/directive.wgsl.expected.glsl

@@ -0,0 +1,25 @@
+diagnostic_filtering/directive.wgsl:9:9 warning: 'textureSample' must only be called from uniform control flow
+    _ = textureSample(t, s, vec2(0, 0));
+        ^^^^^^^^^^^^^
+
+diagnostic_filtering/directive.wgsl:8:3 note: control flow depends on possibly non-uniform value
+  if (x > 0) {
+  ^^
+
+diagnostic_filtering/directive.wgsl:8:7 note: user-defined input 'x' of 'main' may be non-uniform
+  if (x > 0) {
+      ^
+
+#version 310 es
+precision mediump float;
+
+layout(location = 0) in float x_1;
+void tint_symbol(float x) {
+  if ((x > 0.0f)) {
+  }
+}
+
+void main() {
+  tint_symbol(x_1);
+  return;
+}

test/tint/diagnostic_filtering/directive.wgsl.expected.msl

@@ -0,0 +1,29 @@
+diagnostic_filtering/directive.wgsl:9:9 warning: 'textureSample' must only be called from uniform control flow
+    _ = textureSample(t, s, vec2(0, 0));
+        ^^^^^^^^^^^^^
+
+diagnostic_filtering/directive.wgsl:8:3 note: control flow depends on possibly non-uniform value
+  if (x > 0) {
+  ^^
+
+diagnostic_filtering/directive.wgsl:8:7 note: user-defined input 'x' of 'main' may be non-uniform
+  if (x > 0) {
+      ^
+
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_2 {
+  float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x) {
+  if ((x > 0.0f)) {
+  }
+}
+
+fragment void tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+  tint_symbol_inner(tint_symbol_1.x);
+  return;
+}
+

test/tint/diagnostic_filtering/directive.wgsl.expected.spvasm

@@ -0,0 +1,63 @@
+diagnostic_filtering/directive.wgsl:9:9 warning: 'textureSample' must only be called from uniform control flow
+    _ = textureSample(t, s, vec2(0, 0));
+        ^^^^^^^^^^^^^
+
+diagnostic_filtering/directive.wgsl:8:3 note: control flow depends on possibly non-uniform value
+  if (x > 0) {
+  ^^
+
+diagnostic_filtering/directive.wgsl:8:7 note: user-defined input 'x' of 'main' may be non-uniform
+  if (x > 0) {
+      ^
+
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main" %x_1
+               OpExecutionMode %main OriginUpperLeft
+               OpName %x_1 "x_1"
+               OpName %t "t"
+               OpName %s "s"
+               OpName %main_inner "main_inner"
+               OpName %x "x"
+               OpName %main "main"
+               OpDecorate %x_1 Location 0
+               OpDecorate %t DescriptorSet 0
+               OpDecorate %t Binding 1
+               OpDecorate %s DescriptorSet 0
+               OpDecorate %s Binding 2
+      %float = OpTypeFloat 32
+%_ptr_Input_float = OpTypePointer Input %float
+        %x_1 = OpVariable %_ptr_Input_float Input
+          %6 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_6 = OpTypePointer UniformConstant %6
+          %t = OpVariable %_ptr_UniformConstant_6 UniformConstant
+          %9 = OpTypeSampler
+%_ptr_UniformConstant_9 = OpTypePointer UniformConstant %9
+          %s = OpVariable %_ptr_UniformConstant_9 UniformConstant
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void %float
+         %15 = OpConstantNull %float
+       %bool = OpTypeBool
+         %20 = OpTypeFunction %void
+ %main_inner = OpFunction %void None %10
+          %x = OpFunctionParameter %float
+         %14 = OpLabel
+         %16 = OpFOrdGreaterThan %bool %x %15
+               OpSelectionMerge %18 None
+               OpBranchConditional %16 %19 %18
+         %19 = OpLabel
+               OpBranch %18
+         %18 = OpLabel
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %20
+         %22 = OpLabel
+         %24 = OpLoad %float %x_1
+         %23 = OpFunctionCall %void %main_inner %24
+               OpReturn
+               OpFunctionEnd

test/tint/diagnostic_filtering/directive.wgsl.expected.wgsl

@@ -0,0 +1,24 @@
+diagnostic_filtering/directive.wgsl:9:9 warning: 'textureSample' must only be called from uniform control flow
+    _ = textureSample(t, s, vec2(0, 0));
+        ^^^^^^^^^^^^^
+
+diagnostic_filtering/directive.wgsl:8:3 note: control flow depends on possibly non-uniform value
+  if (x > 0) {
+  ^^
+
+diagnostic_filtering/directive.wgsl:8:7 note: user-defined input 'x' of 'main' may be non-uniform
+  if (x > 0) {
+      ^
+
+diagnostic(warning, derivative_uniformity);
+
+@group(0) @binding(1) var t : texture_2d<f32>;
+
+@group(0) @binding(2) var s : sampler;
+
+@fragment
+fn main(@location(0) x : f32) {
+  if ((x > 0)) {
+    _ = textureSample(t, s, vec2(0, 0));
+  }
+}

test/tint/diagnostic_filtering/function_attribute.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(1) var t : texture_2d<f32>;
+@group(0) @binding(2) var s : sampler;
+
+@fragment @diagnostic(warning, derivative_uniformity)
+fn main(@location(0) x : f32) {
+  if (x > 0) {
+    _ = textureSample(t, s, vec2(0, 0));
+  }
+}

test/tint/diagnostic_filtering/function_attribute.wgsl.expected.dxc.hlsl

@@ -0,0 +1,28 @@
+diagnostic_filtering/function_attribute.wgsl:7:9 warning: 'textureSample' must only be called from uniform control flow
+    _ = textureSample(t, s, vec2(0, 0));
+        ^^^^^^^^^^^^^
+
+diagnostic_filtering/function_attribute.wgsl:6:3 note: control flow depends on possibly non-uniform value
+  if (x > 0) {
+  ^^
+
+diagnostic_filtering/function_attribute.wgsl:6:7 note: user-defined input 'x' of 'main' may be non-uniform
+  if (x > 0) {
+      ^
+
+Texture2D<float4> t : register(t1, space0);
+SamplerState s : register(s2, space0);
+
+struct tint_symbol_1 {
+  float x : TEXCOORD0;
+};
+
+void main_inner(float x) {
+  if ((x > 0.0f)) {
+  }
+}
+
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.x);
+  return;
+}

test/tint/diagnostic_filtering/function_attribute.wgsl.expected.fxc.hlsl

@@ -0,0 +1,28 @@
+diagnostic_filtering/function_attribute.wgsl:7:9 warning: 'textureSample' must only be called from uniform control flow
+    _ = textureSample(t, s, vec2(0, 0));
+        ^^^^^^^^^^^^^
+
+diagnostic_filtering/function_attribute.wgsl:6:3 note: control flow depends on possibly non-uniform value
+  if (x > 0) {
+  ^^
+
+diagnostic_filtering/function_attribute.wgsl:6:7 note: user-defined input 'x' of 'main' may be non-uniform
+  if (x > 0) {
+      ^
+
+Texture2D<float4> t : register(t1, space0);
+SamplerState s : register(s2, space0);
+
+struct tint_symbol_1 {
+  float x : TEXCOORD0;
+};
+
+void main_inner(float x) {
+  if ((x > 0.0f)) {
+  }
+}
+
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.x);
+  return;
+}

test/tint/diagnostic_filtering/function_attribute.wgsl.expected.glsl

@@ -0,0 +1,25 @@
+diagnostic_filtering/function_attribute.wgsl:7:9 warning: 'textureSample' must only be called from uniform control flow
+    _ = textureSample(t, s, vec2(0, 0));
+        ^^^^^^^^^^^^^
+
+diagnostic_filtering/function_attribute.wgsl:6:3 note: control flow depends on possibly non-uniform value
+  if (x > 0) {
+  ^^
+
+diagnostic_filtering/function_attribute.wgsl:6:7 note: user-defined input 'x' of 'main' may be non-uniform
+  if (x > 0) {
+      ^
+
+#version 310 es
+precision mediump float;
+
+layout(location = 0) in float x_1;
+void tint_symbol(float x) {
+  if ((x > 0.0f)) {
+  }
+}
+
+void main() {
+  tint_symbol(x_1);
+  return;
+}

test/tint/diagnostic_filtering/function_attribute.wgsl.expected.msl

@@ -0,0 +1,29 @@
+diagnostic_filtering/function_attribute.wgsl:7:9 warning: 'textureSample' must only be called from uniform control flow
+    _ = textureSample(t, s, vec2(0, 0));
+        ^^^^^^^^^^^^^
+
+diagnostic_filtering/function_attribute.wgsl:6:3 note: control flow depends on possibly non-uniform value
+  if (x > 0) {
+  ^^
+
+diagnostic_filtering/function_attribute.wgsl:6:7 note: user-defined input 'x' of 'main' may be non-uniform
+  if (x > 0) {
+      ^
+
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_2 {
+  float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x) {
+  if ((x > 0.0f)) {
+  }
+}
+
+fragment void tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+  tint_symbol_inner(tint_symbol_1.x);
+  return;
+}
+

test/tint/diagnostic_filtering/function_attribute.wgsl.expected.spvasm

@@ -0,0 +1,63 @@
+diagnostic_filtering/function_attribute.wgsl:7:9 warning: 'textureSample' must only be called from uniform control flow
+    _ = textureSample(t, s, vec2(0, 0));
+        ^^^^^^^^^^^^^
+
+diagnostic_filtering/function_attribute.wgsl:6:3 note: control flow depends on possibly non-uniform value
+  if (x > 0) {
+  ^^
+
+diagnostic_filtering/function_attribute.wgsl:6:7 note: user-defined input 'x' of 'main' may be non-uniform
+  if (x > 0) {
+      ^
+
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main" %x_1
+               OpExecutionMode %main OriginUpperLeft
+               OpName %x_1 "x_1"
+               OpName %t "t"
+               OpName %s "s"
+               OpName %main_inner "main_inner"
+               OpName %x "x"
+               OpName %main "main"
+               OpDecorate %x_1 Location 0
+               OpDecorate %t DescriptorSet 0
+               OpDecorate %t Binding 1
+               OpDecorate %s DescriptorSet 0
+               OpDecorate %s Binding 2
+      %float = OpTypeFloat 32
+%_ptr_Input_float = OpTypePointer Input %float
+        %x_1 = OpVariable %_ptr_Input_float Input
+          %6 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_6 = OpTypePointer UniformConstant %6
+          %t = OpVariable %_ptr_UniformConstant_6 UniformConstant
+          %9 = OpTypeSampler
+%_ptr_UniformConstant_9 = OpTypePointer UniformConstant %9
+          %s = OpVariable %_ptr_UniformConstant_9 UniformConstant
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void %float
+         %15 = OpConstantNull %float
+       %bool = OpTypeBool
+         %20 = OpTypeFunction %void
+ %main_inner = OpFunction %void None %10
+          %x = OpFunctionParameter %float
+         %14 = OpLabel
+         %16 = OpFOrdGreaterThan %bool %x %15
+               OpSelectionMerge %18 None
+               OpBranchConditional %16 %19 %18
+         %19 = OpLabel
+               OpBranch %18
+         %18 = OpLabel
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %20
+         %22 = OpLabel
+         %24 = OpLoad %float %x_1
+         %23 = OpFunctionCall %void %main_inner %24
+               OpReturn
+               OpFunctionEnd

test/tint/diagnostic_filtering/function_attribute.wgsl.expected.wgsl

@@ -0,0 +1,22 @@
+diagnostic_filtering/function_attribute.wgsl:7:9 warning: 'textureSample' must only be called from uniform control flow
+    _ = textureSample(t, s, vec2(0, 0));
+        ^^^^^^^^^^^^^
+
+diagnostic_filtering/function_attribute.wgsl:6:3 note: control flow depends on possibly non-uniform value
+  if (x > 0) {
+  ^^
+
+diagnostic_filtering/function_attribute.wgsl:6:7 note: user-defined input 'x' of 'main' may be non-uniform
+  if (x > 0) {
+      ^
+
+@group(0) @binding(1) var t : texture_2d<f32>;
+
+@group(0) @binding(2) var s : sampler;
+
+@fragment @diagnostic(warning, derivative_uniformity)
+fn main(@location(0) x : f32) {
+  if ((x > 0)) {
+    _ = textureSample(t, s, vec2(0, 0));
+  }
+}