Patch sync insertion for redundant predicated writes

[shmsong]

This PR is a quick patch for redundant predicate sync insertion.

A sync is needed for redundant parallel type unless all use chain of the redundantly written value in smem/gmem arrive at redundant write consumers of the same parallel type.

This PR patches the insertion so that all redundant writes are sync'ed to avoid race conditions that may happen in devel TOT.

The detection for the cases where sync is not needed for redundant types will be in a follow up.

[shmsong]

Unrelated formatting.

[naoyam]

Are these meant to do something?

[shmsong]

Not really. Just making sure all the CA parameter has a value. Vaguely remember we didn't have a default behavior without any CA setting but it was a long while ago.

[naoyam]

What I mentioned in the MMA PR was that when we have a chain of redundant exprs, I was wondering if each would be synchronized. I added a variation of the test to see what happens, and here's the generated code:

__global__ void kernel1(Tensor<float, 1> T0, Tensor<float, 2> T1, Tensor<float, 2> T3) {
  alignas(16) extern __shared__ char array[];
  unsigned offset = 0;
  offset = alignBufferSize(offset, 16);
  float* T4 = reinterpret_cast<float*>(array + offset);
  offset += (32 * sizeof(float));
  // Alias Allocation - shared
  auto& T2 = T4;
  if ((((nvfuser_index_t)threadIdx.y) == 0)) {
    T4[((nvfuser_index_t)threadIdx.x)]
       = T0[(((nvfuser_index_t)threadIdx.x) * T0.stride[0])];
  }
  __barrier_sync(0);
  if ((((nvfuser_index_t)threadIdx.y) == 0)) {
    T2[((nvfuser_index_t)threadIdx.x)]
       = T4[((nvfuser_index_t)threadIdx.x)];
  }
  __barrier_sync(0);
  T3[(((nvfuser_index_t)threadIdx.y) * 32) + ((nvfuser_index_t)threadIdx.x)]
    = T2[((nvfuser_index_t)threadIdx.x)]
    + T1[(((nvfuser_index_t)threadIdx.y) * T1.stride[0]) + (((nvfuser_index_t)threadIdx.x) * T1.stride[1])];
}

The point I was making is the first sync is redundant.

I guess this is not common, so I think it's fine with this for now, but I wanted to clarify my concern for the future optimization.

[naoyam]

Please remove or merge the added test as you'd like. I just wanted to demonstrate the case.

[naoyam]

Thanks for the fix.

[shmsong]

What I mentioned in the MMA PR was that when we have a chain of redundant exprs, I was wondering if each would be synchronized. I added a variation of the test to see what happens, and here's the generated code:

__global__ void kernel1(Tensor<float, 1> T0, Tensor<float, 2> T1, Tensor<float, 2> T3) {
  alignas(16) extern __shared__ char array[];
  unsigned offset = 0;
  offset = alignBufferSize(offset, 16);
  float* T4 = reinterpret_cast<float*>(array + offset);
  offset += (32 * sizeof(float));
  // Alias Allocation - shared
  auto& T2 = T4;
  if ((((nvfuser_index_t)threadIdx.y) == 0)) {
    T4[((nvfuser_index_t)threadIdx.x)]
       = T0[(((nvfuser_index_t)threadIdx.x) * T0.stride[0])];
  }
  __barrier_sync(0);
  if ((((nvfuser_index_t)threadIdx.y) == 0)) {
    T2[((nvfuser_index_t)threadIdx.x)]
       = T4[((nvfuser_index_t)threadIdx.x)];
  }
  __barrier_sync(0);
  T3[(((nvfuser_index_t)threadIdx.y) * 32) + ((nvfuser_index_t)threadIdx.x)]
    = T2[((nvfuser_index_t)threadIdx.x)]
    + T1[(((nvfuser_index_t)threadIdx.y) * T1.stride[0]) + (((nvfuser_index_t)threadIdx.x) * T1.stride[1])];
}

The point I was making is the first sync is redundant.

I guess this is not common, so I think it's fine with this for now, but I wanted to clarify my concern for the future optimization.

Yes. I was planning on handling this in a follow up, i.e. a redundant write has a use-chain that has other redundant writes. So T0-> T2 is a redundant chain and T2->T3 isn't.

I think these vertical redundant chains seem to be not too bad to handle. Would need to think a bit more about if there're pathological horizontal patterns, probably worst case we arrive at sub-optimal code without re-considering expr ordering.

[shmsong]

Please remove or merge the added test as you'd like. I just wanted to demonstrate the case.

@naoyam Thanks for the repro. The new test case moved to #1687 and made it a fail for the redundant sync insertion.