encounter/dawn-cmake
1
0
Fork 0
mirror of https://github.com/encounter/dawn-cmake.git synced 2025-12-16 00:17:03 +00:00
Code Issues Packages Projects Releases Wiki Activity

resolver: Remove rule that call statements to functions must return void

Browse Source 
Fixed: tint:1256
Change-Id: Ibff78a1009d57afd7af8867952a9444daaf13a9c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67201
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
This commit is contained in:
Ben Clayton
2021-10-21 20:36:04 +00:00
committed by Tint LUCI CQ
parent 38ed53ce8f
commit 72789de9f5
42 changed files with 318 additions and 353 deletions
Download Patch File Download Diff File Expand all files Collapse all files

15
src/transform/calculate_array_length.cc
View File

@@ -193,14 +193,13 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) {
ast::StorageClass::kNone, ctx.dst->Expr(0u)));
// Call storage_buffer.GetDimensions(&buffer_size_result)
auto* call_get_dims =
ctx.dst->create<ast::CallStatement>(ctx.dst->Call(
// BufferSizeIntrinsic(X, ARGS...) is
// translated to:
// X.GetDimensions(ARGS..) by the writer
buffer_size, ctx.Clone(storage_buffer_expr),
ctx.dst->AddressOf(ctx.dst->Expr(
buffer_size_result->variable->symbol))));
auto* call_get_dims = ctx.dst->CallStmt(ctx.dst->Call(
// BufferSizeIntrinsic(X, ARGS...) is
// translated to:
// X.GetDimensions(ARGS..) by the writer
buffer_size, ctx.Clone(storage_buffer_expr),
ctx.dst->AddressOf(
ctx.dst->Expr(buffer_size_result->variable->symbol))));
// Calculate actual array length
// total_storage_buffer_size - array_offset

2
src/transform/canonicalize_entry_point_io.cc
View File

@@ -497,7 +497,7 @@ struct CanonicalizeEntryPointIO::State {
};
if (func_sem->ReturnType()->Is<sem::Void>()) {
// The function call is just a statement with no result.
wrapper_body.push_back(ctx.dst->create<ast::CallStatement>(call_inner));
wrapper_body.push_back(ctx.dst->CallStmt(call_inner));
} else {
// Capture the result of calling the original function.
auto* inner_result = ctx.dst->Const(

10
src/transform/decompose_memory_access.cc
View File

@@ -606,8 +606,8 @@ struct DecomposeMemoryAccess::State {
auto* arr_el = b.IndexAccessor(array, i);
auto* el_offset =
b.Add(b.Expr("offset"), b.Mul(i, arr_ty->Stride()));
auto* store_stmt = b.create<ast::CallStatement>(
b.Call(store, "buffer", el_offset, arr_el));
auto* store_stmt =
b.CallStmt(b.Call(store, "buffer", el_offset, arr_el));
auto* for_loop =
b.For(for_init, for_cond, for_cont, b.Block(store_stmt));
@@ -619,7 +619,7 @@ struct DecomposeMemoryAccess::State {
auto* offset = b.Add("offset", i * mat_ty->ColumnStride());
auto* access = b.IndexAccessor("value", i);
auto* call = b.Call(store, "buffer", offset, access);
body.emplace_back(b.create<ast::CallStatement>(call));
body.emplace_back(b.CallStmt(call));
}
} else if (auto* str = el_ty->As<sem::Struct>()) {
for (auto* member : str->Members()) {
@@ -629,7 +629,7 @@ struct DecomposeMemoryAccess::State {
Symbol store =
StoreFunc(buf_ty, member->Type()->UnwrapRef(), var_user);
auto* call = b.Call(store, "buffer", offset, access);
body.emplace_back(b.create<ast::CallStatement>(call));
body.emplace_back(b.CallStmt(call));
}
}
b.Func(name, params, b.ty.void_(), body);
@@ -995,7 +995,7 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) {
store.target.var->As<sem::VariableUser>());
auto* call = ctx.dst->Call(func, ctx.CloneWithoutTransform(buf), offset,
ctx.Clone(value));
return ctx.dst->create<ast::CallStatement>(call);
return ctx.dst->CallStmt(call);
});
}

120
src/transform/decompose_memory_access_test.cc
View File

