[glsl][ir] Fix `bitfieldInsert` GLSL def entry.
Fix the `bitfieldInsert` def entry to have the correct name. Fixes using
the vector overloads for insert.
Bug: 42251044
Change-Id: I4075085c689b1309465ab05bbfe857fadd92eb5d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/207214
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/glsl/glsl.def b/src/tint/lang/glsl/glsl.def
index 53d4d61..eab309b 100644
--- a/src/tint/lang/glsl/glsl.def
+++ b/src/tint/lang/glsl/glsl.def
@@ -118,7 +118,7 @@
implicit(T: iu32, N: num) fn bitfieldExtract(value: vec<N, T>, offset: i32, bits: i32) -> vec<N, T>
implicit(T: iu32) fn bitfieldInsert(base: T, insert: T, offset: i32, bits: i32) -> T
-implicit(T: iu32, N: num) fn bitfieldExtract(base: vec<N, T>,
+implicit(T: iu32, N: num) fn bitfieldInsert(base: vec<N, T>,
insert: vec<N, T>,
offset: i32,
bits: i32) -> vec<N, T>
diff --git a/src/tint/lang/glsl/intrinsic/data.cc b/src/tint/lang/glsl/intrinsic/data.cc
index 44bfd68..f4aa29a 100644
--- a/src/tint/lang/glsl/intrinsic/data.cc
+++ b/src/tint/lang/glsl/intrinsic/data.cc
@@ -766,12 +766,12 @@
{
/* [0] */
/* usage */ core::ParameterUsage::kBase,
- /* matcher_indices */ MatcherIndicesIndex(22),
+ /* matcher_indices */ MatcherIndicesIndex(3),
},
{
/* [1] */
/* usage */ core::ParameterUsage::kInsert,
- /* matcher_indices */ MatcherIndicesIndex(22),
+ /* matcher_indices */ MatcherIndicesIndex(3),
},
{
/* [2] */
@@ -786,12 +786,12 @@
{
/* [4] */
/* usage */ core::ParameterUsage::kBase,
- /* matcher_indices */ MatcherIndicesIndex(3),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [5] */
/* usage */ core::ParameterUsage::kInsert,
- /* matcher_indices */ MatcherIndicesIndex(3),
+ /* matcher_indices */ MatcherIndicesIndex(22),
},
{
/* [6] */
@@ -1251,39 +1251,6 @@
{
/* [16] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 3,
- /* num_explicit_templates */ 0,
- /* num_templates */ 1,
- /* templates */ TemplateIndex(0),
- /* parameters */ ParameterIndex(10),
- /* return_matcher_indices */ MatcherIndicesIndex(3),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [17] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 3,
- /* num_explicit_templates */ 0,
- /* num_templates */ 2,
- /* templates */ TemplateIndex(2),
- /* parameters */ ParameterIndex(13),
- /* return_matcher_indices */ MatcherIndicesIndex(22),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [18] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 4,
- /* num_explicit_templates */ 0,
- /* num_templates */ 2,
- /* templates */ TemplateIndex(2),
- /* parameters */ ParameterIndex(0),
- /* return_matcher_indices */ MatcherIndicesIndex(22),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [19] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
/* num_templates */ 0,
@@ -1293,7 +1260,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [20] */
+ /* [17] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -1304,7 +1271,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [21] */
+ /* [18] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -1315,7 +1282,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [22] */
+ /* [19] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -1326,7 +1293,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [23] */
+ /* [20] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -1337,7 +1304,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [24] */
+ /* [21] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -1348,7 +1315,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [25] */
+ /* [22] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -1359,7 +1326,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [26] */
+ /* [23] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -1370,7 +1337,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [27] */
+ /* [24] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -1381,7 +1348,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [28] */
+ /* [25] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -1392,7 +1359,51 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
+ /* [26] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 3,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(0),
+ /* parameters */ ParameterIndex(10),
+ /* return_matcher_indices */ MatcherIndicesIndex(3),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [27] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 3,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 2,
+ /* templates */ TemplateIndex(2),
+ /* parameters */ ParameterIndex(13),
+ /* return_matcher_indices */ MatcherIndicesIndex(22),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [28] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 4,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(0),
+ /* parameters */ ParameterIndex(0),
+ /* return_matcher_indices */ MatcherIndicesIndex(3),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
/* [29] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 4,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 2,
+ /* templates */ TemplateIndex(2),
+ /* parameters */ ParameterIndex(4),
+ /* return_matcher_indices */ MatcherIndicesIndex(22),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [30] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 0,
/* num_explicit_templates */ 0,
@@ -1403,7 +1414,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [30] */
+ /* [31] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -1414,7 +1425,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [31] */
+ /* [32] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -1425,17 +1436,6 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [32] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 4,
- /* num_explicit_templates */ 0,
- /* num_templates */ 1,
- /* templates */ TemplateIndex(0),
- /* parameters */ ParameterIndex(4),
- /* return_matcher_indices */ MatcherIndicesIndex(3),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
/* [33] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
@@ -1467,80 +1467,80 @@
/* [0] */
/* fn barrier() */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(29),
+ /* overloads */ OverloadIndex(30),
},
{
/* [1] */
/* fn memoryBarrierBuffer() */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(29),
+ /* overloads */ OverloadIndex(30),
},
{
/* [2] */
/* fn memoryBarrierImage() */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(29),
+ /* overloads */ OverloadIndex(30),
},
{
/* [3] */
/* fn atomicCompSwap[T : iu32](ptr<workgroup_or_storage, atomic<T>, read_write>, compare_value: T, value: T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(30),
+ /* overloads */ OverloadIndex(31),
},
{
/* [4] */
/* fn atomicSub[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(31),
+ /* overloads */ OverloadIndex(32),
},
{
/* [5] */
/* fn floatBitsToInt(value: f32) -> i32 */
/* fn floatBitsToInt[N : num](value: vec<N, f32>) -> vec<N, i32> */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(19),
+ /* overloads */ OverloadIndex(16),
},
{
/* [6] */
/* fn floatBitsToUint(value: f32) -> u32 */
/* fn floatBitsToUint[N : num](value: vec<N, f32>) -> vec<N, u32> */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(21),
+ /* overloads */ OverloadIndex(18),
},
{
/* [7] */
/* fn intBitsToFloat(value: i32) -> f32 */
/* fn intBitsToFloat[N : num](value: vec<N, i32>) -> vec<N, f32> */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(23),
+ /* overloads */ OverloadIndex(20),
},
{
/* [8] */
/* fn uintBitsToFloat(value: u32) -> f32 */
/* fn uintBitsToFloat[N : num](value: vec<N, u32>) -> vec<N, f32> */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(25),
+ /* overloads */ OverloadIndex(22),
},
{
/* [9] */
/* fn bitCount[T : iu32](value: T) -> i32 */
/* fn bitCount[T : iu32, N : num](value: vec<N, T>) -> vec<N, i32> */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(27),
+ /* overloads */ OverloadIndex(24),
},
{
/* [10] */
/* fn bitfieldExtract[T : iu32](value: T, offset: i32, bits: i32) -> T */
/* fn bitfieldExtract[T : iu32, N : num](value: vec<N, T>, offset: i32, bits: i32) -> vec<N, T> */
- /* fn bitfieldExtract[T : iu32, N : num](base: vec<N, T>, insert: vec<N, T>, offset: i32, bits: i32) -> vec<N, T> */
- /* num overloads */ 3,
- /* overloads */ OverloadIndex(16),
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(26),
},
{
/* [11] */
/* fn bitfieldInsert[T : iu32](base: T, insert: T, offset: i32, bits: i32) -> T */
- /* num overloads */ 1,
- /* overloads */ OverloadIndex(32),
+ /* fn bitfieldInsert[T : iu32, N : num](base: vec<N, T>, insert: vec<N, T>, offset: i32, bits: i32) -> vec<N, T> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(28),
},
{
/* [12] */
diff --git a/test/tint/builtins/gen/var/insertBits/3c7ba5.wgsl.expected.ir.glsl b/test/tint/builtins/gen/var/insertBits/3c7ba5.wgsl.expected.ir.glsl
index f7bcd30..b39716c 100644
--- a/test/tint/builtins/gen/var/insertBits/3c7ba5.wgsl.expected.ir.glsl
+++ b/test/tint/builtins/gen/var/insertBits/3c7ba5.wgsl.expected.ir.glsl
@@ -1,234 +1,86 @@
-SKIP: FAILED
+#version 310 es
+precision highp float;
+precision highp int;
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec2<u32>;
-
-fn insertBits_3c7ba5() -> vec2<u32> {
- var arg_0 = vec2<u32>(1u);
- var arg_1 = vec2<u32>(1u);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec2<u32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(binding = 0, std430)
+buffer tint_symbol_1_1_ssbo {
+ uvec2 tint_symbol;
+} v;
+uvec2 insertBits_3c7ba5() {
+ uvec2 arg_0 = uvec2(1u);
+ uvec2 arg_1 = uvec2(1u);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ uvec2 v_1 = arg_0;
+ uvec2 v_2 = arg_1;
+ uint v_3 = arg_3;
+ uint v_4 = min(arg_2, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ uvec2 res = bitfieldInsert(v_1, v_2, v_6, int(v_5));
return res;
}
-
-struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec2<u32>,
+void main() {
+ v.tint_symbol = insertBits_3c7ba5();
}
+#version 310 es
-@fragment
-fn fragment_main() {
- prevent_dce = insertBits_3c7ba5();
-}
-
-Failed to generate: :24:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec2<u32>, vec2<u32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %16:vec2<u32> = glsl.bitfieldInsert %7, %8, %14, %15
- ^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
- $B2: {
- ^^^
-
-note: # Disassembly
-tint_symbol_1 = struct @align(8), @block {
- tint_symbol:vec2<u32> @offset(0)
-}
-
-$B1: { # root
- %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
-}
-
-%insertBits_3c7ba5 = func():vec2<u32> {
- $B2: {
- %arg_0:ptr<function, vec2<u32>, read_write> = var, vec2<u32>(1u)
- %arg_1:ptr<function, vec2<u32>, read_write> = var, vec2<u32>(1u)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:vec2<u32> = load %arg_0
- %8:vec2<u32> = load %arg_1
- %9:u32 = load %arg_2
- %10:u32 = load %arg_3
- %11:u32 = min %9, 32u
- %12:u32 = sub 32u, %11
- %13:u32 = min %10, %12
- %14:i32 = convert %11
- %15:i32 = convert %13
- %16:vec2<u32> = glsl.bitfieldInsert %7, %8, %14, %15
- %res:ptr<function, vec2<u32>, read_write> = var, %16
- %18:vec2<u32> = load %res
- ret %18
- }
-}
-%fragment_main = @fragment func():void {
- $B3: {
- %20:vec2<u32> = call %insertBits_3c7ba5
- %21:ptr<storage, vec2<u32>, read_write> = access %1, 0u
- store %21, %20
- ret
- }
-}
-
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec2<u32>;
-
-fn insertBits_3c7ba5() -> vec2<u32> {
- var arg_0 = vec2<u32>(1u);
- var arg_1 = vec2<u32>(1u);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec2<u32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(binding = 0, std430)
+buffer tint_symbol_1_1_ssbo {
+ uvec2 tint_symbol;
+} v;
+uvec2 insertBits_3c7ba5() {
+ uvec2 arg_0 = uvec2(1u);
+ uvec2 arg_1 = uvec2(1u);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ uvec2 v_1 = arg_0;
+ uvec2 v_2 = arg_1;
+ uint v_3 = arg_3;
+ uint v_4 = min(arg_2, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ uvec2 res = bitfieldInsert(v_1, v_2, v_6, int(v_5));
return res;
}
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ v.tint_symbol = insertBits_3c7ba5();
+}
+#version 310 es
+
struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec2<u32>,
-}
+ vec4 pos;
+ uvec2 prevent_dce;
+};
-@compute @workgroup_size(1)
-fn compute_main() {
- prevent_dce = insertBits_3c7ba5();
-}
-
-Failed to generate: :24:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec2<u32>, vec2<u32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %16:vec2<u32> = glsl.bitfieldInsert %7, %8, %14, %15
- ^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
- $B2: {
- ^^^
-
-note: # Disassembly
-tint_symbol_1 = struct @align(8), @block {
- tint_symbol:vec2<u32> @offset(0)
-}
-
-$B1: { # root
- %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
-}
-
-%insertBits_3c7ba5 = func():vec2<u32> {
- $B2: {
- %arg_0:ptr<function, vec2<u32>, read_write> = var, vec2<u32>(1u)
- %arg_1:ptr<function, vec2<u32>, read_write> = var, vec2<u32>(1u)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:vec2<u32> = load %arg_0
- %8:vec2<u32> = load %arg_1
- %9:u32 = load %arg_2
- %10:u32 = load %arg_3
- %11:u32 = min %9, 32u
- %12:u32 = sub 32u, %11
- %13:u32 = min %10, %12
- %14:i32 = convert %11
- %15:i32 = convert %13
- %16:vec2<u32> = glsl.bitfieldInsert %7, %8, %14, %15
- %res:ptr<function, vec2<u32>, read_write> = var, %16
- %18:vec2<u32> = load %res
- ret %18
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B3: {
- %20:vec2<u32> = call %insertBits_3c7ba5
- %21:ptr<storage, vec2<u32>, read_write> = access %1, 0u
- store %21, %20
- ret
- }
-}
-
-
-fn insertBits_3c7ba5() -> vec2<u32> {
- var arg_0 = vec2<u32>(1u);
- var arg_1 = vec2<u32>(1u);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec2<u32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(location = 0) flat out uvec2 vertex_main_loc0_Output;
+uvec2 insertBits_3c7ba5() {
+ uvec2 arg_0 = uvec2(1u);
+ uvec2 arg_1 = uvec2(1u);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ uvec2 v = arg_0;
+ uvec2 v_1 = arg_1;
+ uint v_2 = arg_3;
+ uint v_3 = min(arg_2, 32u);
+ uint v_4 = min(v_2, (32u - v_3));
+ int v_5 = int(v_3);
+ uvec2 res = bitfieldInsert(v, v_1, v_5, int(v_4));
return res;
}
-
-struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec2<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
- var tint_symbol : VertexOutput;
- tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+ VertexOutput tint_symbol = VertexOutput(vec4(0.0f), uvec2(0u));
+ tint_symbol.pos = vec4(0.0f);
tint_symbol.prevent_dce = insertBits_3c7ba5();
return tint_symbol;
}
-
-Failed to generate: :21:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec2<u32>, vec2<u32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %15:vec2<u32> = glsl.bitfieldInsert %6, %7, %13, %14
- ^^^^^^^^^^^^^^^^^^^
-
-:7:3 note: in block
- $B1: {
- ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
- pos:vec4<f32> @offset(0), @builtin(position)
- prevent_dce:vec2<u32> @offset(16), @location(0), @interpolate(flat)
+void main() {
+ VertexOutput v_6 = vertex_main_inner();
+ gl_Position = v_6.pos;
+ gl_Position[1u] = -(gl_Position.y);
+ gl_Position[2u] = ((2.0f * gl_Position.z) - gl_Position.w);
+ vertex_main_loc0_Output = v_6.prevent_dce;
+ gl_PointSize = 1.0f;
}
-
-%insertBits_3c7ba5 = func():vec2<u32> {
- $B1: {
- %arg_0:ptr<function, vec2<u32>, read_write> = var, vec2<u32>(1u)
- %arg_1:ptr<function, vec2<u32>, read_write> = var, vec2<u32>(1u)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %6:vec2<u32> = load %arg_0
- %7:vec2<u32> = load %arg_1
- %8:u32 = load %arg_2
- %9:u32 = load %arg_3
- %10:u32 = min %8, 32u
- %11:u32 = sub 32u, %10
- %12:u32 = min %9, %11
- %13:i32 = convert %10
- %14:i32 = convert %12
- %15:vec2<u32> = glsl.bitfieldInsert %6, %7, %13, %14
- %res:ptr<function, vec2<u32>, read_write> = var, %15
- %17:vec2<u32> = load %res
- ret %17
- }
-}
-%vertex_main = @vertex func():VertexOutput {
- $B2: {
- %tint_symbol:ptr<function, VertexOutput, read_write> = var
- %20:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
- store %20, vec4<f32>(0.0f)
- %21:ptr<function, vec2<u32>, read_write> = access %tint_symbol, 1u
- %22:vec2<u32> = call %insertBits_3c7ba5
- store %21, %22
- %23:VertexOutput = load %tint_symbol
- ret %23
- }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/insertBits/428b0b.wgsl.expected.ir.glsl b/test/tint/builtins/gen/var/insertBits/428b0b.wgsl.expected.ir.glsl
index 8e50bee..4312432 100644
--- a/test/tint/builtins/gen/var/insertBits/428b0b.wgsl.expected.ir.glsl
+++ b/test/tint/builtins/gen/var/insertBits/428b0b.wgsl.expected.ir.glsl
@@ -1,234 +1,86 @@
-SKIP: FAILED
+#version 310 es
+precision highp float;
+precision highp int;
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<i32>;
-
-fn insertBits_428b0b() -> vec3<i32> {
- var arg_0 = vec3<i32>(1i);
- var arg_1 = vec3<i32>(1i);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec3<i32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(binding = 0, std430)
+buffer tint_symbol_1_1_ssbo {
+ ivec3 tint_symbol;
+} v;
+ivec3 insertBits_428b0b() {
+ ivec3 arg_0 = ivec3(1);
+ ivec3 arg_1 = ivec3(1);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ ivec3 v_1 = arg_0;
+ ivec3 v_2 = arg_1;
+ uint v_3 = arg_3;
+ uint v_4 = min(arg_2, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ ivec3 res = bitfieldInsert(v_1, v_2, v_6, int(v_5));
return res;
}
-
-struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec3<i32>,
+void main() {
+ v.tint_symbol = insertBits_428b0b();
}
+#version 310 es
-@fragment
-fn fragment_main() {
- prevent_dce = insertBits_428b0b();
-}
-
-Failed to generate: :24:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec3<i32>, vec3<i32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %16:vec3<i32> = glsl.bitfieldInsert %7, %8, %14, %15
- ^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
- $B2: {
- ^^^
-
-note: # Disassembly
-tint_symbol_1 = struct @align(16), @block {
- tint_symbol:vec3<i32> @offset(0)
-}
-
-$B1: { # root
- %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
-}
-
-%insertBits_428b0b = func():vec3<i32> {
- $B2: {
- %arg_0:ptr<function, vec3<i32>, read_write> = var, vec3<i32>(1i)
- %arg_1:ptr<function, vec3<i32>, read_write> = var, vec3<i32>(1i)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:vec3<i32> = load %arg_0
- %8:vec3<i32> = load %arg_1
- %9:u32 = load %arg_2
- %10:u32 = load %arg_3
- %11:u32 = min %9, 32u
- %12:u32 = sub 32u, %11
- %13:u32 = min %10, %12
- %14:i32 = convert %11
- %15:i32 = convert %13
- %16:vec3<i32> = glsl.bitfieldInsert %7, %8, %14, %15
- %res:ptr<function, vec3<i32>, read_write> = var, %16
- %18:vec3<i32> = load %res
- ret %18
- }
-}
-%fragment_main = @fragment func():void {
- $B3: {
- %20:vec3<i32> = call %insertBits_428b0b
- %21:ptr<storage, vec3<i32>, read_write> = access %1, 0u
- store %21, %20
- ret
- }
-}
-
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<i32>;
-
-fn insertBits_428b0b() -> vec3<i32> {
- var arg_0 = vec3<i32>(1i);
- var arg_1 = vec3<i32>(1i);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec3<i32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(binding = 0, std430)
+buffer tint_symbol_1_1_ssbo {
+ ivec3 tint_symbol;
+} v;
+ivec3 insertBits_428b0b() {
+ ivec3 arg_0 = ivec3(1);
+ ivec3 arg_1 = ivec3(1);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ ivec3 v_1 = arg_0;
+ ivec3 v_2 = arg_1;
+ uint v_3 = arg_3;
+ uint v_4 = min(arg_2, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ ivec3 res = bitfieldInsert(v_1, v_2, v_6, int(v_5));
return res;
}
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ v.tint_symbol = insertBits_428b0b();
+}
+#version 310 es
+
struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec3<i32>,
-}
+ vec4 pos;
+ ivec3 prevent_dce;
+};
-@compute @workgroup_size(1)
-fn compute_main() {
- prevent_dce = insertBits_428b0b();
-}
-
-Failed to generate: :24:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec3<i32>, vec3<i32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %16:vec3<i32> = glsl.bitfieldInsert %7, %8, %14, %15
- ^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
- $B2: {
- ^^^
-
-note: # Disassembly
-tint_symbol_1 = struct @align(16), @block {
- tint_symbol:vec3<i32> @offset(0)
-}
-
-$B1: { # root
- %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
-}
-
-%insertBits_428b0b = func():vec3<i32> {
- $B2: {
- %arg_0:ptr<function, vec3<i32>, read_write> = var, vec3<i32>(1i)
- %arg_1:ptr<function, vec3<i32>, read_write> = var, vec3<i32>(1i)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:vec3<i32> = load %arg_0
- %8:vec3<i32> = load %arg_1
- %9:u32 = load %arg_2
- %10:u32 = load %arg_3
- %11:u32 = min %9, 32u
- %12:u32 = sub 32u, %11
- %13:u32 = min %10, %12
- %14:i32 = convert %11
- %15:i32 = convert %13
- %16:vec3<i32> = glsl.bitfieldInsert %7, %8, %14, %15
- %res:ptr<function, vec3<i32>, read_write> = var, %16
- %18:vec3<i32> = load %res
- ret %18
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B3: {
- %20:vec3<i32> = call %insertBits_428b0b
- %21:ptr<storage, vec3<i32>, read_write> = access %1, 0u
- store %21, %20
- ret
- }
-}
-
-
-fn insertBits_428b0b() -> vec3<i32> {
- var arg_0 = vec3<i32>(1i);
- var arg_1 = vec3<i32>(1i);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec3<i32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(location = 0) flat out ivec3 vertex_main_loc0_Output;
+ivec3 insertBits_428b0b() {
+ ivec3 arg_0 = ivec3(1);
+ ivec3 arg_1 = ivec3(1);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ ivec3 v = arg_0;
+ ivec3 v_1 = arg_1;
+ uint v_2 = arg_3;
+ uint v_3 = min(arg_2, 32u);
+ uint v_4 = min(v_2, (32u - v_3));
+ int v_5 = int(v_3);
+ ivec3 res = bitfieldInsert(v, v_1, v_5, int(v_4));
return res;
}
-
-struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec3<i32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
- var tint_symbol : VertexOutput;
- tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+ VertexOutput tint_symbol = VertexOutput(vec4(0.0f), ivec3(0));
+ tint_symbol.pos = vec4(0.0f);
tint_symbol.prevent_dce = insertBits_428b0b();
return tint_symbol;
}
-
-Failed to generate: :21:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec3<i32>, vec3<i32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %15:vec3<i32> = glsl.bitfieldInsert %6, %7, %13, %14
- ^^^^^^^^^^^^^^^^^^^
-
-:7:3 note: in block
- $B1: {
- ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
- pos:vec4<f32> @offset(0), @builtin(position)
- prevent_dce:vec3<i32> @offset(16), @location(0), @interpolate(flat)
+void main() {
+ VertexOutput v_6 = vertex_main_inner();
+ gl_Position = v_6.pos;
+ gl_Position[1u] = -(gl_Position.y);
+ gl_Position[2u] = ((2.0f * gl_Position.z) - gl_Position.w);
+ vertex_main_loc0_Output = v_6.prevent_dce;
+ gl_PointSize = 1.0f;
}
-
-%insertBits_428b0b = func():vec3<i32> {
- $B1: {
- %arg_0:ptr<function, vec3<i32>, read_write> = var, vec3<i32>(1i)
- %arg_1:ptr<function, vec3<i32>, read_write> = var, vec3<i32>(1i)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %6:vec3<i32> = load %arg_0
- %7:vec3<i32> = load %arg_1
- %8:u32 = load %arg_2
- %9:u32 = load %arg_3
- %10:u32 = min %8, 32u
- %11:u32 = sub 32u, %10
- %12:u32 = min %9, %11
- %13:i32 = convert %10
- %14:i32 = convert %12
- %15:vec3<i32> = glsl.bitfieldInsert %6, %7, %13, %14
- %res:ptr<function, vec3<i32>, read_write> = var, %15
- %17:vec3<i32> = load %res
- ret %17
- }
-}
-%vertex_main = @vertex func():VertexOutput {
- $B2: {
- %tint_symbol:ptr<function, VertexOutput, read_write> = var
- %20:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
- store %20, vec4<f32>(0.0f)
- %21:ptr<function, vec3<i32>, read_write> = access %tint_symbol, 1u
- %22:vec3<i32> = call %insertBits_428b0b
- store %21, %22
- %23:VertexOutput = load %tint_symbol
- ret %23
- }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/insertBits/51ede1.wgsl.expected.ir.glsl b/test/tint/builtins/gen/var/insertBits/51ede1.wgsl.expected.ir.glsl
index a46202e..3a11080 100644
--- a/test/tint/builtins/gen/var/insertBits/51ede1.wgsl.expected.ir.glsl
+++ b/test/tint/builtins/gen/var/insertBits/51ede1.wgsl.expected.ir.glsl
@@ -1,234 +1,86 @@
-SKIP: FAILED
+#version 310 es
+precision highp float;
+precision highp int;
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
-
-fn insertBits_51ede1() -> vec4<u32> {
- var arg_0 = vec4<u32>(1u);
- var arg_1 = vec4<u32>(1u);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec4<u32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(binding = 0, std430)
+buffer tint_symbol_1_1_ssbo {
+ uvec4 tint_symbol;
+} v;
+uvec4 insertBits_51ede1() {
+ uvec4 arg_0 = uvec4(1u);
+ uvec4 arg_1 = uvec4(1u);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ uvec4 v_1 = arg_0;
+ uvec4 v_2 = arg_1;
+ uint v_3 = arg_3;
+ uint v_4 = min(arg_2, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ uvec4 res = bitfieldInsert(v_1, v_2, v_6, int(v_5));
return res;
}
-
-struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec4<u32>,
+void main() {
+ v.tint_symbol = insertBits_51ede1();
}
+#version 310 es
-@fragment
-fn fragment_main() {
- prevent_dce = insertBits_51ede1();
-}
-
-Failed to generate: :24:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec4<u32>, vec4<u32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %16:vec4<u32> = glsl.bitfieldInsert %7, %8, %14, %15
- ^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
- $B2: {
- ^^^
-
-note: # Disassembly
-tint_symbol_1 = struct @align(16), @block {
- tint_symbol:vec4<u32> @offset(0)
-}
-
-$B1: { # root
- %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
-}
-
-%insertBits_51ede1 = func():vec4<u32> {
- $B2: {
- %arg_0:ptr<function, vec4<u32>, read_write> = var, vec4<u32>(1u)
- %arg_1:ptr<function, vec4<u32>, read_write> = var, vec4<u32>(1u)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:vec4<u32> = load %arg_0
- %8:vec4<u32> = load %arg_1
- %9:u32 = load %arg_2
- %10:u32 = load %arg_3
- %11:u32 = min %9, 32u
- %12:u32 = sub 32u, %11
- %13:u32 = min %10, %12
- %14:i32 = convert %11
- %15:i32 = convert %13
- %16:vec4<u32> = glsl.bitfieldInsert %7, %8, %14, %15
- %res:ptr<function, vec4<u32>, read_write> = var, %16
- %18:vec4<u32> = load %res
- ret %18
- }
-}
-%fragment_main = @fragment func():void {
- $B3: {
- %20:vec4<u32> = call %insertBits_51ede1
- %21:ptr<storage, vec4<u32>, read_write> = access %1, 0u
- store %21, %20
- ret
- }
-}
-
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec4<u32>;
-
-fn insertBits_51ede1() -> vec4<u32> {
- var arg_0 = vec4<u32>(1u);
- var arg_1 = vec4<u32>(1u);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec4<u32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(binding = 0, std430)
+buffer tint_symbol_1_1_ssbo {
+ uvec4 tint_symbol;
+} v;
+uvec4 insertBits_51ede1() {
+ uvec4 arg_0 = uvec4(1u);
+ uvec4 arg_1 = uvec4(1u);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ uvec4 v_1 = arg_0;
+ uvec4 v_2 = arg_1;
+ uint v_3 = arg_3;
+ uint v_4 = min(arg_2, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ uvec4 res = bitfieldInsert(v_1, v_2, v_6, int(v_5));
return res;
}
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ v.tint_symbol = insertBits_51ede1();
+}
+#version 310 es
+
struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec4<u32>,
-}
+ vec4 pos;
+ uvec4 prevent_dce;
+};
-@compute @workgroup_size(1)
-fn compute_main() {
- prevent_dce = insertBits_51ede1();
-}
-
-Failed to generate: :24:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec4<u32>, vec4<u32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %16:vec4<u32> = glsl.bitfieldInsert %7, %8, %14, %15
- ^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
- $B2: {
- ^^^
-
-note: # Disassembly
-tint_symbol_1 = struct @align(16), @block {
- tint_symbol:vec4<u32> @offset(0)
-}
-
-$B1: { # root
- %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
-}
-
-%insertBits_51ede1 = func():vec4<u32> {
- $B2: {
- %arg_0:ptr<function, vec4<u32>, read_write> = var, vec4<u32>(1u)
- %arg_1:ptr<function, vec4<u32>, read_write> = var, vec4<u32>(1u)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:vec4<u32> = load %arg_0
- %8:vec4<u32> = load %arg_1
- %9:u32 = load %arg_2
- %10:u32 = load %arg_3
- %11:u32 = min %9, 32u
- %12:u32 = sub 32u, %11
- %13:u32 = min %10, %12
- %14:i32 = convert %11
- %15:i32 = convert %13
- %16:vec4<u32> = glsl.bitfieldInsert %7, %8, %14, %15
- %res:ptr<function, vec4<u32>, read_write> = var, %16
- %18:vec4<u32> = load %res
- ret %18
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B3: {
- %20:vec4<u32> = call %insertBits_51ede1
- %21:ptr<storage, vec4<u32>, read_write> = access %1, 0u
- store %21, %20
- ret
- }
-}
-
-
-fn insertBits_51ede1() -> vec4<u32> {
- var arg_0 = vec4<u32>(1u);
- var arg_1 = vec4<u32>(1u);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec4<u32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(location = 0) flat out uvec4 vertex_main_loc0_Output;
+uvec4 insertBits_51ede1() {
+ uvec4 arg_0 = uvec4(1u);
+ uvec4 arg_1 = uvec4(1u);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ uvec4 v = arg_0;
+ uvec4 v_1 = arg_1;
+ uint v_2 = arg_3;
+ uint v_3 = min(arg_2, 32u);
+ uint v_4 = min(v_2, (32u - v_3));
+ int v_5 = int(v_3);
+ uvec4 res = bitfieldInsert(v, v_1, v_5, int(v_4));
return res;
}
-
-struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec4<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
- var tint_symbol : VertexOutput;
- tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+ VertexOutput tint_symbol = VertexOutput(vec4(0.0f), uvec4(0u));
+ tint_symbol.pos = vec4(0.0f);
tint_symbol.prevent_dce = insertBits_51ede1();
return tint_symbol;
}
-
-Failed to generate: :21:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec4<u32>, vec4<u32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %15:vec4<u32> = glsl.bitfieldInsert %6, %7, %13, %14
- ^^^^^^^^^^^^^^^^^^^
-
-:7:3 note: in block
- $B1: {
- ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
- pos:vec4<f32> @offset(0), @builtin(position)
- prevent_dce:vec4<u32> @offset(16), @location(0), @interpolate(flat)
+void main() {
+ VertexOutput v_6 = vertex_main_inner();
+ gl_Position = v_6.pos;
+ gl_Position[1u] = -(gl_Position.y);
+ gl_Position[2u] = ((2.0f * gl_Position.z) - gl_Position.w);
+ vertex_main_loc0_Output = v_6.prevent_dce;
+ gl_PointSize = 1.0f;
}
-
-%insertBits_51ede1 = func():vec4<u32> {
- $B1: {
- %arg_0:ptr<function, vec4<u32>, read_write> = var, vec4<u32>(1u)
- %arg_1:ptr<function, vec4<u32>, read_write> = var, vec4<u32>(1u)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %6:vec4<u32> = load %arg_0
- %7:vec4<u32> = load %arg_1
- %8:u32 = load %arg_2
- %9:u32 = load %arg_3
- %10:u32 = min %8, 32u
- %11:u32 = sub 32u, %10
- %12:u32 = min %9, %11
- %13:i32 = convert %10
- %14:i32 = convert %12
- %15:vec4<u32> = glsl.bitfieldInsert %6, %7, %13, %14
- %res:ptr<function, vec4<u32>, read_write> = var, %15
- %17:vec4<u32> = load %res
- ret %17
- }
-}
-%vertex_main = @vertex func():VertexOutput {
- $B2: {
- %tint_symbol:ptr<function, VertexOutput, read_write> = var
- %20:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
- store %20, vec4<f32>(0.0f)
- %21:ptr<function, vec4<u32>, read_write> = access %tint_symbol, 1u
- %22:vec4<u32> = call %insertBits_51ede1
- store %21, %22
- %23:VertexOutput = load %tint_symbol
- ret %23
- }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/insertBits/87826b.wgsl.expected.ir.glsl b/test/tint/builtins/gen/var/insertBits/87826b.wgsl.expected.ir.glsl
index 6e3a061..9200f44 100644
--- a/test/tint/builtins/gen/var/insertBits/87826b.wgsl.expected.ir.glsl
+++ b/test/tint/builtins/gen/var/insertBits/87826b.wgsl.expected.ir.glsl
@@ -1,234 +1,86 @@
-SKIP: FAILED
+#version 310 es
+precision highp float;
+precision highp int;
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-fn insertBits_87826b() -> vec3<u32> {
- var arg_0 = vec3<u32>(1u);
- var arg_1 = vec3<u32>(1u);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec3<u32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(binding = 0, std430)
+buffer tint_symbol_1_1_ssbo {
+ uvec3 tint_symbol;
+} v;
+uvec3 insertBits_87826b() {
+ uvec3 arg_0 = uvec3(1u);
+ uvec3 arg_1 = uvec3(1u);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ uvec3 v_1 = arg_0;
+ uvec3 v_2 = arg_1;
+ uint v_3 = arg_3;
+ uint v_4 = min(arg_2, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ uvec3 res = bitfieldInsert(v_1, v_2, v_6, int(v_5));
return res;
}
-
-struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec3<u32>,
+void main() {
+ v.tint_symbol = insertBits_87826b();
}
+#version 310 es
-@fragment
-fn fragment_main() {
- prevent_dce = insertBits_87826b();
-}
-
-Failed to generate: :24:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec3<u32>, vec3<u32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %16:vec3<u32> = glsl.bitfieldInsert %7, %8, %14, %15
- ^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
- $B2: {
- ^^^
-
-note: # Disassembly
-tint_symbol_1 = struct @align(16), @block {
- tint_symbol:vec3<u32> @offset(0)
-}
-
-$B1: { # root
- %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
-}
-
-%insertBits_87826b = func():vec3<u32> {
- $B2: {
- %arg_0:ptr<function, vec3<u32>, read_write> = var, vec3<u32>(1u)
- %arg_1:ptr<function, vec3<u32>, read_write> = var, vec3<u32>(1u)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:vec3<u32> = load %arg_0
- %8:vec3<u32> = load %arg_1
- %9:u32 = load %arg_2
- %10:u32 = load %arg_3
- %11:u32 = min %9, 32u
- %12:u32 = sub 32u, %11
- %13:u32 = min %10, %12
- %14:i32 = convert %11
- %15:i32 = convert %13
- %16:vec3<u32> = glsl.bitfieldInsert %7, %8, %14, %15
- %res:ptr<function, vec3<u32>, read_write> = var, %16
- %18:vec3<u32> = load %res
- ret %18
- }
-}
-%fragment_main = @fragment func():void {
- $B3: {
- %20:vec3<u32> = call %insertBits_87826b
- %21:ptr<storage, vec3<u32>, read_write> = access %1, 0u
- store %21, %20
- ret
- }
-}
-
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-fn insertBits_87826b() -> vec3<u32> {
- var arg_0 = vec3<u32>(1u);
- var arg_1 = vec3<u32>(1u);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec3<u32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(binding = 0, std430)
+buffer tint_symbol_1_1_ssbo {
+ uvec3 tint_symbol;
+} v;
+uvec3 insertBits_87826b() {
+ uvec3 arg_0 = uvec3(1u);
+ uvec3 arg_1 = uvec3(1u);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ uvec3 v_1 = arg_0;
+ uvec3 v_2 = arg_1;
+ uint v_3 = arg_3;
+ uint v_4 = min(arg_2, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ uvec3 res = bitfieldInsert(v_1, v_2, v_6, int(v_5));
return res;
}
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ v.tint_symbol = insertBits_87826b();
+}
+#version 310 es
+
struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec3<u32>,
-}
+ vec4 pos;
+ uvec3 prevent_dce;
+};
-@compute @workgroup_size(1)
-fn compute_main() {
- prevent_dce = insertBits_87826b();
-}
-
-Failed to generate: :24:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec3<u32>, vec3<u32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %16:vec3<u32> = glsl.bitfieldInsert %7, %8, %14, %15
- ^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
- $B2: {
- ^^^
-
-note: # Disassembly
-tint_symbol_1 = struct @align(16), @block {
- tint_symbol:vec3<u32> @offset(0)
-}
-
-$B1: { # root
- %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
-}
-
-%insertBits_87826b = func():vec3<u32> {
- $B2: {
- %arg_0:ptr<function, vec3<u32>, read_write> = var, vec3<u32>(1u)
- %arg_1:ptr<function, vec3<u32>, read_write> = var, vec3<u32>(1u)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:vec3<u32> = load %arg_0
- %8:vec3<u32> = load %arg_1
- %9:u32 = load %arg_2
- %10:u32 = load %arg_3
- %11:u32 = min %9, 32u
- %12:u32 = sub 32u, %11
- %13:u32 = min %10, %12
- %14:i32 = convert %11
- %15:i32 = convert %13
- %16:vec3<u32> = glsl.bitfieldInsert %7, %8, %14, %15
- %res:ptr<function, vec3<u32>, read_write> = var, %16
- %18:vec3<u32> = load %res
- ret %18
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B3: {
- %20:vec3<u32> = call %insertBits_87826b
- %21:ptr<storage, vec3<u32>, read_write> = access %1, 0u
- store %21, %20
- ret
- }
-}
-
-
-fn insertBits_87826b() -> vec3<u32> {
- var arg_0 = vec3<u32>(1u);
- var arg_1 = vec3<u32>(1u);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec3<u32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(location = 0) flat out uvec3 vertex_main_loc0_Output;
+uvec3 insertBits_87826b() {
+ uvec3 arg_0 = uvec3(1u);
+ uvec3 arg_1 = uvec3(1u);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ uvec3 v = arg_0;
+ uvec3 v_1 = arg_1;
+ uint v_2 = arg_3;
+ uint v_3 = min(arg_2, 32u);
+ uint v_4 = min(v_2, (32u - v_3));
+ int v_5 = int(v_3);
+ uvec3 res = bitfieldInsert(v, v_1, v_5, int(v_4));
return res;
}
-
-struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
- var tint_symbol : VertexOutput;
- tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+ VertexOutput tint_symbol = VertexOutput(vec4(0.0f), uvec3(0u));
+ tint_symbol.pos = vec4(0.0f);
tint_symbol.prevent_dce = insertBits_87826b();
return tint_symbol;
}
-
-Failed to generate: :21:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec3<u32>, vec3<u32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %15:vec3<u32> = glsl.bitfieldInsert %6, %7, %13, %14
- ^^^^^^^^^^^^^^^^^^^
-
-:7:3 note: in block
- $B1: {
- ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
- pos:vec4<f32> @offset(0), @builtin(position)
- prevent_dce:vec3<u32> @offset(16), @location(0), @interpolate(flat)
+void main() {
+ VertexOutput v_6 = vertex_main_inner();
+ gl_Position = v_6.pos;
+ gl_Position[1u] = -(gl_Position.y);
+ gl_Position[2u] = ((2.0f * gl_Position.z) - gl_Position.w);
+ vertex_main_loc0_Output = v_6.prevent_dce;
+ gl_PointSize = 1.0f;
}
-
-%insertBits_87826b = func():vec3<u32> {
- $B1: {
- %arg_0:ptr<function, vec3<u32>, read_write> = var, vec3<u32>(1u)
- %arg_1:ptr<function, vec3<u32>, read_write> = var, vec3<u32>(1u)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %6:vec3<u32> = load %arg_0
- %7:vec3<u32> = load %arg_1
- %8:u32 = load %arg_2
- %9:u32 = load %arg_3
- %10:u32 = min %8, 32u
- %11:u32 = sub 32u, %10
- %12:u32 = min %9, %11
- %13:i32 = convert %10
- %14:i32 = convert %12
- %15:vec3<u32> = glsl.bitfieldInsert %6, %7, %13, %14
- %res:ptr<function, vec3<u32>, read_write> = var, %15
- %17:vec3<u32> = load %res
- ret %17
- }
-}
-%vertex_main = @vertex func():VertexOutput {
- $B2: {
- %tint_symbol:ptr<function, VertexOutput, read_write> = var
- %20:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
- store %20, vec4<f32>(0.0f)
- %21:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
- %22:vec3<u32> = call %insertBits_87826b
- store %21, %22
- %23:VertexOutput = load %tint_symbol
- ret %23
- }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/insertBits/d86978.wgsl.expected.ir.glsl b/test/tint/builtins/gen/var/insertBits/d86978.wgsl.expected.ir.glsl
index a53551c..7aab0bd 100644
--- a/test/tint/builtins/gen/var/insertBits/d86978.wgsl.expected.ir.glsl
+++ b/test/tint/builtins/gen/var/insertBits/d86978.wgsl.expected.ir.glsl
@@ -1,234 +1,86 @@
-SKIP: FAILED
+#version 310 es
+precision highp float;
+precision highp int;
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
-
-fn insertBits_d86978() -> vec4<i32> {
- var arg_0 = vec4<i32>(1i);
- var arg_1 = vec4<i32>(1i);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec4<i32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(binding = 0, std430)
+buffer tint_symbol_1_1_ssbo {
+ ivec4 tint_symbol;
+} v;
+ivec4 insertBits_d86978() {
+ ivec4 arg_0 = ivec4(1);
+ ivec4 arg_1 = ivec4(1);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ ivec4 v_1 = arg_0;
+ ivec4 v_2 = arg_1;
+ uint v_3 = arg_3;
+ uint v_4 = min(arg_2, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ ivec4 res = bitfieldInsert(v_1, v_2, v_6, int(v_5));
return res;
}
-
-struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec4<i32>,
+void main() {
+ v.tint_symbol = insertBits_d86978();
}
+#version 310 es
-@fragment
-fn fragment_main() {
- prevent_dce = insertBits_d86978();
-}
-
-Failed to generate: :24:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec4<i32>, vec4<i32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %16:vec4<i32> = glsl.bitfieldInsert %7, %8, %14, %15
- ^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
- $B2: {
- ^^^
-
-note: # Disassembly
-tint_symbol_1 = struct @align(16), @block {
- tint_symbol:vec4<i32> @offset(0)
-}
-
-$B1: { # root
- %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
-}
-
-%insertBits_d86978 = func():vec4<i32> {
- $B2: {
- %arg_0:ptr<function, vec4<i32>, read_write> = var, vec4<i32>(1i)
- %arg_1:ptr<function, vec4<i32>, read_write> = var, vec4<i32>(1i)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:vec4<i32> = load %arg_0
- %8:vec4<i32> = load %arg_1
- %9:u32 = load %arg_2
- %10:u32 = load %arg_3
- %11:u32 = min %9, 32u
- %12:u32 = sub 32u, %11
- %13:u32 = min %10, %12
- %14:i32 = convert %11
- %15:i32 = convert %13
- %16:vec4<i32> = glsl.bitfieldInsert %7, %8, %14, %15
- %res:ptr<function, vec4<i32>, read_write> = var, %16
- %18:vec4<i32> = load %res
- ret %18
- }
-}
-%fragment_main = @fragment func():void {
- $B3: {
- %20:vec4<i32> = call %insertBits_d86978
- %21:ptr<storage, vec4<i32>, read_write> = access %1, 0u
- store %21, %20
- ret
- }
-}
-
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec4<i32>;
-
-fn insertBits_d86978() -> vec4<i32> {
- var arg_0 = vec4<i32>(1i);
- var arg_1 = vec4<i32>(1i);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec4<i32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(binding = 0, std430)
+buffer tint_symbol_1_1_ssbo {
+ ivec4 tint_symbol;
+} v;
+ivec4 insertBits_d86978() {
+ ivec4 arg_0 = ivec4(1);
+ ivec4 arg_1 = ivec4(1);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ ivec4 v_1 = arg_0;
+ ivec4 v_2 = arg_1;
+ uint v_3 = arg_3;
+ uint v_4 = min(arg_2, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ ivec4 res = bitfieldInsert(v_1, v_2, v_6, int(v_5));
return res;
}
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ v.tint_symbol = insertBits_d86978();
+}
+#version 310 es
+
struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec4<i32>,
-}
+ vec4 pos;
+ ivec4 prevent_dce;
+};
-@compute @workgroup_size(1)
-fn compute_main() {
- prevent_dce = insertBits_d86978();
-}
-
-Failed to generate: :24:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec4<i32>, vec4<i32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %16:vec4<i32> = glsl.bitfieldInsert %7, %8, %14, %15
- ^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
- $B2: {
- ^^^
-
-note: # Disassembly
-tint_symbol_1 = struct @align(16), @block {
- tint_symbol:vec4<i32> @offset(0)
-}
-
-$B1: { # root
- %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
-}
-
-%insertBits_d86978 = func():vec4<i32> {
- $B2: {
- %arg_0:ptr<function, vec4<i32>, read_write> = var, vec4<i32>(1i)
- %arg_1:ptr<function, vec4<i32>, read_write> = var, vec4<i32>(1i)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:vec4<i32> = load %arg_0
- %8:vec4<i32> = load %arg_1
- %9:u32 = load %arg_2
- %10:u32 = load %arg_3
- %11:u32 = min %9, 32u
- %12:u32 = sub 32u, %11
- %13:u32 = min %10, %12
- %14:i32 = convert %11
- %15:i32 = convert %13
- %16:vec4<i32> = glsl.bitfieldInsert %7, %8, %14, %15
- %res:ptr<function, vec4<i32>, read_write> = var, %16
- %18:vec4<i32> = load %res
- ret %18
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B3: {
- %20:vec4<i32> = call %insertBits_d86978
- %21:ptr<storage, vec4<i32>, read_write> = access %1, 0u
- store %21, %20
- ret
- }
-}
-
-
-fn insertBits_d86978() -> vec4<i32> {
- var arg_0 = vec4<i32>(1i);
- var arg_1 = vec4<i32>(1i);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec4<i32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(location = 0) flat out ivec4 vertex_main_loc0_Output;
+ivec4 insertBits_d86978() {
+ ivec4 arg_0 = ivec4(1);
+ ivec4 arg_1 = ivec4(1);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ ivec4 v = arg_0;
+ ivec4 v_1 = arg_1;
+ uint v_2 = arg_3;
+ uint v_3 = min(arg_2, 32u);
+ uint v_4 = min(v_2, (32u - v_3));
+ int v_5 = int(v_3);
+ ivec4 res = bitfieldInsert(v, v_1, v_5, int(v_4));
return res;
}
-
-struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec4<i32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
- var tint_symbol : VertexOutput;
- tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+ VertexOutput tint_symbol = VertexOutput(vec4(0.0f), ivec4(0));
+ tint_symbol.pos = vec4(0.0f);
tint_symbol.prevent_dce = insertBits_d86978();
return tint_symbol;
}
-
-Failed to generate: :21:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec4<i32>, vec4<i32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %15:vec4<i32> = glsl.bitfieldInsert %6, %7, %13, %14
- ^^^^^^^^^^^^^^^^^^^
-
-:7:3 note: in block
- $B1: {
- ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
- pos:vec4<f32> @offset(0), @builtin(position)
- prevent_dce:vec4<i32> @offset(16), @location(0), @interpolate(flat)
+void main() {
+ VertexOutput v_6 = vertex_main_inner();
+ gl_Position = v_6.pos;
+ gl_Position[1u] = -(gl_Position.y);
+ gl_Position[2u] = ((2.0f * gl_Position.z) - gl_Position.w);
+ vertex_main_loc0_Output = v_6.prevent_dce;
+ gl_PointSize = 1.0f;
}
-
-%insertBits_d86978 = func():vec4<i32> {
- $B1: {
- %arg_0:ptr<function, vec4<i32>, read_write> = var, vec4<i32>(1i)
- %arg_1:ptr<function, vec4<i32>, read_write> = var, vec4<i32>(1i)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %6:vec4<i32> = load %arg_0
- %7:vec4<i32> = load %arg_1
- %8:u32 = load %arg_2
- %9:u32 = load %arg_3
- %10:u32 = min %8, 32u
- %11:u32 = sub 32u, %10
- %12:u32 = min %9, %11
- %13:i32 = convert %10
- %14:i32 = convert %12
- %15:vec4<i32> = glsl.bitfieldInsert %6, %7, %13, %14
- %res:ptr<function, vec4<i32>, read_write> = var, %15
- %17:vec4<i32> = load %res
- ret %17
- }
-}
-%vertex_main = @vertex func():VertexOutput {
- $B2: {
- %tint_symbol:ptr<function, VertexOutput, read_write> = var
- %20:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
- store %20, vec4<f32>(0.0f)
- %21:ptr<function, vec4<i32>, read_write> = access %tint_symbol, 1u
- %22:vec4<i32> = call %insertBits_d86978
- store %21, %22
- %23:VertexOutput = load %tint_symbol
- ret %23
- }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/insertBits/fe6ba6.wgsl.expected.ir.glsl b/test/tint/builtins/gen/var/insertBits/fe6ba6.wgsl.expected.ir.glsl
index 101251a..a872bf5 100644
--- a/test/tint/builtins/gen/var/insertBits/fe6ba6.wgsl.expected.ir.glsl
+++ b/test/tint/builtins/gen/var/insertBits/fe6ba6.wgsl.expected.ir.glsl
@@ -1,234 +1,86 @@
-SKIP: FAILED
+#version 310 es
+precision highp float;
+precision highp int;
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec2<i32>;
-
-fn insertBits_fe6ba6() -> vec2<i32> {
- var arg_0 = vec2<i32>(1i);
- var arg_1 = vec2<i32>(1i);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec2<i32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(binding = 0, std430)
+buffer tint_symbol_1_1_ssbo {
+ ivec2 tint_symbol;
+} v;
+ivec2 insertBits_fe6ba6() {
+ ivec2 arg_0 = ivec2(1);
+ ivec2 arg_1 = ivec2(1);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ ivec2 v_1 = arg_0;
+ ivec2 v_2 = arg_1;
+ uint v_3 = arg_3;
+ uint v_4 = min(arg_2, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ ivec2 res = bitfieldInsert(v_1, v_2, v_6, int(v_5));
return res;
}
-
-struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec2<i32>,
+void main() {
+ v.tint_symbol = insertBits_fe6ba6();
}
+#version 310 es
-@fragment
-fn fragment_main() {
- prevent_dce = insertBits_fe6ba6();
-}
-
-Failed to generate: :24:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec2<i32>, vec2<i32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %16:vec2<i32> = glsl.bitfieldInsert %7, %8, %14, %15
- ^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
- $B2: {
- ^^^
-
-note: # Disassembly
-tint_symbol_1 = struct @align(8), @block {
- tint_symbol:vec2<i32> @offset(0)
-}
-
-$B1: { # root
- %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
-}
-
-%insertBits_fe6ba6 = func():vec2<i32> {
- $B2: {
- %arg_0:ptr<function, vec2<i32>, read_write> = var, vec2<i32>(1i)
- %arg_1:ptr<function, vec2<i32>, read_write> = var, vec2<i32>(1i)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:vec2<i32> = load %arg_0
- %8:vec2<i32> = load %arg_1
- %9:u32 = load %arg_2
- %10:u32 = load %arg_3
- %11:u32 = min %9, 32u
- %12:u32 = sub 32u, %11
- %13:u32 = min %10, %12
- %14:i32 = convert %11
- %15:i32 = convert %13
- %16:vec2<i32> = glsl.bitfieldInsert %7, %8, %14, %15
- %res:ptr<function, vec2<i32>, read_write> = var, %16
- %18:vec2<i32> = load %res
- ret %18
- }
-}
-%fragment_main = @fragment func():void {
- $B3: {
- %20:vec2<i32> = call %insertBits_fe6ba6
- %21:ptr<storage, vec2<i32>, read_write> = access %1, 0u
- store %21, %20
- ret
- }
-}
-
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec2<i32>;
-
-fn insertBits_fe6ba6() -> vec2<i32> {
- var arg_0 = vec2<i32>(1i);
- var arg_1 = vec2<i32>(1i);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec2<i32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(binding = 0, std430)
+buffer tint_symbol_1_1_ssbo {
+ ivec2 tint_symbol;
+} v;
+ivec2 insertBits_fe6ba6() {
+ ivec2 arg_0 = ivec2(1);
+ ivec2 arg_1 = ivec2(1);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ ivec2 v_1 = arg_0;
+ ivec2 v_2 = arg_1;
+ uint v_3 = arg_3;
+ uint v_4 = min(arg_2, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ ivec2 res = bitfieldInsert(v_1, v_2, v_6, int(v_5));
return res;
}
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ v.tint_symbol = insertBits_fe6ba6();
+}
+#version 310 es
+
struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec2<i32>,
-}
+ vec4 pos;
+ ivec2 prevent_dce;
+};
-@compute @workgroup_size(1)
-fn compute_main() {
- prevent_dce = insertBits_fe6ba6();
-}
-
-Failed to generate: :24:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec2<i32>, vec2<i32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %16:vec2<i32> = glsl.bitfieldInsert %7, %8, %14, %15
- ^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
- $B2: {
- ^^^
-
-note: # Disassembly
-tint_symbol_1 = struct @align(8), @block {
- tint_symbol:vec2<i32> @offset(0)
-}
-
-$B1: { # root
- %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
-}
-
-%insertBits_fe6ba6 = func():vec2<i32> {
- $B2: {
- %arg_0:ptr<function, vec2<i32>, read_write> = var, vec2<i32>(1i)
- %arg_1:ptr<function, vec2<i32>, read_write> = var, vec2<i32>(1i)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %7:vec2<i32> = load %arg_0
- %8:vec2<i32> = load %arg_1
- %9:u32 = load %arg_2
- %10:u32 = load %arg_3
- %11:u32 = min %9, 32u
- %12:u32 = sub 32u, %11
- %13:u32 = min %10, %12
- %14:i32 = convert %11
- %15:i32 = convert %13
- %16:vec2<i32> = glsl.bitfieldInsert %7, %8, %14, %15
- %res:ptr<function, vec2<i32>, read_write> = var, %16
- %18:vec2<i32> = load %res
- ret %18
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B3: {
- %20:vec2<i32> = call %insertBits_fe6ba6
- %21:ptr<storage, vec2<i32>, read_write> = access %1, 0u
- store %21, %20
- ret
- }
-}
-
-
-fn insertBits_fe6ba6() -> vec2<i32> {
- var arg_0 = vec2<i32>(1i);
- var arg_1 = vec2<i32>(1i);
- var arg_2 = 1u;
- var arg_3 = 1u;
- var res : vec2<i32> = insertBits(arg_0, arg_1, arg_2, arg_3);
+layout(location = 0) flat out ivec2 vertex_main_loc0_Output;
+ivec2 insertBits_fe6ba6() {
+ ivec2 arg_0 = ivec2(1);
+ ivec2 arg_1 = ivec2(1);
+ uint arg_2 = 1u;
+ uint arg_3 = 1u;
+ ivec2 v = arg_0;
+ ivec2 v_1 = arg_1;
+ uint v_2 = arg_3;
+ uint v_3 = min(arg_2, 32u);
+ uint v_4 = min(v_2, (32u - v_3));
+ int v_5 = int(v_3);
+ ivec2 res = bitfieldInsert(v, v_1, v_5, int(v_4));
return res;
}
-
-struct VertexOutput {
- @builtin(position)
- pos : vec4<f32>,
- @location(0) @interpolate(flat)
- prevent_dce : vec2<i32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
- var tint_symbol : VertexOutput;
- tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+ VertexOutput tint_symbol = VertexOutput(vec4(0.0f), ivec2(0));
+ tint_symbol.pos = vec4(0.0f);
tint_symbol.prevent_dce = insertBits_fe6ba6();
return tint_symbol;
}
-
-Failed to generate: :21:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec2<i32>, vec2<i32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %15:vec2<i32> = glsl.bitfieldInsert %6, %7, %13, %14
- ^^^^^^^^^^^^^^^^^^^
-
-:7:3 note: in block
- $B1: {
- ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
- pos:vec4<f32> @offset(0), @builtin(position)
- prevent_dce:vec2<i32> @offset(16), @location(0), @interpolate(flat)
+void main() {
+ VertexOutput v_6 = vertex_main_inner();
+ gl_Position = v_6.pos;
+ gl_Position[1u] = -(gl_Position.y);
+ gl_Position[2u] = ((2.0f * gl_Position.z) - gl_Position.w);
+ vertex_main_loc0_Output = v_6.prevent_dce;
+ gl_PointSize = 1.0f;
}
-
-%insertBits_fe6ba6 = func():vec2<i32> {
- $B1: {
- %arg_0:ptr<function, vec2<i32>, read_write> = var, vec2<i32>(1i)
- %arg_1:ptr<function, vec2<i32>, read_write> = var, vec2<i32>(1i)
- %arg_2:ptr<function, u32, read_write> = var, 1u
- %arg_3:ptr<function, u32, read_write> = var, 1u
- %6:vec2<i32> = load %arg_0
- %7:vec2<i32> = load %arg_1
- %8:u32 = load %arg_2
- %9:u32 = load %arg_3
- %10:u32 = min %8, 32u
- %11:u32 = sub 32u, %10
- %12:u32 = min %9, %11
- %13:i32 = convert %10
- %14:i32 = convert %12
- %15:vec2<i32> = glsl.bitfieldInsert %6, %7, %13, %14
- %res:ptr<function, vec2<i32>, read_write> = var, %15
- %17:vec2<i32> = load %res
- ret %17
- }
-}
-%vertex_main = @vertex func():VertexOutput {
- $B2: {
- %tint_symbol:ptr<function, VertexOutput, read_write> = var
- %20:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
- store %20, vec4<f32>(0.0f)
- %21:ptr<function, vec2<i32>, read_write> = access %tint_symbol, 1u
- %22:vec2<i32> = call %insertBits_fe6ba6
- store %21, %22
- %23:VertexOutput = load %tint_symbol
- ret %23
- }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/insertBits/vec3/i32.spvasm.expected.ir.glsl b/test/tint/builtins/insertBits/vec3/i32.spvasm.expected.ir.glsl
index 1d3e6c4..7b5ab20 100644
--- a/test/tint/builtins/insertBits/vec3/i32.spvasm.expected.ir.glsl
+++ b/test/tint/builtins/insertBits/vec3/i32.spvasm.expected.ir.glsl
@@ -1,60 +1,19 @@
-SKIP: FAILED
+#version 310 es
-
-fn f_1() {
- var v = vec3i();
- var n = vec3i();
- var offset_1 = 0u;
- var count = 0u;
- let x_16 = insertBits(v, n, offset_1, count);
- return;
+void f_1() {
+ ivec3 v = ivec3(0);
+ ivec3 n = ivec3(0);
+ uint offset_1 = 0u;
+ uint count = 0u;
+ ivec3 v_1 = v;
+ ivec3 v_2 = n;
+ uint v_3 = count;
+ uint v_4 = min(offset_1, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ ivec3 x_16 = bitfieldInsert(v_1, v_2, v_6, int(v_5));
}
-
-@compute @workgroup_size(1i, 1i, 1i)
-fn f() {
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
f_1();
}
-
-Failed to generate: :16:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec3<i32>, vec3<i32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %15:vec3<i32> = glsl.bitfieldInsert %6, %7, %13, %14
- ^^^^^^^^^^^^^^^^^^^
-
-:2:3 note: in block
- $B1: {
- ^^^
-
-note: # Disassembly
-%f_1 = func():void {
- $B1: {
- %v:ptr<function, vec3<i32>, read_write> = var, vec3<i32>(0i)
- %n:ptr<function, vec3<i32>, read_write> = var, vec3<i32>(0i)
- %offset_1:ptr<function, u32, read_write> = var, 0u
- %count:ptr<function, u32, read_write> = var, 0u
- %6:vec3<i32> = load %v
- %7:vec3<i32> = load %n
- %8:u32 = load %offset_1
- %9:u32 = load %count
- %10:u32 = min %8, 32u
- %11:u32 = sub 32u, %10
- %12:u32 = min %9, %11
- %13:i32 = convert %10
- %14:i32 = convert %12
- %15:vec3<i32> = glsl.bitfieldInsert %6, %7, %13, %14
- %x_16:vec3<i32> = let %15
- ret
- }
-}
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %18:void = call %f_1
- ret
- }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/insertBits/vec3/u32.spvasm.expected.ir.glsl b/test/tint/builtins/insertBits/vec3/u32.spvasm.expected.ir.glsl
index 34fdd67..b737b57 100644
--- a/test/tint/builtins/insertBits/vec3/u32.spvasm.expected.ir.glsl
+++ b/test/tint/builtins/insertBits/vec3/u32.spvasm.expected.ir.glsl
@@ -1,60 +1,19 @@
-SKIP: FAILED
+#version 310 es
-
-fn f_1() {
- var v = vec3u();
- var n = vec3u();
- var offset_1 = 0u;
- var count = 0u;
- let x_15 = insertBits(v, n, offset_1, count);
- return;
+void f_1() {
+ uvec3 v = uvec3(0u);
+ uvec3 n = uvec3(0u);
+ uint offset_1 = 0u;
+ uint count = 0u;
+ uvec3 v_1 = v;
+ uvec3 v_2 = n;
+ uint v_3 = count;
+ uint v_4 = min(offset_1, 32u);
+ uint v_5 = min(v_3, (32u - v_4));
+ int v_6 = int(v_4);
+ uvec3 x_15 = bitfieldInsert(v_1, v_2, v_6, int(v_5));
}
-
-@compute @workgroup_size(1i, 1i, 1i)
-fn f() {
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
f_1();
}
-
-Failed to generate: :16:21 error: glsl.bitfieldInsert: no matching call to 'glsl.bitfieldInsert(vec3<u32>, vec3<u32>, i32, i32)'
-
-1 candidate function:
- • 'glsl.bitfieldInsert(base: T ✗ , insert: T ✗ , offset: i32 ✓ , bits: i32 ✓ ) -> T' where:
- ✗ 'T' is 'i32' or 'u32'
-
- %15:vec3<u32> = glsl.bitfieldInsert %6, %7, %13, %14
- ^^^^^^^^^^^^^^^^^^^
-
-:2:3 note: in block
- $B1: {
- ^^^
-
-note: # Disassembly
-%f_1 = func():void {
- $B1: {
- %v:ptr<function, vec3<u32>, read_write> = var, vec3<u32>(0u)
- %n:ptr<function, vec3<u32>, read_write> = var, vec3<u32>(0u)
- %offset_1:ptr<function, u32, read_write> = var, 0u
- %count:ptr<function, u32, read_write> = var, 0u
- %6:vec3<u32> = load %v
- %7:vec3<u32> = load %n
- %8:u32 = load %offset_1
- %9:u32 = load %count
- %10:u32 = min %8, 32u
- %11:u32 = sub 32u, %10
- %12:u32 = min %9, %11
- %13:i32 = convert %10
- %14:i32 = convert %12
- %15:vec3<u32> = glsl.bitfieldInsert %6, %7, %13, %14
- %x_15:vec3<u32> = let %15
- ret
- }
-}
-%f = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %18:void = call %f_1
- ret
- }
-}
-
-
-tint executable returned error: exit status 1
