Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

More control barrier tests #579

Merged
merged 6 commits into from
Nov 30, 2023
Merged

More control barrier tests #579

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
f4de543
Add arrive control barrier to PTX
hernan-poncedeleon Nov 28, 2023
e722344
Add new barrier tests for vulkan
hernan-poncedeleon Nov 28, 2023
34e144a
Add new barrier tests for ptx
hernan-poncedeleon Nov 28, 2023
8be3663
Fix ACQ+REL = SEMSC0 + SEMC1 check
hernan-poncedeleon Nov 30, 2023
dcd0c4d
Fix benchmarks with ACQ+REL = SEMSC0 + SEMC1 problem
hernan-poncedeleon Nov 30, 2023
948f7a3
Regenerate CADP tests guaranteeing ACQ+REL = SEMSC0 + SEMC1
hernan-poncedeleon Nov 30, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
The table of contents is too big for display.
Diff view
Diff view
  •  
  •  
  •  
4 changes: 4 additions & 0 deletions cat/spirv-check.cat
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -193,6 +193,10 @@ let acqIsSem = ([ACQ] \ [SEMSC0 | SEMSC1])
flag ~empty acqIsSem as checkAcqIsSem
let relIsSem = ([REL] \ [SEMSC0 | SEMSC1])
flag ~empty relIsSem as checkRelIsSem
let semcs0IsAcqRel = ([SEMSC0] \ [ACQ | REL])
flag ~empty semcs0IsAcqRel as checkSemcs0IsAcqRel
let semcs1IsAcqRel = ([SEMSC1] \ [ACQ | REL])
flag ~empty semcs1IsAcqRel as checkSemcs1IsAcqRel

// no pair of the same control barrier instance can be in the same thread
let scbarinstIsPo = (syncbar & (po | po^-1))
Expand Down
8 changes: 7 additions & 1 deletion dartagnan/src/main/antlr4/LitmusPTX.g4
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -151,7 +151,12 @@ fenceAlias
;

barrier
: Barrier Period CTA Period Sync value
: Barrier Period CTA Period barrierMode value
;

barrierMode
: Sync
| Arrive
;

atomInstruction
Expand Down Expand Up @@ -301,6 +306,7 @@ Fence : 'fence';
Barrier : 'bar';

Sync : 'sync';
Arrive : 'arrive';

