intrinsics: Remove ignore()

This has been deprecated since M97.

Fixed: tint:1214
Change-Id: I970898f2ae8baf1916e2f8d43230d8b724641730
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/78785
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
This commit is contained in:
Ben Clayton
2022-02-01 17:21:52 +00:00
committed by Tint LUCI CQ
parent e1159c7180
commit b80e2f3b6e
114 changed files with 1617 additions and 5688 deletions

File diff suppressed because it is too large Load Diff

View File

@@ -326,7 +326,6 @@ fn frexp<N: num>(vec<N, f32>) -> __frexp_result_vec<N>
[[stage("fragment")]] fn fwidthCoarse<N: num>(vec<N, f32>) -> vec<N, f32>
[[stage("fragment")]] fn fwidthFine(f32) -> f32
[[stage("fragment")]] fn fwidthFine<N: num>(vec<N, f32>) -> vec<N, f32>
[[deprecated]] fn ignore<T>(T)
fn inverseSqrt(f32) -> f32
fn inverseSqrt<N: num>(vec<N, f32>) -> vec<N, f32>
[[deprecated]] fn isFinite(f32) -> bool

View File

@@ -2880,7 +2880,7 @@ TEST_F(ResolverTypeConstructorValidationTest, Expr_Constructor_Struct_Empty) {
TEST_F(ResolverTypeConstructorValidationTest, NonConstructibleType_Atomic) {
WrapInFunction(
Call("ignore", Construct(Source{{12, 34}}, ty.atomic(ty.i32()))));
Assign(Phony(), Construct(Source{{12, 34}}, ty.atomic(ty.i32()))));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(), "12:34 error: type is not constructible");
@@ -2888,8 +2888,8 @@ TEST_F(ResolverTypeConstructorValidationTest, NonConstructibleType_Atomic) {
TEST_F(ResolverTypeConstructorValidationTest,
NonConstructibleType_AtomicArray) {
WrapInFunction(Call(
"ignore", Construct(Source{{12, 34}}, ty.array(ty.atomic(ty.i32()), 4))));
WrapInFunction(Assign(
Phony(), Construct(Source{{12, 34}}, ty.array(ty.atomic(ty.i32()), 4))));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(
@@ -2900,7 +2900,7 @@ TEST_F(ResolverTypeConstructorValidationTest,
TEST_F(ResolverTypeConstructorValidationTest,
NonConstructibleType_AtomicStructMember) {
auto* str = Structure("S", {Member("a", ty.atomic(ty.i32()))});
WrapInFunction(Call("ignore", Construct(Source{{12, 34}}, ty.Of(str))));
WrapInFunction(Assign(Phony(), Construct(Source{{12, 34}}, ty.Of(str))));
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
@@ -2908,8 +2908,8 @@ TEST_F(ResolverTypeConstructorValidationTest,
}
TEST_F(ResolverTypeConstructorValidationTest, NonConstructibleType_Sampler) {
WrapInFunction(Call(
"ignore",
WrapInFunction(Assign(
Phony(),
Construct(Source{{12, 34}}, ty.sampler(ast::SamplerKind::kSampler))));
EXPECT_FALSE(r()->Resolve());

View File

@@ -132,9 +132,6 @@ IntrinsicType ParseIntrinsicType(const std::string& name) {
if (name == "fwidthFine") {
return IntrinsicType::kFwidthFine;
}
if (name == "ignore") {
return IntrinsicType::kIgnore;
}
if (name == "inverseSqrt") {
return IntrinsicType::kInverseSqrt;
}
@@ -414,8 +411,6 @@ const char* str(IntrinsicType i) {
return "fwidthCoarse";
case IntrinsicType::kFwidthFine:
return "fwidthFine";
case IntrinsicType::kIgnore:
return "ignore";
case IntrinsicType::kInverseSqrt:
return "inverseSqrt";
case IntrinsicType::kIsFinite:

View File

@@ -68,7 +68,6 @@ enum class IntrinsicType {
kFwidth,
kFwidthCoarse,
kFwidthFine,
kIgnore,
kInverseSqrt,
kIsFinite,
kIsInf,

View File

@@ -454,7 +454,7 @@ struct SB {
@stage(compute) @workgroup_size(1)
fn main() {
ignore(&(sb.arr));
_ = &(sb.arr);
}
)";

View File

@@ -922,19 +922,6 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx,
if (auto* call_expr = node->As<ast::CallExpression>()) {
auto* call = sem.Get(call_expr);
if (auto* intrinsic = call->Target()->As<sem::Intrinsic>()) {
if (intrinsic->Type() == sem::IntrinsicType::kIgnore) { // [DEPRECATED]
// ignore(X)
// If X is an memory access, don't transform it into a load, as it
// may refer to a structure holding a runtime array, which cannot be
// loaded. Instead replace X with the underlying storage / uniform
// buffer variable.
if (auto access = state.TakeAccess(call_expr->args[0])) {
ctx.Replace(call_expr->args[0], [=, &ctx] {
return ctx.CloneWithoutTransform(access.var->Declaration());
});
}
continue;
}
if (intrinsic->Type() == sem::IntrinsicType::kArrayLength) {
// arrayLength(X)
// Don't convert X into a load, this intrinsic actually requires the

View File

@@ -59,7 +59,7 @@ fn f() {
break;
}
ignore(123);
_ = 123;
continuing {
i = i + 1;
@@ -73,7 +73,7 @@ fn f() {
var i : i32;
i = 0;
for(; !((i > 15)); i = (i + 1)) {
ignore(123);
_ = 123;
}
}
)";
@@ -94,7 +94,7 @@ fn f() {
break;
}
ignore(123);
_ = 123;
continuing {
i = i + 1;
@@ -108,7 +108,7 @@ fn f() {
var i : i32;
i = 0;
for(; (i < 15); i = (i + 1)) {
ignore(123);
_ = 123;
}
}
)";
@@ -135,8 +135,8 @@ fn f() {
break;
}
ignore(i);
ignore(j);
_ = i;
_ = j;
continuing {
j = (j + 1u);
@@ -160,8 +160,8 @@ fn f() {
{
var j : u32 = 0u;
for(; !((j >= N)); j = (j + 1u)) {
ignore(i);
ignore(j);
_ = i;
_ = j;
}
}
}
@@ -180,10 +180,10 @@ fn f() {
i = 0;
loop {
if ((i < 15)) {
ignore(i);
_ = i;
break;
}
ignore(123);
_ = 123;
continuing {
i = (i + 1);
@@ -207,10 +207,10 @@ fn f() {
loop {
if ((i < 15)) {
} else {
ignore(i);
_ = i;
break;
}
ignore(123);
_ = 123;
continuing {
i = (i + 1);
@@ -235,7 +235,7 @@ fn f() {
if ((i < 15)) {
break;
}
ignore(123);
_ = 123;
continuing {
if (false) {
@@ -261,11 +261,11 @@ fn f() {
if ((i < 15)) {
break;
}
ignore(123);
_ = 123;
continuing {
i = (i + 1);
ignore(i);
_ = i;
}
}
}

View File

@@ -49,7 +49,7 @@ fn test(vert_idx : u32) -> u32 {
@stage(vertex)
fn entry(@builtin(vertex_index) vert_idx : u32
) -> @builtin(position) vec4<f32> {
ignore(test(vert_idx));
_ = test(vert_idx);
return vec4<f32>();
}
)";
@@ -61,7 +61,7 @@ fn tint_symbol(tint_symbol_1 : u32) -> u32 {
@stage(vertex)
fn tint_symbol_2(@builtin(vertex_index) tint_symbol_1 : u32) -> @builtin(position) vec4<f32> {
ignore(tint_symbol(tint_symbol_1));
_ = tint_symbol(tint_symbol_1);
return vec4<f32>();
}
)";

View File

@@ -99,7 +99,7 @@ var<workgroup> v : i32;
@stage(compute) @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
ignore(v); // Initialization should be inserted above this statement
_ = v; // Initialization should be inserted above this statement
}
)";
auto* expect = R"(
@@ -111,7 +111,7 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
v = i32();
}
workgroupBarrier();
ignore(v);
_ = v;
}
)";
@@ -131,7 +131,7 @@ struct Params {
@stage(compute) @workgroup_size(1)
fn f(params : Params) {
ignore(v); // Initialization should be inserted above this statement
_ = v; // Initialization should be inserted above this statement
}
)";
auto* expect = R"(
@@ -148,7 +148,7 @@ fn f(params : Params) {
v = i32();
}
workgroupBarrier();
ignore(v);
_ = v;
}
)";
@@ -163,7 +163,7 @@ var<workgroup> v : i32;
@stage(compute) @workgroup_size(1)
fn f() {
ignore(v); // Initialization should be inserted above this statement
_ = v; // Initialization should be inserted above this statement
}
)";
auto* expect = R"(
@@ -175,7 +175,7 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
v = i32();
}
workgroupBarrier();
ignore(v);
_ = v;
}
)";
@@ -200,9 +200,9 @@ var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
ignore(a); // Initialization should be inserted above this statement
ignore(b);
ignore(c);
_ = a; // Initialization should be inserted above this statement
_ = b;
_ = c;
}
)";
auto* expect = R"(
@@ -237,9 +237,9 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
c[i_2].y[i] = i32();
}
workgroupBarrier();
ignore(a);
ignore(b);
ignore(c);
_ = a;
_ = b;
_ = c;
}
)";
@@ -264,9 +264,9 @@ var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(2, 3)
fn f(@builtin(local_invocation_index) local_idx : u32) {
ignore(a); // Initialization should be inserted above this statement
ignore(b);
ignore(c);
_ = a; // Initialization should be inserted above this statement
_ = b;
_ = c;
}
)";
auto* expect = R"(
@@ -301,9 +301,9 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
c[i_2].y[i] = i32();
}
workgroupBarrier();
ignore(a);
ignore(b);
ignore(c);
_ = a;
_ = b;
_ = c;
}
)";
@@ -330,9 +330,9 @@ var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(2, 3, X)
fn f(@builtin(local_invocation_index) local_idx : u32) {
ignore(a); // Initialization should be inserted above this statement
ignore(b);
ignore(c);
_ = a; // Initialization should be inserted above this statement
_ = b;
_ = c;
}
)";
auto* expect =
@@ -370,9 +370,9 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
c[i_2].y[i] = i32();
}
workgroupBarrier();
ignore(a);
ignore(b);
ignore(c);
_ = a;
_ = b;
_ = c;
}
)";
@@ -400,9 +400,9 @@ var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(5u, X, 10u)
fn f(@builtin(local_invocation_index) local_idx : u32) {
ignore(a); // Initialization should be inserted above this statement
ignore(b);
ignore(c);
_ = a; // Initialization should be inserted above this statement
_ = b;
_ = c;
}
)";
auto* expect =
@@ -460,9 +460,9 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
c[i_5].z[i_2][i][i_1] = i32();
}
workgroupBarrier();
ignore(a);
ignore(b);
ignore(c);
_ = a;
_ = b;
_ = c;
}
)";
@@ -486,9 +486,9 @@ var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(1)
fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>) {
ignore(a); // Initialization should be inserted above this statement
ignore(b);
ignore(c);
_ = a; // Initialization should be inserted above this statement
_ = b;
_ = c;
}
)";
auto* expect = R"(
@@ -523,9 +523,9 @@ fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(loc
c[i_2].y[i] = i32();
}
workgroupBarrier();
ignore(a);
ignore(b);
ignore(c);
_ = a;
_ = b;
_ = c;
}
)";
@@ -549,19 +549,19 @@ var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(1)
fn f1() {
ignore(a); // Initialization should be inserted above this statement
ignore(c);
_ = a; // Initialization should be inserted above this statement
_ = c;
}
@stage(compute) @workgroup_size(1, 2, 3)
fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>) {
ignore(b); // Initialization should be inserted above this statement
_ = b; // Initialization should be inserted above this statement
}
@stage(compute) @workgroup_size(4, 5, 6)
fn f3() {
ignore(c); // Initialization should be inserted above this statement
ignore(a);
_ = c; // Initialization should be inserted above this statement
_ = a;
}
)";
auto* expect = R"(
@@ -591,8 +591,8 @@ fn f1(@builtin(local_invocation_index) local_invocation_index : u32) {
c[i_1].y[i_2] = i32();
}
workgroupBarrier();
ignore(a);
ignore(c);
_ = a;
_ = c;
}
@stage(compute) @workgroup_size(1, 2, 3)
@@ -605,7 +605,7 @@ fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(lo
b.y[i_3] = i32();
}
workgroupBarrier();
ignore(b);
_ = b;
}
@stage(compute) @workgroup_size(4, 5, 6)
@@ -623,8 +623,8 @@ fn f3(@builtin(local_invocation_index) local_invocation_index_2 : u32) {
c[i_5].y[i_6] = i32();
}
workgroupBarrier();
ignore(c);
ignore(a);
_ = c;
_ = a;
}
)";
@@ -638,7 +638,7 @@ TEST_F(ZeroInitWorkgroupMemoryTest, TransitiveUsage) {
var<workgroup> v : i32;
fn use_v() {
ignore(v);
_ = v;
}
fn call_use_v() {
@@ -654,7 +654,7 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
var<workgroup> v : i32;
fn use_v() {
ignore(v);
_ = v;
}
fn call_use_v() {
@@ -683,8 +683,8 @@ var<workgroup> u : atomic<u32>;
@stage(compute) @workgroup_size(1)
fn f() {
ignore(i); // Initialization should be inserted above this statement
ignore(u);
atomicLoad(&(i)); // Initialization should be inserted above this statement
atomicLoad(&(u));
}
)";
auto* expect = R"(
@@ -699,8 +699,8 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
atomicStore(&(u), u32());
}
workgroupBarrier();
ignore(i);
ignore(u);
atomicLoad(&(i));
atomicLoad(&(u));
}
)";
@@ -723,7 +723,7 @@ var<workgroup> w : S;
@stage(compute) @workgroup_size(1)
fn f() {
ignore(w); // Initialization should be inserted above this statement
_ = w.a; // Initialization should be inserted above this statement
}
)";
auto* expect = R"(
@@ -747,7 +747,7 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
w.c = u32();
}
workgroupBarrier();
ignore(w);
_ = w.a;
}
)";
@@ -762,7 +762,7 @@ var<workgroup> w : array<atomic<u32>, 4>;
@stage(compute) @workgroup_size(1)
fn f() {
ignore(w); // Initialization should be inserted above this statement
atomicLoad(&w[0]); // Initialization should be inserted above this statement
}
)";
auto* expect = R"(
@@ -775,7 +775,7 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
atomicStore(&(w[i]), u32());
}
workgroupBarrier();
ignore(w);
atomicLoad(&(w[0]));
}
)";
@@ -798,7 +798,7 @@ var<workgroup> w : array<S, 4>;
@stage(compute) @workgroup_size(1)
fn f() {
ignore(w); // Initialization should be inserted above this statement
_ = w[0].a; // Initialization should be inserted above this statement
}
)";
auto* expect = R"(
@@ -823,7 +823,7 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
w[i_1].c = u32();
}
workgroupBarrier();
ignore(w);
_ = w[0].a;
}
)";

