Skip to content

Commit 1b75d8b

Browse files
jrpriceDawn LUCI CQ
jrprice
authored and
Dawn LUCI CQ
committed
[msl] Fix PackedVec3 for atomic builtins
Only unwrap pointers before load instructions, instead of before calling `UpdateUsage()`, so that the `packed == unpacked` check at the start of `UpdateUsage()` can correctly determine when the target type is not in fact packed. Fixed: 366314931 Change-Id: Ie8e3c0bf87a0afd0d13199227fc9687e57fe0809 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/206395 Auto-Submit: James Price <[email protected]> Reviewed-by: Antonio Maiorano <[email protected]> Commit-Queue: Antonio Maiorano <[email protected]>
1 parent edaec92 commit 1b75d8b

13 files changed

+491
-4
lines changed

‎src/tint/lang/msl/writer/raise/packed_vec3.cc‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -291,12 +291,12 @@ struct State {
291291
auto* packed_result_type = RewriteType(unpacked_result_type);
292292
let->Result(0)->SetType(packed_result_type);
293293
let->Result(0)->ForEachUseSorted([&](core::ir::Usage let_use) { //
294-
UpdateUsage(let_use, unpacked_result_type->UnwrapPtr(), packed_result_type);
294+
UpdateUsage(let_use, unpacked_result_type, packed_result_type);
295295
});
296296
},
297297
[&](core::ir::Load* load) {
298298
b.InsertAfter(load, [&] {
299-
auto* result = LoadPackedToUnpacked(unpacked_type, load->From());
299+
auto* result = LoadPackedToUnpacked(unpacked_type->UnwrapPtr(), load->From());
300300
load->Result(0)->ReplaceAllUsesWith(result);
301301
});
302302
load->Destroy();
@@ -327,7 +327,7 @@ struct State {
327327
// Rebuild the indices of the access instruction.
328328
// Walk through the intermediate types that the access chain will be traversing, and
329329
// check for packed vectors that would be wrapped in structures.
330-
auto* obj_type = unpacked_type;
330+
auto* obj_type = unpacked_type->UnwrapPtr();
331331
Vector<core::ir::Value*, 4> operands;
332332
operands.Push(access->Object());
333333
for (auto* idx : access->Indices()) {
@@ -354,7 +354,7 @@ struct State {
354354
access->SetOperands(std::move(operands));
355355
access->Result(0)->SetType(packed_result_type);
356356
access->Result(0)->ForEachUseSorted([&](core::ir::Usage access_use) { //
357-
UpdateUsage(access_use, unpacked_result_type->UnwrapPtr(), packed_result_type);
357+
UpdateUsage(access_use, unpacked_result_type, packed_result_type);
358358
});
359359
}
360360

‎src/tint/lang/msl/writer/raise/packed_vec3_test.cc‎

Lines changed: 132 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3692,5 +3692,137 @@ S_packed_vec3 = struct @align(16) {
36923692
EXPECT_EQ(expect, str());
36933693
}
36943694

3695+
TEST_F(MslWriter_PackedVec3Test, AtomicOnPackedStructMember) {
3696+
auto* s = ty.Struct(mod.symbols.New("S"), {
3697+
{mod.symbols.Register("vec"), ty.vec3<u32>()},
3698+
{mod.symbols.Register("u"), ty.atomic<u32>()},
3699+
});
3700+
3701+
auto* var = b.Var("v", ty.ptr<workgroup>(s));
3702+
mod.root_block->Append(var);
3703+
3704+
auto* func = b.Function("foo", ty.u32());
3705+
b.Append(func->Block(), [&] { //
3706+
auto* p = b.Access<ptr<workgroup, atomic<u32>>>(var, 1_u);
3707+
auto* result = b.Call<u32>(core::BuiltinFn::kAtomicLoad, p);
3708+
b.Return(func, result);
3709+
});
3710+
3711+
auto* src = R"(
3712+
S = struct @align(16) {
3713+
vec:vec3<u32> @offset(0)
3714+
u:atomic<u32> @offset(12)
3715+
}
3716+
3717+
$B1: { # root
3718+
%v:ptr<workgroup, S, read_write> = var
3719+
}
3720+
3721+
%foo = func():u32 {
3722+
$B2: {
3723+
%3:ptr<workgroup, atomic<u32>, read_write> = access %v, 1u
3724+
%4:u32 = atomicLoad %3
3725+
ret %4
3726+
}
3727+
}
3728+
)";
3729+
EXPECT_EQ(src, str());
3730+
3731+
auto* expect = R"(
3732+
S = struct @align(16) {
3733+
vec:vec3<u32> @offset(0)
3734+
u:atomic<u32> @offset(12)
3735+
}
3736+
3737+
S_packed_vec3 = struct @align(16) {
3738+
vec:__packed_vec3<u32> @offset(0)
3739+
u:atomic<u32> @offset(12)
3740+
}
3741+
3742+
$B1: { # root
3743+
%v:ptr<workgroup, S_packed_vec3, read_write> = var
3744+
}
3745+
3746+
%foo = func():u32 {
3747+
$B2: {
3748+
%3:ptr<workgroup, atomic<u32>, read_write> = access %v, 1u
3749+
%4:u32 = atomicLoad %3
3750+
ret %4
3751+
}
3752+
}
3753+
)";
3754+
3755+
Run(PackedVec3);
3756+
3757+
EXPECT_EQ(expect, str());
3758+
}
3759+
3760+
TEST_F(MslWriter_PackedVec3Test, AtomicOnPackedStructMember_ViaLet) {
3761+
auto* s = ty.Struct(mod.symbols.New("S"), {
3762+
{mod.symbols.Register("vec"), ty.vec3<u32>()},
3763+
{mod.symbols.Register("u"), ty.atomic<u32>()},
3764+
});
3765+
3766+
auto* var = b.Var("v", ty.ptr<workgroup>(s));
3767+
mod.root_block->Append(var);
3768+
3769+
auto* func = b.Function("foo", ty.u32());
3770+
b.Append(func->Block(), [&] { //
3771+
auto* p = b.Let("p", b.Access<ptr<workgroup, atomic<u32>>>(var, 1_u));
3772+
auto* result = b.Call<u32>(core::BuiltinFn::kAtomicLoad, p);
3773+
b.Return(func, result);
3774+
});
3775+
3776+
auto* src = R"(
3777+
S = struct @align(16) {
3778+
vec:vec3<u32> @offset(0)
3779+
u:atomic<u32> @offset(12)
3780+
}
3781+
3782+
$B1: { # root
3783+
%v:ptr<workgroup, S, read_write> = var
3784+
}
3785+
3786+
%foo = func():u32 {
3787+
$B2: {
3788+
%3:ptr<workgroup, atomic<u32>, read_write> = access %v, 1u
3789+
%p:ptr<workgroup, atomic<u32>, read_write> = let %3
3790+
%5:u32 = atomicLoad %p
3791+
ret %5
3792+
}
3793+
}
3794+
)";
3795+
EXPECT_EQ(src, str());
3796+
3797+
auto* expect = R"(
3798+
S = struct @align(16) {
3799+
vec:vec3<u32> @offset(0)
3800+
u:atomic<u32> @offset(12)
3801+
}
3802+
3803+
S_packed_vec3 = struct @align(16) {
3804+
vec:__packed_vec3<u32> @offset(0)
3805+
u:atomic<u32> @offset(12)
3806+
}
3807+
3808+
$B1: { # root
3809+
%v:ptr<workgroup, S_packed_vec3, read_write> = var
3810+
}
3811+
3812+
%foo = func():u32 {
3813+
$B2: {
3814+
%3:ptr<workgroup, atomic<u32>, read_write> = access %v, 1u
3815+
%p:ptr<workgroup, atomic<u32>, read_write> = let %3
3816+
%5:u32 = atomicLoad %p
3817+
ret %5
3818+
}
3819+
}
3820+
)";
3821+
3822+
Run(PackedVec3);
3823+
3824+
EXPECT_EQ(expect, str());
3825+
}
3826+
36953827
} // namespace
36963828
} // namespace tint::msl::writer::raise

