publications | Publications from Trail of Bits

Q1:

Why so?

This is an exceptional case. The programming guide doesn't give a complete description of the detailed behavior of __shfl_sync() to understand this case (that I know of), although the statements given in the programming guide are correct. To get a detailed behavioral description of the instruction, I suggest looking at the PTX guide:

shfl.sync will cause executing thread to wait until all non-exited threads corresponding to membermask have executed shfl.sync with the same qualifiers and same membermask value before resuming execution.

Careful study of that statement may be sufficient for understanding. But we can unpack it a bit.

• As already stated, this doesn't apply to compute capability less than 7.0. For those compute capabilities, all threads named in member mask must participate in the exact line of code/instruction, and for any warp lane's result to be valid, the source lane must be named in the member mask and must not be excluded from participation due to forced divergence at that line of code
• I would describe __shfl_sync() as "exceptional" in the cc7.0+ case because it causes partial-warp execution to pause at that point of the instruction, and control/scheduling would then be given to other warp fragments. Those other warp fragments would be allowed to proceed (due to Volta ITS) until all threads named in the member mask have arrived at a __shfl_sync() statement that "matches", i.e. has the same member mask and qualifiers. Then the shuffle statement executes. Therefore, in spite of the enforced divergence at this point, the __shfl_sync() operation behaves as if the warp were sufficiently converged at that point to match the member mask.

I would describe that as "unusual" or "exceptional" behavior.

If so, the programming guide also states that "if the target thread is inactive, the retrieved value is undefined" and that "threads can be inactive for a variety of reasons including ... having taken a different branch path than the branch path currently executed by the warp."

In my view, the "if the target thread is inactive, the retrieved value is undefined" statement most directly applies to compute capability less than 7.0. It also applies to compute capability 7.0+ if there is no corresponding/matching shuffle statement elsewhere, that the thread scheduler can use to create an appropriate warp-wide shuffle op. The provided code example only gives sensible results because there is a matching op both in the if portion and the else portion. If we made the else portion an empty statement, the code would not give interesting results for any thread in the warp.

Q2:

On GPUs with current implementation of independent thread scheduling (Volta~Ampere), when the if branch is executed, are inactive threads still doing NOOP? That is, should I still think of warp execution as lockstep?

If we consider the general case, I would suggest that the way to think about inactive threads is that they are inactive. You can call that a NOOP if you like. Warp execution at that point is not "lockstep" across the entire warp, because of the enforced divergence (in my view). I don't wish to argue the semantics here. If you feel an accurate description there is "lockstep execution given that some threads are executing the instruction and some aren't", that is ok. We have now seen, however, that for the specific case of the shuffle sync ops, the Volta+ thread scheduler works around the enforced divergence, combining ops from different execution paths, to satisfy the expectations for that particular instruction.

Q3:

Is synchronization (such as __shfl_sync, __ballot_sync) the only cause for statement interleaving (statements A and B from the if branch interleaved with X and Y from the else branch)?

I don't believe so. Any time you have a conditional if-else construct that causes a division intra-warp, you have the possibility for interleaving. I define Volta+ interleaving (figure 12) as forward progress of one warp fragment, followed by forward progress of another warp fragment, perhaps with continued alternation, prior to reconvergence. This ability to alternate back and forth doesn't only apply to the sync ops. Atomics could be handled this way (that is a particular use-case for the Volta ITS model - e.g. use in a producer/consumer algorithm or for intra-warp negotiation of locks - referred to as "starvation free" in the previously linked article) and we could also imagine that a warp fragment could stall for any number of reasons (e.g. a data dependency, perhaps due to a load instruction) which prevents forward progress of that warp fragment "for a while". I believe the Volta ITS can handle a variety of possible latencies, by alternating forward progress scheduling from one warp fragment to another. This idea is covered in the paper in the introduction ("load-to-use"). Sorry, I won't be able to provide an extended discussion of the paper here.

EDIT: Responding to a question in the comments, paraphrased "Under what circumstances can the scheduler use a subsequent shuffle op to satisfy the needs of a warp fragment that is waiting for shuffle op completion?"

First, let's notice that the PTX description above implies some sort of synchronization. The scheduler has halted execution of the warp fragment that encounters the shuffle op, waiting for other warp fragments to participate (somehow). This is a description of synchronization.

Second, the PTX description makes allowance for exited threads.

What does all this mean? The simplest description is just that a subsequent "matching" shuffle op can/will be "found by the scheduler", if it is possible, to satisfy the shuffle op. let's consider some examples.

Test case 1: As given in the programming guide, we see expected results:

You can remove the related objects using this query:

Unfortunately, I don't think it's possible to do what you're hoping directly with a Prisma Query at the moment. I can suggest two possible workarounds though:

Approach 1: Fetch the most recent publicationStatus along with the publication and check the status inside your node application.

This is what the query would look like:

You can do calculations within a function for the x and y values to construct the ggplot which extends the circle all the way round and gives labels correct heights.

I've adapted a function to work with other datasets. This takes a dataset in a tidy format, with:

• a 'year' column
• one row per 'event'
• a grouping variable (such as country)

I've used Nobel laurate data from here as an example dataset to show the function in practice. Data setup:

Use pd.to_datetime:

Redis Server does not use Java and is therefore not impacted by this vulnerability.

See more here: https://redis.com/security/notice-apache-log4j2-cve-2021-44228/

You are binding all inputs to a single property comment.text. You should create one property for each input and bind input with it.

Unless you are creating a Spring Boot application, the Spring Boot Gradle plugin should not be applied. It is not meant for creating Spring Boot libraries or Spring libraries in general. It is for creating Spring Boot applications. The Spring Boot Gradle plugin reacts to various plugins applied to the project as documented in the documentation.

Second if you want your published artifact to be consumed by another, then you should use the java-library instead.

With the said, your Gradle build should look something like:

components.java should contains the shadow jar and the original jar, so simply remove the artifact shadowJar line will work.