CTA : 'cta';
GPU : 'gpu';
Expand Down
3 changes: 3 additions & 0 deletions dartagnan/src/main/java/com/dat3m/dartagnan/parsers/program/visitors/VisitorLitmusPTX.java
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -311,6 +311,9 @@ public Object visitFenceAlias(LitmusPTXParser.FenceAliasContext ctx) {
public Object visitBarrier(LitmusPTXParser.BarrierContext ctx) {
Expression fenceId = (Expression) ctx.value().accept(this);
Event fence = EventFactory.newFenceWithId(ctx.getText().toLowerCase(), fenceId);
if(ctx.barrierMode().Arrive() != null) {
fence.addTags(Tag.PTX.ARRIVE);
}
return programBuilder.addChild(mainThread, fence);
}

Expand Down
2 changes: 2 additions & 0 deletions dartagnan/src/main/java/com/dat3m/dartagnan/program/event/Tag.java
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -285,6 +285,8 @@ public static final class PTX {
public static final String CON = "CON"; // CONSTANT
// Virtual memory
public static final String ALIAS = "ALIAS";
// Barrier Mode
public static final String ARRIVE = "__ARRIVE";

public static String loadMO(String mo) {
return switch (mo) {
Expand Down
6 changes: 1 addition & 5 deletions dartagnan/src/main/java/com/dat3m/dartagnan/program/processing/compilation/VisitorPTX.java
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,13 +1,9 @@
package com.dat3m.dartagnan.program.processing.compilation;

import com.dat3m.dartagnan.expression.Expression;
import com.dat3m.dartagnan.expression.op.COpBin;
import com.dat3m.dartagnan.program.Register;
import com.dat3m.dartagnan.program.event.Tag;
import com.dat3m.dartagnan.program.event.arch.ptx.PTXAtomCAS;
import com.dat3m.dartagnan.program.event.arch.ptx.PTXAtomExch;
import com.dat3m.dartagnan.program.event.arch.ptx.PTXAtomOp;
import com.dat3m.dartagnan.program.event.arch.ptx.PTXRedOp;
import com.dat3m.dartagnan.program.event.arch.ptx.*;
import com.dat3m.dartagnan.program.event.core.Event;
import com.dat3m.dartagnan.program.event.core.Load;
import com.dat3m.dartagnan.program.event.core.rmw.RMWStore;
Expand Down
18 changes: 18 additions & 0 deletions ...agnan/src/main/java/com/dat3m/dartagnan/program/processing/compilation/VisitorVulkan.java
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -66,6 +66,15 @@ private void propagateTags(Event source, Event target) {
if (source.hasTag(Tag.Vulkan.SEM_VISIBLE)) {
target.addTags(Tag.Vulkan.SEM_VISIBLE);
}
// If a RMW is a release, we do not propagate semscX to the read
if (!(source.hasTag(Tag.Vulkan.ACQUIRE) || source.hasTag(Tag.Vulkan.ACQ_REL))) {
if (target.hasTag(Tag.Vulkan.SEMSC0)) {
target.removeTags(Tag.Vulkan.SEMSC0);
}
if (target.hasTag(Tag.Vulkan.SEMSC1)) {
target.removeTags(Tag.Vulkan.SEMSC1);
}
}
} else if (target instanceof Store) {
// Atomic stores are always available
if (source.hasTag(Tag.Vulkan.ATOM)) {
Expand All @@ -74,6 +83,15 @@ private void propagateTags(Event source, Event target) {
if (source.hasTag(Tag.Vulkan.SEM_AVAILABLE)) {
target.addTags(Tag.Vulkan.SEM_AVAILABLE);
}
// If a RMW is a acquire, we do not propagate semscX to the write
if (!(source.hasTag(Tag.Vulkan.RELEASE) || source.hasTag(Tag.Vulkan.ACQ_REL))) {
if (target.hasTag(Tag.Vulkan.SEMSC0)) {
target.removeTags(Tag.Vulkan.SEMSC0);
}
if (target.hasTag(Tag.Vulkan.SEMSC1)) {
target.removeTags(Tag.Vulkan.SEMSC1);
}
}
}
}
}
4 changes: 3 additions & 1 deletion dartagnan/src/main/java/com/dat3m/dartagnan/wmm/analysis/RelationAnalysis.java
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -961,7 +961,9 @@ public Knowledge visitSyncBarrier(Relation sync_bar) {
List<FenceWithId> fenceEvents = program.getThreadEvents(FenceWithId.class);
for (FenceWithId e1 : fenceEvents) {
for (FenceWithId e2 : fenceEvents) {
if(exec.areMutuallyExclusive(e1, e2)) {
// “A bar.sync or bar.red or bar.arrive operation synchronizes with a bar.sync
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

what is bar.red?

Copy link
Owner Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

// Use bar.red.and to compare results across the entire CTA:
setp.eq.u32 p,r1,r2; // p is True if r1==r2
bar.cta.red.and.pred r3,1,p; // r3=AND(p) forall threads in CTA

// Use bar.red.popc to compute the size of a group of threads
// that have a specific condition True:
setp.eq.u32 p,r1,r2; // p is True if r1==r2
bar.cta.red.popc.u32 r3,1,p; // r3=SUM(p) forall threads in CTA

We don't currently support those instructions, but I just copied the comment above from the documentation.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Okay, I think the current implementation support (bar.sync, bar.sync) and (bar.arrive, bar.sync).
Shall we also update the comments on two ptx cat files about the sync_barrier relation?

Copy link
Owner Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't think we need that level of detail in the cat file. E.g. we do not specify either all the cases covered by sr.

// or bar.red operation executed on the same barrier.”
if(exec.areMutuallyExclusive(e1, e2) || e2.hasTag(PTX.ARRIVE)) {
continue;
}
may.add(e1, e2);
Expand Down
5 changes: 5 additions & 0 deletions dartagnan/src/test/resources/PTXv6_0-expected.csv
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -79,3 +79,8 @@ litmus/PTX/Manual/SL-future-plus.litmus,0
litmus/PTX/Manual/SL-future-minus.litmus,1
litmus/PTX/Manual/Coherence.litmus,0
litmus/PTX/Manual/Coherence-weak.litmus,1
litmus/PTX/Manual/PC-bar-sync-arrive.litmus,1
litmus/PTX/Manual/PC-bar-sync-sync-1.litmus,1
litmus/PTX/Manual/PC-bar-sync-sync-2.litmus,1
litmus/PTX/Manual/PC-bar-sync-sync-3.litmus,1
litmus/PTX/Manual/PC-bar-sync-sync-4.litmus,1
5 changes: 5 additions & 0 deletions dartagnan/src/test/resources/PTXv7_5-expected.csv
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -201,3 +201,8 @@ litmus/PTX/Manual/XF-Barrier-rlx.litmus,1
litmus/PTX/Manual/XF-Barrier-weak.litmus,1
litmus/PTX/Manual/Coherence.litmus,0
litmus/PTX/Manual/Coherence-weak.litmus,1
litmus/PTX/Manual/PC-bar-sync-arrive.litmus,1
litmus/PTX/Manual/PC-bar-sync-sync-1.litmus,1
litmus/PTX/Manual/PC-bar-sync-sync-2.litmus,1
litmus/PTX/Manual/PC-bar-sync-sync-3.litmus,1
litmus/PTX/Manual/PC-bar-sync-sync-4.litmus,1
7 changes: 6 additions & 1 deletion dartagnan/src/test/resources/VULKAN-CK-expected.csv
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -230,4 +230,9 @@ litmus/VULKAN/CADP/3_threads_4_instructions/100_simple.litmus,1
litmus/VULKAN/CADP/3_threads_4_instructions/101_simple.litmus,1
litmus/VULKAN/CADP/3_threads_4_instructions/102_simple.litmus,1
litmus/VULKAN/CADP/3_threads_4_instructions/103_simple.litmus,1
litmus/VULKAN/CADP/3_threads_4_instructions/104_simple.litmus,1
litmus/VULKAN/CADP/3_threads_4_instructions/104_simple.litmus,1
litmus/VULKAN/Manual/PC-bar-acq-rel-atom.litmus,1
litmus/VULKAN/Manual/PC-bar-acq-rel-nonpriv.litmus,1
litmus/VULKAN/Manual/PC-bar-acq-rel-priv.litmus,1
litmus/VULKAN/Manual/PC-bar-atom.litmus,1
litmus/VULKAN/Manual/PC-bar-nonpriv.litmus,1
5 changes: 5 additions & 0 deletions dartagnan/src/test/resources/VULKAN-DR-expected.csv
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -86,3 +86,8 @@ litmus/VULKAN/Manual/XF-Barrier-relacq.litmus,1
litmus/VULKAN/Manual/XF-Barrier-rlx.litmus,0
litmus/VULKAN/Manual/XF-Barrier-weak.litmus,0
litmus/VULKAN/Manual/asmo-weak.litmus,0
litmus/VULKAN/Manual/PC-bar-acq-rel-atom.litmus,1
litmus/VULKAN/Manual/PC-bar-acq-rel-nonpriv.litmus,1
litmus/VULKAN/Manual/PC-bar-acq-rel-priv.litmus,0
litmus/VULKAN/Manual/PC-bar-atom.litmus,1
litmus/VULKAN/Manual/PC-bar-nonpriv.litmus,0
5 changes: 5 additions & 0 deletions dartagnan/src/test/resources/VULKAN-expected.csv
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -89,3 +89,8 @@ litmus/VULKAN/Manual/CoWW-RR.litmus,1
litmus/VULKAN/Manual/XF-Barrier-relacq.litmus,0
litmus/VULKAN/Manual/XF-Barrier-rlx.litmus,1
litmus/VULKAN/Manual/XF-Barrier-weak.litmus,1
litmus/VULKAN/Manual/PC-bar-acq-rel-atom.litmus,1
litmus/VULKAN/Manual/PC-bar-acq-rel-nonpriv.litmus,1
litmus/VULKAN/Manual/PC-bar-acq-rel-priv.litmus,0
litmus/VULKAN/Manual/PC-bar-atom.litmus,0
litmus/VULKAN/Manual/PC-bar-nonpriv.litmus,0
14 changes: 14 additions & 0 deletions litmus/PTX/Manual/PC-bar-sync-arrive.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,14 @@
PTX PC-bar-sync-arrive
"Producer-Consumer pattern using named barriers.
Arrive does not block P1.
Load and store are not ordered."
{
x=0;
P0:r0=0
}
P0@cta 0,gpu 0 | P1@cta 0,gpu 0 ;
ld.weak r0, x | bar.cta.arrive 1 ;
bar.cta.sync 0 | st.weak x, 1 ;
bar.cta.arrive 1 | bar.cta.sync 0 ;
exists
(P0:r0 == 1)
13 changes: 13 additions & 0 deletions litmus/PTX/Manual/PC-bar-sync-sync-1.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,13 @@
PTX PC-bar-sync-sync-1
"Producer-Consumer pattern using named barriers.
Load is ordered before store due to sync over 0."
{
x=0;
P0:r0=0
}
P0@cta 0,gpu 0 | P1@cta 0,gpu 0 ;
ld.weak r0, x | bar.cta.sync 0 ;
bar.cta.sync 0 | st.weak x, 1 ;
bar.cta.sync 1 | bar.cta.sync 1 ;
~exists
(P0:r0 == 1)
13 changes: 13 additions & 0 deletions litmus/PTX/Manual/PC-bar-sync-sync-2.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,13 @@
PTX PC-bar-sync-sync-2
"Producer-Consumer pattern using named barriers.
Load is ordered before store due to sync over 0."
{
x=0;
P0:r0=0
}
P0@cta 0,gpu 0 | P1@cta 0,gpu 0 ;
ld.weak r0, x | bar.cta.sync 0 ;
bar.cta.sync 0 | st.weak x, 1 ;
bar.cta.sync 1 | bar.cta.sync 1 ;
exists
(P0:r0 == 0)
13 changes: 13 additions & 0 deletions litmus/PTX/Manual/PC-bar-sync-sync-3.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,13 @@
PTX PC-bar-sync-sync-3
"Producer-Consumer pattern using named barriers.
Deadlock."
{
x=0;
P0:r0=0
}
P0@cta 0,gpu 0 | P1@cta 0,gpu 0 ;
ld.weak r0, x | bar.cta.sync 1 ;
bar.cta.sync 0 | st.weak x, 1 ;
bar.cta.sync 1 | bar.cta.sync 0 ;
~exists
(P0:r0 == 1)
13 changes: 13 additions & 0 deletions litmus/PTX/Manual/PC-bar-sync-sync-4.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,13 @@
PTX PC-bar-sync-sync-4
"Producer-Consumer pattern using named barriers.
Deadlock."
{
x=0;
P0:r0=0
}
P0@cta 0,gpu 0 | P1@cta 0,gpu 0 ;
ld.weak r0, x | bar.cta.sync 1 ;
bar.cta.sync 0 | st.weak x, 1 ;
bar.cta.sync 1 | bar.cta.sync 0 ;
~exists
(P0:r0 == 0)
12 changes: 6 additions & 6 deletions litmus/VULKAN/CADP/2_threads_2_instructions/0_simple.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,10 +5,10 @@ Mem0=0;
P0:r0=0;
P1:r1=0;
}
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
rmw.atom.wg.sc0.semsc0 r0, Mem0, 0 | rmw.atom.wg.sc0.semsc0 r1, Mem0, 1 ;
bne r0, 1, LC01 | bne r1, 0, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
rmw.atom.wg.sc0 r0, Mem0, 0 | rmw.atom.wg.sc0 r1, Mem0, 1 ;
bne r0, 1, LC01 | bne r1, 0, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
exists 0==0
12 changes: 6 additions & 6 deletions litmus/VULKAN/CADP/2_threads_2_instructions/1_simple.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,10 +5,10 @@ Mem0=0;
P0:r0=0;
P1:r1=0;
}
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
rmw.atom.wg.sc0.semsc0 r0, Mem0, 1 | rmw.atom.wg.sc0.semsc0 r1, Mem0, 0 ;
bne r0, 0, LC01 | bne r1, 1, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
rmw.atom.wg.sc0 r0, Mem0, 1 | rmw.atom.wg.sc0 r1, Mem0, 0 ;
bne r0, 0, LC01 | bne r1, 1, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
exists 0==0
12 changes: 6 additions & 6 deletions litmus/VULKAN/CADP/2_threads_2_instructions/2_simple.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,10 +5,10 @@ Mem0=0;
P0:r0=0;
P1:r1=0;
}
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
rmw.atom.wg.sc0.semsc0 r0, Mem0, 0 | rmw.atom.wg.sc0.semsc0 r1, Mem0, 1 ;
bne r0, 0, LC01 | bne r1, 0, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
rmw.atom.wg.sc0 r0, Mem0, 0 | rmw.atom.wg.sc0 r1, Mem0, 1 ;
bne r0, 0, LC01 | bne r1, 0, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
exists 0==0
12 changes: 6 additions & 6 deletions litmus/VULKAN/CADP/2_threads_2_instructions/3_simple.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,10 +5,10 @@ Mem0=0;
P0:r0=0;
P1:r1=0;
}
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
rmw.atom.wg.sc0.semsc0 r0, Mem0, 1 | rmw.atom.wg.sc0.semsc0 r1, Mem0, 0 ;
bne r0, 0, LC01 | bne r1, 0, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
rmw.atom.wg.sc0 r0, Mem0, 1 | rmw.atom.wg.sc0 r1, Mem0, 0 ;
bne r0, 0, LC01 | bne r1, 0, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
exists 0==0
12 changes: 6 additions & 6 deletions litmus/VULKAN/CADP/2_threads_2_instructions/4_simple.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,10 +4,10 @@ VULKAN 4_simple
Mem0=0;
P0:r0=0;
}
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
ld.atom.wg.sc0.semsc0 r0, Mem0 | st.atom.wg.sc0 Mem0, 1 ;
bne r0, 0, LC01 | LC11: ;
goto LC00 | ;
LC01: | ;
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
ld.atom.wg.sc0 r0, Mem0 | st.atom.wg.sc0 Mem0, 1 ;
bne r0, 0, LC01 | LC11: ;
goto LC00 | ;
LC01: | ;
exists 0==0
12 changes: 6 additions & 6 deletions litmus/VULKAN/CADP/2_threads_2_instructions/5_simple.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,10 +4,10 @@ VULKAN 5_simple
Mem0=0;
P1:r0=0;
}
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
st.atom.wg.sc0 Mem0, 1 | ld.atom.wg.sc0.semsc0 r0, Mem0 ;
LC01: | bne r0, 0, LC11 ;
| goto LC10 ;
| LC11: ;
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
st.atom.wg.sc0 Mem0, 1 | ld.atom.wg.sc0 r0, Mem0 ;
LC01: | bne r0, 0, LC11 ;
| goto LC10 ;
| LC11: ;
exists 0==0
12 changes: 6 additions & 6 deletions litmus/VULKAN/CADP/2_threads_2_instructions/6_simple.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,10 +5,10 @@ Mem0=0;
P0:r0=0;
P1:r1=0;
}
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
rmw.atom.wg.sc0.semsc0 r0, Mem0, 1 | ld.atom.wg.sc0.semsc0 r1, Mem0 ;
bne r0, 0, LC01 | bne r1, 0, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
rmw.atom.wg.sc0 r0, Mem0, 1 | ld.atom.wg.sc0 r1, Mem0 ;
bne r0, 0, LC01 | bne r1, 0, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
exists 0==0
12 changes: 6 additions & 6 deletions litmus/VULKAN/CADP/2_threads_2_instructions/7_simple.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,10 +5,10 @@ Mem0=0;
P0:r0=0;
P1:r1=0;
}
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
ld.atom.wg.sc0.semsc0 r0, Mem0 | rmw.atom.wg.sc0.semsc0 r1, Mem0, 1 ;
bne r0, 0, LC01 | bne r1, 0, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
ld.atom.wg.sc0 r0, Mem0 | rmw.atom.wg.sc0 r1, Mem0, 1 ;
bne r0, 0, LC01 | bne r1, 0, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
exists 0==0
16 changes: 8 additions & 8 deletions litmus/VULKAN/CADP/2_threads_3_instructions/0_simple.litmus
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,12 +5,12 @@ Mem0=0;
P0:r0=0;
P1:r1=0;
}
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
rmw.atom.wg.sc0.semsc0 r0, Mem0, 1 | rmw.atom.wg.sc0.semsc0 r1, Mem0, 0 ;
bne r0, 0, LC01 | bne r1, 1, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
| st.atom.wg.sc0 Mem0, 0 ;
| LC12: ;
P0@sg 0,wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ;
LC00: | LC10: ;
rmw.atom.wg.sc0 r0, Mem0, 1 | rmw.atom.wg.sc0 r1, Mem0, 0 ;
bne r0, 0, LC01 | bne r1, 1, LC11 ;
goto LC00 | goto LC10 ;
LC01: | LC11: ;
| st.atom.wg.sc0 Mem0, 0 ;
| LC12: ;
exists 0==0
Loading