From 3a431d71be0c46641170635f75fba8ea923d7711 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Sat, 4 Mar 2023 00:43:21 +0000 Subject: [PATCH] tint/intrinsics.def: Fix textureStore overload The texture_storage_3d overload should allow for unsigned coordinates. Change-Id: I6278571fb9dc7bba644a4ba88cce6b8bd7c790bf Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/122521 Kokoro: Kokoro Commit-Queue: Ben Clayton Reviewed-by: James Price --- src/tint/intrinsics.def | 2 +- src/tint/resolver/intrinsic_table.inl | 8 +- .../gen/literal/textureStore/1efc36.wgsl | 44 ++++++++++ .../1efc36.wgsl.expected.dxc.hlsl | 32 +++++++ .../1efc36.wgsl.expected.fxc.hlsl | 32 +++++++ .../textureStore/1efc36.wgsl.expected.glsl | 52 +++++++++++ .../textureStore/1efc36.wgsl.expected.msl | 33 +++++++ .../textureStore/1efc36.wgsl.expected.spvasm | 74 ++++++++++++++++ .../textureStore/1efc36.wgsl.expected.wgsl | 21 +++++ .../gen/literal/textureStore/3fb31f.wgsl | 44 ++++++++++ .../3fb31f.wgsl.expected.dxc.hlsl | 32 +++++++ .../3fb31f.wgsl.expected.fxc.hlsl | 32 +++++++ .../textureStore/3fb31f.wgsl.expected.glsl | 75 ++++++++++++++++ .../textureStore/3fb31f.wgsl.expected.msl | 33 +++++++ .../textureStore/3fb31f.wgsl.expected.spvasm | 75 ++++++++++++++++ .../textureStore/3fb31f.wgsl.expected.wgsl | 21 +++++ .../gen/literal/textureStore/cd6755.wgsl | 44 ++++++++++ .../cd6755.wgsl.expected.dxc.hlsl | 32 +++++++ .../cd6755.wgsl.expected.fxc.hlsl | 32 +++++++ .../textureStore/cd6755.wgsl.expected.glsl | 52 +++++++++++ .../textureStore/cd6755.wgsl.expected.msl | 33 +++++++ .../textureStore/cd6755.wgsl.expected.spvasm | 74 ++++++++++++++++ .../textureStore/cd6755.wgsl.expected.wgsl | 21 +++++ .../gen/literal/textureStore/d2b565.wgsl | 44 ++++++++++ .../d2b565.wgsl.expected.dxc.hlsl | 32 +++++++ .../d2b565.wgsl.expected.fxc.hlsl | 32 +++++++ .../textureStore/d2b565.wgsl.expected.glsl | 52 +++++++++++ .../textureStore/d2b565.wgsl.expected.msl | 33 +++++++ .../textureStore/d2b565.wgsl.expected.spvasm | 74 ++++++++++++++++ .../textureStore/d2b565.wgsl.expected.wgsl | 21 +++++ .../gen/literal/textureStore/d4aa95.wgsl | 44 ++++++++++ .../d4aa95.wgsl.expected.dxc.hlsl | 32 +++++++ .../d4aa95.wgsl.expected.fxc.hlsl | 32 +++++++ .../textureStore/d4aa95.wgsl.expected.glsl | 52 +++++++++++ .../textureStore/d4aa95.wgsl.expected.msl | 33 +++++++ .../textureStore/d4aa95.wgsl.expected.spvasm | 74 ++++++++++++++++ .../textureStore/d4aa95.wgsl.expected.wgsl | 21 +++++ .../builtins/gen/var/textureStore/1efc36.wgsl | 46 ++++++++++ .../1efc36.wgsl.expected.dxc.hlsl | 34 ++++++++ .../1efc36.wgsl.expected.fxc.hlsl | 34 ++++++++ .../textureStore/1efc36.wgsl.expected.glsl | 58 +++++++++++++ .../var/textureStore/1efc36.wgsl.expected.msl | 35 ++++++++ .../textureStore/1efc36.wgsl.expected.spvasm | 86 ++++++++++++++++++ .../textureStore/1efc36.wgsl.expected.wgsl | 23 +++++ .../builtins/gen/var/textureStore/3fb31f.wgsl | 46 ++++++++++ .../3fb31f.wgsl.expected.dxc.hlsl | 34 ++++++++ .../3fb31f.wgsl.expected.fxc.hlsl | 34 ++++++++ .../textureStore/3fb31f.wgsl.expected.glsl | 81 +++++++++++++++++ .../var/textureStore/3fb31f.wgsl.expected.msl | 35 ++++++++ .../textureStore/3fb31f.wgsl.expected.spvasm | 87 +++++++++++++++++++ .../textureStore/3fb31f.wgsl.expected.wgsl | 23 +++++ .../builtins/gen/var/textureStore/cd6755.wgsl | 46 ++++++++++ .../cd6755.wgsl.expected.dxc.hlsl | 34 ++++++++ .../cd6755.wgsl.expected.fxc.hlsl | 34 ++++++++ .../textureStore/cd6755.wgsl.expected.glsl | 58 +++++++++++++ .../var/textureStore/cd6755.wgsl.expected.msl | 35 ++++++++ .../textureStore/cd6755.wgsl.expected.spvasm | 86 ++++++++++++++++++ .../textureStore/cd6755.wgsl.expected.wgsl | 23 +++++ .../builtins/gen/var/textureStore/d2b565.wgsl | 46 ++++++++++ .../d2b565.wgsl.expected.dxc.hlsl | 34 ++++++++ .../d2b565.wgsl.expected.fxc.hlsl | 34 ++++++++ .../textureStore/d2b565.wgsl.expected.glsl | 58 +++++++++++++ .../var/textureStore/d2b565.wgsl.expected.msl | 35 ++++++++ .../textureStore/d2b565.wgsl.expected.spvasm | 86 ++++++++++++++++++ .../textureStore/d2b565.wgsl.expected.wgsl | 23 +++++ .../builtins/gen/var/textureStore/d4aa95.wgsl | 46 ++++++++++ .../d4aa95.wgsl.expected.dxc.hlsl | 34 ++++++++ .../d4aa95.wgsl.expected.fxc.hlsl | 34 ++++++++ .../textureStore/d4aa95.wgsl.expected.glsl | 58 +++++++++++++ .../var/textureStore/d4aa95.wgsl.expected.msl | 35 ++++++++ .../textureStore/d4aa95.wgsl.expected.spvasm | 86 ++++++++++++++++++ .../textureStore/d4aa95.wgsl.expected.wgsl | 23 +++++ 72 files changed, 3073 insertions(+), 5 deletions(-) create mode 100644 test/tint/builtins/gen/literal/textureStore/1efc36.wgsl create mode 100644 test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl create mode 100644 test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/literal/textureStore/cd6755.wgsl create mode 100644 test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/literal/textureStore/d2b565.wgsl create mode 100644 test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl create mode 100644 test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/var/textureStore/1efc36.wgsl create mode 100644 test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/var/textureStore/3fb31f.wgsl create mode 100644 test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/var/textureStore/cd6755.wgsl create mode 100644 test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/var/textureStore/d2b565.wgsl create mode 100644 test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/var/textureStore/d4aa95.wgsl create mode 100644 test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.wgsl diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def index 45d77c8ed6..fae4ee419c 100644 --- a/src/tint/intrinsics.def +++ b/src/tint/intrinsics.def @@ -825,7 +825,7 @@ fn textureStore(texture: texture_storage_3d, c fn textureStore(texture: texture_storage_1d, coords: C, value: vec4) fn textureStore(texture: texture_storage_2d, coords: vec2, value: vec4) fn textureStore(texture: texture_storage_2d_array, coords: vec2, array_index: A, value: vec4) -fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) +fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) @must_use fn textureLoad(texture: texture_1d, coords: C, level: L) -> vec4 @must_use fn textureLoad(texture: texture_2d, coords: vec2, level: L) -> vec4 @must_use fn textureLoad(texture: texture_2d_array, coords: vec2, array_index: A, level: L) -> vec4 diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl index c383265486..895d4893e9 100644 --- a/src/tint/resolver/intrinsic_table.inl +++ b/src/tint/resolver/intrinsic_table.inl @@ -5843,7 +5843,7 @@ constexpr ParameterInfo kParameters[] = { { /* [550] */ /* usage */ ParameterUsage::kCoords, - /* matcher indices */ &kMatcherIndices[148], + /* matcher indices */ &kMatcherIndices[106], }, { /* [551] */ @@ -9540,9 +9540,9 @@ constexpr OverloadInfo kOverloads[] = { { /* [94] */ /* num parameters */ 3, - /* num template types */ 0, + /* num template types */ 1, /* num template numbers */ 0, - /* template types */ &kTemplateTypes[38], + /* template types */ &kTemplateTypes[1], /* template numbers */ &kTemplateNumbers[10], /* parameters */ &kParameters[549], /* return matcher indices */ nullptr, @@ -14862,7 +14862,7 @@ constexpr IntrinsicInfo kBuiltins[] = { /* fn textureStore(texture: texture_storage_1d, coords: C, value: vec4) */ /* fn textureStore(texture: texture_storage_2d, coords: vec2, value: vec4) */ /* fn textureStore(texture: texture_storage_2d_array, coords: vec2, array_index: A, value: vec4) */ - /* fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) */ + /* fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) */ /* num overloads */ 12, /* overloads */ &kOverloads[83], }, diff --git a/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl new file mode 100644 index 0000000000..cdd5a30537 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl @@ -0,0 +1,44 @@ +// Copyright 2023 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 +//////////////////////////////////////////////////////////////////////////////// + +@group(1) @binding(0) var arg_0: texture_storage_3d; + +// fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) +fn textureStore_1efc36() { + textureStore(arg_0, vec3(1u), vec4(1u)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_1efc36(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_1efc36(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_1efc36(); +} diff --git a/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..f04775c5a2 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.dxc.hlsl @@ -0,0 +1,32 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_1efc36() { + arg_0[(1u).xxx] = (1u).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_1efc36(); + 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() { + textureStore_1efc36(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_1efc36(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..f04775c5a2 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.fxc.hlsl @@ -0,0 +1,32 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_1efc36() { + arg_0[(1u).xxx] = (1u).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_1efc36(); + 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() { + textureStore_1efc36(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_1efc36(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.glsl b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.glsl new file mode 100644 index 0000000000..39a7b31ac5 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.glsl @@ -0,0 +1,52 @@ +#version 310 es + +layout(rgba16ui) uniform highp writeonly uimage3D arg_0; +void textureStore_1efc36() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +vec4 vertex_main() { + textureStore_1efc36(); + return vec4(0.0f); +} + +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 mediump float; + +layout(rgba16ui) uniform highp writeonly uimage3D arg_0; +void textureStore_1efc36() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +void fragment_main() { + textureStore_1efc36(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +layout(rgba16ui) uniform highp writeonly uimage3D arg_0; +void textureStore_1efc36() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +void compute_main() { + textureStore_1efc36(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.msl b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.msl new file mode 100644 index 0000000000..0f4313a9af --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; +void textureStore_1efc36(texture3d tint_symbol_1) { + tint_symbol_1.write(uint4(1u), uint3(uint3(1u))); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner(texture3d tint_symbol_2) { + textureStore_1efc36(tint_symbol_2); + return float4(0.0f); +} + +vertex tint_symbol vertex_main(texture3d tint_symbol_3 [[texture(0)]]) { + float4 const inner_result = vertex_main_inner(tint_symbol_3); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main(texture3d tint_symbol_4 [[texture(0)]]) { + textureStore_1efc36(tint_symbol_4); + return; +} + +kernel void compute_main(texture3d tint_symbol_5 [[texture(0)]]) { + textureStore_1efc36(tint_symbol_5); + return; +} + diff --git a/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.spvasm new file mode 100644 index 0000000000..26c6b84178 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.spvasm @@ -0,0 +1,74 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 38 +; Schema: 0 + OpCapability Shader + 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 %arg_0 "arg_0" + OpName %textureStore_1efc36 "textureStore_1efc36" + 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 + OpDecorate %arg_0 NonReadable + OpDecorate %arg_0 DescriptorSet 1 + OpDecorate %arg_0 Binding 0 + %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 + %uint = OpTypeInt 32 0 + %11 = OpTypeImage %uint 3D 0 0 0 2 Rgba16ui +%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11 + %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant + %void = OpTypeVoid + %13 = OpTypeFunction %void + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %21 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 + %v4uint = OpTypeVector %uint 4 + %23 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1 + %24 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%textureStore_1efc36 = OpFunction %void None %13 + %16 = OpLabel + %18 = OpLoad %11 %arg_0 + OpImageWrite %18 %21 %23 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %24 + %26 = OpLabel + %27 = OpFunctionCall %void %textureStore_1efc36 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %13 + %29 = OpLabel + %30 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %30 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %13 + %33 = OpLabel + %34 = OpFunctionCall %void %textureStore_1efc36 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %13 + %36 = OpLabel + %37 = OpFunctionCall %void %textureStore_1efc36 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.wgsl new file mode 100644 index 0000000000..84bbf25de1 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/1efc36.wgsl.expected.wgsl @@ -0,0 +1,21 @@ +@group(1) @binding(0) var arg_0 : texture_storage_3d; + +fn textureStore_1efc36() { + textureStore(arg_0, vec3(1u), vec4(1u)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_1efc36(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_1efc36(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_1efc36(); +} diff --git a/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl new file mode 100644 index 0000000000..eabdb199c7 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl @@ -0,0 +1,44 @@ +// Copyright 2023 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 +//////////////////////////////////////////////////////////////////////////////// + +@group(1) @binding(0) var arg_0: texture_storage_3d; + +// fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) +fn textureStore_3fb31f() { + textureStore(arg_0, vec3(1u), vec4(1u)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_3fb31f(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_3fb31f(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_3fb31f(); +} diff --git a/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..825c1dad7c --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.dxc.hlsl @@ -0,0 +1,32 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_3fb31f() { + arg_0[(1u).xxx] = (1u).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_3fb31f(); + 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() { + textureStore_3fb31f(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_3fb31f(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..825c1dad7c --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.fxc.hlsl @@ -0,0 +1,32 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_3fb31f() { + arg_0[(1u).xxx] = (1u).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_3fb31f(); + 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() { + textureStore_3fb31f(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_3fb31f(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.glsl b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.glsl new file mode 100644 index 0000000000..5efc829765 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.glsl @@ -0,0 +1,75 @@ +SKIP: FAILED + +#version 310 es + +layout(rg32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_3fb31f() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +vec4 vertex_main() { + textureStore_3fb31f(); + return vec4(0.0f); +} + +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; +} +Error parsing GLSL shader: +ERROR: 0:3: 'image load-store format' : not supported with this profile: es +ERROR: 0:3: '' : compilation terminated +ERROR: 2 compilation errors. No code generated. + + + +#version 310 es +precision mediump float; + +layout(rg32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_3fb31f() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +void fragment_main() { + textureStore_3fb31f(); +} + +void main() { + fragment_main(); + return; +} +Error parsing GLSL shader: +ERROR: 0:4: 'image load-store format' : not supported with this profile: es +ERROR: 0:4: '' : compilation terminated +ERROR: 2 compilation errors. No code generated. + + + +#version 310 es + +layout(rg32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_3fb31f() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +void compute_main() { + textureStore_3fb31f(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} +Error parsing GLSL shader: +ERROR: 0:3: 'image load-store format' : not supported with this profile: es +ERROR: 0:3: '' : compilation terminated +ERROR: 2 compilation errors. No code generated. + + + diff --git a/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.msl b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.msl new file mode 100644 index 0000000000..e625cc54d4 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; +void textureStore_3fb31f(texture3d tint_symbol_1) { + tint_symbol_1.write(uint4(1u), uint3(uint3(1u))); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner(texture3d tint_symbol_2) { + textureStore_3fb31f(tint_symbol_2); + return float4(0.0f); +} + +vertex tint_symbol vertex_main(texture3d tint_symbol_3 [[texture(0)]]) { + float4 const inner_result = vertex_main_inner(tint_symbol_3); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main(texture3d tint_symbol_4 [[texture(0)]]) { + textureStore_3fb31f(tint_symbol_4); + return; +} + +kernel void compute_main(texture3d tint_symbol_5 [[texture(0)]]) { + textureStore_3fb31f(tint_symbol_5); + return; +} + diff --git a/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.spvasm new file mode 100644 index 0000000000..34b5172290 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.spvasm @@ -0,0 +1,75 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 38 +; Schema: 0 + OpCapability Shader + OpCapability StorageImageExtendedFormats + 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 %arg_0 "arg_0" + OpName %textureStore_3fb31f "textureStore_3fb31f" + 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 + OpDecorate %arg_0 NonReadable + OpDecorate %arg_0 DescriptorSet 1 + OpDecorate %arg_0 Binding 0 + %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 + %uint = OpTypeInt 32 0 + %11 = OpTypeImage %uint 3D 0 0 0 2 Rg32ui +%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11 + %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant + %void = OpTypeVoid + %13 = OpTypeFunction %void + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %21 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 + %v4uint = OpTypeVector %uint 4 + %23 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1 + %24 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%textureStore_3fb31f = OpFunction %void None %13 + %16 = OpLabel + %18 = OpLoad %11 %arg_0 + OpImageWrite %18 %21 %23 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %24 + %26 = OpLabel + %27 = OpFunctionCall %void %textureStore_3fb31f + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %13 + %29 = OpLabel + %30 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %30 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %13 + %33 = OpLabel + %34 = OpFunctionCall %void %textureStore_3fb31f + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %13 + %36 = OpLabel + %37 = OpFunctionCall %void %textureStore_3fb31f + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.wgsl new file mode 100644 index 0000000000..3251213b59 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/3fb31f.wgsl.expected.wgsl @@ -0,0 +1,21 @@ +@group(1) @binding(0) var arg_0 : texture_storage_3d; + +fn textureStore_3fb31f() { + textureStore(arg_0, vec3(1u), vec4(1u)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_3fb31f(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_3fb31f(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_3fb31f(); +} diff --git a/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl new file mode 100644 index 0000000000..2c7ef4f305 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl @@ -0,0 +1,44 @@ +// Copyright 2023 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 +//////////////////////////////////////////////////////////////////////////////// + +@group(1) @binding(0) var arg_0: texture_storage_3d; + +// fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) +fn textureStore_cd6755() { + textureStore(arg_0, vec3(1u), vec4(1u)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_cd6755(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_cd6755(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_cd6755(); +} diff --git a/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..8a54a6d104 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.dxc.hlsl @@ -0,0 +1,32 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_cd6755() { + arg_0[(1u).xxx] = (1u).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_cd6755(); + 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() { + textureStore_cd6755(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_cd6755(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..8a54a6d104 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.fxc.hlsl @@ -0,0 +1,32 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_cd6755() { + arg_0[(1u).xxx] = (1u).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_cd6755(); + 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() { + textureStore_cd6755(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_cd6755(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.glsl b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.glsl new file mode 100644 index 0000000000..848590cf1d --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.glsl @@ -0,0 +1,52 @@ +#version 310 es + +layout(r32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_cd6755() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +vec4 vertex_main() { + textureStore_cd6755(); + return vec4(0.0f); +} + +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 mediump float; + +layout(r32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_cd6755() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +void fragment_main() { + textureStore_cd6755(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +layout(r32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_cd6755() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +void compute_main() { + textureStore_cd6755(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.msl b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.msl new file mode 100644 index 0000000000..250426bd33 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; +void textureStore_cd6755(texture3d tint_symbol_1) { + tint_symbol_1.write(uint4(1u), uint3(uint3(1u))); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner(texture3d tint_symbol_2) { + textureStore_cd6755(tint_symbol_2); + return float4(0.0f); +} + +vertex tint_symbol vertex_main(texture3d tint_symbol_3 [[texture(0)]]) { + float4 const inner_result = vertex_main_inner(tint_symbol_3); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main(texture3d tint_symbol_4 [[texture(0)]]) { + textureStore_cd6755(tint_symbol_4); + return; +} + +kernel void compute_main(texture3d tint_symbol_5 [[texture(0)]]) { + textureStore_cd6755(tint_symbol_5); + return; +} + diff --git a/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.spvasm new file mode 100644 index 0000000000..a56d0fb204 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.spvasm @@ -0,0 +1,74 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 38 +; Schema: 0 + OpCapability Shader + 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 %arg_0 "arg_0" + OpName %textureStore_cd6755 "textureStore_cd6755" + 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 + OpDecorate %arg_0 NonReadable + OpDecorate %arg_0 DescriptorSet 1 + OpDecorate %arg_0 Binding 0 + %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 + %uint = OpTypeInt 32 0 + %11 = OpTypeImage %uint 3D 0 0 0 2 R32ui +%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11 + %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant + %void = OpTypeVoid + %13 = OpTypeFunction %void + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %21 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 + %v4uint = OpTypeVector %uint 4 + %23 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1 + %24 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%textureStore_cd6755 = OpFunction %void None %13 + %16 = OpLabel + %18 = OpLoad %11 %arg_0 + OpImageWrite %18 %21 %23 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %24 + %26 = OpLabel + %27 = OpFunctionCall %void %textureStore_cd6755 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %13 + %29 = OpLabel + %30 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %30 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %13 + %33 = OpLabel + %34 = OpFunctionCall %void %textureStore_cd6755 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %13 + %36 = OpLabel + %37 = OpFunctionCall %void %textureStore_cd6755 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.wgsl new file mode 100644 index 0000000000..f821741e21 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/cd6755.wgsl.expected.wgsl @@ -0,0 +1,21 @@ +@group(1) @binding(0) var arg_0 : texture_storage_3d; + +fn textureStore_cd6755() { + textureStore(arg_0, vec3(1u), vec4(1u)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_cd6755(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_cd6755(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_cd6755(); +} diff --git a/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl new file mode 100644 index 0000000000..540692309c --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl @@ -0,0 +1,44 @@ +// Copyright 2023 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 +//////////////////////////////////////////////////////////////////////////////// + +@group(1) @binding(0) var arg_0: texture_storage_3d; + +// fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) +fn textureStore_d2b565() { + textureStore(arg_0, vec3(1u), vec4(1u)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_d2b565(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_d2b565(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_d2b565(); +} diff --git a/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..2d99449c2f --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.dxc.hlsl @@ -0,0 +1,32 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_d2b565() { + arg_0[(1u).xxx] = (1u).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_d2b565(); + 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() { + textureStore_d2b565(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_d2b565(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..2d99449c2f --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.fxc.hlsl @@ -0,0 +1,32 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_d2b565() { + arg_0[(1u).xxx] = (1u).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_d2b565(); + 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() { + textureStore_d2b565(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_d2b565(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.glsl b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.glsl new file mode 100644 index 0000000000..eeb67a8a76 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.glsl @@ -0,0 +1,52 @@ +#version 310 es + +layout(rgba8ui) uniform highp writeonly uimage3D arg_0; +void textureStore_d2b565() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +vec4 vertex_main() { + textureStore_d2b565(); + return vec4(0.0f); +} + +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 mediump float; + +layout(rgba8ui) uniform highp writeonly uimage3D arg_0; +void textureStore_d2b565() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +void fragment_main() { + textureStore_d2b565(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +layout(rgba8ui) uniform highp writeonly uimage3D arg_0; +void textureStore_d2b565() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +void compute_main() { + textureStore_d2b565(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.msl b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.msl new file mode 100644 index 0000000000..ab3f3833cc --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; +void textureStore_d2b565(texture3d tint_symbol_1) { + tint_symbol_1.write(uint4(1u), uint3(uint3(1u))); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner(texture3d tint_symbol_2) { + textureStore_d2b565(tint_symbol_2); + return float4(0.0f); +} + +vertex tint_symbol vertex_main(texture3d tint_symbol_3 [[texture(0)]]) { + float4 const inner_result = vertex_main_inner(tint_symbol_3); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main(texture3d tint_symbol_4 [[texture(0)]]) { + textureStore_d2b565(tint_symbol_4); + return; +} + +kernel void compute_main(texture3d tint_symbol_5 [[texture(0)]]) { + textureStore_d2b565(tint_symbol_5); + return; +} + diff --git a/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.spvasm new file mode 100644 index 0000000000..db6ee7e53a --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.spvasm @@ -0,0 +1,74 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 38 +; Schema: 0 + OpCapability Shader + 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 %arg_0 "arg_0" + OpName %textureStore_d2b565 "textureStore_d2b565" + 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 + OpDecorate %arg_0 NonReadable + OpDecorate %arg_0 DescriptorSet 1 + OpDecorate %arg_0 Binding 0 + %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 + %uint = OpTypeInt 32 0 + %11 = OpTypeImage %uint 3D 0 0 0 2 Rgba8ui +%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11 + %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant + %void = OpTypeVoid + %13 = OpTypeFunction %void + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %21 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 + %v4uint = OpTypeVector %uint 4 + %23 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1 + %24 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%textureStore_d2b565 = OpFunction %void None %13 + %16 = OpLabel + %18 = OpLoad %11 %arg_0 + OpImageWrite %18 %21 %23 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %24 + %26 = OpLabel + %27 = OpFunctionCall %void %textureStore_d2b565 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %13 + %29 = OpLabel + %30 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %30 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %13 + %33 = OpLabel + %34 = OpFunctionCall %void %textureStore_d2b565 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %13 + %36 = OpLabel + %37 = OpFunctionCall %void %textureStore_d2b565 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.wgsl new file mode 100644 index 0000000000..9b56e898cd --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d2b565.wgsl.expected.wgsl @@ -0,0 +1,21 @@ +@group(1) @binding(0) var arg_0 : texture_storage_3d; + +fn textureStore_d2b565() { + textureStore(arg_0, vec3(1u), vec4(1u)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_d2b565(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_d2b565(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_d2b565(); +} diff --git a/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl new file mode 100644 index 0000000000..5dadc3f84d --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl @@ -0,0 +1,44 @@ +// Copyright 2023 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 +//////////////////////////////////////////////////////////////////////////////// + +@group(1) @binding(0) var arg_0: texture_storage_3d; + +// fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) +fn textureStore_d4aa95() { + textureStore(arg_0, vec3(1u), vec4(1u)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_d4aa95(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_d4aa95(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_d4aa95(); +} diff --git a/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..0c933db23e --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.dxc.hlsl @@ -0,0 +1,32 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_d4aa95() { + arg_0[(1u).xxx] = (1u).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_d4aa95(); + 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() { + textureStore_d4aa95(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_d4aa95(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..0c933db23e --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.fxc.hlsl @@ -0,0 +1,32 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_d4aa95() { + arg_0[(1u).xxx] = (1u).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_d4aa95(); + 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() { + textureStore_d4aa95(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_d4aa95(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.glsl b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.glsl new file mode 100644 index 0000000000..fc5f8f7c5c --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.glsl @@ -0,0 +1,52 @@ +#version 310 es + +layout(rgba32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_d4aa95() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +vec4 vertex_main() { + textureStore_d4aa95(); + return vec4(0.0f); +} + +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 mediump float; + +layout(rgba32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_d4aa95() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +void fragment_main() { + textureStore_d4aa95(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +layout(rgba32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_d4aa95() { + imageStore(arg_0, ivec3(uvec3(1u)), uvec4(1u)); +} + +void compute_main() { + textureStore_d4aa95(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.msl b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.msl new file mode 100644 index 0000000000..b1596208f4 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; +void textureStore_d4aa95(texture3d tint_symbol_1) { + tint_symbol_1.write(uint4(1u), uint3(uint3(1u))); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner(texture3d tint_symbol_2) { + textureStore_d4aa95(tint_symbol_2); + return float4(0.0f); +} + +vertex tint_symbol vertex_main(texture3d tint_symbol_3 [[texture(0)]]) { + float4 const inner_result = vertex_main_inner(tint_symbol_3); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main(texture3d tint_symbol_4 [[texture(0)]]) { + textureStore_d4aa95(tint_symbol_4); + return; +} + +kernel void compute_main(texture3d tint_symbol_5 [[texture(0)]]) { + textureStore_d4aa95(tint_symbol_5); + return; +} + diff --git a/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.spvasm new file mode 100644 index 0000000000..a19e893ea0 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.spvasm @@ -0,0 +1,74 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 38 +; Schema: 0 + OpCapability Shader + 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 %arg_0 "arg_0" + OpName %textureStore_d4aa95 "textureStore_d4aa95" + 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 + OpDecorate %arg_0 NonReadable + OpDecorate %arg_0 DescriptorSet 1 + OpDecorate %arg_0 Binding 0 + %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 + %uint = OpTypeInt 32 0 + %11 = OpTypeImage %uint 3D 0 0 0 2 Rgba32ui +%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11 + %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant + %void = OpTypeVoid + %13 = OpTypeFunction %void + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %21 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 + %v4uint = OpTypeVector %uint 4 + %23 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1 + %24 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%textureStore_d4aa95 = OpFunction %void None %13 + %16 = OpLabel + %18 = OpLoad %11 %arg_0 + OpImageWrite %18 %21 %23 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %24 + %26 = OpLabel + %27 = OpFunctionCall %void %textureStore_d4aa95 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %13 + %29 = OpLabel + %30 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %30 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %13 + %33 = OpLabel + %34 = OpFunctionCall %void %textureStore_d4aa95 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %13 + %36 = OpLabel + %37 = OpFunctionCall %void %textureStore_d4aa95 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.wgsl new file mode 100644 index 0000000000..f6f3d79ca4 --- /dev/null +++ b/test/tint/builtins/gen/literal/textureStore/d4aa95.wgsl.expected.wgsl @@ -0,0 +1,21 @@ +@group(1) @binding(0) var arg_0 : texture_storage_3d; + +fn textureStore_d4aa95() { + textureStore(arg_0, vec3(1u), vec4(1u)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_d4aa95(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_d4aa95(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_d4aa95(); +} diff --git a/test/tint/builtins/gen/var/textureStore/1efc36.wgsl b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl new file mode 100644 index 0000000000..36c3e7747c --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl @@ -0,0 +1,46 @@ +// Copyright 2023 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 +//////////////////////////////////////////////////////////////////////////////// + +@group(1) @binding(0) var arg_0: texture_storage_3d; + +// fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) +fn textureStore_1efc36() { + var arg_1 = vec3(1u); + var arg_2 = vec4(1u); + textureStore(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_1efc36(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_1efc36(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_1efc36(); +} diff --git a/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..1d12542952 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.dxc.hlsl @@ -0,0 +1,34 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_1efc36() { + uint3 arg_1 = (1u).xxx; + uint4 arg_2 = (1u).xxxx; + arg_0[arg_1] = arg_2; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_1efc36(); + 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() { + textureStore_1efc36(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_1efc36(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..1d12542952 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.fxc.hlsl @@ -0,0 +1,34 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_1efc36() { + uint3 arg_1 = (1u).xxx; + uint4 arg_2 = (1u).xxxx; + arg_0[arg_1] = arg_2; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_1efc36(); + 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() { + textureStore_1efc36(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_1efc36(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.glsl b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.glsl new file mode 100644 index 0000000000..b6ba970a61 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.glsl @@ -0,0 +1,58 @@ +#version 310 es + +layout(rgba16ui) uniform highp writeonly uimage3D arg_0; +void textureStore_1efc36() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +vec4 vertex_main() { + textureStore_1efc36(); + return vec4(0.0f); +} + +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 mediump float; + +layout(rgba16ui) uniform highp writeonly uimage3D arg_0; +void textureStore_1efc36() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +void fragment_main() { + textureStore_1efc36(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +layout(rgba16ui) uniform highp writeonly uimage3D arg_0; +void textureStore_1efc36() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +void compute_main() { + textureStore_1efc36(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.msl b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.msl new file mode 100644 index 0000000000..d1e2e8b6ae --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.msl @@ -0,0 +1,35 @@ +#include + +using namespace metal; +void textureStore_1efc36(texture3d tint_symbol_1) { + uint3 arg_1 = uint3(1u); + uint4 arg_2 = uint4(1u); + tint_symbol_1.write(arg_2, uint3(arg_1)); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner(texture3d tint_symbol_2) { + textureStore_1efc36(tint_symbol_2); + return float4(0.0f); +} + +vertex tint_symbol vertex_main(texture3d tint_symbol_3 [[texture(0)]]) { + float4 const inner_result = vertex_main_inner(tint_symbol_3); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main(texture3d tint_symbol_4 [[texture(0)]]) { + textureStore_1efc36(tint_symbol_4); + return; +} + +kernel void compute_main(texture3d tint_symbol_5 [[texture(0)]]) { + textureStore_1efc36(tint_symbol_5); + return; +} + diff --git a/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.spvasm b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.spvasm new file mode 100644 index 0000000000..09c59f39dc --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.spvasm @@ -0,0 +1,86 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 46 +; Schema: 0 + OpCapability Shader + 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 %arg_0 "arg_0" + OpName %textureStore_1efc36 "textureStore_1efc36" + OpName %arg_1 "arg_1" + OpName %arg_2 "arg_2" + 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 + OpDecorate %arg_0 NonReadable + OpDecorate %arg_0 DescriptorSet 1 + OpDecorate %arg_0 Binding 0 + %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 + %uint = OpTypeInt 32 0 + %11 = OpTypeImage %uint 3D 0 0 0 2 Rgba16ui +%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11 + %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant + %void = OpTypeVoid + %13 = OpTypeFunction %void + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %19 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %22 = OpConstantNull %v3uint + %v4uint = OpTypeVector %uint 4 + %24 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1 +%_ptr_Function_v4uint = OpTypePointer Function %v4uint + %27 = OpConstantNull %v4uint + %32 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%textureStore_1efc36 = OpFunction %void None %13 + %16 = OpLabel + %arg_1 = OpVariable %_ptr_Function_v3uint Function %22 + %arg_2 = OpVariable %_ptr_Function_v4uint Function %27 + OpStore %arg_1 %19 + OpStore %arg_2 %24 + %29 = OpLoad %11 %arg_0 + %30 = OpLoad %v3uint %arg_1 + %31 = OpLoad %v4uint %arg_2 + OpImageWrite %29 %30 %31 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %32 + %34 = OpLabel + %35 = OpFunctionCall %void %textureStore_1efc36 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %13 + %37 = OpLabel + %38 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %38 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %13 + %41 = OpLabel + %42 = OpFunctionCall %void %textureStore_1efc36 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %13 + %44 = OpLabel + %45 = OpFunctionCall %void %textureStore_1efc36 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.wgsl b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.wgsl new file mode 100644 index 0000000000..cccfb03b56 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/1efc36.wgsl.expected.wgsl @@ -0,0 +1,23 @@ +@group(1) @binding(0) var arg_0 : texture_storage_3d; + +fn textureStore_1efc36() { + var arg_1 = vec3(1u); + var arg_2 = vec4(1u); + textureStore(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_1efc36(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_1efc36(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_1efc36(); +} diff --git a/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl new file mode 100644 index 0000000000..534878786e --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl @@ -0,0 +1,46 @@ +// Copyright 2023 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 +//////////////////////////////////////////////////////////////////////////////// + +@group(1) @binding(0) var arg_0: texture_storage_3d; + +// fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) +fn textureStore_3fb31f() { + var arg_1 = vec3(1u); + var arg_2 = vec4(1u); + textureStore(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_3fb31f(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_3fb31f(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_3fb31f(); +} diff --git a/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..0d0410e8e5 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.dxc.hlsl @@ -0,0 +1,34 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_3fb31f() { + uint3 arg_1 = (1u).xxx; + uint4 arg_2 = (1u).xxxx; + arg_0[arg_1] = arg_2; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_3fb31f(); + 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() { + textureStore_3fb31f(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_3fb31f(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..0d0410e8e5 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.fxc.hlsl @@ -0,0 +1,34 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_3fb31f() { + uint3 arg_1 = (1u).xxx; + uint4 arg_2 = (1u).xxxx; + arg_0[arg_1] = arg_2; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_3fb31f(); + 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() { + textureStore_3fb31f(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_3fb31f(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.glsl b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.glsl new file mode 100644 index 0000000000..0b575900b2 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.glsl @@ -0,0 +1,81 @@ +SKIP: FAILED + +#version 310 es + +layout(rg32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_3fb31f() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +vec4 vertex_main() { + textureStore_3fb31f(); + return vec4(0.0f); +} + +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; +} +Error parsing GLSL shader: +ERROR: 0:3: 'image load-store format' : not supported with this profile: es +ERROR: 0:3: '' : compilation terminated +ERROR: 2 compilation errors. No code generated. + + + +#version 310 es +precision mediump float; + +layout(rg32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_3fb31f() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +void fragment_main() { + textureStore_3fb31f(); +} + +void main() { + fragment_main(); + return; +} +Error parsing GLSL shader: +ERROR: 0:4: 'image load-store format' : not supported with this profile: es +ERROR: 0:4: '' : compilation terminated +ERROR: 2 compilation errors. No code generated. + + + +#version 310 es + +layout(rg32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_3fb31f() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +void compute_main() { + textureStore_3fb31f(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} +Error parsing GLSL shader: +ERROR: 0:3: 'image load-store format' : not supported with this profile: es +ERROR: 0:3: '' : compilation terminated +ERROR: 2 compilation errors. No code generated. + + + diff --git a/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.msl b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.msl new file mode 100644 index 0000000000..cf4db1baea --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.msl @@ -0,0 +1,35 @@ +#include + +using namespace metal; +void textureStore_3fb31f(texture3d tint_symbol_1) { + uint3 arg_1 = uint3(1u); + uint4 arg_2 = uint4(1u); + tint_symbol_1.write(arg_2, uint3(arg_1)); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner(texture3d tint_symbol_2) { + textureStore_3fb31f(tint_symbol_2); + return float4(0.0f); +} + +vertex tint_symbol vertex_main(texture3d tint_symbol_3 [[texture(0)]]) { + float4 const inner_result = vertex_main_inner(tint_symbol_3); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main(texture3d tint_symbol_4 [[texture(0)]]) { + textureStore_3fb31f(tint_symbol_4); + return; +} + +kernel void compute_main(texture3d tint_symbol_5 [[texture(0)]]) { + textureStore_3fb31f(tint_symbol_5); + return; +} + diff --git a/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.spvasm b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.spvasm new file mode 100644 index 0000000000..a52e745e58 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.spvasm @@ -0,0 +1,87 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 46 +; Schema: 0 + OpCapability Shader + OpCapability StorageImageExtendedFormats + 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 %arg_0 "arg_0" + OpName %textureStore_3fb31f "textureStore_3fb31f" + OpName %arg_1 "arg_1" + OpName %arg_2 "arg_2" + 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 + OpDecorate %arg_0 NonReadable + OpDecorate %arg_0 DescriptorSet 1 + OpDecorate %arg_0 Binding 0 + %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 + %uint = OpTypeInt 32 0 + %11 = OpTypeImage %uint 3D 0 0 0 2 Rg32ui +%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11 + %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant + %void = OpTypeVoid + %13 = OpTypeFunction %void + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %19 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %22 = OpConstantNull %v3uint + %v4uint = OpTypeVector %uint 4 + %24 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1 +%_ptr_Function_v4uint = OpTypePointer Function %v4uint + %27 = OpConstantNull %v4uint + %32 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%textureStore_3fb31f = OpFunction %void None %13 + %16 = OpLabel + %arg_1 = OpVariable %_ptr_Function_v3uint Function %22 + %arg_2 = OpVariable %_ptr_Function_v4uint Function %27 + OpStore %arg_1 %19 + OpStore %arg_2 %24 + %29 = OpLoad %11 %arg_0 + %30 = OpLoad %v3uint %arg_1 + %31 = OpLoad %v4uint %arg_2 + OpImageWrite %29 %30 %31 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %32 + %34 = OpLabel + %35 = OpFunctionCall %void %textureStore_3fb31f + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %13 + %37 = OpLabel + %38 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %38 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %13 + %41 = OpLabel + %42 = OpFunctionCall %void %textureStore_3fb31f + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %13 + %44 = OpLabel + %45 = OpFunctionCall %void %textureStore_3fb31f + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.wgsl b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.wgsl new file mode 100644 index 0000000000..34d9f7c824 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/3fb31f.wgsl.expected.wgsl @@ -0,0 +1,23 @@ +@group(1) @binding(0) var arg_0 : texture_storage_3d; + +fn textureStore_3fb31f() { + var arg_1 = vec3(1u); + var arg_2 = vec4(1u); + textureStore(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_3fb31f(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_3fb31f(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_3fb31f(); +} diff --git a/test/tint/builtins/gen/var/textureStore/cd6755.wgsl b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl new file mode 100644 index 0000000000..a5a0a9398b --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl @@ -0,0 +1,46 @@ +// Copyright 2023 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 +//////////////////////////////////////////////////////////////////////////////// + +@group(1) @binding(0) var arg_0: texture_storage_3d; + +// fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) +fn textureStore_cd6755() { + var arg_1 = vec3(1u); + var arg_2 = vec4(1u); + textureStore(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_cd6755(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_cd6755(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_cd6755(); +} diff --git a/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..d61380b361 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.dxc.hlsl @@ -0,0 +1,34 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_cd6755() { + uint3 arg_1 = (1u).xxx; + uint4 arg_2 = (1u).xxxx; + arg_0[arg_1] = arg_2; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_cd6755(); + 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() { + textureStore_cd6755(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_cd6755(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..d61380b361 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.fxc.hlsl @@ -0,0 +1,34 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_cd6755() { + uint3 arg_1 = (1u).xxx; + uint4 arg_2 = (1u).xxxx; + arg_0[arg_1] = arg_2; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_cd6755(); + 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() { + textureStore_cd6755(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_cd6755(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.glsl b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.glsl new file mode 100644 index 0000000000..76dd31acd4 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.glsl @@ -0,0 +1,58 @@ +#version 310 es + +layout(r32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_cd6755() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +vec4 vertex_main() { + textureStore_cd6755(); + return vec4(0.0f); +} + +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 mediump float; + +layout(r32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_cd6755() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +void fragment_main() { + textureStore_cd6755(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +layout(r32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_cd6755() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +void compute_main() { + textureStore_cd6755(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.msl b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.msl new file mode 100644 index 0000000000..de0c7a8a9f --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.msl @@ -0,0 +1,35 @@ +#include + +using namespace metal; +void textureStore_cd6755(texture3d tint_symbol_1) { + uint3 arg_1 = uint3(1u); + uint4 arg_2 = uint4(1u); + tint_symbol_1.write(arg_2, uint3(arg_1)); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner(texture3d tint_symbol_2) { + textureStore_cd6755(tint_symbol_2); + return float4(0.0f); +} + +vertex tint_symbol vertex_main(texture3d tint_symbol_3 [[texture(0)]]) { + float4 const inner_result = vertex_main_inner(tint_symbol_3); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main(texture3d tint_symbol_4 [[texture(0)]]) { + textureStore_cd6755(tint_symbol_4); + return; +} + +kernel void compute_main(texture3d tint_symbol_5 [[texture(0)]]) { + textureStore_cd6755(tint_symbol_5); + return; +} + diff --git a/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.spvasm b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.spvasm new file mode 100644 index 0000000000..651a05df77 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.spvasm @@ -0,0 +1,86 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 46 +; Schema: 0 + OpCapability Shader + 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 %arg_0 "arg_0" + OpName %textureStore_cd6755 "textureStore_cd6755" + OpName %arg_1 "arg_1" + OpName %arg_2 "arg_2" + 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 + OpDecorate %arg_0 NonReadable + OpDecorate %arg_0 DescriptorSet 1 + OpDecorate %arg_0 Binding 0 + %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 + %uint = OpTypeInt 32 0 + %11 = OpTypeImage %uint 3D 0 0 0 2 R32ui +%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11 + %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant + %void = OpTypeVoid + %13 = OpTypeFunction %void + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %19 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %22 = OpConstantNull %v3uint + %v4uint = OpTypeVector %uint 4 + %24 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1 +%_ptr_Function_v4uint = OpTypePointer Function %v4uint + %27 = OpConstantNull %v4uint + %32 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%textureStore_cd6755 = OpFunction %void None %13 + %16 = OpLabel + %arg_1 = OpVariable %_ptr_Function_v3uint Function %22 + %arg_2 = OpVariable %_ptr_Function_v4uint Function %27 + OpStore %arg_1 %19 + OpStore %arg_2 %24 + %29 = OpLoad %11 %arg_0 + %30 = OpLoad %v3uint %arg_1 + %31 = OpLoad %v4uint %arg_2 + OpImageWrite %29 %30 %31 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %32 + %34 = OpLabel + %35 = OpFunctionCall %void %textureStore_cd6755 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %13 + %37 = OpLabel + %38 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %38 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %13 + %41 = OpLabel + %42 = OpFunctionCall %void %textureStore_cd6755 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %13 + %44 = OpLabel + %45 = OpFunctionCall %void %textureStore_cd6755 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.wgsl b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.wgsl new file mode 100644 index 0000000000..8c6cbdb4dd --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/cd6755.wgsl.expected.wgsl @@ -0,0 +1,23 @@ +@group(1) @binding(0) var arg_0 : texture_storage_3d; + +fn textureStore_cd6755() { + var arg_1 = vec3(1u); + var arg_2 = vec4(1u); + textureStore(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_cd6755(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_cd6755(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_cd6755(); +} diff --git a/test/tint/builtins/gen/var/textureStore/d2b565.wgsl b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl new file mode 100644 index 0000000000..d78172c746 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl @@ -0,0 +1,46 @@ +// Copyright 2023 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 +//////////////////////////////////////////////////////////////////////////////// + +@group(1) @binding(0) var arg_0: texture_storage_3d; + +// fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) +fn textureStore_d2b565() { + var arg_1 = vec3(1u); + var arg_2 = vec4(1u); + textureStore(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_d2b565(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_d2b565(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_d2b565(); +} diff --git a/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..cbac6995db --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.dxc.hlsl @@ -0,0 +1,34 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_d2b565() { + uint3 arg_1 = (1u).xxx; + uint4 arg_2 = (1u).xxxx; + arg_0[arg_1] = arg_2; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_d2b565(); + 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() { + textureStore_d2b565(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_d2b565(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..cbac6995db --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.fxc.hlsl @@ -0,0 +1,34 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_d2b565() { + uint3 arg_1 = (1u).xxx; + uint4 arg_2 = (1u).xxxx; + arg_0[arg_1] = arg_2; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_d2b565(); + 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() { + textureStore_d2b565(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_d2b565(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.glsl b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.glsl new file mode 100644 index 0000000000..2e4bde5be4 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.glsl @@ -0,0 +1,58 @@ +#version 310 es + +layout(rgba8ui) uniform highp writeonly uimage3D arg_0; +void textureStore_d2b565() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +vec4 vertex_main() { + textureStore_d2b565(); + return vec4(0.0f); +} + +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 mediump float; + +layout(rgba8ui) uniform highp writeonly uimage3D arg_0; +void textureStore_d2b565() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +void fragment_main() { + textureStore_d2b565(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +layout(rgba8ui) uniform highp writeonly uimage3D arg_0; +void textureStore_d2b565() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +void compute_main() { + textureStore_d2b565(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.msl b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.msl new file mode 100644 index 0000000000..1e7ab29de5 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.msl @@ -0,0 +1,35 @@ +#include + +using namespace metal; +void textureStore_d2b565(texture3d tint_symbol_1) { + uint3 arg_1 = uint3(1u); + uint4 arg_2 = uint4(1u); + tint_symbol_1.write(arg_2, uint3(arg_1)); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner(texture3d tint_symbol_2) { + textureStore_d2b565(tint_symbol_2); + return float4(0.0f); +} + +vertex tint_symbol vertex_main(texture3d tint_symbol_3 [[texture(0)]]) { + float4 const inner_result = vertex_main_inner(tint_symbol_3); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main(texture3d tint_symbol_4 [[texture(0)]]) { + textureStore_d2b565(tint_symbol_4); + return; +} + +kernel void compute_main(texture3d tint_symbol_5 [[texture(0)]]) { + textureStore_d2b565(tint_symbol_5); + return; +} + diff --git a/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.spvasm b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.spvasm new file mode 100644 index 0000000000..156622dbe1 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.spvasm @@ -0,0 +1,86 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 46 +; Schema: 0 + OpCapability Shader + 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 %arg_0 "arg_0" + OpName %textureStore_d2b565 "textureStore_d2b565" + OpName %arg_1 "arg_1" + OpName %arg_2 "arg_2" + 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 + OpDecorate %arg_0 NonReadable + OpDecorate %arg_0 DescriptorSet 1 + OpDecorate %arg_0 Binding 0 + %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 + %uint = OpTypeInt 32 0 + %11 = OpTypeImage %uint 3D 0 0 0 2 Rgba8ui +%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11 + %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant + %void = OpTypeVoid + %13 = OpTypeFunction %void + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %19 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %22 = OpConstantNull %v3uint + %v4uint = OpTypeVector %uint 4 + %24 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1 +%_ptr_Function_v4uint = OpTypePointer Function %v4uint + %27 = OpConstantNull %v4uint + %32 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%textureStore_d2b565 = OpFunction %void None %13 + %16 = OpLabel + %arg_1 = OpVariable %_ptr_Function_v3uint Function %22 + %arg_2 = OpVariable %_ptr_Function_v4uint Function %27 + OpStore %arg_1 %19 + OpStore %arg_2 %24 + %29 = OpLoad %11 %arg_0 + %30 = OpLoad %v3uint %arg_1 + %31 = OpLoad %v4uint %arg_2 + OpImageWrite %29 %30 %31 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %32 + %34 = OpLabel + %35 = OpFunctionCall %void %textureStore_d2b565 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %13 + %37 = OpLabel + %38 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %38 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %13 + %41 = OpLabel + %42 = OpFunctionCall %void %textureStore_d2b565 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %13 + %44 = OpLabel + %45 = OpFunctionCall %void %textureStore_d2b565 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.wgsl b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.wgsl new file mode 100644 index 0000000000..fd9e2b3455 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d2b565.wgsl.expected.wgsl @@ -0,0 +1,23 @@ +@group(1) @binding(0) var arg_0 : texture_storage_3d; + +fn textureStore_d2b565() { + var arg_1 = vec3(1u); + var arg_2 = vec4(1u); + textureStore(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_d2b565(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_d2b565(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_d2b565(); +} diff --git a/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl new file mode 100644 index 0000000000..26c39fff7f --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl @@ -0,0 +1,46 @@ +// Copyright 2023 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 +//////////////////////////////////////////////////////////////////////////////// + +@group(1) @binding(0) var arg_0: texture_storage_3d; + +// fn textureStore(texture: texture_storage_3d, coords: vec3, value: vec4) +fn textureStore_d4aa95() { + var arg_1 = vec3(1u); + var arg_2 = vec4(1u); + textureStore(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_d4aa95(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_d4aa95(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_d4aa95(); +} diff --git a/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..9749be0497 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.dxc.hlsl @@ -0,0 +1,34 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_d4aa95() { + uint3 arg_1 = (1u).xxx; + uint4 arg_2 = (1u).xxxx; + arg_0[arg_1] = arg_2; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_d4aa95(); + 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() { + textureStore_d4aa95(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_d4aa95(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..9749be0497 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.fxc.hlsl @@ -0,0 +1,34 @@ +RWTexture3D arg_0 : register(u0, space1); + +void textureStore_d4aa95() { + uint3 arg_1 = (1u).xxx; + uint4 arg_2 = (1u).xxxx; + arg_0[arg_1] = arg_2; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + textureStore_d4aa95(); + 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() { + textureStore_d4aa95(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + textureStore_d4aa95(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.glsl b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.glsl new file mode 100644 index 0000000000..3a743d9caa --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.glsl @@ -0,0 +1,58 @@ +#version 310 es + +layout(rgba32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_d4aa95() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +vec4 vertex_main() { + textureStore_d4aa95(); + return vec4(0.0f); +} + +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 mediump float; + +layout(rgba32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_d4aa95() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +void fragment_main() { + textureStore_d4aa95(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +layout(rgba32ui) uniform highp writeonly uimage3D arg_0; +void textureStore_d4aa95() { + uvec3 arg_1 = uvec3(1u); + uvec4 arg_2 = uvec4(1u); + imageStore(arg_0, ivec3(arg_1), arg_2); +} + +void compute_main() { + textureStore_d4aa95(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.msl b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.msl new file mode 100644 index 0000000000..74abfdbc04 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.msl @@ -0,0 +1,35 @@ +#include + +using namespace metal; +void textureStore_d4aa95(texture3d tint_symbol_1) { + uint3 arg_1 = uint3(1u); + uint4 arg_2 = uint4(1u); + tint_symbol_1.write(arg_2, uint3(arg_1)); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner(texture3d tint_symbol_2) { + textureStore_d4aa95(tint_symbol_2); + return float4(0.0f); +} + +vertex tint_symbol vertex_main(texture3d tint_symbol_3 [[texture(0)]]) { + float4 const inner_result = vertex_main_inner(tint_symbol_3); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main(texture3d tint_symbol_4 [[texture(0)]]) { + textureStore_d4aa95(tint_symbol_4); + return; +} + +kernel void compute_main(texture3d tint_symbol_5 [[texture(0)]]) { + textureStore_d4aa95(tint_symbol_5); + return; +} + diff --git a/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.spvasm b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.spvasm new file mode 100644 index 0000000000..9fc8fd1b8c --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.spvasm @@ -0,0 +1,86 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 46 +; Schema: 0 + OpCapability Shader + 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 %arg_0 "arg_0" + OpName %textureStore_d4aa95 "textureStore_d4aa95" + OpName %arg_1 "arg_1" + OpName %arg_2 "arg_2" + 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 + OpDecorate %arg_0 NonReadable + OpDecorate %arg_0 DescriptorSet 1 + OpDecorate %arg_0 Binding 0 + %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 + %uint = OpTypeInt 32 0 + %11 = OpTypeImage %uint 3D 0 0 0 2 Rgba32ui +%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11 + %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant + %void = OpTypeVoid + %13 = OpTypeFunction %void + %v3uint = OpTypeVector %uint 3 + %uint_1 = OpConstant %uint 1 + %19 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 +%_ptr_Function_v3uint = OpTypePointer Function %v3uint + %22 = OpConstantNull %v3uint + %v4uint = OpTypeVector %uint 4 + %24 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1 +%_ptr_Function_v4uint = OpTypePointer Function %v4uint + %27 = OpConstantNull %v4uint + %32 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 +%textureStore_d4aa95 = OpFunction %void None %13 + %16 = OpLabel + %arg_1 = OpVariable %_ptr_Function_v3uint Function %22 + %arg_2 = OpVariable %_ptr_Function_v4uint Function %27 + OpStore %arg_1 %19 + OpStore %arg_2 %24 + %29 = OpLoad %11 %arg_0 + %30 = OpLoad %v3uint %arg_1 + %31 = OpLoad %v4uint %arg_2 + OpImageWrite %29 %30 %31 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %32 + %34 = OpLabel + %35 = OpFunctionCall %void %textureStore_d4aa95 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %13 + %37 = OpLabel + %38 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %38 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %13 + %41 = OpLabel + %42 = OpFunctionCall %void %textureStore_d4aa95 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %13 + %44 = OpLabel + %45 = OpFunctionCall %void %textureStore_d4aa95 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.wgsl b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.wgsl new file mode 100644 index 0000000000..15d8887eb8 --- /dev/null +++ b/test/tint/builtins/gen/var/textureStore/d4aa95.wgsl.expected.wgsl @@ -0,0 +1,23 @@ +@group(1) @binding(0) var arg_0 : texture_storage_3d; + +fn textureStore_d4aa95() { + var arg_1 = vec3(1u); + var arg_2 = vec4(1u); + textureStore(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + textureStore_d4aa95(); + return vec4(); +} + +@fragment +fn fragment_main() { + textureStore_d4aa95(); +} + +@compute @workgroup_size(1) +fn compute_main() { + textureStore_d4aa95(); +}