View File

@@ -529,9 +529,6 @@ bool GeneratorImpl::EmitIntrinsicCall(std::ostream& out,
if (intrinsic->Type() == sem::IntrinsicType::kRadians) {
return EmitRadiansCall(out, expr, intrinsic);
}
if (intrinsic->Type() == sem::IntrinsicType::kIgnore) {
return EmitExpression(out, expr->args[0]); // [DEPRECATED]
}
if (intrinsic->IsDataPacking()) {
return EmitDataPackingCall(out, expr, intrinsic);
}

View File

@@ -690,31 +690,6 @@ void main() {
)");
}
TEST_F(GlslGeneratorImplTest_Intrinsic, Ignore) {
Func("f", {Param("a", ty.i32()), Param("b", ty.i32()), Param("c", ty.i32())},
ty.i32(), {Return(Mul(Add("a", "b"), "c"))});
Func("main", {}, ty.void_(),
{CallStmt(Call("ignore", Call("f", 1, 2, 3)))},
{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1),
});
GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error();
EXPECT_EQ(gen.result(), R"(int f(int a, int b, int c) {
return ((a + b) * c);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
f(1, 2, 3);
return;
}
)");
}
#endif
TEST_F(GlslGeneratorImplTest_Intrinsic, DotI32) {

View File

@@ -1026,9 +1026,6 @@ bool GeneratorImpl::EmitIntrinsicCall(std::ostream& out,
if (intrinsic->Type() == sem::IntrinsicType::kRadians) {
return EmitRadiansCall(out, expr, intrinsic);
}
if (intrinsic->Type() == sem::IntrinsicType::kIgnore) {
return EmitExpression(out, expr->args[0]); // [DEPRECATED]
}
if (intrinsic->IsDataPacking()) {
return EmitDataPackingCall(out, expr, intrinsic);
}

View File

@@ -791,31 +791,6 @@ void main() {
)");
}
TEST_F(HlslGeneratorImplTest_Intrinsic, Ignore) {
Func("f", {Param("a", ty.i32()), Param("b", ty.i32()), Param("c", ty.i32())},
ty.i32(), {Return(Mul(Add("a", "b"), "c"))});
Func("main", {}, ty.void_(), {CallStmt(Call("ignore", Call("f", 1, 2, 3)))},
{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1),
});
GeneratorImpl& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.error();
EXPECT_EQ(gen.result(), R"(int f(int a, int b, int c) {
return ((a + b) * c);
}
[numthreads(1, 1, 1)]
void main() {
f(1, 2, 3);
return;
}
)");
}
} // namespace
} // namespace hlsl
} // namespace writer

