tint/writer: Disable constant inlining for lets
Once 'const' is introduced, let will no longer resolve to a creation-time constant value. Add temporary code into each of the writers to prevent constants that originate from a 'let' from being inlined. This will reduce the amount of noise in later CLs. Bug: tint:1580 Change-Id: Id3493a43ac09fe9f042ff2d517d04b2ae854d43e Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/94541 Commit-Queue: Ben Clayton <bclayton@google.com> Reviewed-by: David Neto <dneto@google.com> Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: Ben Clayton <bclayton@chromium.org>
This commit is contained in:
Ben Clayton
Dawn LUCI CQ
parent
431c7a4f66
commit
32cb9cf2f8
|
|
@ -1766,6 +1766,11 @@ bool GeneratorImpl::EmitDiscard(const ast::DiscardStatement*) {
|
|||
|
||||
bool GeneratorImpl::EmitExpression(std::ostream& out, const ast::Expression* expr) {
|
||||
if (auto* sem = builder_.Sem().Get(expr)) {
|
||||
if (auto* user = sem->As<sem::VariableUser>();
|
||||
!user || !user->Variable()->Declaration()->Is<ast::Let>()) {
|
||||
// Disable constant inlining if the constant expression is from a 'let' declaration.
|
||||
// TODO(crbug.com/tint/1580): Once 'const' is implemented, 'let' will no longer resolve
|
||||
// to a shader-creation time constant value, and this can be removed.
|
||||
if (auto constant = sem->ConstantValue()) {
|
||||
// We do not want to inline array constants, as this will undo the work of
|
||||
// PromoteInitializersToConstVar, which ensures that arrays are declarated in 'let's
|
||||
|
|
@ -1775,6 +1780,7 @@ bool GeneratorImpl::EmitExpression(std::ostream& out, const ast::Expression* exp
|
|||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
return Switch(
|
||||
expr,
|
||||
[&](const ast::IndexAccessorExpression* a) { //
|
||||
|
|
|
|||
|
|
@ -2611,6 +2611,11 @@ bool GeneratorImpl::EmitDiscard(const ast::DiscardStatement*) {
|
|||
|
||||
bool GeneratorImpl::EmitExpression(std::ostream& out, const ast::Expression* expr) {
|
||||
if (auto* sem = builder_.Sem().Get(expr)) {
|
||||
if (auto* user = sem->As<sem::VariableUser>();
|
||||
!user || !user->Variable()->Declaration()->Is<ast::Let>()) {
|
||||
// Disable constant inlining if the constant expression is from a 'let' declaration.
|
||||
// TODO(crbug.com/tint/1580): Once 'const' is implemented, 'let' will no longer resolve
|
||||
// to a shader-creation time constant value, and this can be removed.
|
||||
if (auto constant = sem->ConstantValue()) {
|
||||
// We do not want to inline array constants, as this will undo the work of
|
||||
// PromoteInitializersToConstVar, which ensures that arrays are declarated in 'let's
|
||||
|
|
@ -2620,6 +2625,7 @@ bool GeneratorImpl::EmitExpression(std::ostream& out, const ast::Expression* exp
|
|||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
return Switch(
|
||||
expr,
|
||||
[&](const ast::IndexAccessorExpression* a) { //
|
||||
|
|
|
|||
|
|
@ -58,7 +58,7 @@ TEST_F(HlslGeneratorImplTest_Assign, Emit_Vector_Assign_ConstantIndex) {
|
|||
float3 lhs = float3(0.0f, 0.0f, 0.0f);
|
||||
float rhs = 0.0f;
|
||||
const uint index = 0u;
|
||||
lhs[0u] = rhs;
|
||||
lhs[index] = rhs;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
|
@ -106,7 +106,7 @@ TEST_F(HlslGeneratorImplTest_Assign, Emit_Matrix_Assign_Vector_ConstantIndex) {
|
|||
float4x2 lhs = float4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
|
||||
float2 rhs = float2(0.0f, 0.0f);
|
||||
const uint index = 0u;
|
||||
lhs[0u] = rhs;
|
||||
lhs[index] = rhs;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
|
@ -159,7 +159,7 @@ TEST_F(HlslGeneratorImplTest_Assign, Emit_Matrix_Assign_Scalar_ConstantIndex) {
|
|||
float4x2 lhs = float4x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
|
||||
float rhs = 0.0f;
|
||||
const uint index = 0u;
|
||||
lhs[0u][0u] = rhs;
|
||||
lhs[index][index] = rhs;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1702,6 +1702,11 @@ bool GeneratorImpl::EmitLiteral(std::ostream& out, const ast::LiteralExpression*
|
|||
|
||||
bool GeneratorImpl::EmitExpression(std::ostream& out, const ast::Expression* expr) {
|
||||
if (auto* sem = builder_.Sem().Get(expr)) {
|
||||
if (auto* user = sem->As<sem::VariableUser>();
|
||||
!user || !user->Variable()->Declaration()->Is<ast::Let>()) {
|
||||
// Disable constant inlining if the constant expression is from a 'let' declaration.
|
||||
// TODO(crbug.com/tint/1580): Once 'const' is implemented, 'let' will no longer resolve
|
||||
// to a shader-creation time constant value, and this can be removed.
|
||||
if (auto constant = sem->ConstantValue()) {
|
||||
// We do not want to inline array constants, as this will undo the work of
|
||||
// PromoteInitializersToConstVar, which ensures that arrays are declarated in 'let's
|
||||
|
|
@ -1711,6 +1716,7 @@ bool GeneratorImpl::EmitExpression(std::ostream& out, const ast::Expression* exp
|
|||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
return Switch(
|
||||
expr,
|
||||
[&](const ast::IndexAccessorExpression* a) { //
|
||||
|
|
|
|||
|
|
@ -132,7 +132,7 @@ TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement_Initializer_Private) {
|
|||
GeneratorImpl& gen = SanitizeAndBuild();
|
||||
|
||||
ASSERT_TRUE(gen.Generate()) << gen.error();
|
||||
EXPECT_THAT(gen.result(), HasSubstr("thread float tint_symbol_1 = 0.0f;\n float const tint_symbol = tint_symbol_1;\n return;\n"));
|
||||
EXPECT_THAT(gen.result(), HasSubstr("thread float tint_symbol_1 = initializer;\n float const tint_symbol = tint_symbol_1;\n return;\n"));
|
||||
}
|
||||
|
||||
TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement_Workgroup) {
|
||||
|
|
|
|||
|
|
@ -2,9 +2,9 @@
|
|||
|
||||
void tint_symbol() {
|
||||
vec3 v = vec3(1.0f, 2.0f, 3.0f);
|
||||
float scalar = vec3(1.0f, 2.0f, 3.0f).y;
|
||||
vec2 swizzle2 = vec3(1.0f, 2.0f, 3.0f).xz;
|
||||
vec3 swizzle3 = vec3(1.0f, 2.0f, 3.0f).xzy;
|
||||
float scalar = v.y;
|
||||
vec2 swizzle2 = v.xz;
|
||||
vec3 swizzle3 = v.xzy;
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -1,8 +1,8 @@
|
|||
[numthreads(1, 1, 1)]
|
||||
void main() {
|
||||
const float3 v = float3(1.0f, 2.0f, 3.0f);
|
||||
const float scalar = float3(1.0f, 2.0f, 3.0f).y;
|
||||
const float2 swizzle2 = float3(1.0f, 2.0f, 3.0f).xz;
|
||||
const float3 swizzle3 = float3(1.0f, 2.0f, 3.0f).xzy;
|
||||
const float scalar = v.y;
|
||||
const float2 swizzle2 = v.xz;
|
||||
const float3 swizzle3 = v.xzy;
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -3,9 +3,9 @@
|
|||
using namespace metal;
|
||||
kernel void tint_symbol() {
|
||||
float3 const v = float3(1.0f, 2.0f, 3.0f);
|
||||
float const scalar = float3(1.0f, 2.0f, 3.0f)[1];
|
||||
float2 const swizzle2 = float3(float3(1.0f, 2.0f, 3.0f)).xz;
|
||||
float3 const swizzle3 = float3(float3(1.0f, 2.0f, 3.0f)).xzy;
|
||||
float const scalar = v[1];
|
||||
float2 const swizzle2 = float3(v).xz;
|
||||
float3 const swizzle3 = float3(v).xzy;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ void tint_symbol() {
|
|||
int x = 42;
|
||||
int empty[4] = int[4](0, 0, 0, 0);
|
||||
int nonempty[4] = int[4](1, 2, 3, 4);
|
||||
int nonempty_with_expr[4] = int[4](1, 42, (42 + 1), 4);
|
||||
int nonempty_with_expr[4] = int[4](1, x, (x + 1), 4);
|
||||
int nested_empty[2][3][4] = int[2][3][4](int[3][4](int[4](0, 0, 0, 0), int[4](0, 0, 0, 0), int[4](0, 0, 0, 0)), int[3][4](int[4](0, 0, 0, 0), int[4](0, 0, 0, 0), int[4](0, 0, 0, 0)));
|
||||
int tint_symbol_1[4] = int[4](1, 2, 3, 4);
|
||||
int tint_symbol_2[4] = int[4](5, 6, 7, 8);
|
||||
|
|
@ -15,7 +15,7 @@ void tint_symbol() {
|
|||
int tint_symbol_7[4] = int[4](21, 22, 23, 24);
|
||||
int tint_symbol_8[3][4] = int[3][4](tint_symbol_5, tint_symbol_6, tint_symbol_7);
|
||||
int nested_nonempty[2][3][4] = int[2][3][4](tint_symbol_4, tint_symbol_8);
|
||||
int tint_symbol_9[4] = int[4](1, 2, 42, (42 + 1));
|
||||
int tint_symbol_9[4] = int[4](1, 2, x, (x + 1));
|
||||
int tint_symbol_10[4] = int[4](5, 6, 3, (4 + 1));
|
||||
int tint_symbol_11[3][4] = int[3][4](tint_symbol_9, tint_symbol_10, nonempty);
|
||||
int nested_nonempty_with_expr[2][3][4] = int[2][3][4](tint_symbol_11, nested_nonempty[1]);
|
||||
|
|
@ -23,7 +23,7 @@ void tint_symbol() {
|
|||
int subexpr_empty = 0;
|
||||
int tint_symbol_13[4] = int[4](1, 2, 3, 4);
|
||||
int subexpr_nonempty = 3;
|
||||
int tint_symbol_14[4] = int[4](1, 42, (42 + 1), 4);
|
||||
int tint_symbol_14[4] = int[4](1, x, (x + 1), 4);
|
||||
int subexpr_nonempty_with_expr = tint_symbol_14[2];
|
||||
int tint_symbol_15[2][4] = int[2][4](int[4](0, 0, 0, 0), int[4](0, 0, 0, 0));
|
||||
int subexpr_nested_empty[4] = tint_symbol_15[1];
|
||||
|
|
@ -31,7 +31,7 @@ void tint_symbol() {
|
|||
int tint_symbol_17[4] = int[4](5, 6, 7, 8);
|
||||
int tint_symbol_18[2][4] = int[2][4](tint_symbol_16, tint_symbol_17);
|
||||
int subexpr_nested_nonempty[4] = tint_symbol_18[1];
|
||||
int tint_symbol_19[4] = int[4](1, 42, (42 + 1), 4);
|
||||
int tint_symbol_19[4] = int[4](1, x, (x + 1), 4);
|
||||
int tint_symbol_20[2][4] = int[2][4](tint_symbol_19, nested_nonempty[1][2]);
|
||||
int subexpr_nested_nonempty_with_expr[4] = tint_symbol_20[1];
|
||||
}
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@ void main() {
|
|||
const int x = 42;
|
||||
const int empty[4] = (int[4])0;
|
||||
const int nonempty[4] = {1, 2, 3, 4};
|
||||
const int nonempty_with_expr[4] = {1, 42, (42 + 1), 4};
|
||||
const int nonempty_with_expr[4] = {1, x, (x + 1), 4};
|
||||
const int nested_empty[2][3][4] = (int[2][3][4])0;
|
||||
const int tint_symbol[4] = {1, 2, 3, 4};
|
||||
const int tint_symbol_1[4] = {5, 6, 7, 8};
|
||||
|
|
@ -14,7 +14,7 @@ void main() {
|
|||
const int tint_symbol_6[4] = {21, 22, 23, 24};
|
||||
const int tint_symbol_7[3][4] = {tint_symbol_4, tint_symbol_5, tint_symbol_6};
|
||||
const int nested_nonempty[2][3][4] = {tint_symbol_3, tint_symbol_7};
|
||||
const int tint_symbol_8[4] = {1, 2, 42, (42 + 1)};
|
||||
const int tint_symbol_8[4] = {1, 2, x, (x + 1)};
|
||||
const int tint_symbol_9[4] = {5, 6, 3, (4 + 1)};
|
||||
const int tint_symbol_10[3][4] = {tint_symbol_8, tint_symbol_9, nonempty};
|
||||
const int nested_nonempty_with_expr[2][3][4] = {tint_symbol_10, nested_nonempty[1]};
|
||||
|
|
@ -22,7 +22,7 @@ void main() {
|
|||
const int subexpr_empty = 0;
|
||||
const int tint_symbol_12[4] = {1, 2, 3, 4};
|
||||
const int subexpr_nonempty = 3;
|
||||
const int tint_symbol_13[4] = {1, 42, (42 + 1), 4};
|
||||
const int tint_symbol_13[4] = {1, x, (x + 1), 4};
|
||||
const int subexpr_nonempty_with_expr = tint_symbol_13[2];
|
||||
const int tint_symbol_14[2][4] = (int[2][4])0;
|
||||
const int subexpr_nested_empty[4] = tint_symbol_14[1];
|
||||
|
|
@ -30,7 +30,7 @@ void main() {
|
|||
const int tint_symbol_16[4] = {5, 6, 7, 8};
|
||||
const int tint_symbol_17[2][4] = {tint_symbol_15, tint_symbol_16};
|
||||
const int subexpr_nested_nonempty[4] = tint_symbol_17[1];
|
||||
const int tint_symbol_18[4] = {1, 42, (42 + 1), 4};
|
||||
const int tint_symbol_18[4] = {1, x, (x + 1), 4};
|
||||
const int tint_symbol_19[2][4] = {tint_symbol_18, nested_nonempty[1][2]};
|
||||
const int subexpr_nested_nonempty_with_expr[4] = tint_symbol_19[1];
|
||||
return;
|
||||
|
|
|
|||
|
|
@ -18,7 +18,7 @@ kernel void tint_symbol() {
|
|||
int const x = 42;
|
||||
tint_array<int, 4> const empty = tint_array<int, 4>{};
|
||||
tint_array<int, 4> const nonempty = tint_array<int, 4>{1, 2, 3, 4};
|
||||
tint_array<int, 4> const nonempty_with_expr = tint_array<int, 4>{1, 42, as_type<int>((as_type<uint>(42) + as_type<uint>(1))), 4};
|
||||
tint_array<int, 4> const nonempty_with_expr = tint_array<int, 4>{1, x, as_type<int>((as_type<uint>(x) + as_type<uint>(1))), 4};
|
||||
tint_array<tint_array<tint_array<int, 4>, 3>, 2> const nested_empty = tint_array<tint_array<tint_array<int, 4>, 3>, 2>{};
|
||||
tint_array<int, 4> const tint_symbol_1 = tint_array<int, 4>{1, 2, 3, 4};
|
||||
tint_array<int, 4> const tint_symbol_2 = tint_array<int, 4>{5, 6, 7, 8};
|
||||
|
|
@ -29,7 +29,7 @@ kernel void tint_symbol() {
|
|||
tint_array<int, 4> const tint_symbol_7 = tint_array<int, 4>{21, 22, 23, 24};
|
||||
tint_array<tint_array<int, 4>, 3> const tint_symbol_8 = tint_array<tint_array<int, 4>, 3>{tint_symbol_5, tint_symbol_6, tint_symbol_7};
|
||||
tint_array<tint_array<tint_array<int, 4>, 3>, 2> const nested_nonempty = tint_array<tint_array<tint_array<int, 4>, 3>, 2>{tint_symbol_4, tint_symbol_8};
|
||||
tint_array<int, 4> const tint_symbol_9 = tint_array<int, 4>{1, 2, 42, as_type<int>((as_type<uint>(42) + as_type<uint>(1)))};
|
||||
tint_array<int, 4> const tint_symbol_9 = tint_array<int, 4>{1, 2, x, as_type<int>((as_type<uint>(x) + as_type<uint>(1)))};
|
||||
tint_array<int, 4> const tint_symbol_10 = tint_array<int, 4>{5, 6, 3, as_type<int>((as_type<uint>(4) + as_type<uint>(1)))};
|
||||
tint_array<tint_array<int, 4>, 3> const tint_symbol_11 = tint_array<tint_array<int, 4>, 3>{tint_symbol_9, tint_symbol_10, nonempty};
|
||||
tint_array<tint_array<tint_array<int, 4>, 3>, 2> const nested_nonempty_with_expr = tint_array<tint_array<tint_array<int, 4>, 3>, 2>{tint_symbol_11, nested_nonempty[1]};
|
||||
|
|
@ -37,7 +37,7 @@ kernel void tint_symbol() {
|
|||
int const subexpr_empty = 0;
|
||||
tint_array<int, 4> const tint_symbol_13 = tint_array<int, 4>{1, 2, 3, 4};
|
||||
int const subexpr_nonempty = 3;
|
||||
tint_array<int, 4> const tint_symbol_14 = tint_array<int, 4>{1, 42, as_type<int>((as_type<uint>(42) + as_type<uint>(1))), 4};
|
||||
tint_array<int, 4> const tint_symbol_14 = tint_array<int, 4>{1, x, as_type<int>((as_type<uint>(x) + as_type<uint>(1))), 4};
|
||||
int const subexpr_nonempty_with_expr = tint_symbol_14[2];
|
||||
tint_array<tint_array<int, 4>, 2> const tint_symbol_15 = tint_array<tint_array<int, 4>, 2>{};
|
||||
tint_array<int, 4> const subexpr_nested_empty = tint_symbol_15[1];
|
||||
|
|
@ -45,7 +45,7 @@ kernel void tint_symbol() {
|
|||
tint_array<int, 4> const tint_symbol_17 = tint_array<int, 4>{5, 6, 7, 8};
|
||||
tint_array<tint_array<int, 4>, 2> const tint_symbol_18 = tint_array<tint_array<int, 4>, 2>{tint_symbol_16, tint_symbol_17};
|
||||
tint_array<int, 4> const subexpr_nested_nonempty = tint_symbol_18[1];
|
||||
tint_array<int, 4> const tint_symbol_19 = tint_array<int, 4>{1, 42, as_type<int>((as_type<uint>(42) + as_type<uint>(1))), 4};
|
||||
tint_array<int, 4> const tint_symbol_19 = tint_array<int, 4>{1, x, as_type<int>((as_type<uint>(x) + as_type<uint>(1))), 4};
|
||||
tint_array<tint_array<int, 4>, 2> const tint_symbol_20 = tint_array<tint_array<int, 4>, 2>{tint_symbol_19, nested_nonempty[1][2]};
|
||||
tint_array<int, 4> const subexpr_nested_nonempty_with_expr = tint_symbol_20[1];
|
||||
return;
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
int const a = 1;
|
||||
int const b = 0;
|
||||
int const c = (1 / 0);
|
||||
int const c = (a / b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -79,7 +79,7 @@ void main_1() {
|
|||
}
|
||||
vec4 x_34 = x_29.vEyePosition;
|
||||
vec3 x_38 = vec3(0.0f);
|
||||
viewDirectionW = normalize((vec3(x_34.x, x_34.y, x_34.z) - vec3(0.0f)));
|
||||
viewDirectionW = normalize((vec3(x_34.x, x_34.y, x_34.z) - x_38));
|
||||
baseColor = vec4(1.0f);
|
||||
vec4 x_52 = x_49.vDiffuseColor;
|
||||
diffuseColor = vec3(x_52.x, x_52.y, x_52.z);
|
||||
|
|
@ -87,11 +87,11 @@ void main_1() {
|
|||
alpha = x_60;
|
||||
vec3 x_62 = vec3(0.0f);
|
||||
vec3 x_64 = vec3(0.0f);
|
||||
normalW = normalize(-(cross(dFdx(vec3(0.0f)), dFdy(vec3(0.0f)))));
|
||||
normalW = normalize(-(cross(dFdx(x_62), dFdy(x_64))));
|
||||
uvOffset = vec2(0.0f);
|
||||
vec4 x_74 = vec4(0.0f);
|
||||
vec4 x_76 = baseColor;
|
||||
vec3 x_78 = (vec3(x_76.x, x_76.y, x_76.z) * vec3(vec4(0.0f).x, vec4(0.0f).y, vec4(0.0f).z));
|
||||
vec3 x_78 = (vec3(x_76.x, x_76.y, x_76.z) * vec3(x_74.x, x_74.y, x_74.z));
|
||||
baseColor = vec4(x_78.x, x_78.y, x_78.z, baseColor.w);
|
||||
baseAmbientColor = vec3(1.0f);
|
||||
glossiness = 0.0f;
|
||||
|
|
|
|||
|
|
@ -51,7 +51,7 @@ void main_1() {
|
|||
}
|
||||
const float4 x_34 = asfloat(x_29[0]);
|
||||
const float3 x_38 = (0.0f).xxx;
|
||||
viewDirectionW = normalize((float3(x_34.x, x_34.y, x_34.z) - (0.0f).xxx));
|
||||
viewDirectionW = normalize((float3(x_34.x, x_34.y, x_34.z) - x_38));
|
||||
baseColor = (1.0f).xxxx;
|
||||
const float4 x_52 = asfloat(x_49[0]);
|
||||
diffuseColor = float3(x_52.x, x_52.y, x_52.z);
|
||||
|
|
@ -59,11 +59,11 @@ void main_1() {
|
|||
alpha = x_60;
|
||||
const float3 x_62 = (0.0f).xxx;
|
||||
const float3 x_64 = (0.0f).xxx;
|
||||
normalW = normalize(-(cross(ddx((0.0f).xxx), ddy((0.0f).xxx))));
|
||||
normalW = normalize(-(cross(ddx(x_62), ddy(x_64))));
|
||||
uvOffset = (0.0f).xx;
|
||||
const float4 x_74 = (0.0f).xxxx;
|
||||
const float4 x_76 = baseColor;
|
||||
const float3 x_78 = (float3(x_76.x, x_76.y, x_76.z) * float3((0.0f).xxxx.x, (0.0f).xxxx.y, (0.0f).xxxx.z));
|
||||
const float3 x_78 = (float3(x_76.x, x_76.y, x_76.z) * float3(x_74.x, x_74.y, x_74.z));
|
||||
baseColor = float4(x_78.x, x_78.y, x_78.z, baseColor.w);
|
||||
baseAmbientColor = (1.0f).xxx;
|
||||
glossiness = 0.0f;
|
||||
|
|
|
|||
|
|
@ -69,7 +69,7 @@ void main_1(thread float* const tint_symbol_7, thread bool* const tint_symbol_8,
|
|||
}
|
||||
float4 const x_34 = (*(tint_symbol_10)).vEyePosition;
|
||||
float3 const x_38 = float3(0.0f);
|
||||
viewDirectionW = normalize((float3(x_34[0], x_34[1], x_34[2]) - float3(0.0f)));
|
||||
viewDirectionW = normalize((float3(x_34[0], x_34[1], x_34[2]) - x_38));
|
||||
baseColor = float4(1.0f);
|
||||
float4 const x_52 = (*(tint_symbol_11)).vDiffuseColor;
|
||||
diffuseColor = float3(x_52[0], x_52[1], x_52[2]);
|
||||
|
|
@ -77,11 +77,11 @@ void main_1(thread float* const tint_symbol_7, thread bool* const tint_symbol_8,
|
|||
alpha = x_60;
|
||||
float3 const x_62 = float3(0.0f);
|
||||
float3 const x_64 = float3(0.0f);
|
||||
normalW = normalize(-(cross(dfdx(float3(0.0f)), dfdy(float3(0.0f)))));
|
||||
normalW = normalize(-(cross(dfdx(x_62), dfdy(x_64))));
|
||||
uvOffset = float2(0.0f);
|
||||
float4 const x_74 = float4(0.0f);
|
||||
float4 const x_76 = baseColor;
|
||||
float3 const x_78 = (float3(x_76[0], x_76[1], x_76[2]) * float3(float4(0.0f)[0], float4(0.0f)[1], float4(0.0f)[2]));
|
||||
float3 const x_78 = (float3(x_76[0], x_76[1], x_76[2]) * float3(x_74[0], x_74[1], x_74[2]));
|
||||
float4 const x_79 = baseColor;
|
||||
baseColor = float4(x_78[0], x_78[1], x_78[2], x_79[3]);
|
||||
baseAmbientColor = float3(1.0f);
|
||||
|
|
|
|||
|
|
@ -81,8 +81,8 @@ void tint_symbol_2(uvec3 GlobalInvocationID) {
|
|||
{
|
||||
for(int y_1 = 0; (y_1 < 2); y_1 = (y_1 + 1)) {
|
||||
{
|
||||
for(int x_1 = 0; (x_1 < 2); x_1 = (x_1 + 1)) {
|
||||
ivec2 tilePixel0Idx = ivec2((x_1 * 16), (y_1 * 16));
|
||||
for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = (x_1 + 1)) {
|
||||
ivec2 tilePixel0Idx = ivec2((x_1 * TILE_SIZE), (y_1 * TILE_SIZE));
|
||||
vec2 floorCoord = (((2.0f * vec2(tilePixel0Idx)) / uniforms.fullScreenSize.xy) - vec2(1.0f));
|
||||
vec2 ceilCoord = (((2.0f * vec2((tilePixel0Idx + ivec2(16)))) / uniforms.fullScreenSize.xy) - vec2(1.0f));
|
||||
vec2 viewFloorCoord = vec2((((-(viewNear) * floorCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
|
||||
|
|
@ -115,7 +115,7 @@ void tint_symbol_2(uvec3 GlobalInvocationID) {
|
|||
}
|
||||
}
|
||||
if ((dp >= 0.0f)) {
|
||||
uint tileId = uint((x_1 + (y_1 * 2)));
|
||||
uint tileId = uint((x_1 + (y_1 * TILE_COUNT_X)));
|
||||
bool tint_tmp = (tileId < 0u);
|
||||
if (!tint_tmp) {
|
||||
tint_tmp = (tileId >= config.numTiles);
|
||||
|
|
|
|||
|
|
@ -55,8 +55,8 @@ void main_inner(uint3 GlobalInvocationID) {
|
|||
{
|
||||
[loop] for(int y_1 = 0; (y_1 < 2); y_1 = (y_1 + 1)) {
|
||||
{
|
||||
[loop] for(int x_1 = 0; (x_1 < 2); x_1 = (x_1 + 1)) {
|
||||
int2 tilePixel0Idx = int2((x_1 * 16), (y_1 * 16));
|
||||
[loop] for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = (x_1 + 1)) {
|
||||
int2 tilePixel0Idx = int2((x_1 * TILE_SIZE), (y_1 * TILE_SIZE));
|
||||
float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / asfloat(uniforms[10]).xy) - (1.0f).xx);
|
||||
float2 ceilCoord = (((2.0f * float2((tilePixel0Idx + (16).xx))) / asfloat(uniforms[10]).xy) - (1.0f).xx);
|
||||
float2 viewFloorCoord = float2((((-(viewNear) * floorCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
|
||||
|
|
@ -89,7 +89,7 @@ void main_inner(uint3 GlobalInvocationID) {
|
|||
}
|
||||
}
|
||||
if ((dp >= 0.0f)) {
|
||||
uint tileId = uint((x_1 + (y_1 * 2)));
|
||||
uint tileId = uint((x_1 + (y_1 * TILE_COUNT_X)));
|
||||
bool tint_tmp = (tileId < 0u);
|
||||
if (!tint_tmp) {
|
||||
tint_tmp = (tileId >= config[0].y);
|
||||
|
|
|
|||
|
|
@ -84,9 +84,9 @@ void tint_symbol_inner(uint3 GlobalInvocationID, const constant Config* const ti
|
|||
int const TILE_SIZE = 16;
|
||||
int const TILE_COUNT_X = 2;
|
||||
int const TILE_COUNT_Y = 2;
|
||||
for(int y_1 = 0; (y_1 < 2); y_1 = as_type<int>((as_type<uint>(y_1) + as_type<uint>(1)))) {
|
||||
for(int x_1 = 0; (x_1 < 2); x_1 = as_type<int>((as_type<uint>(x_1) + as_type<uint>(1)))) {
|
||||
int2 tilePixel0Idx = int2(as_type<int>((as_type<uint>(x_1) * as_type<uint>(16))), as_type<int>((as_type<uint>(y_1) * as_type<uint>(16))));
|
||||
for(int y_1 = 0; (y_1 < TILE_COUNT_Y); y_1 = as_type<int>((as_type<uint>(y_1) + as_type<uint>(1)))) {
|
||||
for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = as_type<int>((as_type<uint>(x_1) + as_type<uint>(1)))) {
|
||||
int2 tilePixel0Idx = int2(as_type<int>((as_type<uint>(x_1) * as_type<uint>(TILE_SIZE))), as_type<int>((as_type<uint>(y_1) * as_type<uint>(TILE_SIZE))));
|
||||
float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
|
||||
float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(tilePixel0Idx) + as_type<uint2>(int2(16)))))) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
|
||||
float2 viewFloorCoord = float2((((-(viewNear) * floorCoord[0]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord[1]) - (M[2][1] * viewNear)) / M[1][1]));
|
||||
|
|
@ -117,7 +117,7 @@ void tint_symbol_inner(uint3 GlobalInvocationID, const constant Config* const ti
|
|||
dp = (dp + fmin(0.0f, dot(p, frustumPlanes[i])));
|
||||
}
|
||||
if ((dp >= 0.0f)) {
|
||||
uint tileId = uint(as_type<int>((as_type<uint>(x_1) + as_type<uint>(as_type<int>((as_type<uint>(y_1) * as_type<uint>(2)))))));
|
||||
uint tileId = uint(as_type<int>((as_type<uint>(x_1) + as_type<uint>(as_type<int>((as_type<uint>(y_1) * as_type<uint>(TILE_COUNT_X)))))));
|
||||
if (((tileId < 0u) || (tileId >= (*(tint_symbol_1)).numTiles))) {
|
||||
continue;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
using namespace metal;
|
||||
kernel void compute_main() {
|
||||
float const a = 1.230000019f;
|
||||
float b = fmax(1.230000019f, 1.17549435e-38f);
|
||||
float b = fmax(a, 1.17549435e-38f);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -8,6 +8,6 @@ void f() {
|
|||
int i = 0;
|
||||
int j = 0;
|
||||
mat2 m = mat2(vec2(1.0f, 2.0f), vec2(3.0f, 4.0f));
|
||||
float f_1 = mat2(vec2(1.0f, 2.0f), vec2(3.0f, 4.0f))[i][j];
|
||||
float f_1 = m[i][j];
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -7,5 +7,5 @@ void f() {
|
|||
int i = 0;
|
||||
int j = 0;
|
||||
const float2x2 m = float2x2(float2(1.0f, 2.0f), float2(3.0f, 4.0f));
|
||||
const float f_1 = float2x2(float2(1.0f, 2.0f), float2(3.0f, 4.0f))[i][j];
|
||||
const float f_1 = m[i][j];
|
||||
}
|
||||
|
|
|
|||
|
|
@ -5,6 +5,6 @@ void f() {
|
|||
int i = 0;
|
||||
int j = 0;
|
||||
float2x2 const m = float2x2(float2(1.0f, 2.0f), float2(3.0f, 4.0f));
|
||||
float const f_1 = float2x2(float2(1.0f, 2.0f), float2(3.0f, 4.0f))[i][j];
|
||||
float const f_1 = m[i][j];
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -6,7 +6,7 @@ layout(binding = 1, std430) buffer Result_1 {
|
|||
} result;
|
||||
uniform highp sampler2D tex_1;
|
||||
void tint_symbol(uvec3 GlobalInvocationId) {
|
||||
result.values[((GlobalInvocationId.y * 128u) + GlobalInvocationId.x)] = texelFetch(tex_1, ivec2(int(GlobalInvocationId.x), int(GlobalInvocationId.y)), 0).x;
|
||||
result.values[((GlobalInvocationId.y * width) + GlobalInvocationId.x)] = texelFetch(tex_1, ivec2(int(GlobalInvocationId.x), int(GlobalInvocationId.y)), 0).x;
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -8,7 +8,7 @@ struct tint_symbol_1 {
|
|||
};
|
||||
|
||||
void main_inner(uint3 GlobalInvocationId) {
|
||||
result.Store((4u * ((GlobalInvocationId.y * 128u) + GlobalInvocationId.x)), asuint(tex.Load(int3(int(GlobalInvocationId.x), int(GlobalInvocationId.y), 0)).x));
|
||||
result.Store((4u * ((GlobalInvocationId.y * width) + GlobalInvocationId.x)), asuint(tex.Load(int3(int(GlobalInvocationId.x), int(GlobalInvocationId.y), 0)).x));
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
|
|
|
|||
|
|
@ -21,7 +21,7 @@ struct Result {
|
|||
constant uint width = 128u;
|
||||
|
||||
void tint_symbol_inner(uint3 GlobalInvocationId, device Result* const tint_symbol_1, depth2d<float, access::sample> tint_symbol_2) {
|
||||
(*(tint_symbol_1)).values[((GlobalInvocationId[1] * 128u) + GlobalInvocationId[0])] = tint_symbol_2.read(uint2(int2(int(GlobalInvocationId[0]), int(GlobalInvocationId[1]))), 0);
|
||||
(*(tint_symbol_1)).values[((GlobalInvocationId[1] * width) + GlobalInvocationId[0])] = tint_symbol_2.read(uint2(int2(int(GlobalInvocationId[0]), int(GlobalInvocationId[1]))), 0);
|
||||
}
|
||||
|
||||
kernel void tint_symbol(device Result* tint_symbol_3 [[buffer(0)]], depth2d<float, access::sample> tint_symbol_4 [[texture(0)]], uint3 GlobalInvocationId [[thread_position_in_grid]]) {
|
||||
|
|
|
|||
|
|
@ -46,7 +46,7 @@ void tint_symbol_1(uvec3 GlobalInvocationID) {
|
|||
if ((tint_tmp)) {
|
||||
bool tint_tmp_3 = success;
|
||||
if (tint_tmp_3) {
|
||||
tint_tmp_3 = all(equal(texelFetch(dst_1, ivec2(dstTexCoord), 0), vec4(0.0f, 1.0f, 0.0f, 1.0f)));
|
||||
tint_tmp_3 = all(equal(texelFetch(dst_1, ivec2(dstTexCoord), 0), nonCoveredColor));
|
||||
}
|
||||
success = (tint_tmp_3);
|
||||
} else {
|
||||
|
|
|
|||
|
|
@ -38,7 +38,7 @@ void main_inner(uint3 GlobalInvocationID) {
|
|||
if ((tint_tmp_2)) {
|
||||
bool tint_tmp_5 = success;
|
||||
if (tint_tmp_5) {
|
||||
tint_tmp_5 = all((tint_symbol.Load(int3(int2(dstTexCoord), 0)) == float4(0.0f, 1.0f, 0.0f, 1.0f)));
|
||||
tint_tmp_5 = all((tint_symbol.Load(int3(int2(dstTexCoord), 0)) == nonCoveredColor));
|
||||
}
|
||||
success = (tint_tmp_5);
|
||||
} else {
|
||||
|
|
|
|||
|
|
@ -37,7 +37,7 @@ void tint_symbol_inner(uint3 GlobalInvocationID, texture2d<float, access::sample
|
|||
float4 const nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f);
|
||||
bool success = true;
|
||||
if (((((dstTexCoord[0] < (*(tint_symbol_9)).dstCopyOrigin[0]) || (dstTexCoord[1] < (*(tint_symbol_9)).dstCopyOrigin[1])) || (dstTexCoord[0] >= ((*(tint_symbol_9)).dstCopyOrigin[0] + (*(tint_symbol_9)).copySize[0]))) || (dstTexCoord[1] >= ((*(tint_symbol_9)).dstCopyOrigin[1] + (*(tint_symbol_9)).copySize[1])))) {
|
||||
success = (success && all((tint_symbol_8.read(uint2(int2(dstTexCoord)), 0) == float4(0.0f, 1.0f, 0.0f, 1.0f))));
|
||||
success = (success && all((tint_symbol_8.read(uint2(int2(dstTexCoord)), 0) == nonCoveredColor)));
|
||||
} else {
|
||||
uint2 srcTexCoord = ((dstTexCoord - (*(tint_symbol_9)).dstCopyOrigin) + (*(tint_symbol_9)).srcCopyOrigin);
|
||||
if (((*(tint_symbol_9)).dstTextureFlipY == 1u)) {
|
||||
|
|
|
|||
|
|
@ -71,32 +71,32 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
|
|||
}
|
||||
}
|
||||
barrier();
|
||||
uint tileRow = (local_id.y * 4u);
|
||||
uint tileCol = (local_id.x * 4u);
|
||||
uint globalRow = (global_id.y * 4u);
|
||||
uint globalCol = (global_id.x * 4u);
|
||||
uint numTiles = (((uniforms.dimInner - 1u) / 64u) + 1u);
|
||||
uint tileRow = (local_id.y * RowPerThread);
|
||||
uint tileCol = (local_id.x * ColPerThread);
|
||||
uint globalRow = (global_id.y * RowPerThread);
|
||||
uint globalCol = (global_id.x * ColPerThread);
|
||||
uint numTiles = (((uniforms.dimInner - 1u) / TileInner) + 1u);
|
||||
float acc[16] = float[16](0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
|
||||
float ACached = 0.0f;
|
||||
float BCached[4] = float[4](0.0f, 0.0f, 0.0f, 0.0f);
|
||||
{
|
||||
for(uint index = 0u; (index < (4u * 4u)); index = (index + 1u)) {
|
||||
for(uint index = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
|
||||
acc[index] = 0.0f;
|
||||
}
|
||||
}
|
||||
uint ColPerThreadA = (64u / 16u);
|
||||
uint ColPerThreadA = (TileInner / 16u);
|
||||
uint tileColA = (local_id.x * ColPerThreadA);
|
||||
uint RowPerThreadB = (64u / 16u);
|
||||
uint RowPerThreadB = (TileInner / 16u);
|
||||
uint tileRowB = (local_id.y * RowPerThreadB);
|
||||
{
|
||||
for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
|
||||
{
|
||||
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
{
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
|
||||
uint inputRow = (tileRow + innerRow);
|
||||
uint inputCol = (tileColA + innerCol);
|
||||
float tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
|
||||
float tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
|
||||
mm_Asub[inputRow][inputCol] = tint_symbol_1;
|
||||
}
|
||||
}
|
||||
|
|
@ -105,10 +105,10 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
|
|||
{
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
|
||||
{
|
||||
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
uint inputRow = (tileRowB + innerRow);
|
||||
uint inputCol = (tileCol + innerCol);
|
||||
float tint_symbol_2 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
|
||||
float tint_symbol_2 = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
|
||||
mm_Bsub[innerCol][inputCol] = tint_symbol_2;
|
||||
}
|
||||
}
|
||||
|
|
@ -116,18 +116,18 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
|
|||
}
|
||||
barrier();
|
||||
{
|
||||
for(uint k = 0u; (k < 64u); k = (k + 1u)) {
|
||||
for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
|
||||
{
|
||||
for(uint inner = 0u; (inner < 4u); inner = (inner + 1u)) {
|
||||
for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
|
||||
BCached[inner] = mm_Bsub[k][(tileCol + inner)];
|
||||
}
|
||||
}
|
||||
{
|
||||
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
ACached = mm_Asub[(tileRow + innerRow)][k];
|
||||
{
|
||||
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
uint index = ((innerRow * 4u) + innerCol);
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
uint index = ((innerRow * ColPerThread) + innerCol);
|
||||
acc[index] = (acc[index] + (ACached * BCached[innerCol]));
|
||||
}
|
||||
}
|
||||
|
|
@ -139,10 +139,10 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
|
|||
}
|
||||
}
|
||||
{
|
||||
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
{
|
||||
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
uint index = ((innerRow * 4u) + innerCol);
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
uint index = ((innerRow * ColPerThread) + innerCol);
|
||||
mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index]);
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -65,32 +65,32 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
|
|||
}
|
||||
}
|
||||
GroupMemoryBarrierWithGroupSync();
|
||||
const uint tileRow = (local_id.y * 4u);
|
||||
const uint tileCol = (local_id.x * 4u);
|
||||
const uint globalRow = (global_id.y * 4u);
|
||||
const uint globalCol = (global_id.x * 4u);
|
||||
const uint numTiles = (((uniforms[0].y - 1u) / 64u) + 1u);
|
||||
const uint tileRow = (local_id.y * RowPerThread);
|
||||
const uint tileCol = (local_id.x * ColPerThread);
|
||||
const uint globalRow = (global_id.y * RowPerThread);
|
||||
const uint globalCol = (global_id.x * ColPerThread);
|
||||
const uint numTiles = (((uniforms[0].y - 1u) / TileInner) + 1u);
|
||||
float acc[16] = (float[16])0;
|
||||
float ACached = 0.0f;
|
||||
float BCached[4] = (float[4])0;
|
||||
{
|
||||
[loop] for(uint index = 0u; (index < (4u * 4u)); index = (index + 1u)) {
|
||||
[loop] for(uint index = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
|
||||
acc[index] = 0.0f;
|
||||
}
|
||||
}
|
||||
const uint ColPerThreadA = (64u / 16u);
|
||||
const uint ColPerThreadA = (TileInner / 16u);
|
||||
const uint tileColA = (local_id.x * ColPerThreadA);
|
||||
const uint RowPerThreadB = (64u / 16u);
|
||||
const uint RowPerThreadB = (TileInner / 16u);
|
||||
const uint tileRowB = (local_id.y * RowPerThreadB);
|
||||
{
|
||||
[loop] for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
|
||||
{
|
||||
[loop] for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
[loop] for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
{
|
||||
[loop] for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
|
||||
const uint inputRow = (tileRow + innerRow);
|
||||
const uint inputCol = (tileColA + innerCol);
|
||||
const float tint_symbol_2 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
|
||||
const float tint_symbol_2 = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
|
||||
mm_Asub[inputRow][inputCol] = tint_symbol_2;
|
||||
}
|
||||
}
|
||||
|
|
@ -99,10 +99,10 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
|
|||
{
|
||||
[loop] for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
|
||||
{
|
||||
[loop] for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
[loop] for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
const uint inputRow = (tileRowB + innerRow);
|
||||
const uint inputCol = (tileCol + innerCol);
|
||||
const float tint_symbol_3 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
|
||||
const float tint_symbol_3 = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
|
||||
mm_Bsub[innerCol][inputCol] = tint_symbol_3;
|
||||
}
|
||||
}
|
||||
|
|
@ -110,18 +110,18 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
|
|||
}
|
||||
GroupMemoryBarrierWithGroupSync();
|
||||
{
|
||||
[loop] for(uint k = 0u; (k < 64u); k = (k + 1u)) {
|
||||
[loop] for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
|
||||
{
|
||||
[loop] for(uint inner = 0u; (inner < 4u); inner = (inner + 1u)) {
|
||||
[loop] for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
|
||||
BCached[inner] = mm_Bsub[k][(tileCol + inner)];
|
||||
}
|
||||
}
|
||||
{
|
||||
[loop] for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
[loop] for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
ACached = mm_Asub[(tileRow + innerRow)][k];
|
||||
{
|
||||
[loop] for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
const uint index = ((innerRow * 4u) + innerCol);
|
||||
[loop] for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
const uint index = ((innerRow * ColPerThread) + innerCol);
|
||||
acc[index] = (acc[index] + (ACached * BCached[innerCol]));
|
||||
}
|
||||
}
|
||||
|
|
@ -133,10 +133,10 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
|
|||
}
|
||||
}
|
||||
{
|
||||
[loop] for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
[loop] for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
{
|
||||
[loop] for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
const uint index = ((innerRow * 4u) + innerCol);
|
||||
[loop] for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
const uint index = ((innerRow * ColPerThread) + innerCol);
|
||||
mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index]);
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -65,56 +65,56 @@ void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_in
|
|||
(*(tint_symbol_10))[i][i_1] = 0.0f;
|
||||
}
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
uint const tileRow = (local_id[1] * 4u);
|
||||
uint const tileCol = (local_id[0] * 4u);
|
||||
uint const globalRow = (global_id[1] * 4u);
|
||||
uint const globalCol = (global_id[0] * 4u);
|
||||
uint const numTiles = ((((*(tint_symbol_11)).dimInner - 1u) / 64u) + 1u);
|
||||
uint const tileRow = (local_id[1] * RowPerThread);
|
||||
uint const tileCol = (local_id[0] * ColPerThread);
|
||||
uint const globalRow = (global_id[1] * RowPerThread);
|
||||
uint const globalCol = (global_id[0] * ColPerThread);
|
||||
uint const numTiles = ((((*(tint_symbol_11)).dimInner - 1u) / TileInner) + 1u);
|
||||
tint_array<float, 16> acc = {};
|
||||
float ACached = 0.0f;
|
||||
tint_array<float, 4> BCached = {};
|
||||
for(uint index = 0u; (index < (4u * 4u)); index = (index + 1u)) {
|
||||
for(uint index = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
|
||||
acc[index] = 0.0f;
|
||||
}
|
||||
uint const ColPerThreadA = (64u / 16u);
|
||||
uint const ColPerThreadA = (TileInner / 16u);
|
||||
uint const tileColA = (local_id[0] * ColPerThreadA);
|
||||
uint const RowPerThreadB = (64u / 16u);
|
||||
uint const RowPerThreadB = (TileInner / 16u);
|
||||
uint const tileRowB = (local_id[1] * RowPerThreadB);
|
||||
for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
|
||||
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
|
||||
uint const inputRow = (tileRow + innerRow);
|
||||
uint const inputCol = (tileColA + innerCol);
|
||||
float const tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol), tint_symbol_11, tint_symbol_12);
|
||||
float const tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol), tint_symbol_11, tint_symbol_12);
|
||||
(*(tint_symbol_9))[inputRow][inputCol] = tint_symbol_1;
|
||||
}
|
||||
}
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
uint const inputRow = (tileRowB + innerRow);
|
||||
uint const inputCol = (tileCol + innerCol);
|
||||
float const tint_symbol_2 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol), tint_symbol_11, tint_symbol_13);
|
||||
float const tint_symbol_2 = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol), tint_symbol_11, tint_symbol_13);
|
||||
(*(tint_symbol_10))[innerCol][inputCol] = tint_symbol_2;
|
||||
}
|
||||
}
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
for(uint k = 0u; (k < 64u); k = (k + 1u)) {
|
||||
for(uint inner = 0u; (inner < 4u); inner = (inner + 1u)) {
|
||||
for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
|
||||
for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
|
||||
BCached[inner] = (*(tint_symbol_10))[k][(tileCol + inner)];
|
||||
}
|
||||
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
ACached = (*(tint_symbol_9))[(tileRow + innerRow)][k];
|
||||
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
uint const index = ((innerRow * 4u) + innerCol);
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
uint const index = ((innerRow * ColPerThread) + innerCol);
|
||||
acc[index] = (acc[index] + (ACached * BCached[innerCol]));
|
||||
}
|
||||
}
|
||||
}
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
}
|
||||
for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
|
||||
uint const index = ((innerRow * 4u) + innerCol);
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
uint const index = ((innerRow * ColPerThread) + innerCol);
|
||||
mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index], tint_symbol_11, tint_symbol_14);
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
mat3 a = mat3(vec3(1.0f, 2.0f, 3.0f), vec3(4.0f, 5.0f, 6.0f), vec3(7.0f, 8.0f, 9.0f));
|
||||
mat3 b = mat3(vec3(-1.0f, -2.0f, -3.0f), vec3(-4.0f, -5.0f, -6.0f), vec3(-7.0f, -8.0f, -9.0f));
|
||||
mat3 r = (mat3(vec3(1.0f, 2.0f, 3.0f), vec3(4.0f, 5.0f, 6.0f), vec3(7.0f, 8.0f, 9.0f)) + mat3(vec3(-1.0f, -2.0f, -3.0f), vec3(-4.0f, -5.0f, -6.0f), vec3(-7.0f, -8.0f, -9.0f)));
|
||||
mat3 r = (a + b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const float3x3 a = float3x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f));
|
||||
const float3x3 b = float3x3(float3(-1.0f, -2.0f, -3.0f), float3(-4.0f, -5.0f, -6.0f), float3(-7.0f, -8.0f, -9.0f));
|
||||
const float3x3 r = (float3x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f)) + float3x3(float3(-1.0f, -2.0f, -3.0f), float3(-4.0f, -5.0f, -6.0f), float3(-7.0f, -8.0f, -9.0f)));
|
||||
const float3x3 r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
float3x3 const a = float3x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f));
|
||||
float3x3 const b = float3x3(float3(-1.0f, -2.0f, -3.0f), float3(-4.0f, -5.0f, -6.0f), float3(-7.0f, -8.0f, -9.0f));
|
||||
float3x3 const r = (float3x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f)) + float3x3(float3(-1.0f, -2.0f, -3.0f), float3(-4.0f, -5.0f, -6.0f), float3(-7.0f, -8.0f, -9.0f)));
|
||||
float3x3 const r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
float const a = 1.0f;
|
||||
float const b = 2.0f;
|
||||
float const r = (1.0f + 2.0f);
|
||||
float const r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
int const a = 1;
|
||||
int const b = 2;
|
||||
int const r = as_type<int>((as_type<uint>(1) + as_type<uint>(2)));
|
||||
int const r = as_type<int>((as_type<uint>(a) + as_type<uint>(b)));
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
uint const a = 1u;
|
||||
uint const b = 2u;
|
||||
uint const r = (1u + 2u);
|
||||
uint const r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
float a = 4.0f;
|
||||
vec3 b = vec3(1.0f, 2.0f, 3.0f);
|
||||
vec3 r = (4.0f + vec3(1.0f, 2.0f, 3.0f));
|
||||
vec3 r = (a + b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const float a = 4.0f;
|
||||
const float3 b = float3(1.0f, 2.0f, 3.0f);
|
||||
const float3 r = (4.0f + float3(1.0f, 2.0f, 3.0f));
|
||||
const float3 r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
float const a = 4.0f;
|
||||
float3 const b = float3(1.0f, 2.0f, 3.0f);
|
||||
float3 const r = (4.0f + float3(1.0f, 2.0f, 3.0f));
|
||||
float3 const r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
int a = 4;
|
||||
ivec3 b = ivec3(1, 2, 3);
|
||||
ivec3 r = (4 + ivec3(1, 2, 3));
|
||||
ivec3 r = (a + b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const int a = 4;
|
||||
const int3 b = int3(1, 2, 3);
|
||||
const int3 r = (4 + int3(1, 2, 3));
|
||||
const int3 r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
int const a = 4;
|
||||
int3 const b = int3(1, 2, 3);
|
||||
int3 const r = as_type<int3>((as_type<uint>(4) + as_type<uint3>(int3(1, 2, 3))));
|
||||
int3 const r = as_type<int3>((as_type<uint>(a) + as_type<uint3>(b)));
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
uint a = 4u;
|
||||
uvec3 b = uvec3(1u, 2u, 3u);
|
||||
uvec3 r = (4u + uvec3(1u, 2u, 3u));
|
||||
uvec3 r = (a + b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const uint a = 4u;
|
||||
const uint3 b = uint3(1u, 2u, 3u);
|
||||
const uint3 r = (4u + uint3(1u, 2u, 3u));
|
||||
const uint3 r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
uint const a = 4u;
|
||||
uint3 const b = uint3(1u, 2u, 3u);
|
||||
uint3 const r = (4u + uint3(1u, 2u, 3u));
|
||||
uint3 const r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -2,7 +2,7 @@
|
|||
|
||||
void f() {
|
||||
vec3 a = vec3(1.0f, 2.0f, 3.0f);
|
||||
vec3 r = (vec3(1.0f, 2.0f, 3.0f) + 4.0f);
|
||||
vec3 r = (a + 4.0f);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -1,6 +1,6 @@
|
|||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
const float3 a = float3(1.0f, 2.0f, 3.0f);
|
||||
const float3 r = (float3(1.0f, 2.0f, 3.0f) + 4.0f);
|
||||
const float3 r = (a + 4.0f);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
float3 const a = float3(1.0f, 2.0f, 3.0f);
|
||||
float const b = 4.0f;
|
||||
float3 const r = (float3(1.0f, 2.0f, 3.0f) + 4.0f);
|
||||
float3 const r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -2,7 +2,7 @@
|
|||
|
||||
void f() {
|
||||
ivec3 a = ivec3(1, 2, 3);
|
||||
ivec3 r = (ivec3(1, 2, 3) + 4);
|
||||
ivec3 r = (a + 4);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -1,6 +1,6 @@
|
|||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
const int3 a = int3(1, 2, 3);
|
||||
const int3 r = (int3(1, 2, 3) + 4);
|
||||
const int3 r = (a + 4);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
int3 const a = int3(1, 2, 3);
|
||||
int const b = 4;
|
||||
int3 const r = as_type<int3>((as_type<uint3>(int3(1, 2, 3)) + as_type<uint>(4)));
|
||||
int3 const r = as_type<int3>((as_type<uint3>(a) + as_type<uint>(b)));
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -2,7 +2,7 @@
|
|||
|
||||
void f() {
|
||||
uvec3 a = uvec3(1u, 2u, 3u);
|
||||
uvec3 r = (uvec3(1u, 2u, 3u) + 4u);
|
||||
uvec3 r = (a + 4u);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -1,6 +1,6 @@
|
|||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
const uint3 a = uint3(1u, 2u, 3u);
|
||||
const uint3 r = (uint3(1u, 2u, 3u) + 4u);
|
||||
const uint3 r = (a + 4u);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
uint3 const a = uint3(1u, 2u, 3u);
|
||||
uint const b = 4u;
|
||||
uint3 const r = (uint3(1u, 2u, 3u) + 4u);
|
||||
uint3 const r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
vec3 a = vec3(1.0f, 2.0f, 3.0f);
|
||||
vec3 b = vec3(4.0f, 5.0f, 6.0f);
|
||||
vec3 r = (vec3(1.0f, 2.0f, 3.0f) + vec3(4.0f, 5.0f, 6.0f));
|
||||
vec3 r = (a + b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const float3 a = float3(1.0f, 2.0f, 3.0f);
|
||||
const float3 b = float3(4.0f, 5.0f, 6.0f);
|
||||
const float3 r = (float3(1.0f, 2.0f, 3.0f) + float3(4.0f, 5.0f, 6.0f));
|
||||
const float3 r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
float3 const a = float3(1.0f, 2.0f, 3.0f);
|
||||
float3 const b = float3(4.0f, 5.0f, 6.0f);
|
||||
float3 const r = (float3(1.0f, 2.0f, 3.0f) + float3(4.0f, 5.0f, 6.0f));
|
||||
float3 const r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
ivec3 a = ivec3(1, 2, 3);
|
||||
ivec3 b = ivec3(4, 5, 6);
|
||||
ivec3 r = (ivec3(1, 2, 3) + ivec3(4, 5, 6));
|
||||
ivec3 r = (a + b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const int3 a = int3(1, 2, 3);
|
||||
const int3 b = int3(4, 5, 6);
|
||||
const int3 r = (int3(1, 2, 3) + int3(4, 5, 6));
|
||||
const int3 r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
int3 const a = int3(1, 2, 3);
|
||||
int3 const b = int3(4, 5, 6);
|
||||
int3 const r = as_type<int3>((as_type<uint3>(int3(1, 2, 3)) + as_type<uint3>(int3(4, 5, 6))));
|
||||
int3 const r = as_type<int3>((as_type<uint3>(a) + as_type<uint3>(b)));
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
uvec3 a = uvec3(1u, 2u, 3u);
|
||||
uvec3 b = uvec3(4u, 5u, 6u);
|
||||
uvec3 r = (uvec3(1u, 2u, 3u) + uvec3(4u, 5u, 6u));
|
||||
uvec3 r = (a + b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const uint3 a = uint3(1u, 2u, 3u);
|
||||
const uint3 b = uint3(4u, 5u, 6u);
|
||||
const uint3 r = (uint3(1u, 2u, 3u) + uint3(4u, 5u, 6u));
|
||||
const uint3 r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
uint3 const a = uint3(1u, 2u, 3u);
|
||||
uint3 const b = uint3(4u, 5u, 6u);
|
||||
uint3 const r = (uint3(1u, 2u, 3u) + uint3(4u, 5u, 6u));
|
||||
uint3 const r = (a + b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
bool const a = true;
|
||||
bool const b = false;
|
||||
bool const r = bool(true & false);
|
||||
bool const r = bool(a & b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
int const a = 1;
|
||||
int const b = 2;
|
||||
int const r = (1 & 2);
|
||||
int const r = (a & b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
uint const a = 1u;
|
||||
uint const b = 2u;
|
||||
uint const r = (1u & 2u);
|
||||
uint const r = (a & b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
bvec3 a = bvec3(true, true, false);
|
||||
bvec3 b = bvec3(true, false, true);
|
||||
bvec3 r = bvec3(uvec3(bvec3(true, true, false)) & uvec3(bvec3(true, false, true)));
|
||||
bvec3 r = bvec3(uvec3(a) & uvec3(b));
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const bool3 a = bool3(true, true, false);
|
||||
const bool3 b = bool3(true, false, true);
|
||||
const bool3 r = (bool3(true, true, false) & bool3(true, false, true));
|
||||
const bool3 r = (a & b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
bool3 const a = bool3(true, true, false);
|
||||
bool3 const b = bool3(true, false, true);
|
||||
bool3 const r = (bool3(true, true, false) & bool3(true, false, true));
|
||||
bool3 const r = (a & b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
ivec3 a = ivec3(1, 2, 3);
|
||||
ivec3 b = ivec3(4, 5, 6);
|
||||
ivec3 r = (ivec3(1, 2, 3) & ivec3(4, 5, 6));
|
||||
ivec3 r = (a & b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const int3 a = int3(1, 2, 3);
|
||||
const int3 b = int3(4, 5, 6);
|
||||
const int3 r = (int3(1, 2, 3) & int3(4, 5, 6));
|
||||
const int3 r = (a & b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
int3 const a = int3(1, 2, 3);
|
||||
int3 const b = int3(4, 5, 6);
|
||||
int3 const r = (int3(1, 2, 3) & int3(4, 5, 6));
|
||||
int3 const r = (a & b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
uvec3 a = uvec3(1u, 2u, 3u);
|
||||
uvec3 b = uvec3(4u, 5u, 6u);
|
||||
uvec3 r = (uvec3(1u, 2u, 3u) & uvec3(4u, 5u, 6u));
|
||||
uvec3 r = (a & b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const uint3 a = uint3(1u, 2u, 3u);
|
||||
const uint3 b = uint3(4u, 5u, 6u);
|
||||
const uint3 r = (uint3(1u, 2u, 3u) & uint3(4u, 5u, 6u));
|
||||
const uint3 r = (a & b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
uint3 const a = uint3(1u, 2u, 3u);
|
||||
uint3 const b = uint3(4u, 5u, 6u);
|
||||
uint3 const r = (uint3(1u, 2u, 3u) & uint3(4u, 5u, 6u));
|
||||
uint3 const r = (a & b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
int const a = 1;
|
||||
int const b = 2;
|
||||
int const r = (1 | 2);
|
||||
int const r = (a | b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
uint const a = 1u;
|
||||
uint const b = 2u;
|
||||
uint const r = (1u | 2u);
|
||||
uint const r = (a | b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
ivec3 a = ivec3(1, 2, 3);
|
||||
ivec3 b = ivec3(4, 5, 6);
|
||||
ivec3 r = (ivec3(1, 2, 3) | ivec3(4, 5, 6));
|
||||
ivec3 r = (a | b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const int3 a = int3(1, 2, 3);
|
||||
const int3 b = int3(4, 5, 6);
|
||||
const int3 r = (int3(1, 2, 3) | int3(4, 5, 6));
|
||||
const int3 r = (a | b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
int3 const a = int3(1, 2, 3);
|
||||
int3 const b = int3(4, 5, 6);
|
||||
int3 const r = (int3(1, 2, 3) | int3(4, 5, 6));
|
||||
int3 const r = (a | b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
uvec3 a = uvec3(1u, 2u, 3u);
|
||||
uvec3 b = uvec3(4u, 5u, 6u);
|
||||
uvec3 r = (uvec3(1u, 2u, 3u) | uvec3(4u, 5u, 6u));
|
||||
uvec3 r = (a | b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const uint3 a = uint3(1u, 2u, 3u);
|
||||
const uint3 b = uint3(4u, 5u, 6u);
|
||||
const uint3 r = (uint3(1u, 2u, 3u) | uint3(4u, 5u, 6u));
|
||||
const uint3 r = (a | b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
uint3 const a = uint3(1u, 2u, 3u);
|
||||
uint3 const b = uint3(4u, 5u, 6u);
|
||||
uint3 const r = (uint3(1u, 2u, 3u) | uint3(4u, 5u, 6u));
|
||||
uint3 const r = (a | b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
int const a = 1;
|
||||
int const b = 2;
|
||||
int const r = (1 ^ 2);
|
||||
int const r = (a ^ b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
uint const a = 1u;
|
||||
uint const b = 2u;
|
||||
uint const r = (1u ^ 2u);
|
||||
uint const r = (a ^ b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
ivec3 a = ivec3(1, 2, 3);
|
||||
ivec3 b = ivec3(4, 5, 6);
|
||||
ivec3 r = (ivec3(1, 2, 3) ^ ivec3(4, 5, 6));
|
||||
ivec3 r = (a ^ b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const int3 a = int3(1, 2, 3);
|
||||
const int3 b = int3(4, 5, 6);
|
||||
const int3 r = (int3(1, 2, 3) ^ int3(4, 5, 6));
|
||||
const int3 r = (a ^ b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
int3 const a = int3(1, 2, 3);
|
||||
int3 const b = int3(4, 5, 6);
|
||||
int3 const r = (int3(1, 2, 3) ^ int3(4, 5, 6));
|
||||
int3 const r = (a ^ b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
uvec3 a = uvec3(1u, 2u, 3u);
|
||||
uvec3 b = uvec3(4u, 5u, 6u);
|
||||
uvec3 r = (uvec3(1u, 2u, 3u) ^ uvec3(4u, 5u, 6u));
|
||||
uvec3 r = (a ^ b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const uint3 a = uint3(1u, 2u, 3u);
|
||||
const uint3 b = uint3(4u, 5u, 6u);
|
||||
const uint3 r = (uint3(1u, 2u, 3u) ^ uint3(4u, 5u, 6u));
|
||||
const uint3 r = (a ^ b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
uint3 const a = uint3(1u, 2u, 3u);
|
||||
uint3 const b = uint3(4u, 5u, 6u);
|
||||
uint3 const r = (uint3(1u, 2u, 3u) ^ uint3(4u, 5u, 6u));
|
||||
uint3 const r = (a ^ b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
float const a = 1.0f;
|
||||
float const b = 2.0f;
|
||||
float const r = (1.0f / 2.0f);
|
||||
float const r = (a / b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
int const a = 1;
|
||||
int const b = 2;
|
||||
int const r = (1 / 2);
|
||||
int const r = (a / b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
uint const a = 1u;
|
||||
uint const b = 2u;
|
||||
uint const r = (1u / 2u);
|
||||
uint const r = (a / b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
float a = 4.0f;
|
||||
vec3 b = vec3(1.0f, 2.0f, 3.0f);
|
||||
vec3 r = (4.0f / vec3(1.0f, 2.0f, 3.0f));
|
||||
vec3 r = (a / b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const float a = 4.0f;
|
||||
const float3 b = float3(1.0f, 2.0f, 3.0f);
|
||||
const float3 r = (4.0f / float3(1.0f, 2.0f, 3.0f));
|
||||
const float3 r = (a / b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -4,7 +4,7 @@ using namespace metal;
|
|||
kernel void f() {
|
||||
float const a = 4.0f;
|
||||
float3 const b = float3(1.0f, 2.0f, 3.0f);
|
||||
float3 const r = (4.0f / float3(1.0f, 2.0f, 3.0f));
|
||||
float3 const r = (a / b);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
void f() {
|
||||
int a = 4;
|
||||
ivec3 b = ivec3(1, 2, 3);
|
||||
ivec3 r = (4 / ivec3(1, 2, 3));
|
||||
ivec3 r = (a / b);
|
||||
}
|
||||
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
|
|
|||
|
|
@ -2,6 +2,6 @@
|
|||
void f() {
|
||||
const int a = 4;
|
||||
const int3 b = int3(1, 2, 3);
|
||||
const int3 r = (4 / int3(1, 2, 3));
|
||||
const int3 r = (a / b);
|
||||
return;
|
||||
}
|
||||
|
|
|
|||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in New Issue