encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

builtins: Add firstLeadingBit 

Currently polyfilled for all backends.
HLSL should be able to map this to 'firstbithigh', but there might need
to be some special case handling for 0 (undocumented behavior). For now
just polyfill.

CTS tests: https://github.com/gpuweb/cts/pull/1004

Bug: tint:1367
Bug: tint:1449
Change-Id: I9c9a08ea93d1c4a602e0ab763e95e2eea336fb0d
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/81503
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
This commit is contained in:
Ben Clayton 2022-02-23 18:20:30 +00:00 committed by Tint LUCI CQ
parent 2680d1f846
commit 8169693136
60 changed files with 5366 additions and 1900 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3817
src/tint/builtin_table.inl
View File

File diff suppressed because it is too large Load Diff

2
src/tint/builtins.def
View File

@ -316,6 +316,8 @@ fn exp<N: num>(vec<N, f32>) -> vec<N, f32>
fn exp2(f32) -> f32
fn exp2<N: num>(vec<N, f32>) -> vec<N, f32>
fn faceForward<N: num>(vec<N, f32>, vec<N, f32>, vec<N, f32>) -> vec<N, f32>
fn firstLeadingBit<T: iu32>(T) -> T
fn firstLeadingBit<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
fn firstTrailingBit<T: iu32>(T) -> T
fn firstTrailingBit<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
fn floor(f32) -> f32

9
src/tint/program_builder.h
View File

@ -1596,6 +1596,15 @@ class ProgramBuilder {
Expr(std::forward<EXPR>(expr)));
}
/// @param expr the expression to perform a unary complement on
/// @return an ast::UnaryOpExpression that is the unary complement of the
/// input expression
template <typename EXPR>
const ast::UnaryOpExpression* Complement(EXPR&& expr) {
return create<ast::UnaryOpExpression>(ast::UnaryOp::kComplement,
Expr(std::forward<EXPR>(expr)));
}
/// @param source the source information
/// @param func the function name
/// @param args the function call arguments

5
src/tint/sem/builtin_type.cc
View File

