From f8fb94d9e8f5b8a96359d0001493209e16e33058 Mon Sep 17 00:00:00 2001 From: Zhaoming Jiang Date: Wed, 27 Jul 2022 13:27:56 +0000 Subject: [PATCH] tint: Emitting enable directive for DP4A End-to-End test This patch modify the test/tint/builtins/gen/gen.wgsl.tmpl to emit enable directive for dot4I8Packed and dot4U8Packed built-in function. The expectaion files are added. Bug: tint:1497 Change-Id: I53331695fe2e6609858e94bc261383ba3028d77c Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/96640 Reviewed-by: Ben Clayton Kokoro: Kokoro Commit-Queue: Zhaoming Jiang --- test/tint/builtins/gen/gen.wgsl.tmpl | 20 ++++- .../gen/literal/dot4I8Packed/881e62.wgsl | 45 +++++++++++ .../881e62.wgsl.expected.dxc.hlsl | 35 ++++++++ .../881e62.wgsl.expected.fxc.hlsl | 1 + .../dot4I8Packed/881e62.wgsl.expected.glsl | 1 + .../dot4I8Packed/881e62.wgsl.expected.msl | 1 + .../dot4I8Packed/881e62.wgsl.expected.spvasm | 71 ++++++++++++++++ .../dot4I8Packed/881e62.wgsl.expected.wgsl | 21 +++++ .../gen/literal/dot4U8Packed/fbed7b.wgsl | 45 +++++++++++ .../fbed7b.wgsl.expected.dxc.hlsl | 35 ++++++++ .../fbed7b.wgsl.expected.fxc.hlsl | 1 + .../dot4U8Packed/fbed7b.wgsl.expected.glsl | 1 + .../dot4U8Packed/fbed7b.wgsl.expected.msl | 1 + .../dot4U8Packed/fbed7b.wgsl.expected.spvasm | 70 ++++++++++++++++ .../dot4U8Packed/fbed7b.wgsl.expected.wgsl | 21 +++++ .../builtins/gen/var/dot4I8Packed/881e62.wgsl | 47 +++++++++++ .../881e62.wgsl.expected.dxc.hlsl | 37 +++++++++ .../881e62.wgsl.expected.fxc.hlsl | 1 + .../dot4I8Packed/881e62.wgsl.expected.glsl | 1 + .../var/dot4I8Packed/881e62.wgsl.expected.msl | 1 + .../dot4I8Packed/881e62.wgsl.expected.spvasm | 81 +++++++++++++++++++ .../dot4I8Packed/881e62.wgsl.expected.wgsl | 23 ++++++ .../builtins/gen/var/dot4U8Packed/fbed7b.wgsl | 47 +++++++++++ .../fbed7b.wgsl.expected.dxc.hlsl | 37 +++++++++ .../fbed7b.wgsl.expected.fxc.hlsl | 1 + .../dot4U8Packed/fbed7b.wgsl.expected.glsl | 1 + .../var/dot4U8Packed/fbed7b.wgsl.expected.msl | 1 + .../dot4U8Packed/fbed7b.wgsl.expected.spvasm | 78 ++++++++++++++++++ .../dot4U8Packed/fbed7b.wgsl.expected.wgsl | 23 ++++++ tools/src/tint/intrinsic/resolver/resolve.go | 2 +- 30 files changed, 746 insertions(+), 4 deletions(-) create mode 100644 test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl create mode 100644 test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl create mode 100644 test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl create mode 100644 test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl create mode 100644 test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.wgsl diff --git a/test/tint/builtins/gen/gen.wgsl.tmpl b/test/tint/builtins/gen/gen.wgsl.tmpl index ce0f5cfccd..0c80325b15 100644 --- a/test/tint/builtins/gen/gen.wgsl.tmpl +++ b/test/tint/builtins/gen/gen.wgsl.tmpl @@ -11,8 +11,6 @@ See: {{- /* For each permutation of each overload of each function... */ -}} {{- range Sem.Builtins -}} -{{- /* TODO(crbug.com/tint/1536): Remove the bodge below after we support [[extension("extension_name")]] in intrinsics.def */}} -{{- if not (or (eq .Name "dot4I8Packed") (eq .Name "dot4U8Packed"))}} {{- range .Overloads -}} {{- range Permute . -}} {{- /* Generate a ./literal//.wgsl file using @@ -28,7 +26,6 @@ See: {{- end }} {{- end }} {{- end }} -{{- end }} {{- /* ------------------------------------------------------------------ */ -}} {{- define "Permutation" -}} @@ -40,6 +37,9 @@ See: {{- $permutation := printf "%v_%v" $builtin $overload.Hash -}} {{- $args := Map -}} +{{- /* Generate enable directives */ -}} +{{- template "EnableDirectives" $overload -}} + {{- /* Generate RW storage buffer parameters */ -}} {{- $sb_rw_fields := Eval "EmitBufferFields" "overload" $overload "var_name" "sb_rw" @@ -170,6 +170,20 @@ fn compute_main() { {{- end -}} +{{- /* ------------------------------------------------------------------ */ -}} +{{- define "EnableDirectives" -}} +{{- /* Emits the required enable directives for a given overload */ -}} +{{- /* ------------------------------------------------------------------ */ -}} +{{- $overload := . -}} +{{- $builtin_name := $overload.Intrinsic.Name -}} + +{{- /* Check and emit chromium_experimental_dp4a */ -}} +{{- if or (eq "dot4I8Packed" $builtin_name) (eq "dot4U8Packed" $builtin_name)}} +enable chromium_experimental_dp4a; +{{ end -}} + +{{- end -}} + {{- /* ------------------------------------------------------------------ */ -}} {{- define "EmitBufferFields" -}} diff --git a/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl new file mode 100644 index 0000000000..dafb8b6fc9 --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl @@ -0,0 +1,45 @@ +// Copyright 2021 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +//////////////////////////////////////////////////////////////////////////////// +// File generated by tools/src/cmd/gen +// using the template: +// test/tint/builtins/gen/gen.wgsl.tmpl +// +// Do not modify this file directly +//////////////////////////////////////////////////////////////////////////////// + + +enable chromium_experimental_dp4a; + +// fn dot4I8Packed(u32, u32) -> i32 +fn dot4I8Packed_881e62() { + var res: i32 = dot4I8Packed(1u, 1u); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + dot4I8Packed_881e62(); + return vec4(); +} + +@fragment +fn fragment_main() { + dot4I8Packed_881e62(); +} + +@compute @workgroup_size(1) +fn compute_main() { + dot4I8Packed_881e62(); +} diff --git a/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..3ba498587a --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.dxc.hlsl @@ -0,0 +1,35 @@ +int tint_dot4I8Packed(uint param_0, uint param_1) { + int accumulator = 0; + return dot4add_i8packed(param_0, param_1, accumulator); +} + +void dot4I8Packed_881e62() { + int res = tint_dot4I8Packed(1u, 1u); +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + dot4I8Packed_881e62(); + return (0.0f).xxxx; +} + +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() { + dot4I8Packed_881e62(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + dot4I8Packed_881e62(); + return; +} diff --git a/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..9f06383b3c --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.fxc.hlsl @@ -0,0 +1 @@ +SKIP \ No newline at end of file diff --git a/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.glsl b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.glsl new file mode 100644 index 0000000000..9f06383b3c --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.glsl @@ -0,0 +1 @@ +SKIP \ No newline at end of file diff --git a/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.msl b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.msl new file mode 100644 index 0000000000..9f06383b3c --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.msl @@ -0,0 +1 @@ +SKIP \ No newline at end of file diff --git a/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.spvasm new file mode 100644 index 0000000000..d71f31ef3d --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.spvasm @@ -0,0 +1,71 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 34 +; Schema: 0 + OpCapability Shader + OpCapability DotProduct + OpCapability DotProductInput4x8BitPacked + OpExtension "SPV_KHR_integer_dot_product" + 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 %dot4I8Packed_881e62 "dot4I8Packed_881e62" + 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 + %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 + %void = OpTypeVoid + %9 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %uint = OpTypeInt 32 0 + %uint_1 = OpConstant %uint 1 +%_ptr_Function_int = OpTypePointer Function %int + %19 = OpConstantNull %int + %20 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%dot4I8Packed_881e62 = OpFunction %void None %9 + %12 = OpLabel + %res = OpVariable %_ptr_Function_int Function %19 + %13 = OpSDot %int %uint_1 %uint_1 PackedVectorFormat4x8Bit + OpStore %res %13 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %20 + %22 = OpLabel + %23 = OpFunctionCall %void %dot4I8Packed_881e62 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %9 + %25 = OpLabel + %26 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %26 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %9 + %29 = OpLabel + %30 = OpFunctionCall %void %dot4I8Packed_881e62 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %9 + %32 = OpLabel + %33 = OpFunctionCall %void %dot4I8Packed_881e62 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.wgsl new file mode 100644 index 0000000000..8057884b8e --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.wgsl @@ -0,0 +1,21 @@ +enable chromium_experimental_dp4a; + +fn dot4I8Packed_881e62() { + var res : i32 = dot4I8Packed(1u, 1u); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + dot4I8Packed_881e62(); + return vec4(); +} + +@fragment +fn fragment_main() { + dot4I8Packed_881e62(); +} + +@compute @workgroup_size(1) +fn compute_main() { + dot4I8Packed_881e62(); +} diff --git a/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl new file mode 100644 index 0000000000..da99ad289f --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl @@ -0,0 +1,45 @@ +// Copyright 2021 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +//////////////////////////////////////////////////////////////////////////////// +// File generated by tools/src/cmd/gen +// using the template: +// test/tint/builtins/gen/gen.wgsl.tmpl +// +// Do not modify this file directly +//////////////////////////////////////////////////////////////////////////////// + + +enable chromium_experimental_dp4a; + +// fn dot4U8Packed(u32, u32) -> u32 +fn dot4U8Packed_fbed7b() { + var res: u32 = dot4U8Packed(1u, 1u); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + dot4U8Packed_fbed7b(); + return vec4(); +} + +@fragment +fn fragment_main() { + dot4U8Packed_fbed7b(); +} + +@compute @workgroup_size(1) +fn compute_main() { + dot4U8Packed_fbed7b(); +} diff --git a/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..ffd98789e1 --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.dxc.hlsl @@ -0,0 +1,35 @@ +uint tint_dot4U8Packed(uint param_0, uint param_1) { + uint accumulator = 0u; + return dot4add_u8packed(param_0, param_1, accumulator); +} + +void dot4U8Packed_fbed7b() { + uint res = tint_dot4U8Packed(1u, 1u); +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + dot4U8Packed_fbed7b(); + return (0.0f).xxxx; +} + +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() { + dot4U8Packed_fbed7b(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + dot4U8Packed_fbed7b(); + return; +} diff --git a/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..9f06383b3c --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.fxc.hlsl @@ -0,0 +1 @@ +SKIP \ No newline at end of file diff --git a/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.glsl b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.glsl new file mode 100644 index 0000000000..9f06383b3c --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.glsl @@ -0,0 +1 @@ +SKIP \ No newline at end of file diff --git a/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.msl b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.msl new file mode 100644 index 0000000000..9f06383b3c --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.msl @@ -0,0 +1 @@ +SKIP \ No newline at end of file diff --git a/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.spvasm new file mode 100644 index 0000000000..855730c0d1 --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.spvasm @@ -0,0 +1,70 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 33 +; Schema: 0 + OpCapability Shader + OpCapability DotProduct + OpCapability DotProductInput4x8BitPacked + OpExtension "SPV_KHR_integer_dot_product" + 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 %dot4U8Packed_fbed7b "dot4U8Packed_fbed7b" + 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 + %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 + %void = OpTypeVoid + %9 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_1 = OpConstant %uint 1 +%_ptr_Function_uint = OpTypePointer Function %uint + %18 = OpConstantNull %uint + %19 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%dot4U8Packed_fbed7b = OpFunction %void None %9 + %12 = OpLabel + %res = OpVariable %_ptr_Function_uint Function %18 + %13 = OpUDot %uint %uint_1 %uint_1 PackedVectorFormat4x8Bit + OpStore %res %13 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %19 + %21 = OpLabel + %22 = OpFunctionCall %void %dot4U8Packed_fbed7b + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %9 + %24 = OpLabel + %25 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %25 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %9 + %28 = OpLabel + %29 = OpFunctionCall %void %dot4U8Packed_fbed7b + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %9 + %31 = OpLabel + %32 = OpFunctionCall %void %dot4U8Packed_fbed7b + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.wgsl new file mode 100644 index 0000000000..f261301ee1 --- /dev/null +++ b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.wgsl @@ -0,0 +1,21 @@ +enable chromium_experimental_dp4a; + +fn dot4U8Packed_fbed7b() { + var res : u32 = dot4U8Packed(1u, 1u); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + dot4U8Packed_fbed7b(); + return vec4(); +} + +@fragment +fn fragment_main() { + dot4U8Packed_fbed7b(); +} + +@compute @workgroup_size(1) +fn compute_main() { + dot4U8Packed_fbed7b(); +} diff --git a/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl new file mode 100644 index 0000000000..73fd181ff1 --- /dev/null +++ b/test/tint/builtins/gen/var/dot4I8Packed/881e62.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/src/cmd/gen +// using the template: +// test/tint/builtins/gen/gen.wgsl.tmpl +// +// Do not modify this file directly +//////////////////////////////////////////////////////////////////////////////// + + +enable chromium_experimental_dp4a; + +// fn dot4I8Packed(u32, u32) -> i32 +fn dot4I8Packed_881e62() { + var arg_0 = 1u; + var arg_1 = 1u; + var res: i32 = dot4I8Packed(arg_0, arg_1); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + dot4I8Packed_881e62(); + return vec4(); +} + +@fragment +fn fragment_main() { + dot4I8Packed_881e62(); +} + +@compute @workgroup_size(1) +fn compute_main() { + dot4I8Packed_881e62(); +} diff --git a/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..16555b7136 --- /dev/null +++ b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.dxc.hlsl @@ -0,0 +1,37 @@ +int tint_dot4I8Packed(uint param_0, uint param_1) { + int accumulator = 0; + return dot4add_i8packed(param_0, param_1, accumulator); +} + +void dot4I8Packed_881e62() { + uint arg_0 = 1u; + uint arg_1 = 1u; + int res = tint_dot4I8Packed(arg_0, arg_1); +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + dot4I8Packed_881e62(); + return (0.0f).xxxx; +} + +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() { + dot4I8Packed_881e62(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + dot4I8Packed_881e62(); + return; +} diff --git a/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..9f06383b3c --- /dev/null +++ b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.fxc.hlsl @@ -0,0 +1 @@ +SKIP \ No newline at end of file diff --git a/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.glsl b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.glsl new file mode 100644 index 0000000000..9f06383b3c --- /dev/null +++ b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.glsl @@ -0,0 +1 @@ +SKIP \ No newline at end of file diff --git a/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.msl b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.msl new file mode 100644 index 0000000000..9f06383b3c --- /dev/null +++ b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.msl @@ -0,0 +1 @@ +SKIP \ No newline at end of file diff --git a/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.spvasm b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.spvasm new file mode 100644 index 0000000000..075bf9bf0f --- /dev/null +++ b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.spvasm @@ -0,0 +1,81 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 40 +; Schema: 0 + OpCapability Shader + OpCapability DotProduct + OpCapability DotProductInput4x8BitPacked + OpExtension "SPV_KHR_integer_dot_product" + 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 %dot4I8Packed_881e62 "dot4I8Packed_881e62" + OpName %arg_0 "arg_0" + OpName %arg_1 "arg_1" + 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 + %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 + %void = OpTypeVoid + %9 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_1 = OpConstant %uint 1 +%_ptr_Function_uint = OpTypePointer Function %uint + %17 = OpConstantNull %uint + %int = OpTypeInt 32 1 +%_ptr_Function_int = OpTypePointer Function %int + %25 = OpConstantNull %int + %26 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%dot4I8Packed_881e62 = OpFunction %void None %9 + %12 = OpLabel + %arg_0 = OpVariable %_ptr_Function_uint Function %17 + %arg_1 = OpVariable %_ptr_Function_uint Function %17 + %res = OpVariable %_ptr_Function_int Function %25 + OpStore %arg_0 %uint_1 + OpStore %arg_1 %uint_1 + %21 = OpLoad %uint %arg_0 + %22 = OpLoad %uint %arg_1 + %19 = OpSDot %int %21 %22 PackedVectorFormat4x8Bit + OpStore %res %19 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %26 + %28 = OpLabel + %29 = OpFunctionCall %void %dot4I8Packed_881e62 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %9 + %31 = OpLabel + %32 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %32 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %9 + %35 = OpLabel + %36 = OpFunctionCall %void %dot4I8Packed_881e62 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %9 + %38 = OpLabel + %39 = OpFunctionCall %void %dot4I8Packed_881e62 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.wgsl b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.wgsl new file mode 100644 index 0000000000..ba5f0a753f --- /dev/null +++ b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.wgsl @@ -0,0 +1,23 @@ +enable chromium_experimental_dp4a; + +fn dot4I8Packed_881e62() { + var arg_0 = 1u; + var arg_1 = 1u; + var res : i32 = dot4I8Packed(arg_0, arg_1); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + dot4I8Packed_881e62(); + return vec4(); +} + +@fragment +fn fragment_main() { + dot4I8Packed_881e62(); +} + +@compute @workgroup_size(1) +fn compute_main() { + dot4I8Packed_881e62(); +} diff --git a/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl new file mode 100644 index 0000000000..8c90e74c5e --- /dev/null +++ b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.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/src/cmd/gen +// using the template: +// test/tint/builtins/gen/gen.wgsl.tmpl +// +// Do not modify this file directly +//////////////////////////////////////////////////////////////////////////////// + + +enable chromium_experimental_dp4a; + +// fn dot4U8Packed(u32, u32) -> u32 +fn dot4U8Packed_fbed7b() { + var arg_0 = 1u; + var arg_1 = 1u; + var res: u32 = dot4U8Packed(arg_0, arg_1); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + dot4U8Packed_fbed7b(); + return vec4(); +} + +@fragment +fn fragment_main() { + dot4U8Packed_fbed7b(); +} + +@compute @workgroup_size(1) +fn compute_main() { + dot4U8Packed_fbed7b(); +} diff --git a/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..4b66aa9fbb --- /dev/null +++ b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.dxc.hlsl @@ -0,0 +1,37 @@ +uint tint_dot4U8Packed(uint param_0, uint param_1) { + uint accumulator = 0u; + return dot4add_u8packed(param_0, param_1, accumulator); +} + +void dot4U8Packed_fbed7b() { + uint arg_0 = 1u; + uint arg_1 = 1u; + uint res = tint_dot4U8Packed(arg_0, arg_1); +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + dot4U8Packed_fbed7b(); + return (0.0f).xxxx; +} + +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() { + dot4U8Packed_fbed7b(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + dot4U8Packed_fbed7b(); + return; +} diff --git a/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..9f06383b3c --- /dev/null +++ b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.fxc.hlsl @@ -0,0 +1 @@ +SKIP \ No newline at end of file diff --git a/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.glsl b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.glsl new file mode 100644 index 0000000000..9f06383b3c --- /dev/null +++ b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.glsl @@ -0,0 +1 @@ +SKIP \ No newline at end of file diff --git a/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.msl b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.msl new file mode 100644 index 0000000000..9f06383b3c --- /dev/null +++ b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.msl @@ -0,0 +1 @@ +SKIP \ No newline at end of file diff --git a/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.spvasm b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.spvasm new file mode 100644 index 0000000000..b0a4f48f0e --- /dev/null +++ b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.spvasm @@ -0,0 +1,78 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 37 +; Schema: 0 + OpCapability Shader + OpCapability DotProduct + OpCapability DotProductInput4x8BitPacked + OpExtension "SPV_KHR_integer_dot_product" + 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 %dot4U8Packed_fbed7b "dot4U8Packed_fbed7b" + OpName %arg_0 "arg_0" + OpName %arg_1 "arg_1" + 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 + %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 + %void = OpTypeVoid + %9 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_1 = OpConstant %uint 1 +%_ptr_Function_uint = OpTypePointer Function %uint + %17 = OpConstantNull %uint + %23 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%dot4U8Packed_fbed7b = OpFunction %void None %9 + %12 = OpLabel + %arg_0 = OpVariable %_ptr_Function_uint Function %17 + %arg_1 = OpVariable %_ptr_Function_uint Function %17 + %res = OpVariable %_ptr_Function_uint Function %17 + OpStore %arg_0 %uint_1 + OpStore %arg_1 %uint_1 + %20 = OpLoad %uint %arg_0 + %21 = OpLoad %uint %arg_1 + %19 = OpUDot %uint %20 %21 PackedVectorFormat4x8Bit + OpStore %res %19 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %23 + %25 = OpLabel + %26 = OpFunctionCall %void %dot4U8Packed_fbed7b + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %9 + %28 = OpLabel + %29 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %29 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %9 + %32 = OpLabel + %33 = OpFunctionCall %void %dot4U8Packed_fbed7b + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %9 + %35 = OpLabel + %36 = OpFunctionCall %void %dot4U8Packed_fbed7b + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.wgsl b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.wgsl new file mode 100644 index 0000000000..dc9897f162 --- /dev/null +++ b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.wgsl @@ -0,0 +1,23 @@ +enable chromium_experimental_dp4a; + +fn dot4U8Packed_fbed7b() { + var arg_0 = 1u; + var arg_1 = 1u; + var res : u32 = dot4U8Packed(arg_0, arg_1); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + dot4U8Packed_fbed7b(); + return vec4(); +} + +@fragment +fn fragment_main() { + dot4U8Packed_fbed7b(); +} + +@compute @workgroup_size(1) +fn compute_main() { + dot4U8Packed_fbed7b(); +} diff --git a/tools/src/tint/intrinsic/resolver/resolve.go b/tools/src/tint/intrinsic/resolver/resolve.go index 385ce4298e..da03428029 100644 --- a/tools/src/tint/intrinsic/resolver/resolve.go +++ b/tools/src/tint/intrinsic/resolver/resolve.go @@ -537,7 +537,7 @@ func (r *resolver) lookupNamed(s *scope, a ast.TemplatedName) (sem.Named, error) case sem.Named: return target, nil default: - panic(fmt.Errorf("Unknown resolved type %T", target)) + panic(fmt.Errorf("unknown resolved type %T", target)) } // ... and that something takes template parameters // Check the number of templated name template arguments match the number of