More control barrier tests #579
Conversation
| @@ -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 | |||
There was a problem hiding this comment.
// 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.
There was a problem hiding this comment.
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?
There was a problem hiding this comment.
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.
| } | ||
| P0@sg 0, wg 0, qf 0 | P1@sg 0,wg 0, qf 0 ; | ||
| ld.atom.wg.sc0 r0, x | cbar.wg.semsc0 0 ; | ||
| cbar.wg.semsc0 0 | st.atom.wg.sc0 x, 1 ; |
There was a problem hiding this comment.
Good catch. The problem is that the alloy model specifies that ACQ + REL = SEMSC0 + SEMSC1, but our spriv-check.cat only checked half of the equality (i.e. one inclusion).
I fixed the cat file (which also required fixing the propagation of tags in the Vulkan visitor because we should not propagate semscX to the read event coming from a rmw.rel instruction).
There are also a lot of CADP test failing but it might be easier to fix the generation script than going one by one over the tests. @tonghaining can you do this?
|
LGTM |
hernanponcedeleon
force-pushed
the
ptx.arrive
branch
from
7e69acc to
948f7a3
Compare
@tonghaining can you take a look?