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 <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Zhaoming Jiang <zhaoming.jiang@intel.com>
diff --git a/test/tint/builtins/gen/gen.wgsl.tmpl b/test/tint/builtins/gen/gen.wgsl.tmpl
index ce0f5cf..0c80325 100644
--- a/test/tint/builtins/gen/gen.wgsl.tmpl
+++ b/test/tint/builtins/gen/gen.wgsl.tmpl
@@ -11,8 +11,6 @@
{{- /* 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/<function>/<permuataion-hash>.wgsl file using
@@ -28,7 +26,6 @@
{{- end }}
{{- end }}
{{- end }}
-{{- end }}
{{- /* ------------------------------------------------------------------ */ -}}
{{- define "Permutation" -}}
@@ -40,6 +37,9 @@
{{- $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 @@
{{- 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 0000000..dafb8b6
--- /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<f32> {
+ dot4I8Packed_881e62();
+ return vec4<f32>();
+}
+
+@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 0000000..3ba4985
--- /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 0000000..9f06383
--- /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 0000000..9f06383
--- /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 0000000..9f06383
--- /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 0000000..d71f31e
--- /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 0000000..8057884
--- /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<f32> {
+ dot4I8Packed_881e62();
+ return vec4<f32>();
+}
+
+@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 0000000..da99ad2
--- /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<f32> {
+ dot4U8Packed_fbed7b();
+ return vec4<f32>();
+}
+
+@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 0000000..ffd9878
--- /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 0000000..9f06383
--- /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 0000000..9f06383
--- /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 0000000..9f06383
--- /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 0000000..855730c
--- /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 0000000..f261301
--- /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<f32> {
+ dot4U8Packed_fbed7b();
+ return vec4<f32>();
+}
+
+@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 0000000..73fd181
--- /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<f32> {
+ dot4I8Packed_881e62();
+ return vec4<f32>();
+}
+
+@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 0000000..16555b7
--- /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 0000000..9f06383
--- /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 0000000..9f06383
--- /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 0000000..9f06383
--- /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 0000000..075bf9b
--- /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 0000000..ba5f0a7
--- /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<f32> {
+ dot4I8Packed_881e62();
+ return vec4<f32>();
+}
+
+@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 0000000..8c90e74
--- /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<f32> {
+ dot4U8Packed_fbed7b();
+ return vec4<f32>();
+}
+
+@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 0000000..4b66aa9
--- /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 0000000..9f06383
--- /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 0000000..9f06383
--- /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 0000000..9f06383
--- /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 0000000..b0a4f48
--- /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 0000000..dc9897f
--- /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<f32> {
+ dot4U8Packed_fbed7b();
+ return vec4<f32>();
+}
+
+@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 385ce42..da03428 100644
--- a/tools/src/tint/intrinsic/resolver/resolve.go
+++ b/tools/src/tint/intrinsic/resolver/resolve.go
@@ -537,7 +537,7 @@
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
