diff --git a/src/tint/transform/single_entry_point.cc b/src/tint/transform/single_entry_point.cc index f76eda9ef2..698b2e9f33 100644 --- a/src/tint/transform/single_entry_point.cc +++ b/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; diff --git a/src/tint/transform/single_entry_point_test.cc b/src/tint/transform/single_entry_point_test.cc index b8f79f6bba..d0c3dd4178 100644 --- a/src/tint/transform/single_entry_point_test.cc +++ b/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(cfg); + auto got = Run(src, data); + + EXPECT_EQ(src, str(got)); +} + } // namespace } // namespace tint::transform diff --git a/test/tint/diagnostic_filtering/directive.wgsl b/test/tint/diagnostic_filtering/directive.wgsl new file mode 100644 index 0000000000..f720e61fdd --- /dev/null +++ b/test/tint/diagnostic_filtering/directive.wgsl @@ -0,0 +1,11 @@ +diagnostic(warning, derivative_uniformity); + +@group(0) @binding(1) var t : texture_2d; +@group(0) @binding(2) var s : sampler; + +@fragment +fn main(@location(0) x : f32) { + if (x > 0) { + _ = textureSample(t, s, vec2(0, 0)); + } +} diff --git a/test/tint/diagnostic_filtering/directive.wgsl.expected.dxc.hlsl b/test/tint/diagnostic_filtering/directive.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..fe56ba4723 --- /dev/null +++ b/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 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; +} diff --git a/test/tint/diagnostic_filtering/directive.wgsl.expected.fxc.hlsl b/test/tint/diagnostic_filtering/directive.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..fe56ba4723 --- /dev/null +++ b/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 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; +} diff --git a/test/tint/diagnostic_filtering/directive.wgsl.expected.glsl b/test/tint/diagnostic_filtering/directive.wgsl.expected.glsl new file mode 100644 index 0000000000..6b4086e196 --- /dev/null +++ b/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; +} diff --git a/test/tint/diagnostic_filtering/directive.wgsl.expected.msl b/test/tint/diagnostic_filtering/directive.wgsl.expected.msl new file mode 100644 index 0000000000..d9392672a7 --- /dev/null +++ b/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 + +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; +} + diff --git a/test/tint/diagnostic_filtering/directive.wgsl.expected.spvasm b/test/tint/diagnostic_filtering/directive.wgsl.expected.spvasm new file mode 100644 index 0000000000..ecb0141509 --- /dev/null +++ b/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 diff --git a/test/tint/diagnostic_filtering/directive.wgsl.expected.wgsl b/test/tint/diagnostic_filtering/directive.wgsl.expected.wgsl new file mode 100644 index 0000000000..b6eb1d784a --- /dev/null +++ b/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; + +@group(0) @binding(2) var s : sampler; + +@fragment +fn main(@location(0) x : f32) { + if ((x > 0)) { + _ = textureSample(t, s, vec2(0, 0)); + } +} diff --git a/test/tint/diagnostic_filtering/function_attribute.wgsl b/test/tint/diagnostic_filtering/function_attribute.wgsl new file mode 100644 index 0000000000..e52aa39120 --- /dev/null +++ b/test/tint/diagnostic_filtering/function_attribute.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(1) var t : texture_2d; +@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)); + } +} diff --git a/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.dxc.hlsl b/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..dfe948a904 --- /dev/null +++ b/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 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; +} diff --git a/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.fxc.hlsl b/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..dfe948a904 --- /dev/null +++ b/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 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; +} diff --git a/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.glsl b/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.glsl new file mode 100644 index 0000000000..8cd0f5ac70 --- /dev/null +++ b/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; +} diff --git a/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.msl new file mode 100644 index 0000000000..cee8902f51 --- /dev/null +++ b/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 + +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; +} + diff --git a/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.spvasm b/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.spvasm new file mode 100644 index 0000000000..1ba99cae0a --- /dev/null +++ b/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 diff --git a/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.wgsl b/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.wgsl new file mode 100644 index 0000000000..8b3760d486 --- /dev/null +++ b/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; + +@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)); + } +}