‎test/tint/bug/tint/366314931.wgsl‎

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,15 @@
1+
struct S {
2+
v : vec3u,
3+
u : atomic<u32>,
4+
}
5+
6+
var<workgroup> wgvar: S;
7+
8+
@group(0) @binding(0)
9+
var<storage, read_write> output: S;
10+
11+
@compute @workgroup_size(1,1,1)
12+
fn main() {
13+
let x = atomicLoad(&wgvar.u);
14+
atomicStore(&output.u, x);
15+
}

‎test/tint/bug/tint/366314931.wgsl.expected.dxc.hlsl‎

Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,41 @@
1+
struct S {
2+
uint3 v;
3+
uint u;
4+
};
5+
6+
groupshared S wgvar;
7+
8+
void tint_zero_workgroup_memory(uint local_idx) {
9+
if ((local_idx < 1u)) {
10+
wgvar.v = (0u).xxx;
11+
uint atomic_result = 0u;
12+
InterlockedExchange(wgvar.u, 0u, atomic_result);
13+
}
14+
GroupMemoryBarrierWithGroupSync();
15+
}
16+
17+
RWByteAddressBuffer output : register(u0);
18+
19+
struct tint_symbol_1 {
20+
uint local_invocation_index : SV_GroupIndex;
21+
};
22+
23+
void outputatomicStore(uint offset, uint value) {
24+
uint ignored;
25+
output.InterlockedExchange(offset, value, ignored);
26+
}
27+
28+
29+
void main_inner(uint local_invocation_index) {
30+
tint_zero_workgroup_memory(local_invocation_index);
31+
uint atomic_result_1 = 0u;
32+
InterlockedOr(wgvar.u, 0, atomic_result_1);
33+
uint x = atomic_result_1;
34+
outputatomicStore(12u, x);
35+
}
36+
37+
[numthreads(1, 1, 1)]
38+
void main(tint_symbol_1 tint_symbol) {
39+
main_inner(tint_symbol.local_invocation_index);
40+
return;
41+
}

