interleaving | python library

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:

This sounds like a homework question, and it would be nice to see some effort on your own part, but it's pretty straightforward to do this in R.

Suppose your vector looks like this:

I am wondering if there is a race condition for the i variable

Yes, most definitely. The parent thread writes to i, which is a non-atomic variable, and the child threads read it, without any intervening synchronization. That's the exact definition of a data race in C++.

and if so, how does it interleave?

Data races in C++ cause undefined behavior, and any behavior you may observe does not have to be explainable by interleaving.

I tried to compile it and I constantly got the Thread ID printout as 3, but I was surprised because I thought the variable had to be global in order to be accessed by the various new threads?

No, it doesn't have to be global. Threads can access variables which are local to other threads if they are somehow passed a pointer or reference to such a variable.

This is what I thought would happen: thread 1 is created, Fun starts to run in thread 1 with myid = 0, main thread continues running and increments i, 2nd thread is created and the myid for that would be myid=1... and so on. And so the printout would be the myID in increments i/e 1,2,3

Well, nothing at all in your program forces those events to occur (or become observable) in that order, so there is really no basis for expecting that they will. It's entirely possible, for instance, that the three threads all get started, but don't get a chance to actually run until after the loop in main has completed, at which point i has the value 3. (Or rather, the memory where i used to be located, as it is now out of scope and its lifetime has ended - it's a separate bug that you don't prevent that from happening.)

I haven't yet found an explanation of the cufft output, but I can get the behaviour I want with cupyx.scipy.fft.rfft, which might be useful if anyone else finds the same problem.

This is how I understand it. For example,TrackingExecutor is shutting down before CrawlTask exit, this task may be also recorded as a taskCancelledAtShutdown, because if (isShutdown() && Thread.currentThread().isInterrupted()) in TrackingExecutor#execute may be true , but in fact this task has completed.

Instructions can't be reordered if they violate the sequential semantics of a program.

Simple example (assuming a=b=0):

This one worked for me:

No sequential-across-lanes pack until AVX-512, unfortunately. (And even then only for 1 register, or not with saturation.)

The in-lane behaviour of shuffles like vpacksswd and vpalignr is one of the major warts of AVX2 that make the 256-bit versions of those shuffles less useful than their __m128i versions. But on Intel, and Zen2 CPUs, it is often still best to use __m256i vectors with a vpermq at the end, if you need the elements in a specific order. (Or vpermd with a vector constant after 2 levels of packing: How do I efficiently reorder bytes of a __m256i vector (convert int32_t to uint8_t)?)

If your 32-bit elements came from unpacking narrower elements, and you don't care about order of the wider elements, you can widen with in-lane unpacks, which sets you up to pack back into the original order.

This is cheap for zero-extending unpacks: _mm256_unpacklo/hi_epi16 (with _mm256_setzero_si256()). That's as cheap as vpmovzxwd (_mm256_cvtepu16_epi32), and is actually better because you can do 256-bit loads of your source data and unpack two ways, instead of narrow loads to feed vpmovzx... which only works on data at the bottom of an input register. (And memory-source vpmovzx... ymm, [mem] can't micro-fuse the load with a YMM destination, only for the 128-bit XMM version, on Intel CPUs, so the front-end cost is the same as separate load and shuffle instructions.)

But that trick doesn't work work quite as nicely for data you need to sign-extend. vpcmpgtw to get high halves for vpunpckl/hwd does work, but vpermq when re-packing is about as good, just different execution-port pressure. So vpmovsxwd is simpler there.

Slicing up your data into odd/even instead of low/high can also work, e.g. to get 16 bit elements zero-extended into 32-bit elements:

The java.util Date-Time API and their formatting API, SimpleDateFormat are outdated and error-prone. It is recommended to stop using them completely and switch to the modern Date-Time API^*.

Solution using java.time, the modern API: