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 countTrailingZeros 

Requires polyfilling for all but the MSL backend.

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

Bug: tint:1367
Change-Id: I0cf56b74c01f30436f9ad00595a554a4042587e4
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/81501
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
This commit is contained in:
Ben Clayton 2022-02-22 23:16:39 +00:00 committed by Tint LUCI CQ
parent 27aa57ccac
commit f8672d8c35
60 changed files with 5748 additions and 2449 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4939
src/tint/builtin_table.inl
View File

File diff suppressed because it is too large Load Diff

2
src/tint/builtins.def
View File

@ -290,6 +290,8 @@ fn countLeadingZeros<T: iu32>(T) -> T
fn countLeadingZeros<N: num, T: iu32>(vec<N, T>) -> vec<N, T> fn countLeadingZeros<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
fn countOneBits<T: iu32>(T) -> T fn countOneBits<T: iu32>(T) -> T
fn countOneBits<N: num, T: iu32>(vec<N, T>) -> vec<N, T> fn countOneBits<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
fn countTrailingZeros<T: iu32>(T) -> T
fn countTrailingZeros<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
fn cross(vec3<f32>, vec3<f32>) -> vec3<f32> fn cross(vec3<f32>, vec3<f32>) -> vec3<f32>
fn degrees(f32) -> f32 fn degrees(f32) -> f32
fn degrees<N: num>(vec<N, f32>) -> vec<N, f32> fn degrees<N: num>(vec<N, f32>) -> vec<N, f32>

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

@ -72,6 +72,9 @@ BuiltinType ParseBuiltinType(const std::string& name) {
if (name == "countOneBits") { if (name == "countOneBits") {
return BuiltinType::kCountOneBits; return BuiltinType::kCountOneBits;
} }
if (name == "countTrailingZeros") {
return BuiltinType::kCountTrailingZeros;
}
if (name == "cross") { if (name == "cross") {
return BuiltinType::kCross; return BuiltinType::kCross;
} }
@ -374,6 +377,8 @@ const char* str(BuiltinType i) {
return "countLeadingZeros"; return "countLeadingZeros";
case BuiltinType::kCountOneBits: case BuiltinType::kCountOneBits:
return "countOneBits"; return "countOneBits";
case BuiltinType::kCountTrailingZeros:
return "countTrailingZeros";
case BuiltinType::kCross: case BuiltinType::kCross:
return "cross"; return "cross";
case BuiltinType::kDegrees: case BuiltinType::kDegrees:

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

@ -48,6 +48,7 @@ enum class BuiltinType {
kCosh, kCosh,
kCountLeadingZeros, kCountLeadingZeros,
kCountOneBits, kCountOneBits,
kCountTrailingZeros,
kCross, kCross,
kDegrees, kDegrees,
kDeterminant, kDeterminant,

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

@ -108,6 +108,80 @@ struct BuiltinPolyfill::State {
return name; return name;
} }
/// Builds the polyfill function for the `countTrailingZeros` builtin
/// @param ty the parameter and return type for the function
/// @return the polyfill function name
Symbol countTrailingZeros(const sem::Type* ty) {
auto name = b.Symbols().New("tint_count_trailing_zeros");
uint32_t width = 1;
if (auto* v = ty->As<sem::Vector>()) {
width = v->Width();
}
// 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* {
if (width == 1) {
return b.Expr(value);
}
return b.Construct(b.ty.vec<ProgramBuilder::u32>(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);
};
b.Func(
name, {b.Param("v", T(ty))}, T(ty),
{
// var x = U(v);
b.Decl(b.Var("x", nullptr, b.Construct(U(), b.Expr("v")))),
// let b16 = select(16, 0, bool(x & 0x0000ffff));
b.Decl(b.Const(
"b16", nullptr,
b.Call("select", V(16), V(0), B(b.And("x", V(0x0000ffff)))))),
// x = x >> b16;
b.Assign("x", b.Shr("x", "b16")),
// let b8 = select(8, 0, bool(x & 0x000000ff));
b.Decl(b.Const(
"b8", nullptr,
b.Call("select", V(8), V(0), B(b.And("x", V(0x000000ff)))))),
// x = x >> b8;
b.Assign("x", b.Shr("x", "b8")),
// let b4 = select(4, 0, bool(x & 0x0000000f));
b.Decl(b.Const(
"b4", nullptr,
b.Call("select", V(4), V(0), B(b.And("x", V(0x0000000f)))))),
// x = x >> b4;
b.Assign("x", b.Shr("x", "b4")),
// let b2 = select(2, 0, bool(x & 0x00000003));
b.Decl(b.Const(
"b2", nullptr,
b.Call("select", V(2), V(0), B(b.And("x", V(0x00000003)))))),
// x = x >> b2;
b.Assign("x", b.Shr("x", "b2")),
// let b1 = select(1, 0, bool(x & 0x00000001));
b.Decl(b.Const(
"b1", nullptr,
b.Call("select", V(1), V(0), B(b.And("x", V(0x00000001)))))),
// let is_zero = select(0, 1, x == 0);
b.Decl(b.Const("is_zero", nullptr,
b.Call("select", V(0), V(1), b.Equal("x", V(0))))),
// return R((b16 | b8 | b4 | b2 | b1) + zero);
b.Return(b.Construct(
T(ty),
b.Add(b.Or(b.Or(b.Or(b.Or("b16", "b8"), "b4"), "b2"), "b1"),
"is_zero"))),
});
return name;
}
private: private:
const ast::Type* T(const sem::Type* ty) { return CreateASTTypeFor(ctx, ty); } const ast::Type* T(const sem::Type* ty) { return CreateASTTypeFor(ctx, ty); }
}; };
@ -130,6 +204,11 @@ bool BuiltinPolyfill::ShouldRun(const Program* program,
return true; return true;
} }
break; break;
case sem::BuiltinType::kCountTrailingZeros:
if (builtins.count_trailing_zeros) {
return true;
}
break;
default: default:
break; break;
} }
@ -166,6 +245,13 @@ void BuiltinPolyfill::Run(CloneContext& ctx,
}); });
} }
break; break;
case sem::BuiltinType::kCountTrailingZeros:
if (builtins.count_trailing_zeros) {
polyfill = utils::GetOrCreate(polyfills, builtin, [&] {
return s.countTrailingZeros(builtin->ReturnType());
});
}
break;
default: default:
break; break;
} }

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

@ -32,6 +32,8 @@ class BuiltinPolyfill : public Castable<BuiltinPolyfill, Transform> {
struct Builtins { struct Builtins {
/// Should `countLeadingZeros()` be polyfilled? /// Should `countLeadingZeros()` be polyfilled?
bool count_leading_zeros = false; bool count_leading_zeros = false;
/// Should `countTrailingZeros()` be polyfilled?
bool count_trailing_zeros = false;
}; };
/// Config is consumed by the BuiltinPolyfill transform. /// Config is consumed by the BuiltinPolyfill transform.

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

@ -194,6 +194,160 @@ fn f() {
EXPECT_EQ(expect, str(got)); EXPECT_EQ(expect, str(got));
} }
////////////////////////////////////////////////////////////////////////////////
// countTrailingZeros
////////////////////////////////////////////////////////////////////////////////
DataMap polyfillCountTrailingZeros() {
BuiltinPolyfill::Builtins builtins;
builtins.count_trailing_zeros = true;
DataMap data;
data.Add<BuiltinPolyfill::Config>(builtins);
return data;
}
TEST_F(BuiltinPolyfillTest, ShouldRunCountTrailingZeros) {
auto* src = R"(
fn f() {
countTrailingZeros(0xf);
}
)";
EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src));
EXPECT_TRUE(ShouldRun<BuiltinPolyfill>(src, polyfillCountTrailingZeros()));
}
TEST_F(BuiltinPolyfillTest, CountTrailingZeros_i32) {
auto* src = R"(
fn f() {
let r : i32 = countTrailingZeros(15);
}
)";
auto* expect = R"(
fn tint_count_trailing_zeros(v : i32) -> i32 {
var x = u32(v);
let b16 = select(16u, 0u, bool((x & 65535u)));
x = (x >> b16);
let b8 = select(8u, 0u, bool((x & 255u)));
x = (x >> b8);
let b4 = select(4u, 0u, bool((x & 15u)));
x = (x >> b4);
let b2 = select(2u, 0u, bool((x & 3u)));
x = (x >> b2);
let b1 = select(1u, 0u, bool((x & 1u)));
let is_zero = select(0u, 1u, (x == 0u));
return i32((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
fn f() {
let r : i32 = tint_count_trailing_zeros(15);
}
)";
auto got = Run<BuiltinPolyfill>(src, polyfillCountTrailingZeros());
EXPECT_EQ(expect, str(got));
}
TEST_F(BuiltinPolyfillTest, CountTrailingZeros_u32) {
auto* src = R"(
fn f() {
let r : u32 = countTrailingZeros(15u);
}
)";
auto* expect = R"(
fn tint_count_trailing_zeros(v : u32) -> u32 {
var x = u32(v);
let b16 = select(16u, 0u, bool((x & 65535u)));
x = (x >> b16);
let b8 = select(8u, 0u, bool((x & 255u)));
x = (x >> b8);
let b4 = select(4u, 0u, bool((x & 15u)));
x = (x >> b4);
let b2 = select(2u, 0u, bool((x & 3u)));
x = (x >> b2);
let b1 = select(1u, 0u, bool((x & 1u)));
let is_zero = select(0u, 1u, (x == 0u));
return u32((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
fn f() {
let r : u32 = tint_count_trailing_zeros(15u);
}
)";
auto got = Run<BuiltinPolyfill>(src, polyfillCountTrailingZeros());
EXPECT_EQ(expect, str(got));
}
TEST_F(BuiltinPolyfillTest, CountTrailingZeros_vec3_i32) {
auto* src = R"(
fn f() {
let r : vec3<i32> = countTrailingZeros(vec3<i32>(15));
}
)";
auto* expect = R"(
fn tint_count_trailing_zeros(v : vec3<i32>) -> vec3<i32> {
var x = vec3<u32>(v);
let b16 = select(vec3<u32>(16u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(65535u))));
x = (x >> b16);
let b8 = select(vec3<u32>(8u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(255u))));
x = (x >> b8);
let b4 = select(vec3<u32>(4u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(15u))));
x = (x >> b4);
let b2 = select(vec3<u32>(2u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(3u))));
x = (x >> b2);
let b1 = select(vec3<u32>(1u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(1u))));
let is_zero = select(vec3<u32>(0u), vec3<u32>(1u), (x == vec3<u32>(0u)));
return vec3<i32>((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
fn f() {
let r : vec3<i32> = tint_count_trailing_zeros(vec3<i32>(15));
}
)";
auto got = Run<BuiltinPolyfill>(src, polyfillCountTrailingZeros());
EXPECT_EQ(expect, str(got));
}
TEST_F(BuiltinPolyfillTest, CountTrailingZeros_vec3_u32) {
auto* src = R"(
fn f() {
let r : vec3<u32> = countTrailingZeros(vec3<u32>(15u));
}
)";
auto* expect = R"(
fn tint_count_trailing_zeros(v : vec3<u32>) -> vec3<u32> {
var x = vec3<u32>(v);
let b16 = select(vec3<u32>(16u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(65535u))));
x = (x >> b16);
let b8 = select(vec3<u32>(8u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(255u))));
x = (x >> b8);
let b4 = select(vec3<u32>(4u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(15u))));
x = (x >> b4);
let b2 = select(vec3<u32>(2u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(3u))));
x = (x >> b2);
let b1 = select(vec3<u32>(1u), vec3<u32>(0u), vec3<bool>((x & vec3<u32>(1u))));
let is_zero = select(vec3<u32>(0u), vec3<u32>(1u), (x == vec3<u32>(0u)));
return vec3<u32>((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
fn f() {
let r : vec3<u32> = tint_count_trailing_zeros(vec3<u32>(15u));
}
)";
auto got = Run<BuiltinPolyfill>(src, polyfillCountTrailingZeros());
EXPECT_EQ(expect, str(got));
}
} // namespace } // namespace
} // namespace transform } // namespace transform
} // namespace tint } // namespace tint

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

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

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

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

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

