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>
This commit is contained in:
Zhaoming Jiang
Dawn LUCI CQ
parent
6e9a9bf974
commit
f8fb94d9e8
|
|
@ -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/<function>/<permuataion-hash>.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" -}}
|
||||
|
|
|
|||
|
|
@ -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();
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -0,0 +1 @@
|
|||
SKIP
|
||||
|
|
@ -0,0 +1 @@
|
|||
SKIP
|
||||
|
|
@ -0,0 +1 @@
|
|||
SKIP
|
||||
|
|
@ -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
|
||||
|
|
@ -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();
|
||||
}
|
||||
|
|
@ -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();
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -0,0 +1 @@
|
|||
SKIP
|
||||
|
|
@ -0,0 +1 @@
|
|||
SKIP
|
||||
|
|
@ -0,0 +1 @@
|
|||
SKIP
|
||||
|
|
@ -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
|
||||
|
|
@ -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();
|
||||
}
|
||||
|
|
@ -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();
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -0,0 +1 @@
|
|||
SKIP
|
||||
|
|
@ -0,0 +1 @@
|
|||
SKIP
|
||||
|
|
@ -0,0 +1 @@
|
|||
SKIP
|
||||
|
|
@ -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
|
||||
|
|
@ -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();
|
||||
}
|
||||
|
|
@ -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();
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -0,0 +1 @@
|
|||
SKIP
|
||||
|
|
@ -0,0 +1 @@
|
|||
SKIP
|
||||
|
|
@ -0,0 +1 @@
|
|||
SKIP
|
||||
|
|
@ -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
|
||||
|
|
@ -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();
|
||||
}
|
||||
|
|
@ -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
|
||||
|
|
|
|||
Loading…
Reference in New Issue