@@ -1235,28 +1235,28 @@ struct SB {
[[stage(compute), workgroup_size(1)]]
fn main() {
atomicStore(&sb.a, 123);
ignore(atomicLoad(&sb.a));
ignore(atomicAdd(&sb.a, 123));
ignore(atomicSub(&sb.a, 123));
ignore(atomicMax(&sb.a, 123));
ignore(atomicMin(&sb.a, 123));
ignore(atomicAnd(&sb.a, 123));
ignore(atomicOr(&sb.a, 123));
ignore(atomicXor(&sb.a, 123));
ignore(atomicExchange(&sb.a, 123));
ignore(atomicCompareExchangeWeak(&sb.a, 123, 345));
atomicLoad(&sb.a);
atomicAdd(&sb.a, 123);
atomicSub(&sb.a, 123);
atomicMax(&sb.a, 123);
atomicMin(&sb.a, 123);
atomicAnd(&sb.a, 123);
atomicOr(&sb.a, 123);
atomicXor(&sb.a, 123);
atomicExchange(&sb.a, 123);
atomicCompareExchangeWeak(&sb.a, 123, 345);
atomicStore(&sb.b, 123u);
ignore(atomicLoad(&sb.b));
ignore(atomicAdd(&sb.b, 123u));
ignore(atomicSub(&sb.b, 123u));
ignore(atomicMax(&sb.b, 123u));
ignore(atomicMin(&sb.b, 123u));
ignore(atomicAnd(&sb.b, 123u));
ignore(atomicOr(&sb.b, 123u));
ignore(atomicXor(&sb.b, 123u));
ignore(atomicExchange(&sb.b, 123u));
ignore(atomicCompareExchangeWeak(&sb.b, 123u, 345u));
atomicLoad(&sb.b);
atomicAdd(&sb.b, 123u);
atomicSub(&sb.b, 123u);
atomicMax(&sb.b, 123u);
atomicMin(&sb.b, 123u);
atomicAnd(&sb.b, 123u);
atomicOr(&sb.b, 123u);
atomicXor(&sb.b, 123u);
atomicExchange(&sb.b, 123u);
atomicCompareExchangeWeak(&sb.b, 123u, 345u);
}
)";
@@ -1339,27 +1339,27 @@ fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_p
[[stage(compute), workgroup_size(1)]]
fn main() {
tint_symbol(sb, 16u, 123);
ignore(tint_symbol_1(sb, 16u));
ignore(tint_symbol_2(sb, 16u, 123));
ignore(tint_symbol_3(sb, 16u, 123));
ignore(tint_symbol_4(sb, 16u, 123));
ignore(tint_symbol_5(sb, 16u, 123));
ignore(tint_symbol_6(sb, 16u, 123));
ignore(tint_symbol_7(sb, 16u, 123));
ignore(tint_symbol_8(sb, 16u, 123));
ignore(tint_symbol_9(sb, 16u, 123));
ignore(tint_symbol_10(sb, 16u, 123, 345));
tint_symbol_1(sb, 16u);
tint_symbol_2(sb, 16u, 123);
tint_symbol_3(sb, 16u, 123);
tint_symbol_4(sb, 16u, 123);
tint_symbol_5(sb, 16u, 123);
tint_symbol_6(sb, 16u, 123);
tint_symbol_7(sb, 16u, 123);
tint_symbol_8(sb, 16u, 123);
tint_symbol_9(sb, 16u, 123);
tint_symbol_10(sb, 16u, 123, 345);
tint_symbol_11(sb, 20u, 123u);
ignore(tint_symbol_12(sb, 20u));
ignore(tint_symbol_13(sb, 20u, 123u));
ignore(tint_symbol_14(sb, 20u, 123u));
ignore(tint_symbol_15(sb, 20u, 123u));
ignore(tint_symbol_16(sb, 20u, 123u));
ignore(tint_symbol_17(sb, 20u, 123u));
ignore(tint_symbol_18(sb, 20u, 123u));
ignore(tint_symbol_19(sb, 20u, 123u));
ignore(tint_symbol_20(sb, 20u, 123u));
ignore(tint_symbol_21(sb, 20u, 123u, 345u));
tint_symbol_12(sb, 20u);
tint_symbol_13(sb, 20u, 123u);
tint_symbol_14(sb, 20u, 123u);
tint_symbol_15(sb, 20u, 123u);
tint_symbol_16(sb, 20u, 123u);
tint_symbol_17(sb, 20u, 123u);
tint_symbol_18(sb, 20u, 123u);
tint_symbol_19(sb, 20u, 123u);
tint_symbol_20(sb, 20u, 123u);
tint_symbol_21(sb, 20u, 123u, 345u);
}
)";
@@ -1381,27 +1381,27 @@ var<workgroup> w : S;
[[stage(compute), workgroup_size(1)]]
fn main() {
atomicStore(&(w.a), 123);
ignore(atomicLoad(&(w.a)));
ignore(atomicAdd(&(w.a), 123));
ignore(atomicSub(&(w.a), 123));
ignore(atomicMax(&(w.a), 123));
ignore(atomicMin(&(w.a), 123));
ignore(atomicAnd(&(w.a), 123));
ignore(atomicOr(&(w.a), 123));
ignore(atomicXor(&(w.a), 123));
ignore(atomicExchange(&(w.a), 123));
ignore(atomicCompareExchangeWeak(&(w.a), 123, 345));
atomicLoad(&(w.a));
atomicAdd(&(w.a), 123);
atomicSub(&(w.a), 123);
atomicMax(&(w.a), 123);
atomicMin(&(w.a), 123);
atomicAnd(&(w.a), 123);
atomicOr(&(w.a), 123);
atomicXor(&(w.a), 123);
atomicExchange(&(w.a), 123);
atomicCompareExchangeWeak(&(w.a), 123, 345);
atomicStore(&(w.b), 123u);
ignore(atomicLoad(&(w.b)));
ignore(atomicAdd(&(w.b), 123u));
ignore(atomicSub(&(w.b), 123u));
ignore(atomicMax(&(w.b), 123u));
ignore(atomicMin(&(w.b), 123u));
ignore(atomicAnd(&(w.b), 123u));
ignore(atomicOr(&(w.b), 123u));
ignore(atomicXor(&(w.b), 123u));
ignore(atomicExchange(&(w.b), 123u));
ignore(atomicCompareExchangeWeak(&(w.b), 123u, 345u));
atomicLoad(&(w.b));
atomicAdd(&(w.b), 123u);
atomicSub(&(w.b), 123u);
atomicMax(&(w.b), 123u);
atomicMin(&(w.b), 123u);
atomicAnd(&(w.b), 123u);
atomicOr(&(w.b), 123u);
atomicXor(&(w.b), 123u);
atomicExchange(&(w.b), 123u);
atomicCompareExchangeWeak(&(w.b), 123u, 345u);
}
)";