@ -1351,6 +1351,9 @@ std::string GeneratorImpl::generate_builtin_name(const sem::Builtin* builtin) {
case sem::BuiltinType::kCountOneBits: case sem::BuiltinType::kCountOneBits:
out += "popcount"; out += "popcount";
break; break;
case sem::BuiltinType::kCountTrailingZeros:
out += "ctz";
break;
case sem::BuiltinType::kDpdx: case sem::BuiltinType::kDpdx:
case sem::BuiltinType::kDpdxCoarse: case sem::BuiltinType::kDpdxCoarse:
case sem::BuiltinType::kDpdxFine: case sem::BuiltinType::kDpdxFine:

2
src/tint/writer/msl/generator_impl_builtin_test.cc
View File

@ -126,6 +126,7 @@ const ast::CallExpression* GenerateCall(BuiltinType builtin,
} }
case BuiltinType::kCountLeadingZeros: case BuiltinType::kCountLeadingZeros:
case BuiltinType::kCountOneBits: case BuiltinType::kCountOneBits:
case BuiltinType::kCountTrailingZeros:
case BuiltinType::kReverseBits: case BuiltinType::kReverseBits:
return builder->Call(str.str(), "u2"); return builder->Call(str.str(), "u2");
case BuiltinType::kMax: case BuiltinType::kMax:
@ -215,6 +216,7 @@ INSTANTIATE_TEST_SUITE_P(
BuiltinData{BuiltinType::kCosh, ParamType::kF32, "cosh"}, BuiltinData{BuiltinType::kCosh, ParamType::kF32, "cosh"},
BuiltinData{BuiltinType::kCountLeadingZeros, ParamType::kU32, "clz"}, BuiltinData{BuiltinType::kCountLeadingZeros, ParamType::kU32, "clz"},
BuiltinData{BuiltinType::kCountOneBits, ParamType::kU32, "popcount"}, BuiltinData{BuiltinType::kCountOneBits, ParamType::kU32, "popcount"},
BuiltinData{BuiltinType::kCountTrailingZeros, ParamType::kU32, "ctz"},
BuiltinData{BuiltinType::kCross, ParamType::kF32, "cross"}, BuiltinData{BuiltinType::kCross, ParamType::kF32, "cross"},
BuiltinData{BuiltinType::kDeterminant, ParamType::kF32, "determinant"}, BuiltinData{BuiltinType::kDeterminant, ParamType::kF32, "determinant"},
BuiltinData{BuiltinType::kDistance, ParamType::kF32, "distance"}, BuiltinData{BuiltinType::kDistance, ParamType::kF32, "distance"},

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

@ -262,6 +262,7 @@ SanitizedResult Sanitize(const Program* in,
{ // Builtin polyfills { // Builtin polyfills
transform::BuiltinPolyfill::Builtins polyfills; transform::BuiltinPolyfill::Builtins polyfills;
polyfills.count_leading_zeros = true; polyfills.count_leading_zeros = true;
polyfills.count_trailing_zeros = true;
data.Add<transform::BuiltinPolyfill::Config>(polyfills); data.Add<transform::BuiltinPolyfill::Config>(polyfills);
manager.Add<transform::BuiltinPolyfill>(); manager.Add<transform::BuiltinPolyfill>();
} }

45
test/tint/builtins/gen/countTrailingZeros/1ad138.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 countTrailingZeros(vec<2, u32>) -> vec<2, u32>
fn countTrailingZeros_1ad138() {
var res: vec2<u32> = countTrailingZeros(vec2<u32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
countTrailingZeros_1ad138();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
countTrailingZeros_1ad138();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
countTrailingZeros_1ad138();
}

93
test/tint/builtins/gen/countTrailingZeros/1ad138.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
uvec2 tint_count_trailing_zeros(uvec2 v) {
uvec2 x = uvec2(v);
uvec2 b16 = mix(uvec2(16u), uvec2(0u), bvec2((x & uvec2(65535u))));
x = (x >> b16);
uvec2 b8 = mix(uvec2(8u), uvec2(0u), bvec2((x & uvec2(255u))));
x = (x >> b8);
uvec2 b4 = mix(uvec2(4u), uvec2(0u), bvec2((x & uvec2(15u))));
x = (x >> b4);
uvec2 b2 = mix(uvec2(2u), uvec2(0u), bvec2((x & uvec2(3u))));
x = (x >> b2);
uvec2 b1 = mix(uvec2(1u), uvec2(0u), bvec2((x & uvec2(1u))));
uvec2 is_zero = mix(uvec2(0u), uvec2(1u), equal(x, uvec2(0u)));
return uvec2((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_1ad138() {
uvec2 res = tint_count_trailing_zeros(uvec2(0u, 0u));
}
vec4 vertex_main() {
countTrailingZeros_1ad138();
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_count_trailing_zeros(uvec2 v) {
uvec2 x = uvec2(v);
uvec2 b16 = mix(uvec2(16u), uvec2(0u), bvec2((x & uvec2(65535u))));
x = (x >> b16);
uvec2 b8 = mix(uvec2(8u), uvec2(0u), bvec2((x & uvec2(255u))));
x = (x >> b8);
uvec2 b4 = mix(uvec2(4u), uvec2(0u), bvec2((x & uvec2(15u))));
x = (x >> b4);
uvec2 b2 = mix(uvec2(2u), uvec2(0u), bvec2((x & uvec2(3u))));
x = (x >> b2);
uvec2 b1 = mix(uvec2(1u), uvec2(0u), bvec2((x & uvec2(1u))));
uvec2 is_zero = mix(uvec2(0u), uvec2(1u), equal(x, uvec2(0u)));
return uvec2((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_1ad138() {
uvec2 res = tint_count_trailing_zeros(uvec2(0u, 0u));
}
void fragment_main() {
countTrailingZeros_1ad138();
}
void main() {
fragment_main();
return;
}
#version 310 es
uvec2 tint_count_trailing_zeros(uvec2 v) {
uvec2 x = uvec2(v);
uvec2 b16 = mix(uvec2(16u), uvec2(0u), bvec2((x & uvec2(65535u))));
x = (x >> b16);
uvec2 b8 = mix(uvec2(8u), uvec2(0u), bvec2((x & uvec2(255u))));
x = (x >> b8);
uvec2 b4 = mix(uvec2(4u), uvec2(0u), bvec2((x & uvec2(15u))));
x = (x >> b4);
uvec2 b2 = mix(uvec2(2u), uvec2(0u), bvec2((x & uvec2(3u))));
x = (x >> b2);
uvec2 b1 = mix(uvec2(1u), uvec2(0u), bvec2((x & uvec2(1u))));
uvec2 is_zero = mix(uvec2(0u), uvec2(1u), equal(x, uvec2(0u)));
return uvec2((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_1ad138() {
uvec2 res = tint_count_trailing_zeros(uvec2(0u, 0u));
}
void compute_main() {
countTrailingZeros_1ad138();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/countTrailingZeros/1ad138.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
uint2 tint_count_trailing_zeros(uint2 v) {
uint2 x = uint2(v);
const uint2 b16 = (bool2((x & uint2((65535u).xx))) ? uint2((0u).xx) : uint2((16u).xx));
x = (x >> b16);
const uint2 b8 = (bool2((x & uint2((255u).xx))) ? uint2((0u).xx) : uint2((8u).xx));
x = (x >> b8);
const uint2 b4 = (bool2((x & uint2((15u).xx))) ? uint2((0u).xx) : uint2((4u).xx));
x = (x >> b4);
const uint2 b2 = (bool2((x & uint2((3u).xx))) ? uint2((0u).xx) : uint2((2u).xx));
x = (x >> b2);
const uint2 b1 = (bool2((x & uint2((1u).xx))) ? uint2((0u).xx) : uint2((1u).xx));
const uint2 is_zero = ((x == uint2((0u).xx)) ? uint2((1u).xx) : uint2((0u).xx));
return uint2((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_1ad138() {
uint2 res = tint_count_trailing_zeros(uint2(0u, 0u));
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
countTrailingZeros_1ad138();
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() {
countTrailingZeros_1ad138();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
countTrailingZeros_1ad138();
return;
}

33
test/tint/builtins/gen/countTrailingZeros/1ad138.wgsl.expected.msl Normal file
View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
void countTrailingZeros_1ad138() {
uint2 res = ctz(uint2());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
countTrailingZeros_1ad138();
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() {
countTrailingZeros_1ad138();
return;
}
kernel void compute_main() {
countTrailingZeros_1ad138();
return;
}

140
test/tint/builtins/gen/countTrailingZeros/1ad138.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,140 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 98
; 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_count_trailing_zeros "tint_count_trailing_zeros"
OpName %v "v"
OpName %x "x"
OpName %countTrailingZeros_1ad138 "countTrailingZeros_1ad138"
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
%18 = OpConstantNull %v2uint
%bool = OpTypeBool
%v2bool = OpTypeVector %bool 2
%uint_65535 = OpConstant %uint 65535
%25 = OpConstantComposite %v2uint %uint_65535 %uint_65535
%uint_0 = OpConstant %uint 0
%28 = OpConstantComposite %v2uint %uint_0 %uint_0
%uint_16 = OpConstant %uint 16
%30 = OpConstantComposite %v2uint %uint_16 %uint_16
%uint_255 = OpConstant %uint 255
%37 = OpConstantComposite %v2uint %uint_255 %uint_255
%uint_8 = OpConstant %uint 8
%40 = OpConstantComposite %v2uint %uint_8 %uint_8
%uint_15 = OpConstant %uint 15
%47 = OpConstantComposite %v2uint %uint_15 %uint_15
%uint_4 = OpConstant %uint 4
%50 = OpConstantComposite %v2uint %uint_4 %uint_4
%uint_3 = OpConstant %uint 3
%57 = OpConstantComposite %v2uint %uint_3 %uint_3
%uint_2 = OpConstant %uint 2
%60 = OpConstantComposite %v2uint %uint_2 %uint_2
%uint_1 = OpConstant %uint 1
%67 = OpConstantComposite %v2uint %uint_1 %uint_1
%void = OpTypeVoid
%78 = OpTypeFunction %void
%84 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_count_trailing_zeros = OpFunction %v2uint None %9
%v = OpFunctionParameter %v2uint
%14 = OpLabel
%x = OpVariable %_ptr_Function_v2uint Function %18
OpStore %x %v
%23 = OpLoad %v2uint %x
%26 = OpBitwiseAnd %v2uint %23 %25
%20 = OpINotEqual %v2bool %26 %18
%19 = OpSelect %v2uint %20 %28 %30
%31 = OpLoad %v2uint %x
%32 = OpShiftRightLogical %v2uint %31 %19
OpStore %x %32
%35 = OpLoad %v2uint %x
%38 = OpBitwiseAnd %v2uint %35 %37
%34 = OpINotEqual %v2bool %38 %18
%33 = OpSelect %v2uint %34 %28 %40
%41 = OpLoad %v2uint %x
%42 = OpShiftRightLogical %v2uint %41 %33
OpStore %x %42
%45 = OpLoad %v2uint %x
%48 = OpBitwiseAnd %v2uint %45 %47
%44 = OpINotEqual %v2bool %48 %18
%43 = OpSelect %v2uint %44 %28 %50
%51 = OpLoad %v2uint %x
%52 = OpShiftRightLogical %v2uint %51 %43
OpStore %x %52
%55 = OpLoad %v2uint %x
%58 = OpBitwiseAnd %v2uint %55 %57
%54 = OpINotEqual %v2bool %58 %18
%53 = OpSelect %v2uint %54 %28 %60
%61 = OpLoad %v2uint %x
%62 = OpShiftRightLogical %v2uint %61 %53
OpStore %x %62
%65 = OpLoad %v2uint %x
%68 = OpBitwiseAnd %v2uint %65 %67
%64 = OpINotEqual %v2bool %68 %18
%63 = OpSelect %v2uint %64 %28 %67
%70 = OpLoad %v2uint %x
%71 = OpIEqual %v2bool %70 %28
%69 = OpSelect %v2uint %71 %67 %28
%73 = OpBitwiseOr %v2uint %19 %33
%74 = OpBitwiseOr %v2uint %73 %43
%75 = OpBitwiseOr %v2uint %74 %53
%76 = OpBitwiseOr %v2uint %75 %63
%77 = OpIAdd %v2uint %76 %69
OpReturnValue %77
OpFunctionEnd
%countTrailingZeros_1ad138 = OpFunction %void None %78
%81 = OpLabel
%res = OpVariable %_ptr_Function_v2uint Function %18
%82 = OpFunctionCall %v2uint %tint_count_trailing_zeros %18
OpStore %res %82
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %84
%86 = OpLabel
%87 = OpFunctionCall %void %countTrailingZeros_1ad138
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %78
%89 = OpLabel
%90 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %90
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %78
%93 = OpLabel
%94 = OpFunctionCall %void %countTrailingZeros_1ad138
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %78
%96 = OpLabel
%97 = OpFunctionCall %void %countTrailingZeros_1ad138
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/countTrailingZeros/1ad138.wgsl.expected.wgsl Normal file
View File

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

45
test/tint/builtins/gen/countTrailingZeros/1dc84a.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 countTrailingZeros(vec<4, i32>) -> vec<4, i32>
fn countTrailingZeros_1dc84a() {
var res: vec4<i32> = countTrailingZeros(vec4<i32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
countTrailingZeros_1dc84a();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
countTrailingZeros_1dc84a();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
countTrailingZeros_1dc84a();
}

93
test/tint/builtins/gen/countTrailingZeros/1dc84a.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
ivec4 tint_count_trailing_zeros(ivec4 v) {
uvec4 x = uvec4(v);
uvec4 b16 = mix(uvec4(16u), uvec4(0u), bvec4((x & uvec4(65535u))));
x = (x >> b16);
uvec4 b8 = mix(uvec4(8u), uvec4(0u), bvec4((x & uvec4(255u))));
x = (x >> b8);
uvec4 b4 = mix(uvec4(4u), uvec4(0u), bvec4((x & uvec4(15u))));
x = (x >> b4);
uvec4 b2 = mix(uvec4(2u), uvec4(0u), bvec4((x & uvec4(3u))));
x = (x >> b2);
uvec4 b1 = mix(uvec4(1u), uvec4(0u), bvec4((x & uvec4(1u))));
uvec4 is_zero = mix(uvec4(0u), uvec4(1u), equal(x, uvec4(0u)));
return ivec4((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_1dc84a() {
ivec4 res = tint_count_trailing_zeros(ivec4(0, 0, 0, 0));
}
vec4 vertex_main() {
countTrailingZeros_1dc84a();
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_count_trailing_zeros(ivec4 v) {
uvec4 x = uvec4(v);
uvec4 b16 = mix(uvec4(16u), uvec4(0u), bvec4((x & uvec4(65535u))));
x = (x >> b16);
uvec4 b8 = mix(uvec4(8u), uvec4(0u), bvec4((x & uvec4(255u))));
x = (x >> b8);
uvec4 b4 = mix(uvec4(4u), uvec4(0u), bvec4((x & uvec4(15u))));
x = (x >> b4);
uvec4 b2 = mix(uvec4(2u), uvec4(0u), bvec4((x & uvec4(3u))));
x = (x >> b2);
uvec4 b1 = mix(uvec4(1u), uvec4(0u), bvec4((x & uvec4(1u))));
uvec4 is_zero = mix(uvec4(0u), uvec4(1u), equal(x, uvec4(0u)));
return ivec4((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_1dc84a() {
ivec4 res = tint_count_trailing_zeros(ivec4(0, 0, 0, 0));
}
void fragment_main() {
countTrailingZeros_1dc84a();
}
void main() {
fragment_main();
return;
}
#version 310 es
ivec4 tint_count_trailing_zeros(ivec4 v) {
uvec4 x = uvec4(v);
uvec4 b16 = mix(uvec4(16u), uvec4(0u), bvec4((x & uvec4(65535u))));
x = (x >> b16);
uvec4 b8 = mix(uvec4(8u), uvec4(0u), bvec4((x & uvec4(255u))));
x = (x >> b8);
uvec4 b4 = mix(uvec4(4u), uvec4(0u), bvec4((x & uvec4(15u))));
x = (x >> b4);
uvec4 b2 = mix(uvec4(2u), uvec4(0u), bvec4((x & uvec4(3u))));
x = (x >> b2);
uvec4 b1 = mix(uvec4(1u), uvec4(0u), bvec4((x & uvec4(1u))));
uvec4 is_zero = mix(uvec4(0u), uvec4(1u), equal(x, uvec4(0u)));
return ivec4((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_1dc84a() {
ivec4 res = tint_count_trailing_zeros(ivec4(0, 0, 0, 0));
}
void compute_main() {
countTrailingZeros_1dc84a();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/countTrailingZeros/1dc84a.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
int4 tint_count_trailing_zeros(int4 v) {
uint4 x = uint4(v);
const uint4 b16 = (bool4((x & uint4((65535u).xxxx))) ? uint4((0u).xxxx) : uint4((16u).xxxx));
x = (x >> b16);
const uint4 b8 = (bool4((x & uint4((255u).xxxx))) ? uint4((0u).xxxx) : uint4((8u).xxxx));
x = (x >> b8);
const uint4 b4 = (bool4((x & uint4((15u).xxxx))) ? uint4((0u).xxxx) : uint4((4u).xxxx));
x = (x >> b4);
const uint4 b2 = (bool4((x & uint4((3u).xxxx))) ? uint4((0u).xxxx) : uint4((2u).xxxx));
x = (x >> b2);
const uint4 b1 = (bool4((x & uint4((1u).xxxx))) ? uint4((0u).xxxx) : uint4((1u).xxxx));
const uint4 is_zero = ((x == uint4((0u).xxxx)) ? uint4((1u).xxxx) : uint4((0u).xxxx));
return int4((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_1dc84a() {
int4 res = tint_count_trailing_zeros(int4(0, 0, 0, 0));
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
countTrailingZeros_1dc84a();
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() {
countTrailingZeros_1dc84a();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
countTrailingZeros_1dc84a();
return;
}

33
test/tint/builtins/gen/countTrailingZeros/1dc84a.wgsl.expected.msl Normal file
View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
void countTrailingZeros_1dc84a() {
int4 res = ctz(int4());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
countTrailingZeros_1dc84a();
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() {
countTrailingZeros_1dc84a();
return;
}
kernel void compute_main() {
countTrailingZeros_1dc84a();
return;
}

146
test/tint/builtins/gen/countTrailingZeros/1dc84a.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,146 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 102
; 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_count_trailing_zeros "tint_count_trailing_zeros"
OpName %v "v"
OpName %x "x"
OpName %countTrailingZeros_1dc84a "countTrailingZeros_1dc84a"
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
%_ptr_Function_v4uint = OpTypePointer Function %v4uint
%20 = OpConstantNull %v4uint
%bool = OpTypeBool
%v4bool = OpTypeVector %bool 4
%uint_65535 = OpConstant %uint 65535
%27 = OpConstantComposite %v4uint %uint_65535 %uint_65535 %uint_65535 %uint_65535
%uint_0 = OpConstant %uint 0
%30 = OpConstantComposite %v4uint %uint_0 %uint_0 %uint_0 %uint_0
%uint_16 = OpConstant %uint 16
%32 = OpConstantComposite %v4uint %uint_16 %uint_16 %uint_16 %uint_16
%uint_255 = OpConstant %uint 255
%39 = OpConstantComposite %v4uint %uint_255 %uint_255 %uint_255 %uint_255
%uint_8 = OpConstant %uint 8
%42 = OpConstantComposite %v4uint %uint_8 %uint_8 %uint_8 %uint_8
%uint_15 = OpConstant %uint 15
%49 = OpConstantComposite %v4uint %uint_15 %uint_15 %uint_15 %uint_15
%uint_4 = OpConstant %uint 4
%52 = OpConstantComposite %v4uint %uint_4 %uint_4 %uint_4 %uint_4
%uint_3 = OpConstant %uint 3
%59 = OpConstantComposite %v4uint %uint_3 %uint_3 %uint_3 %uint_3
%uint_2 = OpConstant %uint 2
%62 = OpConstantComposite %v4uint %uint_2 %uint_2 %uint_2 %uint_2
%uint_1 = OpConstant %uint 1
%69 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1
%void = OpTypeVoid
%80 = OpTypeFunction %void
%85 = OpConstantNull %v4int
%_ptr_Function_v4int = OpTypePointer Function %v4int
%88 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_count_trailing_zeros = OpFunction %v4int None %9
%v = OpFunctionParameter %v4int
%14 = OpLabel
%x = OpVariable %_ptr_Function_v4uint Function %20
%15 = OpBitcast %v4uint %v
OpStore %x %15
%25 = OpLoad %v4uint %x
%28 = OpBitwiseAnd %v4uint %25 %27
%22 = OpINotEqual %v4bool %28 %20
%21 = OpSelect %v4uint %22 %30 %32
%33 = OpLoad %v4uint %x
%34 = OpShiftRightLogical %v4uint %33 %21
OpStore %x %34
%37 = OpLoad %v4uint %x
%40 = OpBitwiseAnd %v4uint %37 %39
%36 = OpINotEqual %v4bool %40 %20
%35 = OpSelect %v4uint %36 %30 %42
%43 = OpLoad %v4uint %x
%44 = OpShiftRightLogical %v4uint %43 %35
OpStore %x %44
%47 = OpLoad %v4uint %x
%50 = OpBitwiseAnd %v4uint %47 %49
%46 = OpINotEqual %v4bool %50 %20
%45 = OpSelect %v4uint %46 %30 %52
%53 = OpLoad %v4uint %x
%54 = OpShiftRightLogical %v4uint %53 %45
OpStore %x %54
%57 = OpLoad %v4uint %x
%60 = OpBitwiseAnd %v4uint %57 %59
%56 = OpINotEqual %v4bool %60 %20
%55 = OpSelect %v4uint %56 %30 %62
%63 = OpLoad %v4uint %x
%64 = OpShiftRightLogical %v4uint %63 %55
OpStore %x %64
%67 = OpLoad %v4uint %x
%70 = OpBitwiseAnd %v4uint %67 %69
%66 = OpINotEqual %v4bool %70 %20
%65 = OpSelect %v4uint %66 %30 %69
%72 = OpLoad %v4uint %x
%73 = OpIEqual %v4bool %72 %30
%71 = OpSelect %v4uint %73 %69 %30
%75 = OpBitwiseOr %v4uint %21 %35
%76 = OpBitwiseOr %v4uint %75 %45
%77 = OpBitwiseOr %v4uint %76 %55
%78 = OpBitwiseOr %v4uint %77 %65
%79 = OpIAdd %v4uint %78 %71
%74 = OpBitcast %v4int %79
OpReturnValue %74
OpFunctionEnd
%countTrailingZeros_1dc84a = OpFunction %void None %80
%83 = OpLabel
%res = OpVariable %_ptr_Function_v4int Function %85
%84 = OpFunctionCall %v4int %tint_count_trailing_zeros %85
OpStore %res %84
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %88
%90 = OpLabel
%91 = OpFunctionCall %void %countTrailingZeros_1dc84a
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %80
%93 = OpLabel
%94 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %94
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %80
%97 = OpLabel
%98 = OpFunctionCall %void %countTrailingZeros_1dc84a
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %80
%100 = OpLabel
%101 = OpFunctionCall %void %countTrailingZeros_1dc84a
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/countTrailingZeros/1dc84a.wgsl.expected.wgsl Normal file
View File

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

45
test/tint/builtins/gen/countTrailingZeros/21e394.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 countTrailingZeros(u32) -> u32
fn countTrailingZeros_21e394() {
var res: u32 = countTrailingZeros(1u);
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
countTrailingZeros_21e394();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
countTrailingZeros_21e394();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
countTrailingZeros_21e394();
}

93
test/tint/builtins/gen/countTrailingZeros/21e394.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
uint tint_count_trailing_zeros(uint v) {
uint x = uint(v);
uint b16 = (bool((x & 65535u)) ? 0u : 16u);
x = (x >> b16);
uint b8 = (bool((x & 255u)) ? 0u : 8u);
x = (x >> b8);
uint b4 = (bool((x & 15u)) ? 0u : 4u);
x = (x >> b4);
uint b2 = (bool((x & 3u)) ? 0u : 2u);
x = (x >> b2);
uint b1 = (bool((x & 1u)) ? 0u : 1u);
uint is_zero = ((x == 0u) ? 1u : 0u);
return uint((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_21e394() {
uint res = tint_count_trailing_zeros(1u);
}
vec4 vertex_main() {
countTrailingZeros_21e394();
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_count_trailing_zeros(uint v) {
uint x = uint(v);
uint b16 = (bool((x & 65535u)) ? 0u : 16u);
x = (x >> b16);
uint b8 = (bool((x & 255u)) ? 0u : 8u);
x = (x >> b8);
uint b4 = (bool((x & 15u)) ? 0u : 4u);
x = (x >> b4);
uint b2 = (bool((x & 3u)) ? 0u : 2u);
x = (x >> b2);
uint b1 = (bool((x & 1u)) ? 0u : 1u);
uint is_zero = ((x == 0u) ? 1u : 0u);
return uint((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_21e394() {
uint res = tint_count_trailing_zeros(1u);
}
void fragment_main() {
countTrailingZeros_21e394();
}
void main() {
fragment_main();
return;
}
#version 310 es
uint tint_count_trailing_zeros(uint v) {
uint x = uint(v);
uint b16 = (bool((x & 65535u)) ? 0u : 16u);
x = (x >> b16);
uint b8 = (bool((x & 255u)) ? 0u : 8u);
x = (x >> b8);
uint b4 = (bool((x & 15u)) ? 0u : 4u);
x = (x >> b4);
uint b2 = (bool((x & 3u)) ? 0u : 2u);
x = (x >> b2);
uint b1 = (bool((x & 1u)) ? 0u : 1u);
uint is_zero = ((x == 0u) ? 1u : 0u);
return uint((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_21e394() {
uint res = tint_count_trailing_zeros(1u);
}
void compute_main() {
countTrailingZeros_21e394();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/countTrailingZeros/21e394.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
uint tint_count_trailing_zeros(uint v) {
uint x = uint(v);
const uint b16 = (bool((x & 65535u)) ? 0u : 16u);
x = (x >> b16);
const uint b8 = (bool((x & 255u)) ? 0u : 8u);
x = (x >> b8);
const uint b4 = (bool((x & 15u)) ? 0u : 4u);
x = (x >> b4);
const uint b2 = (bool((x & 3u)) ? 0u : 2u);
x = (x >> b2);
const uint b1 = (bool((x & 1u)) ? 0u : 1u);
const uint is_zero = ((x == 0u) ? 1u : 0u);
return uint((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_21e394() {
uint res = tint_count_trailing_zeros(1u);
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
countTrailingZeros_21e394();
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() {
countTrailingZeros_21e394();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
countTrailingZeros_21e394();
return;
}

33
test/tint/builtins/gen/countTrailingZeros/21e394.wgsl.expected.msl Normal file
View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
void countTrailingZeros_21e394() {
uint res = ctz(1u);
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
countTrailingZeros_21e394();
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() {
countTrailingZeros_21e394();
return;
}
kernel void compute_main() {
countTrailingZeros_21e394();
return;
}

128
test/tint/builtins/gen/countTrailingZeros/21e394.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,128 @@
; 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_count_trailing_zeros "tint_count_trailing_zeros"
OpName %v "v"
OpName %x "x"
OpName %countTrailingZeros_21e394 "countTrailingZeros_21e394"
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
%17 = OpConstantNull %uint
%bool = OpTypeBool
%uint_65535 = OpConstant %uint 65535
%uint_0 = OpConstant %uint 0
%uint_16 = OpConstant %uint 16
%uint_255 = OpConstant %uint 255
%uint_8 = OpConstant %uint 8
%uint_15 = OpConstant %uint 15
%uint_4 = OpConstant %uint 4
%uint_3 = OpConstant %uint 3
%uint_2 = OpConstant %uint 2
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
%66 = OpTypeFunction %void
%72 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_count_trailing_zeros = OpFunction %uint None %9
%v = OpFunctionParameter %uint
%13 = OpLabel
%x = OpVariable %_ptr_Function_uint Function %17
OpStore %x %v
%21 = OpLoad %uint %x
%23 = OpBitwiseAnd %uint %21 %uint_65535
%19 = OpINotEqual %bool %23 %17
%18 = OpSelect %uint %19 %uint_0 %uint_16
%26 = OpLoad %uint %x
%27 = OpShiftRightLogical %uint %26 %18
OpStore %x %27
%30 = OpLoad %uint %x
%32 = OpBitwiseAnd %uint %30 %uint_255
%29 = OpINotEqual %bool %32 %17
%28 = OpSelect %uint %29 %uint_0 %uint_8
%34 = OpLoad %uint %x
%35 = OpShiftRightLogical %uint %34 %28
OpStore %x %35
%38 = OpLoad %uint %x
%40 = OpBitwiseAnd %uint %38 %uint_15
%37 = OpINotEqual %bool %40 %17
%36 = OpSelect %uint %37 %uint_0 %uint_4
%42 = OpLoad %uint %x
%43 = OpShiftRightLogical %uint %42 %36
OpStore %x %43
%46 = OpLoad %uint %x
%48 = OpBitwiseAnd %uint %46 %uint_3
%45 = OpINotEqual %bool %48 %17
%44 = OpSelect %uint %45 %uint_0 %uint_2
%50 = OpLoad %uint %x
%51 = OpShiftRightLogical %uint %50 %44
OpStore %x %51
%54 = OpLoad %uint %x
%56 = OpBitwiseAnd %uint %54 %uint_1
%53 = OpINotEqual %bool %56 %17
%52 = OpSelect %uint %53 %uint_0 %uint_1
%58 = OpLoad %uint %x
%59 = OpIEqual %bool %58 %uint_0
%57 = OpSelect %uint %59 %uint_1 %uint_0
%61 = OpBitwiseOr %uint %18 %28
%62 = OpBitwiseOr %uint %61 %36
%63 = OpBitwiseOr %uint %62 %44
%64 = OpBitwiseOr %uint %63 %52
%65 = OpIAdd %uint %64 %57
OpReturnValue %65
OpFunctionEnd
%countTrailingZeros_21e394 = OpFunction %void None %66
%69 = OpLabel
%res = OpVariable %_ptr_Function_uint Function %17
%70 = OpFunctionCall %uint %tint_count_trailing_zeros %uint_1
OpStore %res %70
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %72
%74 = OpLabel
%75 = OpFunctionCall %void %countTrailingZeros_21e394
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 %countTrailingZeros_21e394
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %66
%84 = OpLabel
%85 = OpFunctionCall %void %countTrailingZeros_21e394
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/countTrailingZeros/21e394.wgsl.expected.wgsl Normal file
View File

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

45
test/tint/builtins/gen/countTrailingZeros/327c37.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 countTrailingZeros(vec<2, i32>) -> vec<2, i32>
fn countTrailingZeros_327c37() {
var res: vec2<i32> = countTrailingZeros(vec2<i32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
countTrailingZeros_327c37();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
countTrailingZeros_327c37();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
countTrailingZeros_327c37();
}

93
test/tint/builtins/gen/countTrailingZeros/327c37.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
ivec2 tint_count_trailing_zeros(ivec2 v) {
uvec2 x = uvec2(v);
uvec2 b16 = mix(uvec2(16u), uvec2(0u), bvec2((x & uvec2(65535u))));
x = (x >> b16);
uvec2 b8 = mix(uvec2(8u), uvec2(0u), bvec2((x & uvec2(255u))));
x = (x >> b8);
uvec2 b4 = mix(uvec2(4u), uvec2(0u), bvec2((x & uvec2(15u))));
x = (x >> b4);
uvec2 b2 = mix(uvec2(2u), uvec2(0u), bvec2((x & uvec2(3u))));
x = (x >> b2);
uvec2 b1 = mix(uvec2(1u), uvec2(0u), bvec2((x & uvec2(1u))));
uvec2 is_zero = mix(uvec2(0u), uvec2(1u), equal(x, uvec2(0u)));
return ivec2((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_327c37() {
ivec2 res = tint_count_trailing_zeros(ivec2(0, 0));
}
vec4 vertex_main() {
countTrailingZeros_327c37();
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_count_trailing_zeros(ivec2 v) {
uvec2 x = uvec2(v);
uvec2 b16 = mix(uvec2(16u), uvec2(0u), bvec2((x & uvec2(65535u))));
x = (x >> b16);
uvec2 b8 = mix(uvec2(8u), uvec2(0u), bvec2((x & uvec2(255u))));
x = (x >> b8);
uvec2 b4 = mix(uvec2(4u), uvec2(0u), bvec2((x & uvec2(15u))));
x = (x >> b4);
uvec2 b2 = mix(uvec2(2u), uvec2(0u), bvec2((x & uvec2(3u))));
x = (x >> b2);
uvec2 b1 = mix(uvec2(1u), uvec2(0u), bvec2((x & uvec2(1u))));
uvec2 is_zero = mix(uvec2(0u), uvec2(1u), equal(x, uvec2(0u)));
return ivec2((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_327c37() {
ivec2 res = tint_count_trailing_zeros(ivec2(0, 0));
}
void fragment_main() {
countTrailingZeros_327c37();
}
void main() {
fragment_main();
return;
}
#version 310 es
ivec2 tint_count_trailing_zeros(ivec2 v) {
uvec2 x = uvec2(v);
uvec2 b16 = mix(uvec2(16u), uvec2(0u), bvec2((x & uvec2(65535u))));
x = (x >> b16);
uvec2 b8 = mix(uvec2(8u), uvec2(0u), bvec2((x & uvec2(255u))));
x = (x >> b8);
uvec2 b4 = mix(uvec2(4u), uvec2(0u), bvec2((x & uvec2(15u))));
x = (x >> b4);
uvec2 b2 = mix(uvec2(2u), uvec2(0u), bvec2((x & uvec2(3u))));
x = (x >> b2);
uvec2 b1 = mix(uvec2(1u), uvec2(0u), bvec2((x & uvec2(1u))));
uvec2 is_zero = mix(uvec2(0u), uvec2(1u), equal(x, uvec2(0u)));
return ivec2((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_327c37() {
ivec2 res = tint_count_trailing_zeros(ivec2(0, 0));
}
void compute_main() {
countTrailingZeros_327c37();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/countTrailingZeros/327c37.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
int2 tint_count_trailing_zeros(int2 v) {
uint2 x = uint2(v);
const uint2 b16 = (bool2((x & uint2((65535u).xx))) ? uint2((0u).xx) : uint2((16u).xx));
x = (x >> b16);
const uint2 b8 = (bool2((x & uint2((255u).xx))) ? uint2((0u).xx) : uint2((8u).xx));
x = (x >> b8);
const uint2 b4 = (bool2((x & uint2((15u).xx))) ? uint2((0u).xx) : uint2((4u).xx));
x = (x >> b4);
const uint2 b2 = (bool2((x & uint2((3u).xx))) ? uint2((0u).xx) : uint2((2u).xx));
x = (x >> b2);
const uint2 b1 = (bool2((x & uint2((1u).xx))) ? uint2((0u).xx) : uint2((1u).xx));
const uint2 is_zero = ((x == uint2((0u).xx)) ? uint2((1u).xx) : uint2((0u).xx));
return int2((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_327c37() {
int2 res = tint_count_trailing_zeros(int2(0, 0));
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
countTrailingZeros_327c37();
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() {
countTrailingZeros_327c37();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
countTrailingZeros_327c37();
return;
}

33
test/tint/builtins/gen/countTrailingZeros/327c37.wgsl.expected.msl Normal file
View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
void countTrailingZeros_327c37() {
int2 res = ctz(int2());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
countTrailingZeros_327c37();
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() {
countTrailingZeros_327c37();
return;
}
kernel void compute_main() {
countTrailingZeros_327c37();
return;
}

146
test/tint/builtins/gen/countTrailingZeros/327c37.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,146 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 102
; 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_count_trailing_zeros "tint_count_trailing_zeros"
OpName %v "v"
OpName %x "x"
OpName %countTrailingZeros_327c37 "countTrailingZeros_327c37"
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
%_ptr_Function_v2uint = OpTypePointer Function %v2uint
%20 = OpConstantNull %v2uint
%bool = OpTypeBool
%v2bool = OpTypeVector %bool 2
%uint_65535 = OpConstant %uint 65535
%27 = OpConstantComposite %v2uint %uint_65535 %uint_65535
%uint_0 = OpConstant %uint 0
%30 = OpConstantComposite %v2uint %uint_0 %uint_0
%uint_16 = OpConstant %uint 16
%32 = OpConstantComposite %v2uint %uint_16 %uint_16
%uint_255 = OpConstant %uint 255
%39 = OpConstantComposite %v2uint %uint_255 %uint_255
%uint_8 = OpConstant %uint 8
%42 = OpConstantComposite %v2uint %uint_8 %uint_8
%uint_15 = OpConstant %uint 15
%49 = OpConstantComposite %v2uint %uint_15 %uint_15
%uint_4 = OpConstant %uint 4
%52 = OpConstantComposite %v2uint %uint_4 %uint_4
%uint_3 = OpConstant %uint 3
%59 = OpConstantComposite %v2uint %uint_3 %uint_3
%uint_2 = OpConstant %uint 2
%62 = OpConstantComposite %v2uint %uint_2 %uint_2
%uint_1 = OpConstant %uint 1
%69 = OpConstantComposite %v2uint %uint_1 %uint_1
%void = OpTypeVoid
%80 = OpTypeFunction %void
%85 = OpConstantNull %v2int
%_ptr_Function_v2int = OpTypePointer Function %v2int
%88 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_count_trailing_zeros = OpFunction %v2int None %9
%v = OpFunctionParameter %v2int
%14 = OpLabel
%x = OpVariable %_ptr_Function_v2uint Function %20
%15 = OpBitcast %v2uint %v
OpStore %x %15
%25 = OpLoad %v2uint %x
%28 = OpBitwiseAnd %v2uint %25 %27
%22 = OpINotEqual %v2bool %28 %20
%21 = OpSelect %v2uint %22 %30 %32
%33 = OpLoad %v2uint %x
%34 = OpShiftRightLogical %v2uint %33 %21
OpStore %x %34
%37 = OpLoad %v2uint %x
%40 = OpBitwiseAnd %v2uint %37 %39
%36 = OpINotEqual %v2bool %40 %20
%35 = OpSelect %v2uint %36 %30 %42
%43 = OpLoad %v2uint %x
%44 = OpShiftRightLogical %v2uint %43 %35
OpStore %x %44
%47 = OpLoad %v2uint %x
%50 = OpBitwiseAnd %v2uint %47 %49
%46 = OpINotEqual %v2bool %50 %20
%45 = OpSelect %v2uint %46 %30 %52
%53 = OpLoad %v2uint %x
%54 = OpShiftRightLogical %v2uint %53 %45
OpStore %x %54
%57 = OpLoad %v2uint %x
%60 = OpBitwiseAnd %v2uint %57 %59
%56 = OpINotEqual %v2bool %60 %20
%55 = OpSelect %v2uint %56 %30 %62
%63 = OpLoad %v2uint %x
%64 = OpShiftRightLogical %v2uint %63 %55
OpStore %x %64
%67 = OpLoad %v2uint %x
%70 = OpBitwiseAnd %v2uint %67 %69
%66 = OpINotEqual %v2bool %70 %20
%65 = OpSelect %v2uint %66 %30 %69
%72 = OpLoad %v2uint %x
%73 = OpIEqual %v2bool %72 %30
%71 = OpSelect %v2uint %73 %69 %30
%75 = OpBitwiseOr %v2uint %21 %35
%76 = OpBitwiseOr %v2uint %75 %45
%77 = OpBitwiseOr %v2uint %76 %55
%78 = OpBitwiseOr %v2uint %77 %65
%79 = OpIAdd %v2uint %78 %71
%74 = OpBitcast %v2int %79
OpReturnValue %74
OpFunctionEnd
%countTrailingZeros_327c37 = OpFunction %void None %80
%83 = OpLabel
%res = OpVariable %_ptr_Function_v2int Function %85
%84 = OpFunctionCall %v2int %tint_count_trailing_zeros %85
OpStore %res %84
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %88
%90 = OpLabel
%91 = OpFunctionCall %void %countTrailingZeros_327c37
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %80
%93 = OpLabel
%94 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %94
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %80
%97 = OpLabel
%98 = OpFunctionCall %void %countTrailingZeros_327c37
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %80
%100 = OpLabel
%101 = OpFunctionCall %void %countTrailingZeros_327c37
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/countTrailingZeros/327c37.wgsl.expected.wgsl Normal file
View File

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

45
test/tint/builtins/gen/countTrailingZeros/42fed6.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 countTrailingZeros(i32) -> i32
fn countTrailingZeros_42fed6() {
var res: i32 = countTrailingZeros(1);
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
countTrailingZeros_42fed6();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
countTrailingZeros_42fed6();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
countTrailingZeros_42fed6();
}

93
test/tint/builtins/gen/countTrailingZeros/42fed6.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
int tint_count_trailing_zeros(int v) {
uint x = uint(v);
uint b16 = (bool((x & 65535u)) ? 0u : 16u);
x = (x >> b16);
uint b8 = (bool((x & 255u)) ? 0u : 8u);
x = (x >> b8);
uint b4 = (bool((x & 15u)) ? 0u : 4u);
x = (x >> b4);
uint b2 = (bool((x & 3u)) ? 0u : 2u);
x = (x >> b2);
uint b1 = (bool((x & 1u)) ? 0u : 1u);
uint is_zero = ((x == 0u) ? 1u : 0u);
return int((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_42fed6() {
int res = tint_count_trailing_zeros(1);
}
vec4 vertex_main() {
countTrailingZeros_42fed6();
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_count_trailing_zeros(int v) {
uint x = uint(v);
uint b16 = (bool((x & 65535u)) ? 0u : 16u);
x = (x >> b16);
uint b8 = (bool((x & 255u)) ? 0u : 8u);
x = (x >> b8);
uint b4 = (bool((x & 15u)) ? 0u : 4u);
x = (x >> b4);
uint b2 = (bool((x & 3u)) ? 0u : 2u);
x = (x >> b2);
uint b1 = (bool((x & 1u)) ? 0u : 1u);
uint is_zero = ((x == 0u) ? 1u : 0u);
return int((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_42fed6() {
int res = tint_count_trailing_zeros(1);
}
void fragment_main() {
countTrailingZeros_42fed6();
}
void main() {
fragment_main();
return;
}
#version 310 es
int tint_count_trailing_zeros(int v) {
uint x = uint(v);
uint b16 = (bool((x & 65535u)) ? 0u : 16u);
x = (x >> b16);
uint b8 = (bool((x & 255u)) ? 0u : 8u);
x = (x >> b8);
uint b4 = (bool((x & 15u)) ? 0u : 4u);
x = (x >> b4);
uint b2 = (bool((x & 3u)) ? 0u : 2u);
x = (x >> b2);
uint b1 = (bool((x & 1u)) ? 0u : 1u);
uint is_zero = ((x == 0u) ? 1u : 0u);
return int((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_42fed6() {
int res = tint_count_trailing_zeros(1);
}
void compute_main() {
countTrailingZeros_42fed6();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/countTrailingZeros/42fed6.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
int tint_count_trailing_zeros(int v) {
uint x = uint(v);
const uint b16 = (bool((x & 65535u)) ? 0u : 16u);
x = (x >> b16);
const uint b8 = (bool((x & 255u)) ? 0u : 8u);
x = (x >> b8);
const uint b4 = (bool((x & 15u)) ? 0u : 4u);
x = (x >> b4);
const uint b2 = (bool((x & 3u)) ? 0u : 2u);
x = (x >> b2);
const uint b1 = (bool((x & 1u)) ? 0u : 1u);
const uint is_zero = ((x == 0u) ? 1u : 0u);
return int((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_42fed6() {
int res = tint_count_trailing_zeros(1);
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
countTrailingZeros_42fed6();
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() {
countTrailingZeros_42fed6();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
countTrailingZeros_42fed6();
return;
}

33
test/tint/builtins/gen/countTrailingZeros/42fed6.wgsl.expected.msl Normal file
View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
void countTrailingZeros_42fed6() {
int res = ctz(1);
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
countTrailingZeros_42fed6();
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() {
countTrailingZeros_42fed6();
return;
}
kernel void compute_main() {
countTrailingZeros_42fed6();
return;
}

134
test/tint/builtins/gen/countTrailingZeros/42fed6.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,134 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 90
; 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_count_trailing_zeros "tint_count_trailing_zeros"
OpName %v "v"
OpName %x "x"
OpName %countTrailingZeros_42fed6 "countTrailingZeros_42fed6"
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
%_ptr_Function_uint = OpTypePointer Function %uint
%18 = OpConstantNull %uint
%bool = OpTypeBool
%uint_65535 = OpConstant %uint 65535
%uint_0 = OpConstant %uint 0
%uint_16 = OpConstant %uint 16
%uint_255 = OpConstant %uint 255
%uint_8 = OpConstant %uint 8
%uint_15 = OpConstant %uint 15
%uint_4 = OpConstant %uint 4
%uint_3 = OpConstant %uint 3
%uint_2 = OpConstant %uint 2
%uint_1 = OpConstant %uint 1
%void = OpTypeVoid
%67 = OpTypeFunction %void
%int_1 = OpConstant %int 1
%_ptr_Function_int = OpTypePointer Function %int
%75 = OpConstantNull %int
%76 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_count_trailing_zeros = OpFunction %int None %9
%v = OpFunctionParameter %int
%13 = OpLabel
%x = OpVariable %_ptr_Function_uint Function %18
%14 = OpBitcast %uint %v
OpStore %x %14
%22 = OpLoad %uint %x
%24 = OpBitwiseAnd %uint %22 %uint_65535
%20 = OpINotEqual %bool %24 %18
%19 = OpSelect %uint %20 %uint_0 %uint_16
%27 = OpLoad %uint %x
%28 = OpShiftRightLogical %uint %27 %19
OpStore %x %28
%31 = OpLoad %uint %x
%33 = OpBitwiseAnd %uint %31 %uint_255
%30 = OpINotEqual %bool %33 %18
%29 = OpSelect %uint %30 %uint_0 %uint_8
%35 = OpLoad %uint %x
%36 = OpShiftRightLogical %uint %35 %29
OpStore %x %36
%39 = OpLoad %uint %x
%41 = OpBitwiseAnd %uint %39 %uint_15
%38 = OpINotEqual %bool %41 %18
%37 = OpSelect %uint %38 %uint_0 %uint_4
%43 = OpLoad %uint %x
%44 = OpShiftRightLogical %uint %43 %37
OpStore %x %44
%47 = OpLoad %uint %x
%49 = OpBitwiseAnd %uint %47 %uint_3
%46 = OpINotEqual %bool %49 %18
%45 = OpSelect %uint %46 %uint_0 %uint_2
%51 = OpLoad %uint %x
%52 = OpShiftRightLogical %uint %51 %45
OpStore %x %52
%55 = OpLoad %uint %x
%57 = OpBitwiseAnd %uint %55 %uint_1
%54 = OpINotEqual %bool %57 %18
%53 = OpSelect %uint %54 %uint_0 %uint_1
%59 = OpLoad %uint %x
%60 = OpIEqual %bool %59 %uint_0
%58 = OpSelect %uint %60 %uint_1 %uint_0
%62 = OpBitwiseOr %uint %19 %29
%63 = OpBitwiseOr %uint %62 %37
%64 = OpBitwiseOr %uint %63 %45
%65 = OpBitwiseOr %uint %64 %53
%66 = OpIAdd %uint %65 %58
%61 = OpBitcast %int %66
OpReturnValue %61
OpFunctionEnd
%countTrailingZeros_42fed6 = OpFunction %void None %67
%70 = OpLabel
%res = OpVariable %_ptr_Function_int Function %75
%71 = OpFunctionCall %int %tint_count_trailing_zeros %int_1
OpStore %res %71
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %76
%78 = OpLabel
%79 = OpFunctionCall %void %countTrailingZeros_42fed6
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %67
%81 = OpLabel
%82 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %82
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %67
%85 = OpLabel
%86 = OpFunctionCall %void %countTrailingZeros_42fed6
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %67
%88 = OpLabel
%89 = OpFunctionCall %void %countTrailingZeros_42fed6
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/countTrailingZeros/42fed6.wgsl.expected.wgsl Normal file
View File

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

45
test/tint/builtins/gen/countTrailingZeros/8ed26f.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 countTrailingZeros(vec<3, u32>) -> vec<3, u32>
fn countTrailingZeros_8ed26f() {
var res: vec3<u32> = countTrailingZeros(vec3<u32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
countTrailingZeros_8ed26f();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
countTrailingZeros_8ed26f();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
countTrailingZeros_8ed26f();
}

93
test/tint/builtins/gen/countTrailingZeros/8ed26f.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
uvec3 tint_count_trailing_zeros(uvec3 v) {
uvec3 x = uvec3(v);
uvec3 b16 = mix(uvec3(16u), uvec3(0u), bvec3((x & uvec3(65535u))));
x = (x >> b16);
uvec3 b8 = mix(uvec3(8u), uvec3(0u), bvec3((x & uvec3(255u))));
x = (x >> b8);
uvec3 b4 = mix(uvec3(4u), uvec3(0u), bvec3((x & uvec3(15u))));
x = (x >> b4);
uvec3 b2 = mix(uvec3(2u), uvec3(0u), bvec3((x & uvec3(3u))));
x = (x >> b2);
uvec3 b1 = mix(uvec3(1u), uvec3(0u), bvec3((x & uvec3(1u))));
uvec3 is_zero = mix(uvec3(0u), uvec3(1u), equal(x, uvec3(0u)));
return uvec3((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_8ed26f() {
uvec3 res = tint_count_trailing_zeros(uvec3(0u, 0u, 0u));
}
vec4 vertex_main() {
countTrailingZeros_8ed26f();
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_count_trailing_zeros(uvec3 v) {
uvec3 x = uvec3(v);
uvec3 b16 = mix(uvec3(16u), uvec3(0u), bvec3((x & uvec3(65535u))));
x = (x >> b16);
uvec3 b8 = mix(uvec3(8u), uvec3(0u), bvec3((x & uvec3(255u))));
x = (x >> b8);
uvec3 b4 = mix(uvec3(4u), uvec3(0u), bvec3((x & uvec3(15u))));
x = (x >> b4);
uvec3 b2 = mix(uvec3(2u), uvec3(0u), bvec3((x & uvec3(3u))));
x = (x >> b2);
uvec3 b1 = mix(uvec3(1u), uvec3(0u), bvec3((x & uvec3(1u))));
uvec3 is_zero = mix(uvec3(0u), uvec3(1u), equal(x, uvec3(0u)));
return uvec3((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_8ed26f() {
uvec3 res = tint_count_trailing_zeros(uvec3(0u, 0u, 0u));
}
void fragment_main() {
countTrailingZeros_8ed26f();
}
void main() {
fragment_main();
return;
}
#version 310 es
uvec3 tint_count_trailing_zeros(uvec3 v) {
uvec3 x = uvec3(v);
uvec3 b16 = mix(uvec3(16u), uvec3(0u), bvec3((x & uvec3(65535u))));
x = (x >> b16);
uvec3 b8 = mix(uvec3(8u), uvec3(0u), bvec3((x & uvec3(255u))));
x = (x >> b8);
uvec3 b4 = mix(uvec3(4u), uvec3(0u), bvec3((x & uvec3(15u))));
x = (x >> b4);
uvec3 b2 = mix(uvec3(2u), uvec3(0u), bvec3((x & uvec3(3u))));
x = (x >> b2);
uvec3 b1 = mix(uvec3(1u), uvec3(0u), bvec3((x & uvec3(1u))));
uvec3 is_zero = mix(uvec3(0u), uvec3(1u), equal(x, uvec3(0u)));
return uvec3((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_8ed26f() {
uvec3 res = tint_count_trailing_zeros(uvec3(0u, 0u, 0u));
}
void compute_main() {
countTrailingZeros_8ed26f();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/countTrailingZeros/8ed26f.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
uint3 tint_count_trailing_zeros(uint3 v) {
uint3 x = uint3(v);
const uint3 b16 = (bool3((x & uint3((65535u).xxx))) ? uint3((0u).xxx) : uint3((16u).xxx));
x = (x >> b16);
const uint3 b8 = (bool3((x & uint3((255u).xxx))) ? uint3((0u).xxx) : uint3((8u).xxx));
x = (x >> b8);
const uint3 b4 = (bool3((x & uint3((15u).xxx))) ? uint3((0u).xxx) : uint3((4u).xxx));
x = (x >> b4);
const uint3 b2 = (bool3((x & uint3((3u).xxx))) ? uint3((0u).xxx) : uint3((2u).xxx));
x = (x >> b2);
const uint3 b1 = (bool3((x & uint3((1u).xxx))) ? uint3((0u).xxx) : uint3((1u).xxx));
const uint3 is_zero = ((x == uint3((0u).xxx)) ? uint3((1u).xxx) : uint3((0u).xxx));
return uint3((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_8ed26f() {
uint3 res = tint_count_trailing_zeros(uint3(0u, 0u, 0u));
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
countTrailingZeros_8ed26f();
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() {
countTrailingZeros_8ed26f();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
countTrailingZeros_8ed26f();
return;
}

33
test/tint/builtins/gen/countTrailingZeros/8ed26f.wgsl.expected.msl Normal file
View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
void countTrailingZeros_8ed26f() {
uint3 res = ctz(uint3());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
countTrailingZeros_8ed26f();
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() {
countTrailingZeros_8ed26f();
return;
}
kernel void compute_main() {
countTrailingZeros_8ed26f();
return;
}

140
test/tint/builtins/gen/countTrailingZeros/8ed26f.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,140 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 98
; 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_count_trailing_zeros "tint_count_trailing_zeros"
OpName %v "v"
OpName %x "x"
OpName %countTrailingZeros_8ed26f "countTrailingZeros_8ed26f"
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
%18 = OpConstantNull %v3uint
%bool = OpTypeBool
%v3bool = OpTypeVector %bool 3
%uint_65535 = OpConstant %uint 65535
%25 = OpConstantComposite %v3uint %uint_65535 %uint_65535 %uint_65535
%uint_0 = OpConstant %uint 0
%28 = OpConstantComposite %v3uint %uint_0 %uint_0 %uint_0
%uint_16 = OpConstant %uint 16
%30 = OpConstantComposite %v3uint %uint_16 %uint_16 %uint_16
%uint_255 = OpConstant %uint 255
%37 = OpConstantComposite %v3uint %uint_255 %uint_255 %uint_255
%uint_8 = OpConstant %uint 8
%40 = OpConstantComposite %v3uint %uint_8 %uint_8 %uint_8
%uint_15 = OpConstant %uint 15
%47 = OpConstantComposite %v3uint %uint_15 %uint_15 %uint_15
%uint_4 = OpConstant %uint 4
%50 = OpConstantComposite %v3uint %uint_4 %uint_4 %uint_4
%uint_3 = OpConstant %uint 3
%57 = OpConstantComposite %v3uint %uint_3 %uint_3 %uint_3
%uint_2 = OpConstant %uint 2
%60 = OpConstantComposite %v3uint %uint_2 %uint_2 %uint_2
%uint_1 = OpConstant %uint 1
%67 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
%void = OpTypeVoid
%78 = OpTypeFunction %void
%84 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_count_trailing_zeros = OpFunction %v3uint None %9
%v = OpFunctionParameter %v3uint
%14 = OpLabel
%x = OpVariable %_ptr_Function_v3uint Function %18
OpStore %x %v
%23 = OpLoad %v3uint %x
%26 = OpBitwiseAnd %v3uint %23 %25
%20 = OpINotEqual %v3bool %26 %18
%19 = OpSelect %v3uint %20 %28 %30
%31 = OpLoad %v3uint %x
%32 = OpShiftRightLogical %v3uint %31 %19
OpStore %x %32
%35 = OpLoad %v3uint %x
%38 = OpBitwiseAnd %v3uint %35 %37
%34 = OpINotEqual %v3bool %38 %18
%33 = OpSelect %v3uint %34 %28 %40
%41 = OpLoad %v3uint %x
%42 = OpShiftRightLogical %v3uint %41 %33
OpStore %x %42
%45 = OpLoad %v3uint %x
%48 = OpBitwiseAnd %v3uint %45 %47
%44 = OpINotEqual %v3bool %48 %18
%43 = OpSelect %v3uint %44 %28 %50
%51 = OpLoad %v3uint %x
%52 = OpShiftRightLogical %v3uint %51 %43
OpStore %x %52
%55 = OpLoad %v3uint %x
%58 = OpBitwiseAnd %v3uint %55 %57
%54 = OpINotEqual %v3bool %58 %18
%53 = OpSelect %v3uint %54 %28 %60
%61 = OpLoad %v3uint %x
%62 = OpShiftRightLogical %v3uint %61 %53
OpStore %x %62
%65 = OpLoad %v3uint %x
%68 = OpBitwiseAnd %v3uint %65 %67
%64 = OpINotEqual %v3bool %68 %18
%63 = OpSelect %v3uint %64 %28 %67
%70 = OpLoad %v3uint %x
%71 = OpIEqual %v3bool %70 %28
%69 = OpSelect %v3uint %71 %67 %28
%73 = OpBitwiseOr %v3uint %19 %33
%74 = OpBitwiseOr %v3uint %73 %43
%75 = OpBitwiseOr %v3uint %74 %53
%76 = OpBitwiseOr %v3uint %75 %63
%77 = OpIAdd %v3uint %76 %69
OpReturnValue %77
OpFunctionEnd
%countTrailingZeros_8ed26f = OpFunction %void None %78
%81 = OpLabel
%res = OpVariable %_ptr_Function_v3uint Function %18
%82 = OpFunctionCall %v3uint %tint_count_trailing_zeros %18
OpStore %res %82
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %84
%86 = OpLabel
%87 = OpFunctionCall %void %countTrailingZeros_8ed26f
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %78
%89 = OpLabel
%90 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %90
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %78
%93 = OpLabel
%94 = OpFunctionCall %void %countTrailingZeros_8ed26f
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %78
%96 = OpLabel
%97 = OpFunctionCall %void %countTrailingZeros_8ed26f
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/countTrailingZeros/8ed26f.wgsl.expected.wgsl Normal file
View File

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

45
test/tint/builtins/gen/countTrailingZeros/acfacb.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 countTrailingZeros(vec<3, i32>) -> vec<3, i32>
fn countTrailingZeros_acfacb() {
var res: vec3<i32> = countTrailingZeros(vec3<i32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
countTrailingZeros_acfacb();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
countTrailingZeros_acfacb();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
countTrailingZeros_acfacb();
}

93
test/tint/builtins/gen/countTrailingZeros/acfacb.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
ivec3 tint_count_trailing_zeros(ivec3 v) {
uvec3 x = uvec3(v);
uvec3 b16 = mix(uvec3(16u), uvec3(0u), bvec3((x & uvec3(65535u))));
x = (x >> b16);
uvec3 b8 = mix(uvec3(8u), uvec3(0u), bvec3((x & uvec3(255u))));
x = (x >> b8);
uvec3 b4 = mix(uvec3(4u), uvec3(0u), bvec3((x & uvec3(15u))));
x = (x >> b4);
uvec3 b2 = mix(uvec3(2u), uvec3(0u), bvec3((x & uvec3(3u))));
x = (x >> b2);
uvec3 b1 = mix(uvec3(1u), uvec3(0u), bvec3((x & uvec3(1u))));
uvec3 is_zero = mix(uvec3(0u), uvec3(1u), equal(x, uvec3(0u)));
return ivec3((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_acfacb() {
ivec3 res = tint_count_trailing_zeros(ivec3(0, 0, 0));
}
vec4 vertex_main() {
countTrailingZeros_acfacb();
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_count_trailing_zeros(ivec3 v) {
uvec3 x = uvec3(v);
uvec3 b16 = mix(uvec3(16u), uvec3(0u), bvec3((x & uvec3(65535u))));
x = (x >> b16);
uvec3 b8 = mix(uvec3(8u), uvec3(0u), bvec3((x & uvec3(255u))));
x = (x >> b8);
uvec3 b4 = mix(uvec3(4u), uvec3(0u), bvec3((x & uvec3(15u))));
x = (x >> b4);
uvec3 b2 = mix(uvec3(2u), uvec3(0u), bvec3((x & uvec3(3u))));
x = (x >> b2);
uvec3 b1 = mix(uvec3(1u), uvec3(0u), bvec3((x & uvec3(1u))));
uvec3 is_zero = mix(uvec3(0u), uvec3(1u), equal(x, uvec3(0u)));
return ivec3((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_acfacb() {
ivec3 res = tint_count_trailing_zeros(ivec3(0, 0, 0));
}
void fragment_main() {
countTrailingZeros_acfacb();
}
void main() {
fragment_main();
return;
}
#version 310 es
ivec3 tint_count_trailing_zeros(ivec3 v) {
uvec3 x = uvec3(v);
uvec3 b16 = mix(uvec3(16u), uvec3(0u), bvec3((x & uvec3(65535u))));
x = (x >> b16);
uvec3 b8 = mix(uvec3(8u), uvec3(0u), bvec3((x & uvec3(255u))));
x = (x >> b8);
uvec3 b4 = mix(uvec3(4u), uvec3(0u), bvec3((x & uvec3(15u))));
x = (x >> b4);
uvec3 b2 = mix(uvec3(2u), uvec3(0u), bvec3((x & uvec3(3u))));
x = (x >> b2);
uvec3 b1 = mix(uvec3(1u), uvec3(0u), bvec3((x & uvec3(1u))));
uvec3 is_zero = mix(uvec3(0u), uvec3(1u), equal(x, uvec3(0u)));
return ivec3((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_acfacb() {
ivec3 res = tint_count_trailing_zeros(ivec3(0, 0, 0));
}
void compute_main() {
countTrailingZeros_acfacb();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/countTrailingZeros/acfacb.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
int3 tint_count_trailing_zeros(int3 v) {
uint3 x = uint3(v);
const uint3 b16 = (bool3((x & uint3((65535u).xxx))) ? uint3((0u).xxx) : uint3((16u).xxx));
x = (x >> b16);
const uint3 b8 = (bool3((x & uint3((255u).xxx))) ? uint3((0u).xxx) : uint3((8u).xxx));
x = (x >> b8);
const uint3 b4 = (bool3((x & uint3((15u).xxx))) ? uint3((0u).xxx) : uint3((4u).xxx));
x = (x >> b4);
const uint3 b2 = (bool3((x & uint3((3u).xxx))) ? uint3((0u).xxx) : uint3((2u).xxx));
x = (x >> b2);
const uint3 b1 = (bool3((x & uint3((1u).xxx))) ? uint3((0u).xxx) : uint3((1u).xxx));
const uint3 is_zero = ((x == uint3((0u).xxx)) ? uint3((1u).xxx) : uint3((0u).xxx));
return int3((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_acfacb() {
int3 res = tint_count_trailing_zeros(int3(0, 0, 0));
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
countTrailingZeros_acfacb();
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() {
countTrailingZeros_acfacb();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
countTrailingZeros_acfacb();
return;
}

33
test/tint/builtins/gen/countTrailingZeros/acfacb.wgsl.expected.msl Normal file
View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
void countTrailingZeros_acfacb() {
int3 res = ctz(int3());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
countTrailingZeros_acfacb();
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() {
countTrailingZeros_acfacb();
return;
}
kernel void compute_main() {
countTrailingZeros_acfacb();
return;
}

146
test/tint/builtins/gen/countTrailingZeros/acfacb.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,146 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 102
; 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_count_trailing_zeros "tint_count_trailing_zeros"
OpName %v "v"
OpName %x "x"
OpName %countTrailingZeros_acfacb "countTrailingZeros_acfacb"
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
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
%20 = OpConstantNull %v3uint
%bool = OpTypeBool
%v3bool = OpTypeVector %bool 3
%uint_65535 = OpConstant %uint 65535
%27 = OpConstantComposite %v3uint %uint_65535 %uint_65535 %uint_65535
%uint_0 = OpConstant %uint 0
%30 = OpConstantComposite %v3uint %uint_0 %uint_0 %uint_0
%uint_16 = OpConstant %uint 16
%32 = OpConstantComposite %v3uint %uint_16 %uint_16 %uint_16
%uint_255 = OpConstant %uint 255
%39 = OpConstantComposite %v3uint %uint_255 %uint_255 %uint_255
%uint_8 = OpConstant %uint 8
%42 = OpConstantComposite %v3uint %uint_8 %uint_8 %uint_8
%uint_15 = OpConstant %uint 15
%49 = OpConstantComposite %v3uint %uint_15 %uint_15 %uint_15
%uint_4 = OpConstant %uint 4
%52 = OpConstantComposite %v3uint %uint_4 %uint_4 %uint_4
%uint_3 = OpConstant %uint 3
%59 = OpConstantComposite %v3uint %uint_3 %uint_3 %uint_3
%uint_2 = OpConstant %uint 2
%62 = OpConstantComposite %v3uint %uint_2 %uint_2 %uint_2
%uint_1 = OpConstant %uint 1
%69 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
%void = OpTypeVoid
%80 = OpTypeFunction %void
%85 = OpConstantNull %v3int
%_ptr_Function_v3int = OpTypePointer Function %v3int
%88 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_count_trailing_zeros = OpFunction %v3int None %9
%v = OpFunctionParameter %v3int
%14 = OpLabel
%x = OpVariable %_ptr_Function_v3uint Function %20
%15 = OpBitcast %v3uint %v
OpStore %x %15
%25 = OpLoad %v3uint %x
%28 = OpBitwiseAnd %v3uint %25 %27
%22 = OpINotEqual %v3bool %28 %20
%21 = OpSelect %v3uint %22 %30 %32
%33 = OpLoad %v3uint %x
%34 = OpShiftRightLogical %v3uint %33 %21
OpStore %x %34
%37 = OpLoad %v3uint %x
%40 = OpBitwiseAnd %v3uint %37 %39
%36 = OpINotEqual %v3bool %40 %20
%35 = OpSelect %v3uint %36 %30 %42
%43 = OpLoad %v3uint %x
%44 = OpShiftRightLogical %v3uint %43 %35
OpStore %x %44
%47 = OpLoad %v3uint %x
%50 = OpBitwiseAnd %v3uint %47 %49
%46 = OpINotEqual %v3bool %50 %20
%45 = OpSelect %v3uint %46 %30 %52
%53 = OpLoad %v3uint %x
%54 = OpShiftRightLogical %v3uint %53 %45
OpStore %x %54
%57 = OpLoad %v3uint %x
%60 = OpBitwiseAnd %v3uint %57 %59
%56 = OpINotEqual %v3bool %60 %20
%55 = OpSelect %v3uint %56 %30 %62
%63 = OpLoad %v3uint %x
%64 = OpShiftRightLogical %v3uint %63 %55
OpStore %x %64
%67 = OpLoad %v3uint %x
%70 = OpBitwiseAnd %v3uint %67 %69
%66 = OpINotEqual %v3bool %70 %20
%65 = OpSelect %v3uint %66 %30 %69
%72 = OpLoad %v3uint %x
%73 = OpIEqual %v3bool %72 %30
%71 = OpSelect %v3uint %73 %69 %30
%75 = OpBitwiseOr %v3uint %21 %35
%76 = OpBitwiseOr %v3uint %75 %45
%77 = OpBitwiseOr %v3uint %76 %55
%78 = OpBitwiseOr %v3uint %77 %65
%79 = OpIAdd %v3uint %78 %71
%74 = OpBitcast %v3int %79
OpReturnValue %74
OpFunctionEnd
%countTrailingZeros_acfacb = OpFunction %void None %80
%83 = OpLabel
%res = OpVariable %_ptr_Function_v3int Function %85
%84 = OpFunctionCall %v3int %tint_count_trailing_zeros %85
OpStore %res %84
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %88
%90 = OpLabel
%91 = OpFunctionCall %void %countTrailingZeros_acfacb
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %80
%93 = OpLabel
%94 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %94
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %80
%97 = OpLabel
%98 = OpFunctionCall %void %countTrailingZeros_acfacb
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %80
%100 = OpLabel
%101 = OpFunctionCall %void %countTrailingZeros_acfacb
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/countTrailingZeros/acfacb.wgsl.expected.wgsl Normal file
View File

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

45
test/tint/builtins/gen/countTrailingZeros/d2b4a0.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 countTrailingZeros(vec<4, u32>) -> vec<4, u32>
fn countTrailingZeros_d2b4a0() {
var res: vec4<u32> = countTrailingZeros(vec4<u32>());
}
@stage(vertex)
fn vertex_main() -> @builtin(position) vec4<f32> {
countTrailingZeros_d2b4a0();
return vec4<f32>();
}
@stage(fragment)
fn fragment_main() {
countTrailingZeros_d2b4a0();
}
@stage(compute) @workgroup_size(1)
fn compute_main() {
countTrailingZeros_d2b4a0();
}

93
test/tint/builtins/gen/countTrailingZeros/d2b4a0.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,93 @@
#version 310 es
uvec4 tint_count_trailing_zeros(uvec4 v) {
uvec4 x = uvec4(v);
uvec4 b16 = mix(uvec4(16u), uvec4(0u), bvec4((x & uvec4(65535u))));
x = (x >> b16);
uvec4 b8 = mix(uvec4(8u), uvec4(0u), bvec4((x & uvec4(255u))));
x = (x >> b8);
uvec4 b4 = mix(uvec4(4u), uvec4(0u), bvec4((x & uvec4(15u))));
x = (x >> b4);
uvec4 b2 = mix(uvec4(2u), uvec4(0u), bvec4((x & uvec4(3u))));
x = (x >> b2);
uvec4 b1 = mix(uvec4(1u), uvec4(0u), bvec4((x & uvec4(1u))));
uvec4 is_zero = mix(uvec4(0u), uvec4(1u), equal(x, uvec4(0u)));
return uvec4((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_d2b4a0() {
uvec4 res = tint_count_trailing_zeros(uvec4(0u, 0u, 0u, 0u));
}
vec4 vertex_main() {
countTrailingZeros_d2b4a0();
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_count_trailing_zeros(uvec4 v) {
uvec4 x = uvec4(v);
uvec4 b16 = mix(uvec4(16u), uvec4(0u), bvec4((x & uvec4(65535u))));
x = (x >> b16);
uvec4 b8 = mix(uvec4(8u), uvec4(0u), bvec4((x & uvec4(255u))));
x = (x >> b8);
uvec4 b4 = mix(uvec4(4u), uvec4(0u), bvec4((x & uvec4(15u))));
x = (x >> b4);
uvec4 b2 = mix(uvec4(2u), uvec4(0u), bvec4((x & uvec4(3u))));
x = (x >> b2);
uvec4 b1 = mix(uvec4(1u), uvec4(0u), bvec4((x & uvec4(1u))));
uvec4 is_zero = mix(uvec4(0u), uvec4(1u), equal(x, uvec4(0u)));
return uvec4((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_d2b4a0() {
uvec4 res = tint_count_trailing_zeros(uvec4(0u, 0u, 0u, 0u));
}
void fragment_main() {
countTrailingZeros_d2b4a0();
}
void main() {
fragment_main();
return;
}
#version 310 es
uvec4 tint_count_trailing_zeros(uvec4 v) {
uvec4 x = uvec4(v);
uvec4 b16 = mix(uvec4(16u), uvec4(0u), bvec4((x & uvec4(65535u))));
x = (x >> b16);
uvec4 b8 = mix(uvec4(8u), uvec4(0u), bvec4((x & uvec4(255u))));
x = (x >> b8);
uvec4 b4 = mix(uvec4(4u), uvec4(0u), bvec4((x & uvec4(15u))));
x = (x >> b4);
uvec4 b2 = mix(uvec4(2u), uvec4(0u), bvec4((x & uvec4(3u))));
x = (x >> b2);
uvec4 b1 = mix(uvec4(1u), uvec4(0u), bvec4((x & uvec4(1u))));
uvec4 is_zero = mix(uvec4(0u), uvec4(1u), equal(x, uvec4(0u)));
return uvec4((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_d2b4a0() {
uvec4 res = tint_count_trailing_zeros(uvec4(0u, 0u, 0u, 0u));
}
void compute_main() {
countTrailingZeros_d2b4a0();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
compute_main();
return;
}

45
test/tint/builtins/gen/countTrailingZeros/d2b4a0.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,45 @@
uint4 tint_count_trailing_zeros(uint4 v) {
uint4 x = uint4(v);
const uint4 b16 = (bool4((x & uint4((65535u).xxxx))) ? uint4((0u).xxxx) : uint4((16u).xxxx));
x = (x >> b16);
const uint4 b8 = (bool4((x & uint4((255u).xxxx))) ? uint4((0u).xxxx) : uint4((8u).xxxx));
x = (x >> b8);
const uint4 b4 = (bool4((x & uint4((15u).xxxx))) ? uint4((0u).xxxx) : uint4((4u).xxxx));
x = (x >> b4);
const uint4 b2 = (bool4((x & uint4((3u).xxxx))) ? uint4((0u).xxxx) : uint4((2u).xxxx));
x = (x >> b2);
const uint4 b1 = (bool4((x & uint4((1u).xxxx))) ? uint4((0u).xxxx) : uint4((1u).xxxx));
const uint4 is_zero = ((x == uint4((0u).xxxx)) ? uint4((1u).xxxx) : uint4((0u).xxxx));
return uint4((((((b16 | b8) | b4) | b2) | b1) + is_zero));
}
void countTrailingZeros_d2b4a0() {
uint4 res = tint_count_trailing_zeros(uint4(0u, 0u, 0u, 0u));
}
struct tint_symbol {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
countTrailingZeros_d2b4a0();
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() {
countTrailingZeros_d2b4a0();
return;
}
[numthreads(1, 1, 1)]
void compute_main() {
countTrailingZeros_d2b4a0();
return;
}

33
test/tint/builtins/gen/countTrailingZeros/d2b4a0.wgsl.expected.msl Normal file
View File

@ -0,0 +1,33 @@
#include <metal_stdlib>
using namespace metal;
void countTrailingZeros_d2b4a0() {
uint4 res = ctz(uint4());
}
struct tint_symbol {
float4 value [[position]];
};
float4 vertex_main_inner() {
countTrailingZeros_d2b4a0();
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() {
countTrailingZeros_d2b4a0();
return;
}
kernel void compute_main() {
countTrailingZeros_d2b4a0();
return;
}

140
test/tint/builtins/gen/countTrailingZeros/d2b4a0.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,140 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 98
; 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_count_trailing_zeros "tint_count_trailing_zeros"
OpName %v "v"
OpName %x "x"
OpName %countTrailingZeros_d2b4a0 "countTrailingZeros_d2b4a0"
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
%18 = OpConstantNull %v4uint
%bool = OpTypeBool
%v4bool = OpTypeVector %bool 4
%uint_65535 = OpConstant %uint 65535
%25 = OpConstantComposite %v4uint %uint_65535 %uint_65535 %uint_65535 %uint_65535
%uint_0 = OpConstant %uint 0
%28 = OpConstantComposite %v4uint %uint_0 %uint_0 %uint_0 %uint_0
%uint_16 = OpConstant %uint 16
%30 = OpConstantComposite %v4uint %uint_16 %uint_16 %uint_16 %uint_16
%uint_255 = OpConstant %uint 255
%37 = OpConstantComposite %v4uint %uint_255 %uint_255 %uint_255 %uint_255
%uint_8 = OpConstant %uint 8
%40 = OpConstantComposite %v4uint %uint_8 %uint_8 %uint_8 %uint_8
%uint_15 = OpConstant %uint 15
%47 = OpConstantComposite %v4uint %uint_15 %uint_15 %uint_15 %uint_15
%uint_4 = OpConstant %uint 4
%50 = OpConstantComposite %v4uint %uint_4 %uint_4 %uint_4 %uint_4
%uint_3 = OpConstant %uint 3
%57 = OpConstantComposite %v4uint %uint_3 %uint_3 %uint_3 %uint_3
%uint_2 = OpConstant %uint 2
%60 = 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
%void = OpTypeVoid
%78 = OpTypeFunction %void
%84 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%tint_count_trailing_zeros = OpFunction %v4uint None %9
%v = OpFunctionParameter %v4uint
%14 = OpLabel
%x = OpVariable %_ptr_Function_v4uint Function %18
OpStore %x %v
%23 = OpLoad %v4uint %x
%26 = OpBitwiseAnd %v4uint %23 %25
%20 = OpINotEqual %v4bool %26 %18
%19 = OpSelect %v4uint %20 %28 %30
%31 = OpLoad %v4uint %x
%32 = OpShiftRightLogical %v4uint %31 %19
OpStore %x %32
%35 = OpLoad %v4uint %x
%38 = OpBitwiseAnd %v4uint %35 %37
%34 = OpINotEqual %v4bool %38 %18
%33 = OpSelect %v4uint %34 %28 %40
%41 = OpLoad %v4uint %x
%42 = OpShiftRightLogical %v4uint %41 %33
OpStore %x %42
%45 = OpLoad %v4uint %x
%48 = OpBitwiseAnd %v4uint %45 %47
%44 = OpINotEqual %v4bool %48 %18
%43 = OpSelect %v4uint %44 %28 %50
%51 = OpLoad %v4uint %x
%52 = OpShiftRightLogical %v4uint %51 %43
OpStore %x %52
%55 = OpLoad %v4uint %x
%58 = OpBitwiseAnd %v4uint %55 %57
%54 = OpINotEqual %v4bool %58 %18
%53 = OpSelect %v4uint %54 %28 %60
%61 = OpLoad %v4uint %x
%62 = OpShiftRightLogical %v4uint %61 %53
OpStore %x %62
%65 = OpLoad %v4uint %x
%68 = OpBitwiseAnd %v4uint %65 %67
%64 = OpINotEqual %v4bool %68 %18
%63 = OpSelect %v4uint %64 %28 %67
%70 = OpLoad %v4uint %x
%71 = OpIEqual %v4bool %70 %28
%69 = OpSelect %v4uint %71 %67 %28
%73 = OpBitwiseOr %v4uint %19 %33
%74 = OpBitwiseOr %v4uint %73 %43
%75 = OpBitwiseOr %v4uint %74 %53
%76 = OpBitwiseOr %v4uint %75 %63
%77 = OpIAdd %v4uint %76 %69
OpReturnValue %77
OpFunctionEnd
%countTrailingZeros_d2b4a0 = OpFunction %void None %78
%81 = OpLabel
%res = OpVariable %_ptr_Function_v4uint Function %18
%82 = OpFunctionCall %v4uint %tint_count_trailing_zeros %18
OpStore %res %82
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %84
%86 = OpLabel
%87 = OpFunctionCall %void %countTrailingZeros_d2b4a0
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %78
%89 = OpLabel
%90 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value %90
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %78
%93 = OpLabel
%94 = OpFunctionCall %void %countTrailingZeros_d2b4a0
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %78
%96 = OpLabel
%97 = OpFunctionCall %void %countTrailingZeros_d2b4a0
OpReturn
OpFunctionEnd

19
test/tint/builtins/gen/countTrailingZeros/d2b4a0.wgsl.expected.wgsl Normal file
View File

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