@ -117,6 +117,9 @@ BuiltinType ParseBuiltinType(const std::string& name) {
if (name == "faceForward") {
return BuiltinType::kFaceForward;
}
if (name == "firstLeadingBit") {
return BuiltinType::kFirstLeadingBit;
}
if (name == "firstTrailingBit") {
return BuiltinType::kFirstTrailingBit;
}
@ -410,6 +413,8 @@ const char* str(BuiltinType i) {
return "exp2";
case BuiltinType::kFaceForward:
return "faceForward";
case BuiltinType::kFirstLeadingBit:
return "firstLeadingBit";
case BuiltinType::kFirstTrailingBit:
return "firstTrailingBit";
case BuiltinType::kFloor:

1
src/tint/sem/builtin_type.h
View File

@ -63,6 +63,7 @@ enum class BuiltinType {
kExp,
kExp2,
kFaceForward,
kFirstLeadingBit,
kFirstTrailingBit,
kFloor,
kFma,

115
src/tint/transform/builtin_polyfill.cc
View File

@ -55,10 +55,7 @@ struct BuiltinPolyfill::State {
return b.ty.vec<ProgramBuilder::u32>(width);
};
auto V = [&](uint32_t value) -> const ast::Expression* {
if (width == 1) {
return b.Expr(value);
}
return b.Construct(b.ty.vec<ProgramBuilder::u32>(width), value);
return ScalarOrVector(width, value);
};
b.Func(
name, {b.Param("v", T(ty))}, T(ty),
@ -120,10 +117,7 @@ struct BuiltinPolyfill::State {
return b.ty.vec<ProgramBuilder::u32>(width);
};
auto V = [&](uint32_t value) -> const ast::Expression* {
if (width == 1) {
return b.Expr(value);
}
return b.Construct(b.ty.vec<ProgramBuilder::u32>(width), value);
return ScalarOrVector(width, value);
};
auto B = [&](const ast::Expression* value) -> const ast::Expression* {
if (width == 1) {
@ -176,6 +170,87 @@ struct BuiltinPolyfill::State {
return name;
}
/// Builds the polyfill function for the `firstLeadingBit` builtin
/// @param ty the parameter and return type for the function
/// @return the polyfill function name
Symbol firstLeadingBit(const sem::Type* ty) {
auto name = b.Symbols().New("tint_first_leading_bit");
uint32_t width = WidthOf(ty);
// Returns either u32 or vecN<u32>
auto U = [&]() -> const ast::Type* {
if (width == 1) {
return b.ty.u32();
}
return b.ty.vec<ProgramBuilder::u32>(width);
};
auto V = [&](uint32_t value) -> const ast::Expression* {
return ScalarOrVector(width, value);
};
auto B = [&](const ast::Expression* value) -> const ast::Expression* {
if (width == 1) {
return b.Construct<bool>(value);
}
return b.Construct(b.ty.vec<bool>(width), value);
};
const ast::Expression* x = nullptr;
if (ty->is_unsigned_scalar_or_vector()) {
x = b.Expr("v");
} else {
// If ty is signed, then the value is inverted if the sign is negative
x = b.Call("select", //
b.Construct(U(), "v"), //
b.Construct(U(), b.Complement("v")), //
b.LessThan("v", ScalarOrVector(width, 0)));
}
b.Func(name, {b.Param("v", T(ty))}, T(ty),
{
// var x = v; (unsigned)
// var x = select(U(v), ~U(v), v < 0); (signed)
b.Decl(b.Var("x", nullptr, x)),
// let b16 = select(0, 16, bool(x & 0xffff0000));
b.Decl(b.Const("b16", nullptr,
b.Call("select", V(0), V(16),
B(b.And("x", V(0xffff0000)))))),
// x = x >> b16;
b.Assign("x", b.Shr("x", "b16")),
// let b8 = select(0, 8, bool(x & 0x0000ff00));
b.Decl(b.Const(
"b8", nullptr,
b.Call("select", V(0), V(8), B(b.And("x", V(0x0000ff00)))))),
// x = x >> b8;
b.Assign("x", b.Shr("x", "b8")),
// let b4 = select(0, 4, bool(x & 0x000000f0));
b.Decl(b.Const(
"b4", nullptr,
b.Call("select", V(0), V(4), B(b.And("x", V(0x000000f0)))))),
// x = x >> b4;
b.Assign("x", b.Shr("x", "b4")),
// let b2 = select(0, 2, bool(x & 0x0000000c));
b.Decl(b.Const(
"b2", nullptr,
b.Call("select", V(0), V(2), B(b.And("x", V(0x0000000c)))))),
// x = x >> b2;
b.Assign("x", b.Shr("x", "b2")),
// let b1 = select(0, 1, bool(x & 0x00000002));
b.Decl(b.Const(
"b1", nullptr,
b.Call("select", V(0), V(1), B(b.And("x", V(0x00000002)))))),
// let is_zero = select(0, 0xffffffff, x == 0);
b.Decl(b.Const("is_zero", nullptr,
b.Call("select", V(0), V(0xffffffff),
b.Equal("x", V(0))))),
// return R(b16 | b8 | b4 | b2 | b1 | zero);
b.Return(b.Construct(
T(ty),
b.Or(b.Or(b.Or(b.Or(b.Or("b16", "b8"), "b4"), "b2"), "b1"),
"is_zero"))),
});
return name;
}
/// Builds the polyfill function for the `firstTrailingBit` builtin
/// @param ty the parameter and return type for the function
/// @return the polyfill function name
@ -191,10 +266,7 @@ struct BuiltinPolyfill::State {
return b.ty.vec<ProgramBuilder::u32>(width);
};
auto V = [&](uint32_t value) -> const ast::Expression* {
if (width == 1) {
return b.Expr(value);
}
return b.Construct(b.ty.vec<ProgramBuilder::u32>(width), value);
return ScalarOrVector(width, value);
};
auto B = [&](const ast::Expression* value) -> const ast::Expression* {
if (width == 1) {
@ -257,6 +329,13 @@ struct BuiltinPolyfill::State {
}
return 1;
}
template <typename T>
const ast::Expression* ScalarOrVector(uint32_t width, T value) const {
if (width == 1) {
return b.Expr(value);
}
return b.Construct(b.ty.vec<T>(width), value);
}
};
BuiltinPolyfill::BuiltinPolyfill() = default;
@ -282,6 +361,11 @@ bool BuiltinPolyfill::ShouldRun(const Program* program,
return true;
}
break;
case sem::BuiltinType::kFirstLeadingBit:
if (builtins.first_leading_bit) {
return true;
}
break;
case sem::BuiltinType::kFirstTrailingBit:
if (builtins.first_trailing_bit) {
return true;
@ -330,6 +414,13 @@ void BuiltinPolyfill::Run(CloneContext& ctx,
});
}
break;
case sem::BuiltinType::kFirstLeadingBit:
if (builtins.first_leading_bit) {
polyfill = utils::GetOrCreate(polyfills, builtin, [&] {
return s.firstLeadingBit(builtin->ReturnType());
});
}
break;
case sem::BuiltinType::kFirstTrailingBit:
if (builtins.first_trailing_bit) {
polyfill = utils::GetOrCreate(polyfills, builtin, [&] {

2
src/tint/transform/builtin_polyfill.h
View File

@ -34,6 +34,8 @@ class BuiltinPolyfill : public Castable<BuiltinPolyfill, Transform> {
bool count_leading_zeros = false;
/// Should `countTrailingZeros()` be polyfilled?
bool count_trailing_zeros = false;
/// Should `firstLeadingBit()` be polyfilled?
bool first_leading_bit = false;
/// Should `firstTrailingBit()` be polyfilled?
bool first_trailing_bit = false;
};

154
src/tint/transform/builtin_polyfill_test.cc
View File

@ -348,6 +348,160 @@ fn f() {
EXPECT_EQ(expect, str(got));
}
////////////////////////////////////////////////////////////////////////////////
// firstLeadingBit
////////////////////////////////////////////////////////////////////////////////
DataMap polyfillFirstLeadingBit() {
BuiltinPolyfill::Builtins builtins;
builtins.first_leading_bit = true;
DataMap data;
data.Add<BuiltinPolyfill::Config>(builtins);
return data;
}
TEST_F(BuiltinPolyfillTest, ShouldRunFirstLeadingBit) {
auto* src = R"(
fn f() {
firstLeadingBit(0xf);
}
)";
EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src));
EXPECT_TRUE(ShouldRun<BuiltinPolyfill>(src, polyfillFirstLeadingBit()));
}
TEST_F(BuiltinPolyfillTest, FirstLeadingBit_i32) {
auto* src = R"(
fn f() {
let r : i32 = firstLeadingBit(15);
}
)";
auto* expect = R"(
fn tint_first_leading_bit(v : i32) -> i32 {
var x = select(u32(v), u32(~(v)), (v < 0));
let b16 = select(0u, 16u, bool((x & 4294901760u)));
x = (x >> b16);
let b8 = select(0u, 8u, bool((x & 65280u)));
x = (x >> b8);
let b4 = select(0u, 4u, bool((x & 240u)));
x = (x >> b4);
let b2 = select(0u, 2u, bool((x & 12u)));
x = (x >> b2);
let b1 = select(0u, 1u, bool((x & 2u)));
let is_zero = select(0u, 4294967295u, (x == 0u));
return i32((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
fn f() {
let r : i32 = tint_first_leading_bit(15);
}
)";
auto got = Run<BuiltinPolyfill>(src, polyfillFirstLeadingBit());
EXPECT_EQ(expect, str(got));
}
TEST_F(BuiltinPolyfillTest, FirstLeadingBit_u32) {
auto* src = R"(
fn f() {
let r : u32 = firstLeadingBit(15u);
}
)";
auto* expect = R"(
fn tint_first_leading_bit(v : u32) -> u32 {
var x = v;
let b16 = select(0u, 16u, bool((x & 4294901760u)));
x = (x >> b16);
let b8 = select(0u, 8u, bool((x & 65280u)));
x = (x >> b8);
let b4 = select(0u, 4u, bool((x & 240u)));
x = (x >> b4);
let b2 = select(0u, 2u, bool((x & 12u)));
x = (x >> b2);
let b1 = select(0u, 1u, bool((x & 2u)));
let is_zero = select(0u, 4294967295u, (x == 0u));
return u32((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
fn f() {
let r : u32 = tint_first_leading_bit(15u);
}
)";
auto got = Run<BuiltinPolyfill>(src, polyfillFirstLeadingBit());
EXPECT_EQ(expect, str(got));
}
TEST_F(BuiltinPolyfillTest, FirstLeadingBit_vec3_i32) {
auto* src = R"(
fn f() {
let r : vec3<i32> = firstLeadingBit(vec3<i32>(15));
}
)";
auto* expect = R"(
fn tint_first_leading_bit(v : vec3<i32>) -> vec3<i32> {
var x = select(vec3<u32>(v), vec3<u32>(~(v)), (v < vec3<i32>(0)));
let b16 = select(vec3<u32>(0u), vec3<u32>(16u), vec3<bool>((x & vec3<u32>(4294901760u))));
x = (x >> b16);
let b8 = select(vec3<u32>(0u), vec3<u32>(8u), vec3<bool>((x & vec3<u32>(65280u))));
x = (x >> b8);
let b4 = select(vec3<u32>(0u), vec3<u32>(4u), vec3<bool>((x & vec3<u32>(240u))));
x = (x >> b4);
let b2 = select(vec3<u32>(0u), vec3<u32>(2u), vec3<bool>((x & vec3<u32>(12u))));
x = (x >> b2);
let b1 = select(vec3<u32>(0u), vec3<u32>(1u), vec3<bool>((x & vec3<u32>(2u))));
let is_zero = select(vec3<u32>(0u), vec3<u32>(4294967295u), (x == vec3<u32>(0u)));
return vec3<i32>((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
fn f() {
let r : vec3<i32> = tint_first_leading_bit(vec3<i32>(15));
}
)";
auto got = Run<BuiltinPolyfill>(src, polyfillFirstLeadingBit());
EXPECT_EQ(expect, str(got));
}
TEST_F(BuiltinPolyfillTest, FirstLeadingBit_vec3_u32) {
auto* src = R"(
fn f() {
let r : vec3<u32> = firstLeadingBit(vec3<u32>(15u));
}
)";
auto* expect = R"(
fn tint_first_leading_bit(v : vec3<u32>) -> vec3<u32> {
var x = v;
let b16 = select(vec3<u32>(0u), vec3<u32>(16u), vec3<bool>((x & vec3<u32>(4294901760u))));
x = (x >> b16);
let b8 = select(vec3<u32>(0u), vec3<u32>(8u), vec3<bool>((x & vec3<u32>(65280u))));
x = (x >> b8);
let b4 = select(vec3<u32>(0u), vec3<u32>(4u), vec3<bool>((x & vec3<u32>(240u))));
x = (x >> b4);
let b2 = select(vec3<u32>(0u), vec3<u32>(2u), vec3<bool>((x & vec3<u32>(12u))));
x = (x >> b2);
let b1 = select(vec3<u32>(0u), vec3<u32>(1u), vec3<bool>((x & vec3<u32>(2u))));
let is_zero = select(vec3<u32>(0u), vec3<u32>(4294967295u), (x == vec3<u32>(0u)));
return vec3<u32>((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
fn f() {
let r : vec3<u32> = tint_first_leading_bit(vec3<u32>(15u));
}
)";
auto got = Run<BuiltinPolyfill>(src, polyfillFirstLeadingBit());
EXPECT_EQ(expect, str(got));
}
////////////////////////////////////////////////////////////////////////////////
// firstTrailingBit
////////////////////////////////////////////////////////////////////////////////

1
src/tint/transform/glsl.cc
View File

@ -56,6 +56,7 @@ Output Glsl::Run(const Program* in, const DataMap& inputs) const {
BuiltinPolyfill::Builtins polyfills;
polyfills.count_leading_zeros = true;
polyfills.count_trailing_zeros = true;
polyfills.first_leading_bit = true;
polyfills.first_trailing_bit = true;
data.Add<BuiltinPolyfill::Config>(polyfills);
manager.Add<BuiltinPolyfill>();

1
src/tint/writer/hlsl/generator_impl.cc
View File

@ -147,6 +147,7 @@ SanitizedResult Sanitize(
// and `firstbithigh`.
polyfills.count_leading_zeros = true;
polyfills.count_trailing_zeros = true;
polyfills.first_leading_bit = true;
polyfills.first_trailing_bit = true;
data.Add<transform::BuiltinPolyfill::Config>(polyfills);
manager.Add<transform::BuiltinPolyfill>();

1
src/tint/writer/msl/generator_impl.cc
View File

@ -130,6 +130,7 @@ SanitizedResult Sanitize(
{ // Builtin polyfills
transform::BuiltinPolyfill::Builtins polyfills;
polyfills.first_leading_bit = true;
polyfills.first_trailing_bit = true;
data.Add<transform::BuiltinPolyfill::Config>(polyfills);
manager.Add<transform::BuiltinPolyfill>();

1
src/tint/writer/spirv/builder.cc
View File

@ -263,6 +263,7 @@ SanitizedResult Sanitize(const Program* in,
transform::BuiltinPolyfill::Builtins polyfills;
polyfills.count_leading_zeros = true;
polyfills.count_trailing_zeros = true;
polyfills.first_leading_bit = true;
polyfills.first_trailing_bit = true;
data.Add<transform::BuiltinPolyfill::Config>(polyfills);
manager.Add<transform::BuiltinPolyfill>();

45
test/tint/builtins/gen/firstLeadingBit/000ff3.wgsl Normal file
View File

@ -0,0 +1,45 @@
// Copyright 2021 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/builtin-gen
// using the template:
// test/tint/builtins/builtins.wgsl.tmpl
// and the builtin defintion file:
// src/tint/builtins.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn firstLeadingBit(vec<4, u32>) -> vec<4, u32>
fn firstLeadingBit_000ff3() {
var res: vec4<u32> = firstLeadingBit(vec4<u32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_000ff3();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_000ff3();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_000ff3();
}

93
test/tint/builtins/gen/firstLeadingBit/000ff3.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
uvec4 tint_first_leading_bit(uvec4 v) {
uvec4 x = v;
uvec4 b16 = mix(uvec4(0u), uvec4(16u), bvec4((x & uvec4(4294901760u))));
x = (x >> b16);
uvec4 b8 = mix(uvec4(0u), uvec4(8u), bvec4((x & uvec4(65280u))));
x = (x >> b8);
uvec4 b4 = mix(uvec4(0u), uvec4(4u), bvec4((x & uvec4(240u))));
x = (x >> b4);
uvec4 b2 = mix(uvec4(0u), uvec4(2u), bvec4((x & uvec4(12u))));
x = (x >> b2);
uvec4 b1 = mix(uvec4(0u), uvec4(1u), bvec4((x & uvec4(2u))));
uvec4 is_zero = mix(uvec4(0u), uvec4(4294967295u), equal(x, uvec4(0u)));
return uvec4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_000ff3() {
uvec4 res = tint_first_leading_bit(uvec4(0u, 0u, 0u, 0u));
}
vec4 vertex_main() {
firstLeadingBit_000ff3();
return vec4(0.0f, 0.0f, 0.0f, 0.0f);
}
void main() {
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;
uvec4 tint_first_leading_bit(uvec4 v) {
uvec4 x = v;
uvec4 b16 = mix(uvec4(0u), uvec4(16u), bvec4((x & uvec4(4294901760u))));
x = (x >> b16);
uvec4 b8 = mix(uvec4(0u), uvec4(8u), bvec4((x & uvec4(65280u))));
x = (x >> b8);
uvec4 b4 = mix(uvec4(0u), uvec4(4u), bvec4((x & uvec4(240u))));
x = (x >> b4);
uvec4 b2 = mix(uvec4(0u), uvec4(2u), bvec4((x & uvec4(12u))));
x = (x >> b2);
uvec4 b1 = mix(uvec4(0u), uvec4(1u), bvec4((x & uvec4(2u))));
uvec4 is_zero = mix(uvec4(0u), uvec4(4294967295u), equal(x, uvec4(0u)));
return uvec4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_000ff3() {
uvec4 res = tint_first_leading_bit(uvec4(0u, 0u, 0u, 0u));
}
void fragment_main() {
firstLeadingBit_000ff3();
}
void main() {
fragment_main();
return;
}
#version 310 es
uvec4 tint_first_leading_bit(uvec4 v) {
uvec4 x = v;
uvec4 b16 = mix(uvec4(0u), uvec4(16u), bvec4((x & uvec4(4294901760u))));
x = (x >> b16);
uvec4 b8 = mix(uvec4(0u), uvec4(8u), bvec4((x & uvec4(65280u))));
x = (x >> b8);
uvec4 b4 = mix(uvec4(0u), uvec4(4u), bvec4((x & uvec4(240u))));
x = (x >> b4);
uvec4 b2 = mix(uvec4(0u), uvec4(2u), bvec4((x & uvec4(12u))));
x = (x >> b2);
uvec4 b1 = mix(uvec4(0u), uvec4(1u), bvec4((x & uvec4(2u))));
uvec4 is_zero = mix(uvec4(0u), uvec4(4294967295u), equal(x, uvec4(0u)));
return uvec4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_000ff3() {
uvec4 res = tint_first_leading_bit(uvec4(0u, 0u, 0u, 0u));
}
void compute_main() {
firstLeadingBit_000ff3();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/firstLeadingBit/000ff3.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
uint4 tint_first_leading_bit(uint4 v) {
uint4 x = v;
const uint4 b16 = (bool4((x & uint4((4294901760u).xxxx))) ? uint4((16u).xxxx) : uint4((0u).xxxx));
x = (x >> b16);
const uint4 b8 = (bool4((x & uint4((65280u).xxxx))) ? uint4((8u).xxxx) : uint4((0u).xxxx));
x = (x >> b8);
const uint4 b4 = (bool4((x & uint4((240u).xxxx))) ? uint4((4u).xxxx) : uint4((0u).xxxx));
x = (x >> b4);
const uint4 b2 = (bool4((x & uint4((12u).xxxx))) ? uint4((2u).xxxx) : uint4((0u).xxxx));
x = (x >> b2);
const uint4 b1 = (bool4((x & uint4((2u).xxxx))) ? uint4((1u).xxxx) : uint4((0u).xxxx));
const uint4 is_zero = ((x == uint4((0u).xxxx)) ? uint4((4294967295u).xxxx) : uint4((0u).xxxx));
return uint4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_000ff3() {
uint4 res = tint_first_leading_bit(uint4(0u, 0u, 0u, 0u));
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
firstLeadingBit_000ff3();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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() {
firstLeadingBit_000ff3();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
firstLeadingBit_000ff3();
return;
}

48
test/tint/builtins/gen/firstLeadingBit/000ff3.wgsl.expected.msl Normal file
View File

@ -0,0 +1,48 @@
#include <metal_stdlib>
using namespace metal;
uint4 tint_first_leading_bit(uint4 v) {
uint4 x = v;
uint4 const b16 = select(uint4(0u), uint4(16u), bool4((x & uint4(4294901760u))));
x = (x >> b16);
uint4 const b8 = select(uint4(0u), uint4(8u), bool4((x & uint4(65280u))));
x = (x >> b8);
uint4 const b4 = select(uint4(0u), uint4(4u), bool4((x & uint4(240u))));
x = (x >> b4);
uint4 const b2 = select(uint4(0u), uint4(2u), bool4((x & uint4(12u))));
x = (x >> b2);
uint4 const b1 = select(uint4(0u), uint4(1u), bool4((x & uint4(2u))));
uint4 const is_zero = select(uint4(0u), uint4(4294967295u), (x == uint4(0u)));
return uint4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_000ff3() {
uint4 res = tint_first_leading_bit(uint4());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
firstLeadingBit_000ff3();
return float4();
}
vertex tint_symbol vertex_main() {
float4 const inner_result = vertex_main_inner();
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main() {
firstLeadingBit_000ff3();
return;
}
kernel void compute_main() {
firstLeadingBit_000ff3();
return;
}

142
test/tint/builtins/gen/firstLeadingBit/000ff3.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,142 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 99
; 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 %tint_first_leading_bit "tint_first_leading_bit"
OpName %v "v"
OpName %x "x"
OpName %firstLeadingBit_000ff3 "firstLeadingBit_000ff3"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%uint = OpTypeInt 32 0
%v4uint = OpTypeVector %uint 4
%9 = OpTypeFunction %v4uint %v4uint
%_ptr_Function_v4uint = OpTypePointer Function %v4uint
%17 = OpConstantNull %v4uint
%bool = OpTypeBool
%v4bool = OpTypeVector %bool 4
%uint_4294901760 = OpConstant %uint 4294901760
%24 = OpConstantComposite %v4uint %uint_4294901760 %uint_4294901760 %uint_4294901760 %uint_4294901760
%uint_16 = OpConstant %uint 16
%27 = OpConstantComposite %v4uint %uint_16 %uint_16 %uint_16 %uint_16
%uint_0 = OpConstant %uint 0
%29 = OpConstantComposite %v4uint %uint_0 %uint_0 %uint_0 %uint_0
%uint_65280 = OpConstant %uint 65280
%36 = OpConstantComposite %v4uint %uint_65280 %uint_65280 %uint_65280 %uint_65280
%uint_8 = OpConstant %uint 8
%39 = OpConstantComposite %v4uint %uint_8 %uint_8 %uint_8 %uint_8
%uint_240 = OpConstant %uint 240
%46 = OpConstantComposite %v4uint %uint_240 %uint_240 %uint_240 %uint_240
%uint_4 = OpConstant %uint 4
%49 = OpConstantComposite %v4uint %uint_4 %uint_4 %uint_4 %uint_4
%uint_12 = OpConstant %uint 12
%56 = OpConstantComposite %v4uint %uint_12 %uint_12 %uint_12 %uint_12
%uint_2 = OpConstant %uint 2
%59 = OpConstantComposite %v4uint %uint_2 %uint_2 %uint_2 %uint_2
%uint_1 = OpConstant %uint 1
%67 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1
%uint_4294967295 = OpConstant %uint 4294967295
%72 = OpConstantComposite %v4uint %uint_4294967295 %uint_4294967295 %uint_4294967295 %uint_4294967295
%void = OpTypeVoid
%79 = OpTypeFunction %void
%85 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_first_leading_bit = OpFunction %v4uint None %9
%v = OpFunctionParameter %v4uint
%14 = OpLabel
%x = OpVariable %_ptr_Function_v4uint Function %17
OpStore %x %v
%22 = OpLoad %v4uint %x
%25 = OpBitwiseAnd %v4uint %22 %24
%19 = OpINotEqual %v4bool %25 %17
%18 = OpSelect %v4uint %19 %27 %29
%30 = OpLoad %v4uint %x
%31 = OpShiftRightLogical %v4uint %30 %18
OpStore %x %31
%34 = OpLoad %v4uint %x
%37 = OpBitwiseAnd %v4uint %34 %36
%33 = OpINotEqual %v4bool %37 %17
%32 = OpSelect %v4uint %33 %39 %29
%40 = OpLoad %v4uint %x
%41 = OpShiftRightLogical %v4uint %40 %32
OpStore %x %41
%44 = OpLoad %v4uint %x
%47 = OpBitwiseAnd %v4uint %44 %46
%43 = OpINotEqual %v4bool %47 %17
%42 = OpSelect %v4uint %43 %49 %29
%50 = OpLoad %v4uint %x
%51 = OpShiftRightLogical %v4uint %50 %42
OpStore %x %51
%54 = OpLoad %v4uint %x
%57 = OpBitwiseAnd %v4uint %54 %56
%53 = OpINotEqual %v4bool %57 %17
%52 = OpSelect %v4uint %53 %59 %29
%60 = OpLoad %v4uint %x
%61 = OpShiftRightLogical %v4uint %60 %52
OpStore %x %61
%64 = OpLoad %v4uint %x
%65 = OpBitwiseAnd %v4uint %64 %59
%63 = OpINotEqual %v4bool %65 %17
%62 = OpSelect %v4uint %63 %67 %29
%69 = OpLoad %v4uint %x
%70 = OpIEqual %v4bool %69 %29
%68 = OpSelect %v4uint %70 %72 %29
%74 = OpBitwiseOr %v4uint %18 %32
%75 = OpBitwiseOr %v4uint %74 %42
%76 = OpBitwiseOr %v4uint %75 %52
%77 = OpBitwiseOr %v4uint %76 %62
%78 = OpBitwiseOr %v4uint %77 %68
OpReturnValue %78
OpFunctionEnd
%firstLeadingBit_000ff3 = OpFunction %void None %79
%82 = OpLabel
%res = OpVariable %_ptr_Function_v4uint Function %17
%83 = OpFunctionCall %v4uint %tint_first_leading_bit %17
OpStore %res %83
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %85
%87 = OpLabel
%88 = OpFunctionCall %void %firstLeadingBit_000ff3
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %79
%90 = OpLabel
%91 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %91
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %79
%94 = OpLabel
%95 = OpFunctionCall %void %firstLeadingBit_000ff3
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %79
%97 = OpLabel
%98 = OpFunctionCall %void %firstLeadingBit_000ff3
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/firstLeadingBit/000ff3.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,19 @@
fn firstLeadingBit_000ff3() {
var res : vec4<u32> = firstLeadingBit(vec4<u32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_000ff3();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_000ff3();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_000ff3();
}

45
test/tint/builtins/gen/firstLeadingBit/35053e.wgsl Normal file
View File

@ -0,0 +1,45 @@
// Copyright 2021 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/builtin-gen
// using the template:
// test/tint/builtins/builtins.wgsl.tmpl
// and the builtin defintion file:
// src/tint/builtins.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn firstLeadingBit(vec<3, i32>) -> vec<3, i32>
fn firstLeadingBit_35053e() {
var res: vec3<i32> = firstLeadingBit(vec3<i32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_35053e();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_35053e();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_35053e();
}

93
test/tint/builtins/gen/firstLeadingBit/35053e.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
ivec3 tint_first_leading_bit(ivec3 v) {
uvec3 x = mix(uvec3(v), uvec3(~(v)), lessThan(v, ivec3(0)));
uvec3 b16 = mix(uvec3(0u), uvec3(16u), bvec3((x & uvec3(4294901760u))));
x = (x >> b16);
uvec3 b8 = mix(uvec3(0u), uvec3(8u), bvec3((x & uvec3(65280u))));
x = (x >> b8);
uvec3 b4 = mix(uvec3(0u), uvec3(4u), bvec3((x & uvec3(240u))));
x = (x >> b4);
uvec3 b2 = mix(uvec3(0u), uvec3(2u), bvec3((x & uvec3(12u))));
x = (x >> b2);
uvec3 b1 = mix(uvec3(0u), uvec3(1u), bvec3((x & uvec3(2u))));
uvec3 is_zero = mix(uvec3(0u), uvec3(4294967295u), equal(x, uvec3(0u)));
return ivec3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_35053e() {
ivec3 res = tint_first_leading_bit(ivec3(0, 0, 0));
}
vec4 vertex_main() {
firstLeadingBit_35053e();
return vec4(0.0f, 0.0f, 0.0f, 0.0f);
}
void main() {
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;
ivec3 tint_first_leading_bit(ivec3 v) {
uvec3 x = mix(uvec3(v), uvec3(~(v)), lessThan(v, ivec3(0)));
uvec3 b16 = mix(uvec3(0u), uvec3(16u), bvec3((x & uvec3(4294901760u))));
x = (x >> b16);
uvec3 b8 = mix(uvec3(0u), uvec3(8u), bvec3((x & uvec3(65280u))));
x = (x >> b8);
uvec3 b4 = mix(uvec3(0u), uvec3(4u), bvec3((x & uvec3(240u))));
x = (x >> b4);
uvec3 b2 = mix(uvec3(0u), uvec3(2u), bvec3((x & uvec3(12u))));
x = (x >> b2);
uvec3 b1 = mix(uvec3(0u), uvec3(1u), bvec3((x & uvec3(2u))));
uvec3 is_zero = mix(uvec3(0u), uvec3(4294967295u), equal(x, uvec3(0u)));
return ivec3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_35053e() {
ivec3 res = tint_first_leading_bit(ivec3(0, 0, 0));
}
void fragment_main() {
firstLeadingBit_35053e();
}
void main() {
fragment_main();
return;
}
#version 310 es
ivec3 tint_first_leading_bit(ivec3 v) {
uvec3 x = mix(uvec3(v), uvec3(~(v)), lessThan(v, ivec3(0)));
uvec3 b16 = mix(uvec3(0u), uvec3(16u), bvec3((x & uvec3(4294901760u))));
x = (x >> b16);
uvec3 b8 = mix(uvec3(0u), uvec3(8u), bvec3((x & uvec3(65280u))));
x = (x >> b8);
uvec3 b4 = mix(uvec3(0u), uvec3(4u), bvec3((x & uvec3(240u))));
x = (x >> b4);
uvec3 b2 = mix(uvec3(0u), uvec3(2u), bvec3((x & uvec3(12u))));
x = (x >> b2);
uvec3 b1 = mix(uvec3(0u), uvec3(1u), bvec3((x & uvec3(2u))));
uvec3 is_zero = mix(uvec3(0u), uvec3(4294967295u), equal(x, uvec3(0u)));
return ivec3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_35053e() {
ivec3 res = tint_first_leading_bit(ivec3(0, 0, 0));
}
void compute_main() {
firstLeadingBit_35053e();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/firstLeadingBit/35053e.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
int3 tint_first_leading_bit(int3 v) {
uint3 x = ((v < int3((0).xxx)) ? uint3(~(v)) : uint3(v));
const uint3 b16 = (bool3((x & uint3((4294901760u).xxx))) ? uint3((16u).xxx) : uint3((0u).xxx));
x = (x >> b16);
const uint3 b8 = (bool3((x & uint3((65280u).xxx))) ? uint3((8u).xxx) : uint3((0u).xxx));
x = (x >> b8);
const uint3 b4 = (bool3((x & uint3((240u).xxx))) ? uint3((4u).xxx) : uint3((0u).xxx));
x = (x >> b4);
const uint3 b2 = (bool3((x & uint3((12u).xxx))) ? uint3((2u).xxx) : uint3((0u).xxx));
x = (x >> b2);
const uint3 b1 = (bool3((x & uint3((2u).xxx))) ? uint3((1u).xxx) : uint3((0u).xxx));
const uint3 is_zero = ((x == uint3((0u).xxx)) ? uint3((4294967295u).xxx) : uint3((0u).xxx));
return int3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_35053e() {
int3 res = tint_first_leading_bit(int3(0, 0, 0));
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
firstLeadingBit_35053e();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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() {
firstLeadingBit_35053e();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
firstLeadingBit_35053e();
return;
}

48
test/tint/builtins/gen/firstLeadingBit/35053e.wgsl.expected.msl Normal file
View File

@ -0,0 +1,48 @@
#include <metal_stdlib>
using namespace metal;
int3 tint_first_leading_bit(int3 v) {
uint3 x = select(uint3(v), uint3(~(v)), (v < int3(0)));
uint3 const b16 = select(uint3(0u), uint3(16u), bool3((x & uint3(4294901760u))));
x = (x >> b16);
uint3 const b8 = select(uint3(0u), uint3(8u), bool3((x & uint3(65280u))));
x = (x >> b8);
uint3 const b4 = select(uint3(0u), uint3(4u), bool3((x & uint3(240u))));
x = (x >> b4);
uint3 const b2 = select(uint3(0u), uint3(2u), bool3((x & uint3(12u))));
x = (x >> b2);
uint3 const b1 = select(uint3(0u), uint3(1u), bool3((x & uint3(2u))));
uint3 const is_zero = select(uint3(0u), uint3(4294967295u), (x == uint3(0u)));
return int3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_35053e() {
int3 res = tint_first_leading_bit(int3());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
firstLeadingBit_35053e();
return float4();
}
vertex tint_symbol vertex_main() {
float4 const inner_result = vertex_main_inner();
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main() {
firstLeadingBit_35053e();
return;
}
kernel void compute_main() {
firstLeadingBit_35053e();
return;
}

154
test/tint/builtins/gen/firstLeadingBit/35053e.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,154 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 110
; 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 %tint_first_leading_bit "tint_first_leading_bit"
OpName %v "v"
OpName %x "x"
OpName %firstLeadingBit_35053e "firstLeadingBit_35053e"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%int = OpTypeInt 32 1
%v3int = OpTypeVector %int 3
%9 = OpTypeFunction %v3int %v3int
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%int_0 = OpConstant %int 0
%19 = OpConstantComposite %v3int %int_0 %int_0 %int_0
%bool = OpTypeBool
%v3bool = OpTypeVector %bool 3
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
%28 = OpConstantNull %v3uint
%uint_4294901760 = OpConstant %uint 4294901760
%33 = OpConstantComposite %v3uint %uint_4294901760 %uint_4294901760 %uint_4294901760
%uint_16 = OpConstant %uint 16
%36 = OpConstantComposite %v3uint %uint_16 %uint_16 %uint_16
%uint_0 = OpConstant %uint 0
%38 = OpConstantComposite %v3uint %uint_0 %uint_0 %uint_0
%uint_65280 = OpConstant %uint 65280
%45 = OpConstantComposite %v3uint %uint_65280 %uint_65280 %uint_65280
%uint_8 = OpConstant %uint 8
%48 = OpConstantComposite %v3uint %uint_8 %uint_8 %uint_8
%uint_240 = OpConstant %uint 240
%55 = OpConstantComposite %v3uint %uint_240 %uint_240 %uint_240
%uint_4 = OpConstant %uint 4
%58 = OpConstantComposite %v3uint %uint_4 %uint_4 %uint_4
%uint_12 = OpConstant %uint 12
%65 = OpConstantComposite %v3uint %uint_12 %uint_12 %uint_12
%uint_2 = OpConstant %uint 2
%68 = OpConstantComposite %v3uint %uint_2 %uint_2 %uint_2
%uint_1 = OpConstant %uint 1
%76 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
%uint_4294967295 = OpConstant %uint 4294967295
%81 = OpConstantComposite %v3uint %uint_4294967295 %uint_4294967295 %uint_4294967295
%void = OpTypeVoid
%88 = OpTypeFunction %void
%93 = OpConstantNull %v3int
%_ptr_Function_v3int = OpTypePointer Function %v3int
%96 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_first_leading_bit = OpFunction %v3int None %9
%v = OpFunctionParameter %v3int
%14 = OpLabel
%x = OpVariable %_ptr_Function_v3uint Function %28
%20 = OpSLessThan %v3bool %v %19
%24 = OpNot %v3int %v
%23 = OpBitcast %v3uint %24
%25 = OpBitcast %v3uint %v
%15 = OpSelect %v3uint %20 %23 %25
OpStore %x %15
%31 = OpLoad %v3uint %x
%34 = OpBitwiseAnd %v3uint %31 %33
%30 = OpINotEqual %v3bool %34 %28
%29 = OpSelect %v3uint %30 %36 %38
%39 = OpLoad %v3uint %x
%40 = OpShiftRightLogical %v3uint %39 %29
OpStore %x %40
%43 = OpLoad %v3uint %x
%46 = OpBitwiseAnd %v3uint %43 %45
%42 = OpINotEqual %v3bool %46 %28
%41 = OpSelect %v3uint %42 %48 %38
%49 = OpLoad %v3uint %x
%50 = OpShiftRightLogical %v3uint %49 %41
OpStore %x %50
%53 = OpLoad %v3uint %x
%56 = OpBitwiseAnd %v3uint %53 %55
%52 = OpINotEqual %v3bool %56 %28
%51 = OpSelect %v3uint %52 %58 %38
%59 = OpLoad %v3uint %x
%60 = OpShiftRightLogical %v3uint %59 %51
OpStore %x %60
%63 = OpLoad %v3uint %x
%66 = OpBitwiseAnd %v3uint %63 %65
%62 = OpINotEqual %v3bool %66 %28
%61 = OpSelect %v3uint %62 %68 %38
%69 = OpLoad %v3uint %x
%70 = OpShiftRightLogical %v3uint %69 %61
OpStore %x %70
%73 = OpLoad %v3uint %x
%74 = OpBitwiseAnd %v3uint %73 %68
%72 = OpINotEqual %v3bool %74 %28
%71 = OpSelect %v3uint %72 %76 %38
%78 = OpLoad %v3uint %x
%79 = OpIEqual %v3bool %78 %38
%77 = OpSelect %v3uint %79 %81 %38
%83 = OpBitwiseOr %v3uint %29 %41
%84 = OpBitwiseOr %v3uint %83 %51
%85 = OpBitwiseOr %v3uint %84 %61
%86 = OpBitwiseOr %v3uint %85 %71
%87 = OpBitwiseOr %v3uint %86 %77
%82 = OpBitcast %v3int %87
OpReturnValue %82
OpFunctionEnd
%firstLeadingBit_35053e = OpFunction %void None %88
%91 = OpLabel
%res = OpVariable %_ptr_Function_v3int Function %93
%92 = OpFunctionCall %v3int %tint_first_leading_bit %93
OpStore %res %92
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %96
%98 = OpLabel
%99 = OpFunctionCall %void %firstLeadingBit_35053e
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %88
%101 = OpLabel
%102 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %102
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %88
%105 = OpLabel
%106 = OpFunctionCall %void %firstLeadingBit_35053e
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %88
%108 = OpLabel
%109 = OpFunctionCall %void %firstLeadingBit_35053e
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/firstLeadingBit/35053e.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,19 @@
fn firstLeadingBit_35053e() {
var res : vec3<i32> = firstLeadingBit(vec3<i32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_35053e();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_35053e();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_35053e();
}

45
test/tint/builtins/gen/firstLeadingBit/3fd7d0.wgsl Normal file
View File

@ -0,0 +1,45 @@
// Copyright 2021 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/builtin-gen
// using the template:
// test/tint/builtins/builtins.wgsl.tmpl
// and the builtin defintion file:
// src/tint/builtins.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn firstLeadingBit(vec<3, u32>) -> vec<3, u32>
fn firstLeadingBit_3fd7d0() {
var res: vec3<u32> = firstLeadingBit(vec3<u32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_3fd7d0();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_3fd7d0();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_3fd7d0();
}

93
test/tint/builtins/gen/firstLeadingBit/3fd7d0.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
uvec3 tint_first_leading_bit(uvec3 v) {
uvec3 x = v;
uvec3 b16 = mix(uvec3(0u), uvec3(16u), bvec3((x & uvec3(4294901760u))));
x = (x >> b16);
uvec3 b8 = mix(uvec3(0u), uvec3(8u), bvec3((x & uvec3(65280u))));
x = (x >> b8);
uvec3 b4 = mix(uvec3(0u), uvec3(4u), bvec3((x & uvec3(240u))));
x = (x >> b4);
uvec3 b2 = mix(uvec3(0u), uvec3(2u), bvec3((x & uvec3(12u))));
x = (x >> b2);
uvec3 b1 = mix(uvec3(0u), uvec3(1u), bvec3((x & uvec3(2u))));
uvec3 is_zero = mix(uvec3(0u), uvec3(4294967295u), equal(x, uvec3(0u)));
return uvec3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_3fd7d0() {
uvec3 res = tint_first_leading_bit(uvec3(0u, 0u, 0u));
}
vec4 vertex_main() {
firstLeadingBit_3fd7d0();
return vec4(0.0f, 0.0f, 0.0f, 0.0f);
}
void main() {
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;
uvec3 tint_first_leading_bit(uvec3 v) {
uvec3 x = v;
uvec3 b16 = mix(uvec3(0u), uvec3(16u), bvec3((x & uvec3(4294901760u))));
x = (x >> b16);
uvec3 b8 = mix(uvec3(0u), uvec3(8u), bvec3((x & uvec3(65280u))));
x = (x >> b8);
uvec3 b4 = mix(uvec3(0u), uvec3(4u), bvec3((x & uvec3(240u))));
x = (x >> b4);
uvec3 b2 = mix(uvec3(0u), uvec3(2u), bvec3((x & uvec3(12u))));
x = (x >> b2);
uvec3 b1 = mix(uvec3(0u), uvec3(1u), bvec3((x & uvec3(2u))));
uvec3 is_zero = mix(uvec3(0u), uvec3(4294967295u), equal(x, uvec3(0u)));
return uvec3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_3fd7d0() {
uvec3 res = tint_first_leading_bit(uvec3(0u, 0u, 0u));
}
void fragment_main() {
firstLeadingBit_3fd7d0();
}
void main() {
fragment_main();
return;
}
#version 310 es
uvec3 tint_first_leading_bit(uvec3 v) {
uvec3 x = v;
uvec3 b16 = mix(uvec3(0u), uvec3(16u), bvec3((x & uvec3(4294901760u))));
x = (x >> b16);
uvec3 b8 = mix(uvec3(0u), uvec3(8u), bvec3((x & uvec3(65280u))));
x = (x >> b8);
uvec3 b4 = mix(uvec3(0u), uvec3(4u), bvec3((x & uvec3(240u))));
x = (x >> b4);
uvec3 b2 = mix(uvec3(0u), uvec3(2u), bvec3((x & uvec3(12u))));
x = (x >> b2);
uvec3 b1 = mix(uvec3(0u), uvec3(1u), bvec3((x & uvec3(2u))));
uvec3 is_zero = mix(uvec3(0u), uvec3(4294967295u), equal(x, uvec3(0u)));
return uvec3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_3fd7d0() {
uvec3 res = tint_first_leading_bit(uvec3(0u, 0u, 0u));
}
void compute_main() {
firstLeadingBit_3fd7d0();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/firstLeadingBit/3fd7d0.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
uint3 tint_first_leading_bit(uint3 v) {
uint3 x = v;
const uint3 b16 = (bool3((x & uint3((4294901760u).xxx))) ? uint3((16u).xxx) : uint3((0u).xxx));
x = (x >> b16);
const uint3 b8 = (bool3((x & uint3((65280u).xxx))) ? uint3((8u).xxx) : uint3((0u).xxx));
x = (x >> b8);
const uint3 b4 = (bool3((x & uint3((240u).xxx))) ? uint3((4u).xxx) : uint3((0u).xxx));
x = (x >> b4);
const uint3 b2 = (bool3((x & uint3((12u).xxx))) ? uint3((2u).xxx) : uint3((0u).xxx));
x = (x >> b2);
const uint3 b1 = (bool3((x & uint3((2u).xxx))) ? uint3((1u).xxx) : uint3((0u).xxx));
const uint3 is_zero = ((x == uint3((0u).xxx)) ? uint3((4294967295u).xxx) : uint3((0u).xxx));
return uint3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_3fd7d0() {
uint3 res = tint_first_leading_bit(uint3(0u, 0u, 0u));
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
firstLeadingBit_3fd7d0();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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() {
firstLeadingBit_3fd7d0();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
firstLeadingBit_3fd7d0();
return;
}

48
test/tint/builtins/gen/firstLeadingBit/3fd7d0.wgsl.expected.msl Normal file
View File

@ -0,0 +1,48 @@
#include <metal_stdlib>
using namespace metal;
uint3 tint_first_leading_bit(uint3 v) {
uint3 x = v;
uint3 const b16 = select(uint3(0u), uint3(16u), bool3((x & uint3(4294901760u))));
x = (x >> b16);
uint3 const b8 = select(uint3(0u), uint3(8u), bool3((x & uint3(65280u))));
x = (x >> b8);
uint3 const b4 = select(uint3(0u), uint3(4u), bool3((x & uint3(240u))));
x = (x >> b4);
uint3 const b2 = select(uint3(0u), uint3(2u), bool3((x & uint3(12u))));
x = (x >> b2);
uint3 const b1 = select(uint3(0u), uint3(1u), bool3((x & uint3(2u))));
uint3 const is_zero = select(uint3(0u), uint3(4294967295u), (x == uint3(0u)));
return uint3((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_3fd7d0() {
uint3 res = tint_first_leading_bit(uint3());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
firstLeadingBit_3fd7d0();
return float4();
}
vertex tint_symbol vertex_main() {
float4 const inner_result = vertex_main_inner();
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main() {
firstLeadingBit_3fd7d0();
return;
}
kernel void compute_main() {
firstLeadingBit_3fd7d0();
return;
}

142
test/tint/builtins/gen/firstLeadingBit/3fd7d0.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,142 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 99
; 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 %tint_first_leading_bit "tint_first_leading_bit"
OpName %v "v"
OpName %x "x"
OpName %firstLeadingBit_3fd7d0 "firstLeadingBit_3fd7d0"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%9 = OpTypeFunction %v3uint %v3uint
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
%17 = OpConstantNull %v3uint
%bool = OpTypeBool
%v3bool = OpTypeVector %bool 3
%uint_4294901760 = OpConstant %uint 4294901760
%24 = OpConstantComposite %v3uint %uint_4294901760 %uint_4294901760 %uint_4294901760
%uint_16 = OpConstant %uint 16
%27 = OpConstantComposite %v3uint %uint_16 %uint_16 %uint_16
%uint_0 = OpConstant %uint 0
%29 = OpConstantComposite %v3uint %uint_0 %uint_0 %uint_0
%uint_65280 = OpConstant %uint 65280
%36 = OpConstantComposite %v3uint %uint_65280 %uint_65280 %uint_65280
%uint_8 = OpConstant %uint 8
%39 = OpConstantComposite %v3uint %uint_8 %uint_8 %uint_8
%uint_240 = OpConstant %uint 240
%46 = OpConstantComposite %v3uint %uint_240 %uint_240 %uint_240
%uint_4 = OpConstant %uint 4
%49 = OpConstantComposite %v3uint %uint_4 %uint_4 %uint_4
%uint_12 = OpConstant %uint 12
%56 = OpConstantComposite %v3uint %uint_12 %uint_12 %uint_12
%uint_2 = OpConstant %uint 2
%59 = OpConstantComposite %v3uint %uint_2 %uint_2 %uint_2
%uint_1 = OpConstant %uint 1
%67 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
%uint_4294967295 = OpConstant %uint 4294967295
%72 = OpConstantComposite %v3uint %uint_4294967295 %uint_4294967295 %uint_4294967295
%void = OpTypeVoid
%79 = OpTypeFunction %void
%85 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_first_leading_bit = OpFunction %v3uint None %9
%v = OpFunctionParameter %v3uint
%14 = OpLabel
%x = OpVariable %_ptr_Function_v3uint Function %17
OpStore %x %v
%22 = OpLoad %v3uint %x
%25 = OpBitwiseAnd %v3uint %22 %24
%19 = OpINotEqual %v3bool %25 %17
%18 = OpSelect %v3uint %19 %27 %29
%30 = OpLoad %v3uint %x
%31 = OpShiftRightLogical %v3uint %30 %18
OpStore %x %31
%34 = OpLoad %v3uint %x
%37 = OpBitwiseAnd %v3uint %34 %36
%33 = OpINotEqual %v3bool %37 %17
%32 = OpSelect %v3uint %33 %39 %29
%40 = OpLoad %v3uint %x
%41 = OpShiftRightLogical %v3uint %40 %32
OpStore %x %41
%44 = OpLoad %v3uint %x
%47 = OpBitwiseAnd %v3uint %44 %46
%43 = OpINotEqual %v3bool %47 %17
%42 = OpSelect %v3uint %43 %49 %29
%50 = OpLoad %v3uint %x
%51 = OpShiftRightLogical %v3uint %50 %42
OpStore %x %51
%54 = OpLoad %v3uint %x
%57 = OpBitwiseAnd %v3uint %54 %56
%53 = OpINotEqual %v3bool %57 %17
%52 = OpSelect %v3uint %53 %59 %29
%60 = OpLoad %v3uint %x
%61 = OpShiftRightLogical %v3uint %60 %52
OpStore %x %61
%64 = OpLoad %v3uint %x
%65 = OpBitwiseAnd %v3uint %64 %59
%63 = OpINotEqual %v3bool %65 %17
%62 = OpSelect %v3uint %63 %67 %29
%69 = OpLoad %v3uint %x
%70 = OpIEqual %v3bool %69 %29
%68 = OpSelect %v3uint %70 %72 %29
%74 = OpBitwiseOr %v3uint %18 %32
%75 = OpBitwiseOr %v3uint %74 %42
%76 = OpBitwiseOr %v3uint %75 %52
%77 = OpBitwiseOr %v3uint %76 %62
%78 = OpBitwiseOr %v3uint %77 %68
OpReturnValue %78
OpFunctionEnd
%firstLeadingBit_3fd7d0 = OpFunction %void None %79
%82 = OpLabel
%res = OpVariable %_ptr_Function_v3uint Function %17
%83 = OpFunctionCall %v3uint %tint_first_leading_bit %17
OpStore %res %83
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %85
%87 = OpLabel
%88 = OpFunctionCall %void %firstLeadingBit_3fd7d0
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %79
%90 = OpLabel
%91 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %91
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %79
%94 = OpLabel
%95 = OpFunctionCall %void %firstLeadingBit_3fd7d0
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %79
%97 = OpLabel
%98 = OpFunctionCall %void %firstLeadingBit_3fd7d0
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/firstLeadingBit/3fd7d0.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,19 @@
fn firstLeadingBit_3fd7d0() {
var res : vec3<u32> = firstLeadingBit(vec3<u32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_3fd7d0();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_3fd7d0();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_3fd7d0();
}

45
test/tint/builtins/gen/firstLeadingBit/57a1a3.wgsl Normal file
View File

@ -0,0 +1,45 @@
// Copyright 2021 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/builtin-gen
// using the template:
// test/tint/builtins/builtins.wgsl.tmpl
// and the builtin defintion file:
// src/tint/builtins.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn firstLeadingBit(i32) -> i32
fn firstLeadingBit_57a1a3() {
var res: i32 = firstLeadingBit(1);
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_57a1a3();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_57a1a3();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_57a1a3();
}

93
test/tint/builtins/gen/firstLeadingBit/57a1a3.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
int tint_first_leading_bit(int v) {
uint x = ((v < 0) ? uint(~(v)) : uint(v));
uint b16 = (bool((x & 4294901760u)) ? 16u : 0u);
x = (x >> b16);
uint b8 = (bool((x & 65280u)) ? 8u : 0u);
x = (x >> b8);
uint b4 = (bool((x & 240u)) ? 4u : 0u);
x = (x >> b4);
uint b2 = (bool((x & 12u)) ? 2u : 0u);
x = (x >> b2);
uint b1 = (bool((x & 2u)) ? 1u : 0u);
uint is_zero = ((x == 0u) ? 4294967295u : 0u);
return int((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_57a1a3() {
int res = tint_first_leading_bit(1);
}
vec4 vertex_main() {
firstLeadingBit_57a1a3();
return vec4(0.0f, 0.0f, 0.0f, 0.0f);
}
void main() {
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;
int tint_first_leading_bit(int v) {
uint x = ((v < 0) ? uint(~(v)) : uint(v));
uint b16 = (bool((x & 4294901760u)) ? 16u : 0u);
x = (x >> b16);
uint b8 = (bool((x & 65280u)) ? 8u : 0u);
x = (x >> b8);
uint b4 = (bool((x & 240u)) ? 4u : 0u);
x = (x >> b4);
uint b2 = (bool((x & 12u)) ? 2u : 0u);
x = (x >> b2);
uint b1 = (bool((x & 2u)) ? 1u : 0u);
uint is_zero = ((x == 0u) ? 4294967295u : 0u);
return int((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_57a1a3() {
int res = tint_first_leading_bit(1);
}
void fragment_main() {
firstLeadingBit_57a1a3();
}
void main() {
fragment_main();
return;
}
#version 310 es
int tint_first_leading_bit(int v) {
uint x = ((v < 0) ? uint(~(v)) : uint(v));
uint b16 = (bool((x & 4294901760u)) ? 16u : 0u);
x = (x >> b16);
uint b8 = (bool((x & 65280u)) ? 8u : 0u);
x = (x >> b8);
uint b4 = (bool((x & 240u)) ? 4u : 0u);
x = (x >> b4);
uint b2 = (bool((x & 12u)) ? 2u : 0u);
x = (x >> b2);
uint b1 = (bool((x & 2u)) ? 1u : 0u);
uint is_zero = ((x == 0u) ? 4294967295u : 0u);
return int((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_57a1a3() {
int res = tint_first_leading_bit(1);
}
void compute_main() {
firstLeadingBit_57a1a3();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/firstLeadingBit/57a1a3.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
int tint_first_leading_bit(int v) {
uint x = ((v < 0) ? uint(~(v)) : uint(v));
const uint b16 = (bool((x & 4294901760u)) ? 16u : 0u);
x = (x >> b16);
const uint b8 = (bool((x & 65280u)) ? 8u : 0u);
x = (x >> b8);
const uint b4 = (bool((x & 240u)) ? 4u : 0u);
x = (x >> b4);
const uint b2 = (bool((x & 12u)) ? 2u : 0u);
x = (x >> b2);
const uint b1 = (bool((x & 2u)) ? 1u : 0u);
const uint is_zero = ((x == 0u) ? 4294967295u : 0u);
return int((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_57a1a3() {
int res = tint_first_leading_bit(1);
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
firstLeadingBit_57a1a3();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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() {
firstLeadingBit_57a1a3();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
firstLeadingBit_57a1a3();
return;
}

48
test/tint/builtins/gen/firstLeadingBit/57a1a3.wgsl.expected.msl Normal file
View File

@ -0,0 +1,48 @@
#include <metal_stdlib>
using namespace metal;
int tint_first_leading_bit(int v) {
uint x = select(uint(v), uint(~(v)), (v < 0));
uint const b16 = select(0u, 16u, bool((x & 4294901760u)));
x = (x >> b16);
uint const b8 = select(0u, 8u, bool((x & 65280u)));
x = (x >> b8);
uint const b4 = select(0u, 4u, bool((x & 240u)));
x = (x >> b4);
uint const b2 = select(0u, 2u, bool((x & 12u)));
x = (x >> b2);
uint const b1 = select(0u, 1u, bool((x & 2u)));
uint const is_zero = select(0u, 4294967295u, (x == 0u));
return int((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_57a1a3() {
int res = tint_first_leading_bit(1);
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
firstLeadingBit_57a1a3();
return float4();
}
vertex tint_symbol vertex_main() {
float4 const inner_result = vertex_main_inner();
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main() {
firstLeadingBit_57a1a3();
return;
}
kernel void compute_main() {
firstLeadingBit_57a1a3();
return;
}

140
test/tint/builtins/gen/firstLeadingBit/57a1a3.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,140 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 96
; 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 %tint_first_leading_bit "tint_first_leading_bit"
OpName %v "v"
OpName %x "x"
OpName %firstLeadingBit_57a1a3 "firstLeadingBit_57a1a3"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%int = OpTypeInt 32 1
%9 = OpTypeFunction %int %int
%uint = OpTypeInt 32 0
%int_0 = OpConstant %int 0
%bool = OpTypeBool
%_ptr_Function_uint = OpTypePointer Function %uint
%24 = OpConstantNull %uint
%uint_4294901760 = OpConstant %uint 4294901760
%uint_16 = OpConstant %uint 16
%uint_0 = OpConstant %uint 0
%uint_65280 = OpConstant %uint 65280
%uint_8 = OpConstant %uint 8
%uint_240 = OpConstant %uint 240
%uint_4 = OpConstant %uint 4
%uint_12 = OpConstant %uint 12
%uint_2 = OpConstant %uint 2
%uint_1 = OpConstant %uint 1
%uint_4294967295 = OpConstant %uint 4294967295
%void = OpTypeVoid
%73 = OpTypeFunction %void
%int_1 = OpConstant %int 1
%_ptr_Function_int = OpTypePointer Function %int
%81 = OpConstantNull %int
%82 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_first_leading_bit = OpFunction %int None %9
%v = OpFunctionParameter %int
%13 = OpLabel
%x = OpVariable %_ptr_Function_uint Function %24
%17 = OpSLessThan %bool %v %int_0
%20 = OpNot %int %v
%19 = OpBitcast %uint %20
%21 = OpBitcast %uint %v
%14 = OpSelect %uint %17 %19 %21
OpStore %x %14
%27 = OpLoad %uint %x
%29 = OpBitwiseAnd %uint %27 %uint_4294901760
%26 = OpINotEqual %bool %29 %24
%25 = OpSelect %uint %26 %uint_16 %uint_0
%32 = OpLoad %uint %x
%33 = OpShiftRightLogical %uint %32 %25
OpStore %x %33
%36 = OpLoad %uint %x
%38 = OpBitwiseAnd %uint %36 %uint_65280
%35 = OpINotEqual %bool %38 %24
%34 = OpSelect %uint %35 %uint_8 %uint_0
%40 = OpLoad %uint %x
%41 = OpShiftRightLogical %uint %40 %34
OpStore %x %41
%44 = OpLoad %uint %x
%46 = OpBitwiseAnd %uint %44 %uint_240
%43 = OpINotEqual %bool %46 %24
%42 = OpSelect %uint %43 %uint_4 %uint_0
%48 = OpLoad %uint %x
%49 = OpShiftRightLogical %uint %48 %42
OpStore %x %49
%52 = OpLoad %uint %x
%54 = OpBitwiseAnd %uint %52 %uint_12
%51 = OpINotEqual %bool %54 %24
%50 = OpSelect %uint %51 %uint_2 %uint_0
%56 = OpLoad %uint %x
%57 = OpShiftRightLogical %uint %56 %50
OpStore %x %57
%60 = OpLoad %uint %x
%61 = OpBitwiseAnd %uint %60 %uint_2
%59 = OpINotEqual %bool %61 %24
%58 = OpSelect %uint %59 %uint_1 %uint_0
%64 = OpLoad %uint %x
%65 = OpIEqual %bool %64 %uint_0
%63 = OpSelect %uint %65 %uint_4294967295 %uint_0
%68 = OpBitwiseOr %uint %25 %34
%69 = OpBitwiseOr %uint %68 %42
%70 = OpBitwiseOr %uint %69 %50
%71 = OpBitwiseOr %uint %70 %58
%72 = OpBitwiseOr %uint %71 %63
%67 = OpBitcast %int %72
OpReturnValue %67
OpFunctionEnd
%firstLeadingBit_57a1a3 = OpFunction %void None %73
%76 = OpLabel
%res = OpVariable %_ptr_Function_int Function %81
%77 = OpFunctionCall %int %tint_first_leading_bit %int_1
OpStore %res %77
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %82
%84 = OpLabel
%85 = OpFunctionCall %void %firstLeadingBit_57a1a3
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %73
%87 = OpLabel
%88 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %88
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %73
%91 = OpLabel
%92 = OpFunctionCall %void %firstLeadingBit_57a1a3
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %73
%94 = OpLabel
%95 = OpFunctionCall %void %firstLeadingBit_57a1a3
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/firstLeadingBit/57a1a3.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,19 @@
fn firstLeadingBit_57a1a3() {
var res : i32 = firstLeadingBit(1);
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_57a1a3();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_57a1a3();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_57a1a3();
}

45
test/tint/builtins/gen/firstLeadingBit/6fe804.wgsl Normal file
View File

@ -0,0 +1,45 @@
// Copyright 2021 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/builtin-gen
// using the template:
// test/tint/builtins/builtins.wgsl.tmpl
// and the builtin defintion file:
// src/tint/builtins.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn firstLeadingBit(vec<2, u32>) -> vec<2, u32>
fn firstLeadingBit_6fe804() {
var res: vec2<u32> = firstLeadingBit(vec2<u32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_6fe804();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_6fe804();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_6fe804();
}

93
test/tint/builtins/gen/firstLeadingBit/6fe804.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
uvec2 tint_first_leading_bit(uvec2 v) {
uvec2 x = v;
uvec2 b16 = mix(uvec2(0u), uvec2(16u), bvec2((x & uvec2(4294901760u))));
x = (x >> b16);
uvec2 b8 = mix(uvec2(0u), uvec2(8u), bvec2((x & uvec2(65280u))));
x = (x >> b8);
uvec2 b4 = mix(uvec2(0u), uvec2(4u), bvec2((x & uvec2(240u))));
x = (x >> b4);
uvec2 b2 = mix(uvec2(0u), uvec2(2u), bvec2((x & uvec2(12u))));
x = (x >> b2);
uvec2 b1 = mix(uvec2(0u), uvec2(1u), bvec2((x & uvec2(2u))));
uvec2 is_zero = mix(uvec2(0u), uvec2(4294967295u), equal(x, uvec2(0u)));
return uvec2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_6fe804() {
uvec2 res = tint_first_leading_bit(uvec2(0u, 0u));
}
vec4 vertex_main() {
firstLeadingBit_6fe804();
return vec4(0.0f, 0.0f, 0.0f, 0.0f);
}
void main() {
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;
uvec2 tint_first_leading_bit(uvec2 v) {
uvec2 x = v;
uvec2 b16 = mix(uvec2(0u), uvec2(16u), bvec2((x & uvec2(4294901760u))));
x = (x >> b16);
uvec2 b8 = mix(uvec2(0u), uvec2(8u), bvec2((x & uvec2(65280u))));
x = (x >> b8);
uvec2 b4 = mix(uvec2(0u), uvec2(4u), bvec2((x & uvec2(240u))));
x = (x >> b4);
uvec2 b2 = mix(uvec2(0u), uvec2(2u), bvec2((x & uvec2(12u))));
x = (x >> b2);
uvec2 b1 = mix(uvec2(0u), uvec2(1u), bvec2((x & uvec2(2u))));
uvec2 is_zero = mix(uvec2(0u), uvec2(4294967295u), equal(x, uvec2(0u)));
return uvec2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_6fe804() {
uvec2 res = tint_first_leading_bit(uvec2(0u, 0u));
}
void fragment_main() {
firstLeadingBit_6fe804();
}
void main() {
fragment_main();
return;
}
#version 310 es
uvec2 tint_first_leading_bit(uvec2 v) {
uvec2 x = v;
uvec2 b16 = mix(uvec2(0u), uvec2(16u), bvec2((x & uvec2(4294901760u))));
x = (x >> b16);
uvec2 b8 = mix(uvec2(0u), uvec2(8u), bvec2((x & uvec2(65280u))));
x = (x >> b8);
uvec2 b4 = mix(uvec2(0u), uvec2(4u), bvec2((x & uvec2(240u))));
x = (x >> b4);
uvec2 b2 = mix(uvec2(0u), uvec2(2u), bvec2((x & uvec2(12u))));
x = (x >> b2);
uvec2 b1 = mix(uvec2(0u), uvec2(1u), bvec2((x & uvec2(2u))));
uvec2 is_zero = mix(uvec2(0u), uvec2(4294967295u), equal(x, uvec2(0u)));
return uvec2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_6fe804() {
uvec2 res = tint_first_leading_bit(uvec2(0u, 0u));
}
void compute_main() {
firstLeadingBit_6fe804();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/firstLeadingBit/6fe804.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
uint2 tint_first_leading_bit(uint2 v) {
uint2 x = v;
const uint2 b16 = (bool2((x & uint2((4294901760u).xx))) ? uint2((16u).xx) : uint2((0u).xx));
x = (x >> b16);
const uint2 b8 = (bool2((x & uint2((65280u).xx))) ? uint2((8u).xx) : uint2((0u).xx));
x = (x >> b8);
const uint2 b4 = (bool2((x & uint2((240u).xx))) ? uint2((4u).xx) : uint2((0u).xx));
x = (x >> b4);
const uint2 b2 = (bool2((x & uint2((12u).xx))) ? uint2((2u).xx) : uint2((0u).xx));
x = (x >> b2);
const uint2 b1 = (bool2((x & uint2((2u).xx))) ? uint2((1u).xx) : uint2((0u).xx));
const uint2 is_zero = ((x == uint2((0u).xx)) ? uint2((4294967295u).xx) : uint2((0u).xx));
return uint2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_6fe804() {
uint2 res = tint_first_leading_bit(uint2(0u, 0u));
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
firstLeadingBit_6fe804();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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() {
firstLeadingBit_6fe804();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
firstLeadingBit_6fe804();
return;
}

48
test/tint/builtins/gen/firstLeadingBit/6fe804.wgsl.expected.msl Normal file
View File

@ -0,0 +1,48 @@
#include <metal_stdlib>
using namespace metal;
uint2 tint_first_leading_bit(uint2 v) {
uint2 x = v;
uint2 const b16 = select(uint2(0u), uint2(16u), bool2((x & uint2(4294901760u))));
x = (x >> b16);
uint2 const b8 = select(uint2(0u), uint2(8u), bool2((x & uint2(65280u))));
x = (x >> b8);
uint2 const b4 = select(uint2(0u), uint2(4u), bool2((x & uint2(240u))));
x = (x >> b4);
uint2 const b2 = select(uint2(0u), uint2(2u), bool2((x & uint2(12u))));
x = (x >> b2);
uint2 const b1 = select(uint2(0u), uint2(1u), bool2((x & uint2(2u))));
uint2 const is_zero = select(uint2(0u), uint2(4294967295u), (x == uint2(0u)));
return uint2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_6fe804() {
uint2 res = tint_first_leading_bit(uint2());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
firstLeadingBit_6fe804();
return float4();
}
vertex tint_symbol vertex_main() {
float4 const inner_result = vertex_main_inner();
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main() {
firstLeadingBit_6fe804();
return;
}
kernel void compute_main() {
firstLeadingBit_6fe804();
return;
}

142
test/tint/builtins/gen/firstLeadingBit/6fe804.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,142 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 99
; 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 %tint_first_leading_bit "tint_first_leading_bit"
OpName %v "v"
OpName %x "x"
OpName %firstLeadingBit_6fe804 "firstLeadingBit_6fe804"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%uint = OpTypeInt 32 0
%v2uint = OpTypeVector %uint 2
%9 = OpTypeFunction %v2uint %v2uint
%_ptr_Function_v2uint = OpTypePointer Function %v2uint
%17 = OpConstantNull %v2uint
%bool = OpTypeBool
%v2bool = OpTypeVector %bool 2
%uint_4294901760 = OpConstant %uint 4294901760
%24 = OpConstantComposite %v2uint %uint_4294901760 %uint_4294901760
%uint_16 = OpConstant %uint 16
%27 = OpConstantComposite %v2uint %uint_16 %uint_16
%uint_0 = OpConstant %uint 0
%29 = OpConstantComposite %v2uint %uint_0 %uint_0
%uint_65280 = OpConstant %uint 65280
%36 = OpConstantComposite %v2uint %uint_65280 %uint_65280
%uint_8 = OpConstant %uint 8
%39 = OpConstantComposite %v2uint %uint_8 %uint_8
%uint_240 = OpConstant %uint 240
%46 = OpConstantComposite %v2uint %uint_240 %uint_240
%uint_4 = OpConstant %uint 4
%49 = OpConstantComposite %v2uint %uint_4 %uint_4
%uint_12 = OpConstant %uint 12
%56 = OpConstantComposite %v2uint %uint_12 %uint_12
%uint_2 = OpConstant %uint 2
%59 = OpConstantComposite %v2uint %uint_2 %uint_2
%uint_1 = OpConstant %uint 1
%67 = OpConstantComposite %v2uint %uint_1 %uint_1
%uint_4294967295 = OpConstant %uint 4294967295
%72 = OpConstantComposite %v2uint %uint_4294967295 %uint_4294967295
%void = OpTypeVoid
%79 = OpTypeFunction %void
%85 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_first_leading_bit = OpFunction %v2uint None %9
%v = OpFunctionParameter %v2uint
%14 = OpLabel
%x = OpVariable %_ptr_Function_v2uint Function %17
OpStore %x %v
%22 = OpLoad %v2uint %x
%25 = OpBitwiseAnd %v2uint %22 %24
%19 = OpINotEqual %v2bool %25 %17
%18 = OpSelect %v2uint %19 %27 %29
%30 = OpLoad %v2uint %x
%31 = OpShiftRightLogical %v2uint %30 %18
OpStore %x %31
%34 = OpLoad %v2uint %x
%37 = OpBitwiseAnd %v2uint %34 %36
%33 = OpINotEqual %v2bool %37 %17
%32 = OpSelect %v2uint %33 %39 %29
%40 = OpLoad %v2uint %x
%41 = OpShiftRightLogical %v2uint %40 %32
OpStore %x %41
%44 = OpLoad %v2uint %x
%47 = OpBitwiseAnd %v2uint %44 %46
%43 = OpINotEqual %v2bool %47 %17
%42 = OpSelect %v2uint %43 %49 %29
%50 = OpLoad %v2uint %x
%51 = OpShiftRightLogical %v2uint %50 %42
OpStore %x %51
%54 = OpLoad %v2uint %x
%57 = OpBitwiseAnd %v2uint %54 %56
%53 = OpINotEqual %v2bool %57 %17
%52 = OpSelect %v2uint %53 %59 %29
%60 = OpLoad %v2uint %x
%61 = OpShiftRightLogical %v2uint %60 %52
OpStore %x %61
%64 = OpLoad %v2uint %x
%65 = OpBitwiseAnd %v2uint %64 %59
%63 = OpINotEqual %v2bool %65 %17
%62 = OpSelect %v2uint %63 %67 %29
%69 = OpLoad %v2uint %x
%70 = OpIEqual %v2bool %69 %29
%68 = OpSelect %v2uint %70 %72 %29
%74 = OpBitwiseOr %v2uint %18 %32
%75 = OpBitwiseOr %v2uint %74 %42
%76 = OpBitwiseOr %v2uint %75 %52
%77 = OpBitwiseOr %v2uint %76 %62
%78 = OpBitwiseOr %v2uint %77 %68
OpReturnValue %78
OpFunctionEnd
%firstLeadingBit_6fe804 = OpFunction %void None %79
%82 = OpLabel
%res = OpVariable %_ptr_Function_v2uint Function %17
%83 = OpFunctionCall %v2uint %tint_first_leading_bit %17
OpStore %res %83
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %85
%87 = OpLabel
%88 = OpFunctionCall %void %firstLeadingBit_6fe804
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %79
%90 = OpLabel
%91 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %91
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %79
%94 = OpLabel
%95 = OpFunctionCall %void %firstLeadingBit_6fe804
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %79
%97 = OpLabel
%98 = OpFunctionCall %void %firstLeadingBit_6fe804
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/firstLeadingBit/6fe804.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,19 @@
fn firstLeadingBit_6fe804() {
var res : vec2<u32> = firstLeadingBit(vec2<u32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_6fe804();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_6fe804();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_6fe804();
}

45
test/tint/builtins/gen/firstLeadingBit/a622c2.wgsl Normal file
View File

@ -0,0 +1,45 @@
// Copyright 2021 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/builtin-gen
// using the template:
// test/tint/builtins/builtins.wgsl.tmpl
// and the builtin defintion file:
// src/tint/builtins.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn firstLeadingBit(vec<2, i32>) -> vec<2, i32>
fn firstLeadingBit_a622c2() {
var res: vec2<i32> = firstLeadingBit(vec2<i32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_a622c2();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_a622c2();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_a622c2();
}

93
test/tint/builtins/gen/firstLeadingBit/a622c2.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
ivec2 tint_first_leading_bit(ivec2 v) {
uvec2 x = mix(uvec2(v), uvec2(~(v)), lessThan(v, ivec2(0)));
uvec2 b16 = mix(uvec2(0u), uvec2(16u), bvec2((x & uvec2(4294901760u))));
x = (x >> b16);
uvec2 b8 = mix(uvec2(0u), uvec2(8u), bvec2((x & uvec2(65280u))));
x = (x >> b8);
uvec2 b4 = mix(uvec2(0u), uvec2(4u), bvec2((x & uvec2(240u))));
x = (x >> b4);
uvec2 b2 = mix(uvec2(0u), uvec2(2u), bvec2((x & uvec2(12u))));
x = (x >> b2);
uvec2 b1 = mix(uvec2(0u), uvec2(1u), bvec2((x & uvec2(2u))));
uvec2 is_zero = mix(uvec2(0u), uvec2(4294967295u), equal(x, uvec2(0u)));
return ivec2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_a622c2() {
ivec2 res = tint_first_leading_bit(ivec2(0, 0));
}
vec4 vertex_main() {
firstLeadingBit_a622c2();
return vec4(0.0f, 0.0f, 0.0f, 0.0f);
}
void main() {
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;
ivec2 tint_first_leading_bit(ivec2 v) {
uvec2 x = mix(uvec2(v), uvec2(~(v)), lessThan(v, ivec2(0)));
uvec2 b16 = mix(uvec2(0u), uvec2(16u), bvec2((x & uvec2(4294901760u))));
x = (x >> b16);
uvec2 b8 = mix(uvec2(0u), uvec2(8u), bvec2((x & uvec2(65280u))));
x = (x >> b8);
uvec2 b4 = mix(uvec2(0u), uvec2(4u), bvec2((x & uvec2(240u))));
x = (x >> b4);
uvec2 b2 = mix(uvec2(0u), uvec2(2u), bvec2((x & uvec2(12u))));
x = (x >> b2);
uvec2 b1 = mix(uvec2(0u), uvec2(1u), bvec2((x & uvec2(2u))));
uvec2 is_zero = mix(uvec2(0u), uvec2(4294967295u), equal(x, uvec2(0u)));
return ivec2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_a622c2() {
ivec2 res = tint_first_leading_bit(ivec2(0, 0));
}
void fragment_main() {
firstLeadingBit_a622c2();
}
void main() {
fragment_main();
return;
}
#version 310 es
ivec2 tint_first_leading_bit(ivec2 v) {
uvec2 x = mix(uvec2(v), uvec2(~(v)), lessThan(v, ivec2(0)));
uvec2 b16 = mix(uvec2(0u), uvec2(16u), bvec2((x & uvec2(4294901760u))));
x = (x >> b16);
uvec2 b8 = mix(uvec2(0u), uvec2(8u), bvec2((x & uvec2(65280u))));
x = (x >> b8);
uvec2 b4 = mix(uvec2(0u), uvec2(4u), bvec2((x & uvec2(240u))));
x = (x >> b4);
uvec2 b2 = mix(uvec2(0u), uvec2(2u), bvec2((x & uvec2(12u))));
x = (x >> b2);
uvec2 b1 = mix(uvec2(0u), uvec2(1u), bvec2((x & uvec2(2u))));
uvec2 is_zero = mix(uvec2(0u), uvec2(4294967295u), equal(x, uvec2(0u)));
return ivec2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_a622c2() {
ivec2 res = tint_first_leading_bit(ivec2(0, 0));
}
void compute_main() {
firstLeadingBit_a622c2();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/firstLeadingBit/a622c2.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
int2 tint_first_leading_bit(int2 v) {
uint2 x = ((v < int2((0).xx)) ? uint2(~(v)) : uint2(v));
const uint2 b16 = (bool2((x & uint2((4294901760u).xx))) ? uint2((16u).xx) : uint2((0u).xx));
x = (x >> b16);
const uint2 b8 = (bool2((x & uint2((65280u).xx))) ? uint2((8u).xx) : uint2((0u).xx));
x = (x >> b8);
const uint2 b4 = (bool2((x & uint2((240u).xx))) ? uint2((4u).xx) : uint2((0u).xx));
x = (x >> b4);
const uint2 b2 = (bool2((x & uint2((12u).xx))) ? uint2((2u).xx) : uint2((0u).xx));
x = (x >> b2);
const uint2 b1 = (bool2((x & uint2((2u).xx))) ? uint2((1u).xx) : uint2((0u).xx));
const uint2 is_zero = ((x == uint2((0u).xx)) ? uint2((4294967295u).xx) : uint2((0u).xx));
return int2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_a622c2() {
int2 res = tint_first_leading_bit(int2(0, 0));
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
firstLeadingBit_a622c2();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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() {
firstLeadingBit_a622c2();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
firstLeadingBit_a622c2();
return;
}

48
test/tint/builtins/gen/firstLeadingBit/a622c2.wgsl.expected.msl Normal file
View File

@ -0,0 +1,48 @@
#include <metal_stdlib>
using namespace metal;
int2 tint_first_leading_bit(int2 v) {
uint2 x = select(uint2(v), uint2(~(v)), (v < int2(0)));
uint2 const b16 = select(uint2(0u), uint2(16u), bool2((x & uint2(4294901760u))));
x = (x >> b16);
uint2 const b8 = select(uint2(0u), uint2(8u), bool2((x & uint2(65280u))));
x = (x >> b8);
uint2 const b4 = select(uint2(0u), uint2(4u), bool2((x & uint2(240u))));
x = (x >> b4);
uint2 const b2 = select(uint2(0u), uint2(2u), bool2((x & uint2(12u))));
x = (x >> b2);
uint2 const b1 = select(uint2(0u), uint2(1u), bool2((x & uint2(2u))));
uint2 const is_zero = select(uint2(0u), uint2(4294967295u), (x == uint2(0u)));
return int2((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_a622c2() {
int2 res = tint_first_leading_bit(int2());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
firstLeadingBit_a622c2();
return float4();
}
vertex tint_symbol vertex_main() {
float4 const inner_result = vertex_main_inner();
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main() {
firstLeadingBit_a622c2();
return;
}
kernel void compute_main() {
firstLeadingBit_a622c2();
return;
}

154
test/tint/builtins/gen/firstLeadingBit/a622c2.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,154 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 110
; 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 %tint_first_leading_bit "tint_first_leading_bit"
OpName %v "v"
OpName %x "x"
OpName %firstLeadingBit_a622c2 "firstLeadingBit_a622c2"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%int = OpTypeInt 32 1
%v2int = OpTypeVector %int 2
%9 = OpTypeFunction %v2int %v2int
%uint = OpTypeInt 32 0
%v2uint = OpTypeVector %uint 2
%int_0 = OpConstant %int 0
%19 = OpConstantComposite %v2int %int_0 %int_0
%bool = OpTypeBool
%v2bool = OpTypeVector %bool 2
%_ptr_Function_v2uint = OpTypePointer Function %v2uint
%28 = OpConstantNull %v2uint
%uint_4294901760 = OpConstant %uint 4294901760
%33 = OpConstantComposite %v2uint %uint_4294901760 %uint_4294901760
%uint_16 = OpConstant %uint 16
%36 = OpConstantComposite %v2uint %uint_16 %uint_16
%uint_0 = OpConstant %uint 0
%38 = OpConstantComposite %v2uint %uint_0 %uint_0
%uint_65280 = OpConstant %uint 65280
%45 = OpConstantComposite %v2uint %uint_65280 %uint_65280
%uint_8 = OpConstant %uint 8
%48 = OpConstantComposite %v2uint %uint_8 %uint_8
%uint_240 = OpConstant %uint 240
%55 = OpConstantComposite %v2uint %uint_240 %uint_240
%uint_4 = OpConstant %uint 4
%58 = OpConstantComposite %v2uint %uint_4 %uint_4
%uint_12 = OpConstant %uint 12
%65 = OpConstantComposite %v2uint %uint_12 %uint_12
%uint_2 = OpConstant %uint 2
%68 = OpConstantComposite %v2uint %uint_2 %uint_2
%uint_1 = OpConstant %uint 1
%76 = OpConstantComposite %v2uint %uint_1 %uint_1
%uint_4294967295 = OpConstant %uint 4294967295
%81 = OpConstantComposite %v2uint %uint_4294967295 %uint_4294967295
%void = OpTypeVoid
%88 = OpTypeFunction %void
%93 = OpConstantNull %v2int
%_ptr_Function_v2int = OpTypePointer Function %v2int
%96 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_first_leading_bit = OpFunction %v2int None %9
%v = OpFunctionParameter %v2int
%14 = OpLabel
%x = OpVariable %_ptr_Function_v2uint Function %28
%20 = OpSLessThan %v2bool %v %19
%24 = OpNot %v2int %v
%23 = OpBitcast %v2uint %24
%25 = OpBitcast %v2uint %v
%15 = OpSelect %v2uint %20 %23 %25
OpStore %x %15
%31 = OpLoad %v2uint %x
%34 = OpBitwiseAnd %v2uint %31 %33
%30 = OpINotEqual %v2bool %34 %28
%29 = OpSelect %v2uint %30 %36 %38
%39 = OpLoad %v2uint %x
%40 = OpShiftRightLogical %v2uint %39 %29
OpStore %x %40
%43 = OpLoad %v2uint %x
%46 = OpBitwiseAnd %v2uint %43 %45
%42 = OpINotEqual %v2bool %46 %28
%41 = OpSelect %v2uint %42 %48 %38
%49 = OpLoad %v2uint %x
%50 = OpShiftRightLogical %v2uint %49 %41
OpStore %x %50
%53 = OpLoad %v2uint %x
%56 = OpBitwiseAnd %v2uint %53 %55
%52 = OpINotEqual %v2bool %56 %28
%51 = OpSelect %v2uint %52 %58 %38
%59 = OpLoad %v2uint %x
%60 = OpShiftRightLogical %v2uint %59 %51
OpStore %x %60
%63 = OpLoad %v2uint %x
%66 = OpBitwiseAnd %v2uint %63 %65
%62 = OpINotEqual %v2bool %66 %28
%61 = OpSelect %v2uint %62 %68 %38
%69 = OpLoad %v2uint %x
%70 = OpShiftRightLogical %v2uint %69 %61
OpStore %x %70
%73 = OpLoad %v2uint %x
%74 = OpBitwiseAnd %v2uint %73 %68
%72 = OpINotEqual %v2bool %74 %28
%71 = OpSelect %v2uint %72 %76 %38
%78 = OpLoad %v2uint %x
%79 = OpIEqual %v2bool %78 %38
%77 = OpSelect %v2uint %79 %81 %38
%83 = OpBitwiseOr %v2uint %29 %41
%84 = OpBitwiseOr %v2uint %83 %51
%85 = OpBitwiseOr %v2uint %84 %61
%86 = OpBitwiseOr %v2uint %85 %71
%87 = OpBitwiseOr %v2uint %86 %77
%82 = OpBitcast %v2int %87
OpReturnValue %82
OpFunctionEnd
%firstLeadingBit_a622c2 = OpFunction %void None %88
%91 = OpLabel
%res = OpVariable %_ptr_Function_v2int Function %93
%92 = OpFunctionCall %v2int %tint_first_leading_bit %93
OpStore %res %92
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %96
%98 = OpLabel
%99 = OpFunctionCall %void %firstLeadingBit_a622c2
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %88
%101 = OpLabel
%102 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %102
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %88
%105 = OpLabel
%106 = OpFunctionCall %void %firstLeadingBit_a622c2
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %88
%108 = OpLabel
%109 = OpFunctionCall %void %firstLeadingBit_a622c2
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/firstLeadingBit/a622c2.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,19 @@
fn firstLeadingBit_a622c2() {
var res : vec2<i32> = firstLeadingBit(vec2<i32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_a622c2();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_a622c2();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_a622c2();
}

45
test/tint/builtins/gen/firstLeadingBit/c1f940.wgsl Normal file
View File

@ -0,0 +1,45 @@
// Copyright 2021 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/builtin-gen
// using the template:
// test/tint/builtins/builtins.wgsl.tmpl
// and the builtin defintion file:
// src/tint/builtins.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn firstLeadingBit(vec<4, i32>) -> vec<4, i32>
fn firstLeadingBit_c1f940() {
var res: vec4<i32> = firstLeadingBit(vec4<i32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_c1f940();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_c1f940();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_c1f940();
}

93
test/tint/builtins/gen/firstLeadingBit/c1f940.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
ivec4 tint_first_leading_bit(ivec4 v) {
uvec4 x = mix(uvec4(v), uvec4(~(v)), lessThan(v, ivec4(0)));
uvec4 b16 = mix(uvec4(0u), uvec4(16u), bvec4((x & uvec4(4294901760u))));
x = (x >> b16);
uvec4 b8 = mix(uvec4(0u), uvec4(8u), bvec4((x & uvec4(65280u))));
x = (x >> b8);
uvec4 b4 = mix(uvec4(0u), uvec4(4u), bvec4((x & uvec4(240u))));
x = (x >> b4);
uvec4 b2 = mix(uvec4(0u), uvec4(2u), bvec4((x & uvec4(12u))));
x = (x >> b2);
uvec4 b1 = mix(uvec4(0u), uvec4(1u), bvec4((x & uvec4(2u))));
uvec4 is_zero = mix(uvec4(0u), uvec4(4294967295u), equal(x, uvec4(0u)));
return ivec4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_c1f940() {
ivec4 res = tint_first_leading_bit(ivec4(0, 0, 0, 0));
}
vec4 vertex_main() {
firstLeadingBit_c1f940();
return vec4(0.0f, 0.0f, 0.0f, 0.0f);
}
void main() {
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;
ivec4 tint_first_leading_bit(ivec4 v) {
uvec4 x = mix(uvec4(v), uvec4(~(v)), lessThan(v, ivec4(0)));
uvec4 b16 = mix(uvec4(0u), uvec4(16u), bvec4((x & uvec4(4294901760u))));
x = (x >> b16);
uvec4 b8 = mix(uvec4(0u), uvec4(8u), bvec4((x & uvec4(65280u))));
x = (x >> b8);
uvec4 b4 = mix(uvec4(0u), uvec4(4u), bvec4((x & uvec4(240u))));
x = (x >> b4);
uvec4 b2 = mix(uvec4(0u), uvec4(2u), bvec4((x & uvec4(12u))));
x = (x >> b2);
uvec4 b1 = mix(uvec4(0u), uvec4(1u), bvec4((x & uvec4(2u))));
uvec4 is_zero = mix(uvec4(0u), uvec4(4294967295u), equal(x, uvec4(0u)));
return ivec4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_c1f940() {
ivec4 res = tint_first_leading_bit(ivec4(0, 0, 0, 0));
}
void fragment_main() {
firstLeadingBit_c1f940();
}
void main() {
fragment_main();
return;
}
#version 310 es
ivec4 tint_first_leading_bit(ivec4 v) {
uvec4 x = mix(uvec4(v), uvec4(~(v)), lessThan(v, ivec4(0)));
uvec4 b16 = mix(uvec4(0u), uvec4(16u), bvec4((x & uvec4(4294901760u))));
x = (x >> b16);
uvec4 b8 = mix(uvec4(0u), uvec4(8u), bvec4((x & uvec4(65280u))));
x = (x >> b8);
uvec4 b4 = mix(uvec4(0u), uvec4(4u), bvec4((x & uvec4(240u))));
x = (x >> b4);
uvec4 b2 = mix(uvec4(0u), uvec4(2u), bvec4((x & uvec4(12u))));
x = (x >> b2);
uvec4 b1 = mix(uvec4(0u), uvec4(1u), bvec4((x & uvec4(2u))));
uvec4 is_zero = mix(uvec4(0u), uvec4(4294967295u), equal(x, uvec4(0u)));
return ivec4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_c1f940() {
ivec4 res = tint_first_leading_bit(ivec4(0, 0, 0, 0));
}
void compute_main() {
firstLeadingBit_c1f940();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/firstLeadingBit/c1f940.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
int4 tint_first_leading_bit(int4 v) {
uint4 x = ((v < int4((0).xxxx)) ? uint4(~(v)) : uint4(v));
const uint4 b16 = (bool4((x & uint4((4294901760u).xxxx))) ? uint4((16u).xxxx) : uint4((0u).xxxx));
x = (x >> b16);
const uint4 b8 = (bool4((x & uint4((65280u).xxxx))) ? uint4((8u).xxxx) : uint4((0u).xxxx));
x = (x >> b8);
const uint4 b4 = (bool4((x & uint4((240u).xxxx))) ? uint4((4u).xxxx) : uint4((0u).xxxx));
x = (x >> b4);
const uint4 b2 = (bool4((x & uint4((12u).xxxx))) ? uint4((2u).xxxx) : uint4((0u).xxxx));
x = (x >> b2);
const uint4 b1 = (bool4((x & uint4((2u).xxxx))) ? uint4((1u).xxxx) : uint4((0u).xxxx));
const uint4 is_zero = ((x == uint4((0u).xxxx)) ? uint4((4294967295u).xxxx) : uint4((0u).xxxx));
return int4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_c1f940() {
int4 res = tint_first_leading_bit(int4(0, 0, 0, 0));
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
firstLeadingBit_c1f940();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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() {
firstLeadingBit_c1f940();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
firstLeadingBit_c1f940();
return;
}

48
test/tint/builtins/gen/firstLeadingBit/c1f940.wgsl.expected.msl Normal file
View File

@ -0,0 +1,48 @@
#include <metal_stdlib>
using namespace metal;
int4 tint_first_leading_bit(int4 v) {
uint4 x = select(uint4(v), uint4(~(v)), (v < int4(0)));
uint4 const b16 = select(uint4(0u), uint4(16u), bool4((x & uint4(4294901760u))));
x = (x >> b16);
uint4 const b8 = select(uint4(0u), uint4(8u), bool4((x & uint4(65280u))));
x = (x >> b8);
uint4 const b4 = select(uint4(0u), uint4(4u), bool4((x & uint4(240u))));
x = (x >> b4);
uint4 const b2 = select(uint4(0u), uint4(2u), bool4((x & uint4(12u))));
x = (x >> b2);
uint4 const b1 = select(uint4(0u), uint4(1u), bool4((x & uint4(2u))));
uint4 const is_zero = select(uint4(0u), uint4(4294967295u), (x == uint4(0u)));
return int4((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_c1f940() {
int4 res = tint_first_leading_bit(int4());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
firstLeadingBit_c1f940();
return float4();
}
vertex tint_symbol vertex_main() {
float4 const inner_result = vertex_main_inner();
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main() {
firstLeadingBit_c1f940();
return;
}
kernel void compute_main() {
firstLeadingBit_c1f940();
return;
}

154
test/tint/builtins/gen/firstLeadingBit/c1f940.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,154 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 110
; 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 %tint_first_leading_bit "tint_first_leading_bit"
OpName %v "v"
OpName %x "x"
OpName %firstLeadingBit_c1f940 "firstLeadingBit_c1f940"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%int = OpTypeInt 32 1
%v4int = OpTypeVector %int 4
%9 = OpTypeFunction %v4int %v4int
%uint = OpTypeInt 32 0
%v4uint = OpTypeVector %uint 4
%int_0 = OpConstant %int 0
%19 = OpConstantComposite %v4int %int_0 %int_0 %int_0 %int_0
%bool = OpTypeBool
%v4bool = OpTypeVector %bool 4
%_ptr_Function_v4uint = OpTypePointer Function %v4uint
%28 = OpConstantNull %v4uint
%uint_4294901760 = OpConstant %uint 4294901760
%33 = OpConstantComposite %v4uint %uint_4294901760 %uint_4294901760 %uint_4294901760 %uint_4294901760
%uint_16 = OpConstant %uint 16
%36 = OpConstantComposite %v4uint %uint_16 %uint_16 %uint_16 %uint_16
%uint_0 = OpConstant %uint 0
%38 = OpConstantComposite %v4uint %uint_0 %uint_0 %uint_0 %uint_0
%uint_65280 = OpConstant %uint 65280
%45 = OpConstantComposite %v4uint %uint_65280 %uint_65280 %uint_65280 %uint_65280
%uint_8 = OpConstant %uint 8
%48 = OpConstantComposite %v4uint %uint_8 %uint_8 %uint_8 %uint_8
%uint_240 = OpConstant %uint 240
%55 = OpConstantComposite %v4uint %uint_240 %uint_240 %uint_240 %uint_240
%uint_4 = OpConstant %uint 4
%58 = OpConstantComposite %v4uint %uint_4 %uint_4 %uint_4 %uint_4
%uint_12 = OpConstant %uint 12
%65 = OpConstantComposite %v4uint %uint_12 %uint_12 %uint_12 %uint_12
%uint_2 = OpConstant %uint 2
%68 = OpConstantComposite %v4uint %uint_2 %uint_2 %uint_2 %uint_2
%uint_1 = OpConstant %uint 1
%76 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1
%uint_4294967295 = OpConstant %uint 4294967295
%81 = OpConstantComposite %v4uint %uint_4294967295 %uint_4294967295 %uint_4294967295 %uint_4294967295
%void = OpTypeVoid
%88 = OpTypeFunction %void
%93 = OpConstantNull %v4int
%_ptr_Function_v4int = OpTypePointer Function %v4int
%96 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_first_leading_bit = OpFunction %v4int None %9
%v = OpFunctionParameter %v4int
%14 = OpLabel
%x = OpVariable %_ptr_Function_v4uint Function %28
%20 = OpSLessThan %v4bool %v %19
%24 = OpNot %v4int %v
%23 = OpBitcast %v4uint %24
%25 = OpBitcast %v4uint %v
%15 = OpSelect %v4uint %20 %23 %25
OpStore %x %15
%31 = OpLoad %v4uint %x
%34 = OpBitwiseAnd %v4uint %31 %33
%30 = OpINotEqual %v4bool %34 %28
%29 = OpSelect %v4uint %30 %36 %38
%39 = OpLoad %v4uint %x
%40 = OpShiftRightLogical %v4uint %39 %29
OpStore %x %40
%43 = OpLoad %v4uint %x
%46 = OpBitwiseAnd %v4uint %43 %45
%42 = OpINotEqual %v4bool %46 %28
%41 = OpSelect %v4uint %42 %48 %38
%49 = OpLoad %v4uint %x
%50 = OpShiftRightLogical %v4uint %49 %41
OpStore %x %50
%53 = OpLoad %v4uint %x
%56 = OpBitwiseAnd %v4uint %53 %55
%52 = OpINotEqual %v4bool %56 %28
%51 = OpSelect %v4uint %52 %58 %38
%59 = OpLoad %v4uint %x
%60 = OpShiftRightLogical %v4uint %59 %51
OpStore %x %60
%63 = OpLoad %v4uint %x
%66 = OpBitwiseAnd %v4uint %63 %65
%62 = OpINotEqual %v4bool %66 %28
%61 = OpSelect %v4uint %62 %68 %38
%69 = OpLoad %v4uint %x
%70 = OpShiftRightLogical %v4uint %69 %61
OpStore %x %70
%73 = OpLoad %v4uint %x
%74 = OpBitwiseAnd %v4uint %73 %68
%72 = OpINotEqual %v4bool %74 %28
%71 = OpSelect %v4uint %72 %76 %38
%78 = OpLoad %v4uint %x
%79 = OpIEqual %v4bool %78 %38
%77 = OpSelect %v4uint %79 %81 %38
%83 = OpBitwiseOr %v4uint %29 %41
%84 = OpBitwiseOr %v4uint %83 %51
%85 = OpBitwiseOr %v4uint %84 %61
%86 = OpBitwiseOr %v4uint %85 %71
%87 = OpBitwiseOr %v4uint %86 %77
%82 = OpBitcast %v4int %87
OpReturnValue %82
OpFunctionEnd
%firstLeadingBit_c1f940 = OpFunction %void None %88
%91 = OpLabel
%res = OpVariable %_ptr_Function_v4int Function %93
%92 = OpFunctionCall %v4int %tint_first_leading_bit %93
OpStore %res %92
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %96
%98 = OpLabel
%99 = OpFunctionCall %void %firstLeadingBit_c1f940
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %88
%101 = OpLabel
%102 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %102
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %88
%105 = OpLabel
%106 = OpFunctionCall %void %firstLeadingBit_c1f940
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %88
%108 = OpLabel
%109 = OpFunctionCall %void %firstLeadingBit_c1f940
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/firstLeadingBit/c1f940.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,19 @@
fn firstLeadingBit_c1f940() {
var res : vec4<i32> = firstLeadingBit(vec4<i32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_c1f940();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_c1f940();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_c1f940();
}

45
test/tint/builtins/gen/firstLeadingBit/f0779d.wgsl Normal file
View File

@ -0,0 +1,45 @@
// Copyright 2021 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
////////////////////////////////////////////////////////////////////////////////
// File generated by tools/builtin-gen
// using the template:
// test/tint/builtins/builtins.wgsl.tmpl
// and the builtin defintion file:
// src/tint/builtins.def
//
// Do not modify this file directly
////////////////////////////////////////////////////////////////////////////////
// fn firstLeadingBit(u32) -> u32
fn firstLeadingBit_f0779d() {
var res: u32 = firstLeadingBit(1u);
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_f0779d();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_f0779d();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_f0779d();
}

93
test/tint/builtins/gen/firstLeadingBit/f0779d.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
uint tint_first_leading_bit(uint v) {
uint x = v;
uint b16 = (bool((x & 4294901760u)) ? 16u : 0u);
x = (x >> b16);
uint b8 = (bool((x & 65280u)) ? 8u : 0u);
x = (x >> b8);
uint b4 = (bool((x & 240u)) ? 4u : 0u);
x = (x >> b4);
uint b2 = (bool((x & 12u)) ? 2u : 0u);
x = (x >> b2);
uint b1 = (bool((x & 2u)) ? 1u : 0u);
uint is_zero = ((x == 0u) ? 4294967295u : 0u);
return uint((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_f0779d() {
uint res = tint_first_leading_bit(1u);
}
vec4 vertex_main() {
firstLeadingBit_f0779d();
return vec4(0.0f, 0.0f, 0.0f, 0.0f);
}
void main() {
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;
uint tint_first_leading_bit(uint v) {
uint x = v;
uint b16 = (bool((x & 4294901760u)) ? 16u : 0u);
x = (x >> b16);
uint b8 = (bool((x & 65280u)) ? 8u : 0u);
x = (x >> b8);
uint b4 = (bool((x & 240u)) ? 4u : 0u);
x = (x >> b4);
uint b2 = (bool((x & 12u)) ? 2u : 0u);
x = (x >> b2);
uint b1 = (bool((x & 2u)) ? 1u : 0u);
uint is_zero = ((x == 0u) ? 4294967295u : 0u);
return uint((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_f0779d() {
uint res = tint_first_leading_bit(1u);
}
void fragment_main() {
firstLeadingBit_f0779d();
}
void main() {
fragment_main();
return;
}
#version 310 es
uint tint_first_leading_bit(uint v) {
uint x = v;
uint b16 = (bool((x & 4294901760u)) ? 16u : 0u);
x = (x >> b16);
uint b8 = (bool((x & 65280u)) ? 8u : 0u);
x = (x >> b8);
uint b4 = (bool((x & 240u)) ? 4u : 0u);
x = (x >> b4);
uint b2 = (bool((x & 12u)) ? 2u : 0u);
x = (x >> b2);
uint b1 = (bool((x & 2u)) ? 1u : 0u);
uint is_zero = ((x == 0u) ? 4294967295u : 0u);
return uint((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_f0779d() {
uint res = tint_first_leading_bit(1u);
}
void compute_main() {
firstLeadingBit_f0779d();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/firstLeadingBit/f0779d.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
uint tint_first_leading_bit(uint v) {
uint x = v;
const uint b16 = (bool((x & 4294901760u)) ? 16u : 0u);
x = (x >> b16);
const uint b8 = (bool((x & 65280u)) ? 8u : 0u);
x = (x >> b8);
const uint b4 = (bool((x & 240u)) ? 4u : 0u);
x = (x >> b4);
const uint b2 = (bool((x & 12u)) ? 2u : 0u);
x = (x >> b2);
const uint b1 = (bool((x & 2u)) ? 1u : 0u);
const uint is_zero = ((x == 0u) ? 4294967295u : 0u);
return uint((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_f0779d() {
uint res = tint_first_leading_bit(1u);
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
firstLeadingBit_f0779d();
return float4(0.0f, 0.0f, 0.0f, 0.0f);
}
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() {
firstLeadingBit_f0779d();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
firstLeadingBit_f0779d();
return;
}

48
test/tint/builtins/gen/firstLeadingBit/f0779d.wgsl.expected.msl Normal file
View File

@ -0,0 +1,48 @@
#include <metal_stdlib>
using namespace metal;
uint tint_first_leading_bit(uint v) {
uint x = v;
uint const b16 = select(0u, 16u, bool((x & 4294901760u)));
x = (x >> b16);
uint const b8 = select(0u, 8u, bool((x & 65280u)));
x = (x >> b8);
uint const b4 = select(0u, 4u, bool((x & 240u)));
x = (x >> b4);
uint const b2 = select(0u, 2u, bool((x & 12u)));
x = (x >> b2);
uint const b1 = select(0u, 1u, bool((x & 2u)));
uint const is_zero = select(0u, 4294967295u, (x == 0u));
return uint((((((b16 | b8) | b4) | b2) | b1) | is_zero));
}
void firstLeadingBit_f0779d() {
uint res = tint_first_leading_bit(1u);
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
firstLeadingBit_f0779d();
return float4();
}
vertex tint_symbol vertex_main() {
float4 const inner_result = vertex_main_inner();
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main() {
firstLeadingBit_f0779d();
return;
}
kernel void compute_main() {
firstLeadingBit_f0779d();
return;
}

129
test/tint/builtins/gen/firstLeadingBit/f0779d.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,129 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 86
; 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 %tint_first_leading_bit "tint_first_leading_bit"
OpName %v "v"
OpName %x "x"
OpName %firstLeadingBit_f0779d "firstLeadingBit_f0779d"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main"
OpDecorate %value BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%uint = OpTypeInt 32 0
%9 = OpTypeFunction %uint %uint
%_ptr_Function_uint = OpTypePointer Function %uint
%16 = OpConstantNull %uint
%bool = OpTypeBool
%uint_4294901760 = OpConstant %uint 4294901760
%uint_16 = OpConstant %uint 16
%uint_0 = OpConstant %uint 0
%uint_65280 = OpConstant %uint 65280
%uint_8 = OpConstant %uint 8
%uint_240 = OpConstant %uint 240
%uint_4 = OpConstant %uint 4
%uint_12 = OpConstant %uint 12
%uint_2 = OpConstant %uint 2
%uint_1 = OpConstant %uint 1
%uint_4294967295 = OpConstant %uint 4294967295
%void = OpTypeVoid
%66 = OpTypeFunction %void
%72 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_first_leading_bit = OpFunction %uint None %9
%v = OpFunctionParameter %uint
%13 = OpLabel
%x = OpVariable %_ptr_Function_uint Function %16
OpStore %x %v
%20 = OpLoad %uint %x
%22 = OpBitwiseAnd %uint %20 %uint_4294901760
%18 = OpINotEqual %bool %22 %16
%17 = OpSelect %uint %18 %uint_16 %uint_0
%25 = OpLoad %uint %x
%26 = OpShiftRightLogical %uint %25 %17
OpStore %x %26
%29 = OpLoad %uint %x
%31 = OpBitwiseAnd %uint %29 %uint_65280
%28 = OpINotEqual %bool %31 %16
%27 = OpSelect %uint %28 %uint_8 %uint_0
%33 = OpLoad %uint %x
%34 = OpShiftRightLogical %uint %33 %27
OpStore %x %34
%37 = OpLoad %uint %x
%39 = OpBitwiseAnd %uint %37 %uint_240
%36 = OpINotEqual %bool %39 %16
%35 = OpSelect %uint %36 %uint_4 %uint_0
%41 = OpLoad %uint %x
%42 = OpShiftRightLogical %uint %41 %35
OpStore %x %42
%45 = OpLoad %uint %x
%47 = OpBitwiseAnd %uint %45 %uint_12
%44 = OpINotEqual %bool %47 %16
%43 = OpSelect %uint %44 %uint_2 %uint_0
%49 = OpLoad %uint %x
%50 = OpShiftRightLogical %uint %49 %43
OpStore %x %50
%53 = OpLoad %uint %x
%54 = OpBitwiseAnd %uint %53 %uint_2
%52 = OpINotEqual %bool %54 %16
%51 = OpSelect %uint %52 %uint_1 %uint_0
%57 = OpLoad %uint %x
%58 = OpIEqual %bool %57 %uint_0
%56 = OpSelect %uint %58 %uint_4294967295 %uint_0
%61 = OpBitwiseOr %uint %17 %27
%62 = OpBitwiseOr %uint %61 %35
%63 = OpBitwiseOr %uint %62 %43
%64 = OpBitwiseOr %uint %63 %51
%65 = OpBitwiseOr %uint %64 %56
OpReturnValue %65
OpFunctionEnd
%firstLeadingBit_f0779d = OpFunction %void None %66
%69 = OpLabel
%res = OpVariable %_ptr_Function_uint Function %16
%70 = OpFunctionCall %uint %tint_first_leading_bit %uint_1
OpStore %res %70
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %72
%74 = OpLabel
%75 = OpFunctionCall %void %firstLeadingBit_f0779d
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %66
%77 = OpLabel
%78 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %78
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %66
%81 = OpLabel
%82 = OpFunctionCall %void %firstLeadingBit_f0779d
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %66
%84 = OpLabel
%85 = OpFunctionCall %void %firstLeadingBit_f0779d
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/firstLeadingBit/f0779d.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,19 @@
fn firstLeadingBit_f0779d() {
var res : u32 = firstLeadingBit(1u);
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
firstLeadingBit_f0779d();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
firstLeadingBit_f0779d();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
firstLeadingBit_f0779d();
}