28
src/transform/first_index_offset_test.cc
View File

@@ -53,7 +53,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>();
}
)";
@@ -72,7 +72,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 + tint_symbol_1.first_vertex_index)));
test((vert_idx + tint_symbol_1.first_vertex_index));
return vec4<f32>();
}
)";
@@ -100,7 +100,7 @@ fn test(inst_idx : u32) -> u32 {
[[stage(vertex)]]
fn entry([[builtin(instance_index)]] inst_idx : u32) -> [[builtin(position)]] vec4<f32> {
ignore(test(inst_idx));
test(inst_idx);
return vec4<f32>();
}
)";
@@ -119,7 +119,7 @@ fn test(inst_idx : u32) -> u32 {
[[stage(vertex)]]
fn entry([[builtin(instance_index)]] inst_idx : u32) -> [[builtin(position)]] vec4<f32> {
ignore(test((inst_idx + tint_symbol_1.first_instance_index)));
test((inst_idx + tint_symbol_1.first_instance_index));
return vec4<f32>();
}
)";
@@ -152,7 +152,7 @@ struct Inputs {
[[stage(vertex)]]
fn entry(inputs : Inputs) -> [[builtin(position)]] vec4<f32> {
ignore(test(inputs.instance_idx, inputs.vert_idx));
test(inputs.instance_idx, inputs.vert_idx);
return vec4<f32>();
}
)";
@@ -179,7 +179,7 @@ struct Inputs {
[[stage(vertex)]]
fn entry(inputs : Inputs) -> [[builtin(position)]] vec4<f32> {
ignore(test((inputs.instance_idx + tint_symbol_1.first_instance_index), (inputs.vert_idx + tint_symbol_1.first_vertex_index)));
test((inputs.instance_idx + tint_symbol_1.first_instance_index), (inputs.vert_idx + tint_symbol_1.first_vertex_index));
return vec4<f32>();
}
)";
@@ -211,7 +211,7 @@ fn func2(vert_idx : u32) -> u32 {
[[stage(vertex)]]
fn entry([[builtin(vertex_index)]] vert_idx : u32) -> [[builtin(position)]] vec4<f32> {
ignore(func2(vert_idx));
func2(vert_idx);
return vec4<f32>();
}
)";
@@ -234,7 +234,7 @@ fn func2(vert_idx : u32) -> u32 {
[[stage(vertex)]]
fn entry([[builtin(vertex_index)]] vert_idx : u32) -> [[builtin(position)]] vec4<f32> {
ignore(func2((vert_idx + tint_symbol_1.first_vertex_index)));
func2((vert_idx + tint_symbol_1.first_vertex_index));
return vec4<f32>();
}
)";
@@ -262,19 +262,19 @@ fn func(i : u32) -> u32 {
[[stage(vertex)]]
fn entry_a([[builtin(vertex_index)]] vert_idx : u32) -> [[builtin(position)]] vec4<f32> {
ignore(func(vert_idx));
func(vert_idx);
return vec4<f32>();
}
[[stage(vertex)]]
fn entry_b([[builtin(vertex_index)]] vert_idx : u32, [[builtin(instance_index)]] inst_idx : u32) -> [[builtin(position)]] vec4<f32> {
ignore(func(vert_idx + inst_idx));
func(vert_idx + inst_idx);
return vec4<f32>();
}
[[stage(vertex)]]
fn entry_c([[builtin(instance_index)]] inst_idx : u32) -> [[builtin(position)]] vec4<f32> {
ignore(func(inst_idx));
func(inst_idx);
return vec4<f32>();
}
)";
@@ -294,19 +294,19 @@ fn func(i : u32) -> u32 {
[[stage(vertex)]]
fn entry_a([[builtin(vertex_index)]] vert_idx : u32) -> [[builtin(position)]] vec4<f32> {
ignore(func((vert_idx + tint_symbol_1.first_vertex_index)));
func((vert_idx + tint_symbol_1.first_vertex_index));
return vec4<f32>();
}
[[stage(vertex)]]
fn entry_b([[builtin(vertex_index)]] vert_idx : u32, [[builtin(instance_index)]] inst_idx : u32) -> [[builtin(position)]] vec4<f32> {
ignore(func(((vert_idx + tint_symbol_1.first_vertex_index) + (inst_idx + tint_symbol_1.first_instance_index))));
func(((vert_idx + tint_symbol_1.first_vertex_index) + (inst_idx + tint_symbol_1.first_instance_index)));
return vec4<f32>();
}
[[stage(vertex)]]
fn entry_c([[builtin(instance_index)]] inst_idx : u32) -> [[builtin(position)]] vec4<f32> {
ignore(func((inst_idx + tint_symbol_1.first_instance_index)));
func((inst_idx + tint_symbol_1.first_instance_index));
return vec4<f32>();
}
)";

