In the relevant tests, can you also check whether a required syncthreads is actually inserted at its proper location? Result validation may not be able to expose potential race conditions.

I thought about that but didn't come up with a good solution. e.g. check the entire kernel string.
Do you have any ideas?

How about just traversing KIR to find a relevant ForLoop node and check whether it ends with a sync node? Finding relevant ForLoops may not be trivial for complex fusions, though.

Here are two other options:

1. Instead of traversing KIR, flatten KIR and check if sync exists at position x
2. Check if substring at position x is __syncthreads

Do we have access to the KIR from the test? Option 2 is the easiest to implement.

I think (read-only) accesses to KIR should be allowed for verification.

last_op_sync_ seems to be used to suppress inserting syncthreads when the last operation is also syncthreads. It seems that the "last operation" here can mean the last operation in a nested loop body. Something like this:

for i in X
  Write to SMEM
  __syncthreads()
  Read from SMEM
  for j in Y
    do something
     __syncthreads()
  end for
  // Insert __syncthreads() here?
end for

If I read the code correctly, in the above case, the last commented-out syncthreads is NOT inserted. However, since the trip count of the inner loop can be 0, the nested syncthreads may not be executed, so we need to have the last syncthreads.

Am I missing anything?

Don't all for-loops iterate from 0 to some positive integer? If we enforce this constraint, then every for-loop is entered.
Since we are performing this pass at compile time, we have to assume either all for-loops are entered or not.

Don't all for-loops iterate from 0 to some positive integer? If we enforce this constraint, then every for-loop is entered.
Since we are performing this pass at compile time, we have to assume either all for-loops are entered or not.

Are we sure all for loops must have a trip count greater than 0? That may be the case, actually, but not 100% sure.

I ran all the cpp unit tests with this assertion inside the IterDomain constructor and they passed.

  TORCH_INTERNAL_ASSERT(
      _start->isZeroInt(),
      "Cannot create an iter domain with a start that is not zero but received ",
      _extent,
      " .");

  TORCH_INTERNAL_ASSERT(
      !_extent->isZeroInt(),
      "Cannot create an iter domain with a extent that is zero but received ",
      _extent,
      " .");

Thanks for adding the check for inserted syncthreads. Looks very good!