‎test/tint/bug/tint/366314931.wgsl.expected.fxc.hlsl‎

Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,41 @@
1+
struct S {
2+
uint3 v;
3+
uint u;
4+
};
5+
6+
groupshared S wgvar;
7+
8+
void tint_zero_workgroup_memory(uint local_idx) {
9+
if ((local_idx < 1u)) {
10+
wgvar.v = (0u).xxx;
11+
uint atomic_result = 0u;
12+
InterlockedExchange(wgvar.u, 0u, atomic_result);
13+
}
14+
GroupMemoryBarrierWithGroupSync();
15+
}
16+
17+
RWByteAddressBuffer output : register(u0);
18+
19+
struct tint_symbol_1 {
20+
uint local_invocation_index : SV_GroupIndex;
21+
};
22+
23+
void outputatomicStore(uint offset, uint value) {
24+
uint ignored;
25+
output.InterlockedExchange(offset, value, ignored);
26+
}
27+
28+
29+
void main_inner(uint local_invocation_index) {
30+
tint_zero_workgroup_memory(local_invocation_index);
31+
uint atomic_result_1 = 0u;
32+
InterlockedOr(wgvar.u, 0, atomic_result_1);
33+
uint x = atomic_result_1;
34+
outputatomicStore(12u, x);
35+
}
36+
37+
[numthreads(1, 1, 1)]
38+
void main(tint_symbol_1 tint_symbol) {
39+
main_inner(tint_symbol.local_invocation_index);
40+
return;
41+
}