View File

@@ -627,13 +627,6 @@ bool GeneratorImpl::EmitIntrinsicCall(std::ostream& out,
out << "threadgroup_barrier(mem_flags::mem_threadgroup)";
return true;
}
case sem::IntrinsicType::kIgnore: { // [DEPRECATED]
out << "(void) ";
if (!EmitExpression(out, expr->args[0])) {
return false;
}
return true;
}
case sem::IntrinsicType::kLength: {
auto* sem = builder_.Sem().Get(expr->args[0]);

View File

@@ -2460,14 +2460,6 @@ uint32_t Builder::GenerateIntrinsicCall(const sem::Call* call,
case IntrinsicType::kFwidthFine:
op = spv::Op::OpFwidthFine;
break;
case IntrinsicType::kIgnore: // [DEPRECATED]
// Evaluate the single argument, return the non-zero result_id which isn't
// associated with any op (ignore returns void, so this cannot be used in
// an expression).
if (!get_arg_as_value_id(0, false)) {
return 0;
}
return result_id;
case IntrinsicType::kIsInf:
op = spv::Op::OpIsInf;
break;

View File

@@ -2556,45 +2556,6 @@ OpReturn
Validate(b);
}
TEST_F(IntrinsicBuilderTest, Call_Ignore) {
Func("f", {Param("a", ty.i32()), Param("b", ty.i32()), Param("c", ty.i32())},
ty.i32(), {Return(Mul(Add("a", "b"), "c"))});
Func("main", {}, ty.void_(),
{
CallStmt(Call("ignore", Call("f", 1, 2, 3))),
},
{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1),
});
spirv::Builder& b = Build();
ASSERT_TRUE(b.Build()) << b.error();
ASSERT_EQ(b.functions().size(), 2u);
auto* expected_types = R"(%2 = OpTypeInt 32 1
%1 = OpTypeFunction %2 %2 %2 %2
%11 = OpTypeVoid
%10 = OpTypeFunction %11
%16 = OpConstant %2 1
%17 = OpConstant %2 2
%18 = OpConstant %2 3
)";
auto got_types = DumpInstructions(b.types());
EXPECT_EQ(expected_types, got_types);
auto* expected_instructions = R"(%15 = OpFunctionCall %2 %3 %16 %17 %18
OpReturn
)";
auto got_instructions = DumpInstructions(b.functions()[1].instructions());
EXPECT_EQ(expected_instructions, got_instructions);
Validate(b);
}
} // namespace
} // namespace spirv
} // namespace writer