OpenGL: Support dot4I8Packed() and dot4U8Packed()
This patch adds the support of `dot4I8Packed()` and `dot4U8Packed()`
on OpenGL backend with their polyfills.
This patch also removes the toggle `PolyFillPacked4x8DotProduct` on
D3D11 backends as on D3D11 `dot4I8Packed()` and `dot4U8Packed()` must
be polyfilled.
Bug: tint:1497
Test: dawn_end2end_tests
Change-Id: Iae81b4313ca699cbae6209b66c4371e3879af2e2
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/163702
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Jiawei Shao <jiawei.shao@intel.com>
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: Austin Eng <enga@chromium.org>
diff --git a/src/dawn/native/d3d11/PhysicalDeviceD3D11.cpp b/src/dawn/native/d3d11/PhysicalDeviceD3D11.cpp
index 4246a46..456d7e6 100644
--- a/src/dawn/native/d3d11/PhysicalDeviceD3D11.cpp
+++ b/src/dawn/native/d3d11/PhysicalDeviceD3D11.cpp
@@ -307,8 +307,6 @@
// D3D11 can only clear RTV with float values.
deviceToggles->Default(Toggle::ApplyClearBigIntegerColorValueWithDraw, true);
deviceToggles->Default(Toggle::UseBlitForBufferToStencilTextureCopy, true);
- // D3D11 doesn't support shader model 6+ features
- deviceToggles->ForceSet(Toggle::PolyFillPacked4x8DotProduct, true);
}
ResultOrError<Ref<DeviceBase>> PhysicalDevice::CreateDeviceImpl(AdapterBase* adapter,
diff --git a/src/dawn/native/d3d11/ShaderModuleD3D11.cpp b/src/dawn/native/d3d11/ShaderModuleD3D11.cpp
index c0bb9a9..ef9fe91 100644
--- a/src/dawn/native/d3d11/ShaderModuleD3D11.cpp
+++ b/src/dawn/native/d3d11/ShaderModuleD3D11.cpp
@@ -210,7 +210,7 @@
// TODO(dawn:1705): do we need to support it?
req.hlsl.tintOptions.polyfill_reflect_vec2_f32 = false;
- DAWN_ASSERT(device->IsToggleEnabled(Toggle::PolyFillPacked4x8DotProduct));
+ // D3D11 doesn't support shader model 6+ features
req.hlsl.tintOptions.polyfill_dot_4x8_packed = true;
CacheResult<d3d::CompiledShader> compiledShader;
diff --git a/src/dawn/native/opengl/PhysicalDeviceGL.cpp b/src/dawn/native/opengl/PhysicalDeviceGL.cpp
index d78a66f..28c5147 100644
--- a/src/dawn/native/opengl/PhysicalDeviceGL.cpp
+++ b/src/dawn/native/opengl/PhysicalDeviceGL.cpp
@@ -236,6 +236,8 @@
if (mFunctions.IsGLExtensionSupported("GL_EXT_texture_norm16")) {
EnableFeature(Feature::Norm16TextureFormats);
}
+
+ EnableFeature(Feature::ChromiumExperimentalDp4a);
}
namespace {
diff --git a/src/dawn/tests/end2end/ExperimentalDP4aTests.cpp b/src/dawn/tests/end2end/ExperimentalDP4aTests.cpp
index 18d7626..d4df4ea 100644
--- a/src/dawn/tests/end2end/ExperimentalDP4aTests.cpp
+++ b/src/dawn/tests/end2end/ExperimentalDP4aTests.cpp
@@ -134,6 +134,8 @@
D3D12Backend({}, {"use_dxc"}),
D3D12Backend({"polyfill_packed_4x8_dot_product"}),
MetalBackend(),
+ OpenGLBackend(),
+ OpenGLESBackend(),
VulkanBackend(),
},
{true, false});
diff --git a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
index 9694635..cbd2226 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
@@ -189,6 +189,7 @@
polyfills.saturate = true;
polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
polyfills.workgroup_uniform_load = true;
+ polyfills.dot_4x8_packed = true;
data.Add<ast::transform::BuiltinPolyfill::Config>(polyfills);
manager.Add<ast::transform::BuiltinPolyfill>(); // Must come before DirectVariableAccess
}
diff --git a/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.glsl b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.glsl
index ee99c12..3f051af 100644
--- a/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.glsl
@@ -1,83 +1,94 @@
-SKIP: FAILED
+#version 310 es
-
-enable chromium_experimental_dp4a;
-
-fn dot4I8Packed_881e62() {
- var res : i32 = dot4I8Packed(1u, 1u);
- prevent_dce = res;
+int tint_int_dot(ivec4 a, ivec4 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : i32;
+int tint_dot4_i8_packed(uint a, uint b) {
+ ivec4 a_i8 = (ivec4((uvec4(a) << uvec4(24u, 16u, 8u, 0u))) >> uvec4(24u));
+ ivec4 b_i8 = (ivec4((uvec4(b) << uvec4(24u, 16u, 8u, 0u))) >> uvec4(24u));
+ return tint_int_dot(a_i8, b_i8);
+}
-@vertex
-fn vertex_main() -> @builtin(position) vec4<f32> {
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ int inner;
+} prevent_dce;
+
+void dot4I8Packed_881e62() {
+ int res = tint_dot4_i8_packed(1u, 1u);
+ prevent_dce.inner = res;
+}
+
+vec4 vertex_main() {
dot4I8Packed_881e62();
- return vec4<f32>();
+ return vec4(0.0f);
}
-@fragment
-fn fragment_main() {
+void main() {
+ gl_PointSize = 1.0;
+ vec4 inner_result = vertex_main();
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
+#version 310 es
+precision highp float;
+
+int tint_int_dot(ivec4 a, ivec4 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
+}
+
+int tint_dot4_i8_packed(uint a, uint b) {
+ ivec4 a_i8 = (ivec4((uvec4(a) << uvec4(24u, 16u, 8u, 0u))) >> uvec4(24u));
+ ivec4 b_i8 = (ivec4((uvec4(b) << uvec4(24u, 16u, 8u, 0u))) >> uvec4(24u));
+ return tint_int_dot(a_i8, b_i8);
+}
+
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ int inner;
+} prevent_dce;
+
+void dot4I8Packed_881e62() {
+ int res = tint_dot4_i8_packed(1u, 1u);
+ prevent_dce.inner = res;
+}
+
+void fragment_main() {
dot4I8Packed_881e62();
}
-@compute @workgroup_size(1)
-fn compute_main() {
+void main() {
+ fragment_main();
+ return;
+}
+#version 310 es
+
+int tint_int_dot(ivec4 a, ivec4 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
+}
+
+int tint_dot4_i8_packed(uint a, uint b) {
+ ivec4 a_i8 = (ivec4((uvec4(a) << uvec4(24u, 16u, 8u, 0u))) >> uvec4(24u));
+ ivec4 b_i8 = (ivec4((uvec4(b) << uvec4(24u, 16u, 8u, 0u))) >> uvec4(24u));
+ return tint_int_dot(a_i8, b_i8);
+}
+
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ int inner;
+} prevent_dce;
+
+void dot4I8Packed_881e62() {
+ int res = tint_dot4_i8_packed(1u, 1u);
+ prevent_dce.inner = res;
+}
+
+void compute_main() {
dot4I8Packed_881e62();
}
-Failed to generate: error: Unknown builtin method: dot4I8Packed
-
-enable chromium_experimental_dp4a;
-
-fn dot4I8Packed_881e62() {
- var res : i32 = dot4I8Packed(1u, 1u);
- prevent_dce = res;
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main();
+ return;
}
-
-@group(2) @binding(0) var<storage, read_write> prevent_dce : i32;
-
-@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();
-}
-
-Failed to generate: error: Unknown builtin method: dot4I8Packed
-
-enable chromium_experimental_dp4a;
-
-fn dot4I8Packed_881e62() {
- var res : i32 = dot4I8Packed(1u, 1u);
- prevent_dce = res;
-}
-
-@group(2) @binding(0) var<storage, read_write> prevent_dce : i32;
-
-@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();
-}
-
-Failed to generate: error: Unknown builtin method: dot4I8Packed
diff --git a/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.glsl b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.glsl
index 000e824..8d6d992 100644
--- a/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.glsl
@@ -1,83 +1,94 @@
-SKIP: FAILED
+#version 310 es
-
-enable chromium_experimental_dp4a;
-
-fn dot4U8Packed_fbed7b() {
- var res : u32 = dot4U8Packed(1u, 1u);
- prevent_dce = res;
+uint tint_int_dot(uvec4 a, uvec4 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : u32;
+uint tint_dot4_u8_packed(uint a, uint b) {
+ uvec4 a_u8 = ((uvec4(a) >> uvec4(24u, 16u, 8u, 0u)) & uvec4(255u));
+ uvec4 b_u8 = ((uvec4(b) >> uvec4(24u, 16u, 8u, 0u)) & uvec4(255u));
+ return tint_int_dot(a_u8, b_u8);
+}
-@vertex
-fn vertex_main() -> @builtin(position) vec4<f32> {
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ uint inner;
+} prevent_dce;
+
+void dot4U8Packed_fbed7b() {
+ uint res = tint_dot4_u8_packed(1u, 1u);
+ prevent_dce.inner = res;
+}
+
+vec4 vertex_main() {
dot4U8Packed_fbed7b();
- return vec4<f32>();
+ return vec4(0.0f);
}
-@fragment
-fn fragment_main() {
+void main() {
+ gl_PointSize = 1.0;
+ vec4 inner_result = vertex_main();
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
+#version 310 es
+precision highp float;
+
+uint tint_int_dot(uvec4 a, uvec4 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
+}
+
+uint tint_dot4_u8_packed(uint a, uint b) {
+ uvec4 a_u8 = ((uvec4(a) >> uvec4(24u, 16u, 8u, 0u)) & uvec4(255u));
+ uvec4 b_u8 = ((uvec4(b) >> uvec4(24u, 16u, 8u, 0u)) & uvec4(255u));
+ return tint_int_dot(a_u8, b_u8);
+}
+
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ uint inner;
+} prevent_dce;
+
+void dot4U8Packed_fbed7b() {
+ uint res = tint_dot4_u8_packed(1u, 1u);
+ prevent_dce.inner = res;
+}
+
+void fragment_main() {
dot4U8Packed_fbed7b();
}
-@compute @workgroup_size(1)
-fn compute_main() {
+void main() {
+ fragment_main();
+ return;
+}
+#version 310 es
+
+uint tint_int_dot(uvec4 a, uvec4 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
+}
+
+uint tint_dot4_u8_packed(uint a, uint b) {
+ uvec4 a_u8 = ((uvec4(a) >> uvec4(24u, 16u, 8u, 0u)) & uvec4(255u));
+ uvec4 b_u8 = ((uvec4(b) >> uvec4(24u, 16u, 8u, 0u)) & uvec4(255u));
+ return tint_int_dot(a_u8, b_u8);
+}
+
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ uint inner;
+} prevent_dce;
+
+void dot4U8Packed_fbed7b() {
+ uint res = tint_dot4_u8_packed(1u, 1u);
+ prevent_dce.inner = res;
+}
+
+void compute_main() {
dot4U8Packed_fbed7b();
}
-Failed to generate: error: Unknown builtin method: dot4U8Packed
-
-enable chromium_experimental_dp4a;
-
-fn dot4U8Packed_fbed7b() {
- var res : u32 = dot4U8Packed(1u, 1u);
- prevent_dce = res;
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main();
+ return;
}
-
-@group(2) @binding(0) var<storage, read_write> prevent_dce : u32;
-
-@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();
-}
-
-Failed to generate: error: Unknown builtin method: dot4U8Packed
-
-enable chromium_experimental_dp4a;
-
-fn dot4U8Packed_fbed7b() {
- var res : u32 = dot4U8Packed(1u, 1u);
- prevent_dce = res;
-}
-
-@group(2) @binding(0) var<storage, read_write> prevent_dce : u32;
-
-@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();
-}
-
-Failed to generate: error: Unknown builtin method: dot4U8Packed
diff --git a/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.glsl b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.glsl
index bbac611..4a8e535 100644
--- a/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.glsl
@@ -1,89 +1,100 @@
-SKIP: FAILED
+#version 310 es
-
-enable chromium_experimental_dp4a;
-
-fn dot4I8Packed_881e62() {
- var arg_0 = 1u;
- var arg_1 = 1u;
- var res : i32 = dot4I8Packed(arg_0, arg_1);
- prevent_dce = res;
+int tint_int_dot(ivec4 a, ivec4 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : i32;
+int tint_dot4_i8_packed(uint a, uint b) {
+ ivec4 a_i8 = (ivec4((uvec4(a) << uvec4(24u, 16u, 8u, 0u))) >> uvec4(24u));
+ ivec4 b_i8 = (ivec4((uvec4(b) << uvec4(24u, 16u, 8u, 0u))) >> uvec4(24u));
+ return tint_int_dot(a_i8, b_i8);
+}
-@vertex
-fn vertex_main() -> @builtin(position) vec4<f32> {
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ int inner;
+} prevent_dce;
+
+void dot4I8Packed_881e62() {
+ uint arg_0 = 1u;
+ uint arg_1 = 1u;
+ int res = tint_dot4_i8_packed(arg_0, arg_1);
+ prevent_dce.inner = res;
+}
+
+vec4 vertex_main() {
dot4I8Packed_881e62();
- return vec4<f32>();
+ return vec4(0.0f);
}
-@fragment
-fn fragment_main() {
+void main() {
+ gl_PointSize = 1.0;
+ vec4 inner_result = vertex_main();
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
+#version 310 es
+precision highp float;
+
+int tint_int_dot(ivec4 a, ivec4 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
+}
+
+int tint_dot4_i8_packed(uint a, uint b) {
+ ivec4 a_i8 = (ivec4((uvec4(a) << uvec4(24u, 16u, 8u, 0u))) >> uvec4(24u));
+ ivec4 b_i8 = (ivec4((uvec4(b) << uvec4(24u, 16u, 8u, 0u))) >> uvec4(24u));
+ return tint_int_dot(a_i8, b_i8);
+}
+
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ int inner;
+} prevent_dce;
+
+void dot4I8Packed_881e62() {
+ uint arg_0 = 1u;
+ uint arg_1 = 1u;
+ int res = tint_dot4_i8_packed(arg_0, arg_1);
+ prevent_dce.inner = res;
+}
+
+void fragment_main() {
dot4I8Packed_881e62();
}
-@compute @workgroup_size(1)
-fn compute_main() {
+void main() {
+ fragment_main();
+ return;
+}
+#version 310 es
+
+int tint_int_dot(ivec4 a, ivec4 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
+}
+
+int tint_dot4_i8_packed(uint a, uint b) {
+ ivec4 a_i8 = (ivec4((uvec4(a) << uvec4(24u, 16u, 8u, 0u))) >> uvec4(24u));
+ ivec4 b_i8 = (ivec4((uvec4(b) << uvec4(24u, 16u, 8u, 0u))) >> uvec4(24u));
+ return tint_int_dot(a_i8, b_i8);
+}
+
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ int inner;
+} prevent_dce;
+
+void dot4I8Packed_881e62() {
+ uint arg_0 = 1u;
+ uint arg_1 = 1u;
+ int res = tint_dot4_i8_packed(arg_0, arg_1);
+ prevent_dce.inner = res;
+}
+
+void compute_main() {
dot4I8Packed_881e62();
}
-Failed to generate: error: Unknown builtin method: dot4I8Packed
-
-enable chromium_experimental_dp4a;
-
-fn dot4I8Packed_881e62() {
- var arg_0 = 1u;
- var arg_1 = 1u;
- var res : i32 = dot4I8Packed(arg_0, arg_1);
- prevent_dce = res;
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main();
+ return;
}
-
-@group(2) @binding(0) var<storage, read_write> prevent_dce : i32;
-
-@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();
-}
-
-Failed to generate: error: Unknown builtin method: dot4I8Packed
-
-enable chromium_experimental_dp4a;
-
-fn dot4I8Packed_881e62() {
- var arg_0 = 1u;
- var arg_1 = 1u;
- var res : i32 = dot4I8Packed(arg_0, arg_1);
- prevent_dce = res;
-}
-
-@group(2) @binding(0) var<storage, read_write> prevent_dce : i32;
-
-@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();
-}
-
-Failed to generate: error: Unknown builtin method: dot4I8Packed
diff --git a/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.glsl b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.glsl
index e2ab67e..6938c5c 100644
--- a/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.glsl
@@ -1,89 +1,100 @@
-SKIP: FAILED
+#version 310 es
-
-enable chromium_experimental_dp4a;
-
-fn dot4U8Packed_fbed7b() {
- var arg_0 = 1u;
- var arg_1 = 1u;
- var res : u32 = dot4U8Packed(arg_0, arg_1);
- prevent_dce = res;
+uint tint_int_dot(uvec4 a, uvec4 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : u32;
+uint tint_dot4_u8_packed(uint a, uint b) {
+ uvec4 a_u8 = ((uvec4(a) >> uvec4(24u, 16u, 8u, 0u)) & uvec4(255u));
+ uvec4 b_u8 = ((uvec4(b) >> uvec4(24u, 16u, 8u, 0u)) & uvec4(255u));
+ return tint_int_dot(a_u8, b_u8);
+}
-@vertex
-fn vertex_main() -> @builtin(position) vec4<f32> {
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ uint inner;
+} prevent_dce;
+
+void dot4U8Packed_fbed7b() {
+ uint arg_0 = 1u;
+ uint arg_1 = 1u;
+ uint res = tint_dot4_u8_packed(arg_0, arg_1);
+ prevent_dce.inner = res;
+}
+
+vec4 vertex_main() {
dot4U8Packed_fbed7b();
- return vec4<f32>();
+ return vec4(0.0f);
}
-@fragment
-fn fragment_main() {
+void main() {
+ gl_PointSize = 1.0;
+ vec4 inner_result = vertex_main();
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
+#version 310 es
+precision highp float;
+
+uint tint_int_dot(uvec4 a, uvec4 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
+}
+
+uint tint_dot4_u8_packed(uint a, uint b) {
+ uvec4 a_u8 = ((uvec4(a) >> uvec4(24u, 16u, 8u, 0u)) & uvec4(255u));
+ uvec4 b_u8 = ((uvec4(b) >> uvec4(24u, 16u, 8u, 0u)) & uvec4(255u));
+ return tint_int_dot(a_u8, b_u8);
+}
+
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ uint inner;
+} prevent_dce;
+
+void dot4U8Packed_fbed7b() {
+ uint arg_0 = 1u;
+ uint arg_1 = 1u;
+ uint res = tint_dot4_u8_packed(arg_0, arg_1);
+ prevent_dce.inner = res;
+}
+
+void fragment_main() {
dot4U8Packed_fbed7b();
}
-@compute @workgroup_size(1)
-fn compute_main() {
+void main() {
+ fragment_main();
+ return;
+}
+#version 310 es
+
+uint tint_int_dot(uvec4 a, uvec4 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
+}
+
+uint tint_dot4_u8_packed(uint a, uint b) {
+ uvec4 a_u8 = ((uvec4(a) >> uvec4(24u, 16u, 8u, 0u)) & uvec4(255u));
+ uvec4 b_u8 = ((uvec4(b) >> uvec4(24u, 16u, 8u, 0u)) & uvec4(255u));
+ return tint_int_dot(a_u8, b_u8);
+}
+
+layout(binding = 0, std430) buffer prevent_dce_block_ssbo {
+ uint inner;
+} prevent_dce;
+
+void dot4U8Packed_fbed7b() {
+ uint arg_0 = 1u;
+ uint arg_1 = 1u;
+ uint res = tint_dot4_u8_packed(arg_0, arg_1);
+ prevent_dce.inner = res;
+}
+
+void compute_main() {
dot4U8Packed_fbed7b();
}
-Failed to generate: error: Unknown builtin method: dot4U8Packed
-
-enable chromium_experimental_dp4a;
-
-fn dot4U8Packed_fbed7b() {
- var arg_0 = 1u;
- var arg_1 = 1u;
- var res : u32 = dot4U8Packed(arg_0, arg_1);
- prevent_dce = res;
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main();
+ return;
}
-
-@group(2) @binding(0) var<storage, read_write> prevent_dce : u32;
-
-@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();
-}
-
-Failed to generate: error: Unknown builtin method: dot4U8Packed
-
-enable chromium_experimental_dp4a;
-
-fn dot4U8Packed_fbed7b() {
- var arg_0 = 1u;
- var arg_1 = 1u;
- var res : u32 = dot4U8Packed(arg_0, arg_1);
- prevent_dce = res;
-}
-
-@group(2) @binding(0) var<storage, read_write> prevent_dce : u32;
-
-@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();
-}
-
-Failed to generate: error: Unknown builtin method: dot4U8Packed
