warp | Bringing Solidity to Starknet at warp speed | Cryptocurrency library

I'm guessing you have an issue with your initial user state and your PublicRoute component redirecting before the app has actually computed/resolved the user's auth status.

The issue is that your resolved/unauthenticated state is masked by your initial state value. You probably want to use an indeterminant initial state and conditionally render null or some loading indicator until the authentication status has been resolved.

Example:

A dispatch of compute defines with its parameters the global workgroup. The global workgroup has x×y×z invocations.

Each of those invocations are divided into local groups (defined by the shader). A local workgroup also has another set of x×y×z invocations.

A local workgroup is partitioned into subgroups. Its invocations are rearranged into subgroups. A subgroup has (1-dimensional) SubgroupSize amount of invocations, which all need not be assigned a local workgroup invocation. And a subgroup must not span over multiple local workgroups; it can use only invocations from a single local workgroup.

Otherwise how this partitioning is done seems largely unspecified, except that under very specific conditions you are guaranteed full subgroups, which means none of the invocations in a subgroup of SubgroupSize will stay vacant. If those conditions are not fulfilled, then the driver may keep some invocations inactive in the subgroup as it sees fit.

If the local workgroup has in total less invocations than SubgroupSize, then some of the invocations of the subgroup indeed need to stay inactive as there are not enough available local workgroup invocations to fill even one subgroup.

Note that the reported issue is only a problem if compression is used. The websocket-shootout example uses Ws.defaultConnectionOptions which means compression is disabled. As long as you also leave compression disabled, you shouldn't run into any problems with this issue.

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:

1 - COMPRESSION

Please find the solution for the problem of file compression. To be honest, I have already been confronted with the same problem as you and, at the time, I was racking my brains... to finally find the solution which is quite simple (once we know it!): you must not put any spaces (i.e. "COMPRESS=DEFLATE" and not "COMPRESS = DEFLATE")

So, please find below a small reprex.

Reprex

C++ Copy Constructors: must I spell out all member variables in the initializer list?

Yes, if you write a user defined copy constructor, then you must write an initialiser for every sub object - unless you wish to default initialise them, in which case you don't need any initialiser - or if you can use a default member initialiser.

the object contains some member variables which are pointers to its other member variables)

This is a design that should be avoided when possible. Not only does this force you to define custom copy and move assignment operators and constructors, but it is often unnecessarily inefficient.

But, in case that is necessary for some reason - or custom special member functions are needed for any other reason - you can achieve clean code by combining the normally copying parts into a separate dummy class. That way the the user defined constructor has only one sub object to initialise.

Like this:

Thanks to @Stargateur's comment I realised what I was doing wrong. Here's a working prototype should anyone else get stuck at a similar place.

I think that the issue here is that you are sending the message and then immediately aborting the background task:

Your psWarpOptions->hCutline should be a polygon, not a layer.

Also the cutline should be in source pixel/line coordinates.

Check TransformCutlineToSource from gdalwarp_lib.cpp, you can probably simply get the code from there.

This particular GDAL operation, when called from C++, is so full of pitfalls - and there are so many open questions about it here - that I am reproducing a full working example:

Warping (reprojecting) a raster image with a polygon mask (cutline):