‎test/tint/bug/tint/366314931.wgsl.expected.glsl‎

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,31 @@
1+
#version 310 es
2+
3+
struct S {
4+
uvec3 v;
5+
uint u;
6+
};
7+
8+
shared S wgvar;
9+
void tint_zero_workgroup_memory(uint local_idx) {
10+
if ((local_idx < 1u)) {
11+
wgvar.v = uvec3(0u);
12+
atomicExchange(wgvar.u, 0u);
13+
}
14+
barrier();
15+
}
16+
17+
layout(binding = 0, std430) buffer tint_symbol_block_ssbo {
18+
S inner;
19+
} tint_symbol;
20+
21+
void tint_symbol_1(uint local_invocation_index) {
22+
tint_zero_workgroup_memory(local_invocation_index);
23+
uint x = atomicOr(wgvar.u, 0u);
24+
atomicExchange(tint_symbol.inner.u, x);
25+
}
26+
27+
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
28+
void main() {
29+
tint_symbol_1(gl_LocalInvocationIndex);
30+
return;
31+
}

‎test/tint/bug/tint/366314931.wgsl.expected.ir.dxc.hlsl‎

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,31 @@
1+
struct S {
2+
uint3 v;
3+
uint u;
4+
};
5+
6+
struct main_inputs {
7+
uint tint_local_index : SV_GroupIndex;
8+
};
9+
10+
11+
groupshared S wgvar;
12+
RWByteAddressBuffer output : register(u0);
13+
void main_inner(uint tint_local_index) {
14+
if ((tint_local_index == 0u)) {
15+
wgvar.v = (0u).xxx;
16+
uint v_1 = 0u;
17+
InterlockedExchange(wgvar.u, 0u, v_1);
18+
}
19+
GroupMemoryBarrierWithGroupSync();
20+
uint v_2 = 0u;
21+
InterlockedOr(wgvar.u, 0u, v_2);
22+
uint x = v_2;
23+
uint v_3 = 0u;
24+
output.InterlockedExchange(uint(12u), x, v_3);
25+
}
26+
27+
[numthreads(1, 1, 1)]
28+
void main(main_inputs inputs) {
29+
main_inner(inputs.tint_local_index);
30+
}
31+

‎test/tint/bug/tint/366314931.wgsl.expected.ir.fxc.hlsl‎

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,31 @@
1+
struct S {
2+
uint3 v;
3+
uint u;
4+
};
5+
6+
struct main_inputs {
7+
uint tint_local_index : SV_GroupIndex;
8+
};
9+
10+
11+
groupshared S wgvar;
12+
RWByteAddressBuffer output : register(u0);
13+
void main_inner(uint tint_local_index) {
14+
if ((tint_local_index == 0u)) {
15+
wgvar.v = (0u).xxx;
16+
uint v_1 = 0u;
17+
InterlockedExchange(wgvar.u, 0u, v_1);
18+
}
19+
GroupMemoryBarrierWithGroupSync();
20+
uint v_2 = 0u;
21+
InterlockedOr(wgvar.u, 0u, v_2);
22+
uint x = v_2;
23+
uint v_3 = 0u;
24+
output.InterlockedExchange(uint(12u), x, v_3);
25+
}
26+
27+
[numthreads(1, 1, 1)]
28+
void main(main_inputs inputs) {
29+
main_inner(inputs.tint_local_index);
30+
}
31+

‎test/tint/bug/tint/366314931.wgsl.expected.ir.glsl‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,11 @@
1+
SKIP: FAILED
2+
3+
../../src/tint/lang/glsl/writer/printer/printer.cc:1394 internal compiler error: TINT_UNREACHABLE unhandled core builtin: atomicStore
4+
********************************************************************
5+
* The tint shader compiler has encountered an unexpected error. *
6+
* *
7+
* Please help us fix this issue by submitting a bug report at *
8+
* crbug.com/tint with the source program that triggered the bug. *
9+
********************************************************************
10+
11+
tint executable returned error: signal: trace/BPT trap

0 commit comments

Comments
 (0)