Add Circular-Buffer TMA Support for Pointwise and Normalization Kernels

[rdspring1]

Pipelining - (Multiple mbarriers per TensorView)

Launch multiple TMA operations simultaneously but process each stage as they become available.

Motivation

Overlap data movement with computation

Pseudo-code

 for each stage of producer TV:
   launch TMA operation for stage
 end for

 for each stage of consumer:
   wait for corresponding TMA stage to become available
 end for

Example: Synchronous TMA

      uint64_t* T6 = reinterpret_cast<uint64_t*>(array + smem_offset + 8208LL);
      mbarrier::init(toSmem(T6), 1U);
      __syncthreads();
      if (b12) {
        uint64_t i26;
        i26 = mbarrier::arriveExpectTX(toSmem(T6), 4096U);
        Hopper::cpAsyncBulkTensorTileG2S((Hopper::CpAsyncBulkTensorTileG2SIndex<2>{ ptr4, a23, toSmem(T6) }), i6);
        mbarrier::wait(toSmem(T6), i26);
      }
      __syncthreads();
      mbarrier::inval(toSmem(T6));
      NVFUSER_UPDATE_MAGIC_ZERO;
      uint64_t* T7 = reinterpret_cast<uint64_t*>(array + smem_offset + 8192LL);
      mbarrier::init(toSmem(T7), 1U);
      __syncthreads();
      if (b12) {
        uint64_t i27;
        i27 = mbarrier::arriveExpectTX(toSmem(T7), 4096U);
        Hopper::cpAsyncBulkTensorTileG2S((Hopper::CpAsyncBulkTensorTileG2SIndex<2>{ ptr7, a23, toSmem(T7) }), i8);
        mbarrier::wait(toSmem(T7), i27);
      }
      __syncthreads();
      mbarrier::inval(toSmem(T7));
      NVFUSER_UPDATE_MAGIC_ZERO;
      
      // compute

[kevinstephano]

We decided not to do this.