32
src/transform/robustness_test.cc
View File

@@ -670,14 +670,14 @@ fn f() {
var level_idx : i32;
var sample_idx : i32;
ignore(textureLoad(tex_1d, 1, level_idx));
ignore(textureLoad(tex_2d, vec2<i32>(1, 2), level_idx));
ignore(textureLoad(tex_2d_arr, vec2<i32>(1, 2), array_idx, level_idx));
ignore(textureLoad(tex_3d, vec3<i32>(1, 2, 3), level_idx));
ignore(textureLoad(tex_ms_2d, vec2<i32>(1, 2), sample_idx));
ignore(textureLoad(tex_depth_2d, vec2<i32>(1, 2), level_idx));
ignore(textureLoad(tex_depth_2d_arr, vec2<i32>(1, 2), array_idx, level_idx));
ignore(textureLoad(tex_external, vec2<i32>(1, 2)));
textureLoad(tex_1d, 1, level_idx);
textureLoad(tex_2d, vec2<i32>(1, 2), level_idx);
textureLoad(tex_2d_arr, vec2<i32>(1, 2), array_idx, level_idx);
textureLoad(tex_3d, vec3<i32>(1, 2, 3), level_idx);
textureLoad(tex_ms_2d, vec2<i32>(1, 2), sample_idx);
textureLoad(tex_depth_2d, vec2<i32>(1, 2), level_idx);
textureLoad(tex_depth_2d_arr, vec2<i32>(1, 2), array_idx, level_idx);
textureLoad(tex_external, vec2<i32>(1, 2));
}
)";
@@ -703,14 +703,14 @@ fn f() {
var array_idx : i32;
var level_idx : i32;
var sample_idx : i32;
ignore(textureLoad(tex_1d, clamp(1, i32(), (textureDimensions(tex_1d, clamp(level_idx, 0, (textureNumLevels(tex_1d) - 1))) - i32(1))), clamp(level_idx, 0, (textureNumLevels(tex_1d) - 1))));
ignore(textureLoad(tex_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_2d, clamp(level_idx, 0, (textureNumLevels(tex_2d) - 1))) - vec2<i32>(1))), clamp(level_idx, 0, (textureNumLevels(tex_2d) - 1))));
ignore(textureLoad(tex_2d_arr, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_2d_arr, clamp(level_idx, 0, (textureNumLevels(tex_2d_arr) - 1))) - vec2<i32>(1))), clamp(array_idx, 0, (textureNumLayers(tex_2d_arr) - 1)), clamp(level_idx, 0, (textureNumLevels(tex_2d_arr) - 1))));
ignore(textureLoad(tex_3d, clamp(vec3<i32>(1, 2, 3), vec3<i32>(), (textureDimensions(tex_3d, clamp(level_idx, 0, (textureNumLevels(tex_3d) - 1))) - vec3<i32>(1))), clamp(level_idx, 0, (textureNumLevels(tex_3d) - 1))));
ignore(textureLoad(tex_ms_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_ms_2d) - vec2<i32>(1))), sample_idx));
ignore(textureLoad(tex_depth_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_depth_2d, clamp(level_idx, 0, (textureNumLevels(tex_depth_2d) - 1))) - vec2<i32>(1))), clamp(level_idx, 0, (textureNumLevels(tex_depth_2d) - 1))));
ignore(textureLoad(tex_depth_2d_arr, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_depth_2d_arr, clamp(level_idx, 0, (textureNumLevels(tex_depth_2d_arr) - 1))) - vec2<i32>(1))), clamp(array_idx, 0, (textureNumLayers(tex_depth_2d_arr) - 1)), clamp(level_idx, 0, (textureNumLevels(tex_depth_2d_arr) - 1))));
ignore(textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_external) - vec2<i32>(1)))));
textureLoad(tex_1d, clamp(1, i32(), (textureDimensions(tex_1d, clamp(level_idx, 0, (textureNumLevels(tex_1d) - 1))) - i32(1))), clamp(level_idx, 0, (textureNumLevels(tex_1d) - 1)));
textureLoad(tex_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_2d, clamp(level_idx, 0, (textureNumLevels(tex_2d) - 1))) - vec2<i32>(1))), clamp(level_idx, 0, (textureNumLevels(tex_2d) - 1)));
textureLoad(tex_2d_arr, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_2d_arr, clamp(level_idx, 0, (textureNumLevels(tex_2d_arr) - 1))) - vec2<i32>(1))), clamp(array_idx, 0, (textureNumLayers(tex_2d_arr) - 1)), clamp(level_idx, 0, (textureNumLevels(tex_2d_arr) - 1)));
textureLoad(tex_3d, clamp(vec3<i32>(1, 2, 3), vec3<i32>(), (textureDimensions(tex_3d, clamp(level_idx, 0, (textureNumLevels(tex_3d) - 1))) - vec3<i32>(1))), clamp(level_idx, 0, (textureNumLevels(tex_3d) - 1)));
textureLoad(tex_ms_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_ms_2d) - vec2<i32>(1))), sample_idx);
textureLoad(tex_depth_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_depth_2d, clamp(level_idx, 0, (textureNumLevels(tex_depth_2d) - 1))) - vec2<i32>(1))), clamp(level_idx, 0, (textureNumLevels(tex_depth_2d) - 1)));
textureLoad(tex_depth_2d_arr, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_depth_2d_arr, clamp(level_idx, 0, (textureNumLevels(tex_depth_2d_arr) - 1))) - vec2<i32>(1))), clamp(array_idx, 0, (textureNumLayers(tex_depth_2d_arr) - 1)), clamp(level_idx, 0, (textureNumLevels(tex_depth_2d_arr) - 1)));
textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_external) - vec2<i32>(1))));
}
)";

5
src/transform/zero_init_workgroup_memory.cc
View File

@@ -260,7 +260,7 @@ struct ZeroInitWorkgroupMemory::State {
// Append a single workgroup barrier after the zero initialization.
ctx.InsertFront(fn->body->statements,
b.create<ast::CallStatement>(b.Call("workgroupBarrier")));
b.CallStmt(b.Call("workgroupBarrier")));
}
/// BuildZeroingExpr is a function that builds a sub-expression used to zero
@@ -286,8 +286,7 @@ struct ZeroInitWorkgroupMemory::State {
auto* zero_init = b.Construct(CreateASTTypeFor(ctx, atomic->Type()));
auto expr = get_expr(1u);
auto* store = b.Call("atomicStore", b.AddressOf(expr.expr), zero_init);
statements.emplace_back(Statement{b.create<ast::CallStatement>(store),
expr.num_iterations,
statements.emplace_back(Statement{b.CallStmt(store), expr.num_iterations,
expr.array_indices});
return;
}