Fix SIGABRT for 1D circular-buffer split pointer (ptr[i % N])

Extend rewriteSplitPtr to handle rank-1 split pointers produced by
PtrAnalysis::visitOperandRem. Previously the assert required rank == 2.

The failing kernel was triton_poi_fused_mul_4, the RMSNorm weight-broadcast
multiply (weight[i % 2048] * input[i]), when compiled via
torch.compile/inductor with --enable-triton:

  @triton_heuristics.pointwise(
      size_hints={'x': 65536},
      filename=__file__,
      triton_meta={'signature': {'in_out_ptr0': '*fp32', 'in_ptr0': '*fp32',
                                 'xnumel': 'i32', 'XBLOCK': 'constexpr'},
                   'device': DeviceProperties(type='qaic', index=0,
                       multi_processor_count=16, cc=None, major=None,
                       regs_per_multiprocessor=None,
                       max_threads_per_multi_processor=None,
                       max_threads_per_block=1024, warp_size=32),
                   'constants': {}, 'native_matmul': False,
                   'configs': [{(0,): [['tt.divisibility', 16]],
                                (2,): [['tt.divisibility', 16]]}],
                   'enable_fp_fusion': True},
      inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(),
                     'kernel_name': 'triton_poi_fused_mul_4',
                     'mutated_arg_names': ['in_out_ptr0'],
                     'optimize_mem': True, 'no_x_dim': False, ...},
      min_elem_per_thread=0
  )
  @triton.jit
  def triton_poi_fused_mul_4(in_out_ptr0, in_ptr0, xnumel, XBLOCK : tl.constexpr):
      xoffset = tl.program_id(0) * XBLOCK
      xindex = xoffset + tl.arange(0, XBLOCK)[:]
      xmask = xindex < xnumel
      x0 = (xindex % 2048)
      x2 = xindex
      tmp0 = tl.load(in_ptr0 + (x0), xmask, eviction_policy='evict_last')
      tmp1 = tl.load(in_out_ptr0 + (x2), xmask)
      tmp2 = tmp0 * tmp1
      tl.store(in_out_ptr0 + (x2), tmp2, xmask)

The load of in_ptr0 + (x0) — where x0 = xindex % 2048 — is a purely 1D
circular-buffer access. PtrAnalysis::visitOperandRem sets state.shape[0] = N
(here 2048), making isSplitPtr() return true with parentShape.size() == 1.
rewriteSplitPtr previously asserted rank == 2, causing the SIGABRT.

Add create1DCastOps (two ReinterpretCastOps for the two contiguous
circular-buffer segments), create1DCopies, get1DSubviews, and a WRAP_1D
store path to StoreConverter. Add wraparound_1d.mlir FileCheck test.

Author: dpathikonda

lib/Conversion/StructuredToMemref/StructuredToMemref.cpp

@@ -44,6 +44,7 @@ using namespace mlir;
 
 static const std::string WRAP_SIDE_BY_SIDE = "wrap_side_by_side";
 static const std::string WRAP_STACKED = "wrap_stacked";
+static const std::string WRAP_1D = "wrap_1d";
 
 static memref::SubViewOp getSubview(int rank, ArrayRef<OpFoldResult> dims,
                                     Value source, Location loc, OpBuilder &b) {
@@ -456,33 +457,108 @@ struct MakeTensorPtrConverter
     return {cast1, cast2};
   }
 
+  // Handles 1D circular-buffer wrap-around: ptr[i % N] where the tile
+  // [startOffset, startOffset+XBLOCK) may straddle the modulo boundary N.
+  //
+  //   x             = startOffset % N  (position within circular buffer)
+  //   d1            = min(x + XBLOCK, N) - x  (elements before wrap)
+  //   d2            = XBLOCK - d1              (elements after wrap, from 0)
+  //
+  //   chunk1: base[startOffset .. startOffset + d1)   stride 1
+  //   chunk2: base[startOffset - x .. startOffset - x + d2)  stride 1
+  //           (i.e. from the start of the row in flat memory)
+  //
+  // When d2 == 0 (no wrap, common case) chunk2 is a zero-size memref; the
+  // downstream CopyOp over it is a no-op.
+  std::pair<memref::ReinterpretCastOp, memref::ReinterpretCastOp>
+  create1DCastOps(tts::MakeTensorPtrOp op, OpAdaptor adaptor,
+                  ConversionPatternRewriter &rewriter) const {
+    auto loc = op->getLoc();
+
+    auto resultType = getResultMemrefType(
+        op, /* offset */ ShapedType::kDynamic,
+        /* staticStrides */ SmallVector<int64_t>{ShapedType::kDynamic},
+        /* resultShape */ SmallVector<int64_t>{ShapedType::kDynamic});
+
+    auto targetOffset = ofrToIndexValue(
+        accumulateTargetOffset(op.getLoc(), op.getMixedOffsets(), rewriter),
+        loc, rewriter);
+
+    // N = modulo bound stored in shape[0] by PtrAnalysis::visitOperandRem
+    Value modN = ofrToIndexValue(op.getMixedShape()[0], loc, rewriter);
+
+    Value blockSz = arith::ConstantOp::create(
+        rewriter, loc, rewriter.getIndexAttr(op.getSizes()[0]));
+
+    Value stride1 = arith::ConstantOp::create(
+        rewriter, loc, rewriter.getIndexAttr(1));
+
+    // x = targetOffset % N
+    Value x = arith::RemSIOp::create(rewriter, loc, targetOffset, modN);
+
+    // wrappedBase = targetOffset - x  (start of the circular-buffer row)
+    Value wrappedBase = arith::SubIOp::create(rewriter, loc, targetOffset, x);
+
+    // d1 = min(x + blockSz, N) - x
+    Value nextOff    = arith::AddIOp::create(rewriter, loc, x, blockSz);
+    Value clampedOff = arith::MinSIOp::create(rewriter, loc, nextOff, modN);
+    Value d1         = arith::SubIOp::create(rewriter, loc, clampedOff, x);
+
+    // chunk1: [targetOffset, targetOffset + d1)
+    SmallVector<Value> sizes1{d1};
+    SmallVector<Value> strides1{stride1};
+    auto cast1 = memref::ReinterpretCastOp::create(
+        rewriter, loc, resultType, adaptor.getBase(),
+        targetOffset, sizes1, strides1);
+
+    // d2 = blockSz - d1  (wrap-around part, may be zero)
+    Value d2 = arith::SubIOp::create(rewriter, loc, blockSz, d1);
+
+    // chunk2: [wrappedBase, wrappedBase + d2)
+    SmallVector<Value> sizes2{d2};
+    auto cast2 = memref::ReinterpretCastOp::create(
+        rewriter, loc, resultType, adaptor.getBase(),
+        wrappedBase, sizes2, strides1);
+
+    return {cast1, cast2};
+  }
+
   LogicalResult rewriteSplitPtr(tts::MakeTensorPtrOp op, OpAdaptor adaptor,
                                 ConversionPatternRewriter &rewriter) const {
     auto parentShape = op.getStaticShape();
-    assert(parentShape.size() == 2 &&
-           "Only support split pointer for 2D tensors only");
     SmallVector<Value> casts;
     StringRef wrapType;
 
-    // For split pointers, a split dimension is either a dynamic or a non-zero
-    // value. The other dimension must be zero.
-    auto isSplitDimension = [](int64_t dim) {
-      return dim == ShapedType::kDynamic || dim != 0;
-    };
-
-    if (isSplitDimension(parentShape[0])) {
-      // Stacked case
-      assert(parentShape[1] == 0);
-      auto [cast1, cast2] = createStackedCastOps(op, adaptor, rewriter);
-      casts = {cast1.getResult(), cast2.getResult()};
-      wrapType = WRAP_STACKED;
-    } else if (isSplitDimension(parentShape[1])) {
-      assert(parentShape[0] == 0);
-      auto [cast1, cast2] = createSideBySideCastOps(op, adaptor, rewriter);
+    if (parentShape.size() == 1) {
+      // 1D circular-buffer wrap-around: ptr[i % N]
+      // shape[0] carries N (set by PtrAnalysis::visitOperandRem for rank-1).
+      auto [cast1, cast2] = create1DCastOps(op, adaptor, rewriter);
       casts = {cast1.getResult(), cast2.getResult()};
-      wrapType = WRAP_SIDE_BY_SIDE;
+      wrapType = WRAP_1D;
     } else {
-      llvm_unreachable("Unexpected split pointer shape");
+      assert(parentShape.size() == 2 &&
+             "Only support split pointer for 1D and 2D tensors");
+
+      // For split pointers, a split dimension is either a dynamic or a non-zero
+      // value. The other dimension must be zero.
+      auto isSplitDimension = [](int64_t dim) {
+        return dim == ShapedType::kDynamic || dim != 0;
+      };
+
+      if (isSplitDimension(parentShape[0])) {
+        // Stacked case
+        assert(parentShape[1] == 0);
+        auto [cast1, cast2] = createStackedCastOps(op, adaptor, rewriter);
+        casts = {cast1.getResult(), cast2.getResult()};
+        wrapType = WRAP_STACKED;
+      } else if (isSplitDimension(parentShape[1])) {
+        assert(parentShape[0] == 0);
+        auto [cast1, cast2] = createSideBySideCastOps(op, adaptor, rewriter);
+        casts = {cast1.getResult(), cast2.getResult()};
+        wrapType = WRAP_SIDE_BY_SIDE;
+      } else {
+        llvm_unreachable("Unexpected split pointer shape");
+      }
     }
 
     auto combinedCast = UnrealizedConversionCastOp::create(
@@ -643,6 +719,34 @@ struct LoadConverter : public OpConversionPattern<tts::LoadOp> {
     memref::CopyOp::create(rewriter, loc, block2, block2Dst);
   }
 
+  // 1D wrap copy: block1 = [startOffset, startOffset+d1), fills dst[0..d1).
+  //               block2 = wrapped part [wrappedBase, wrappedBase+d2), fills dst[d1..XBLOCK).
+  // When d2 == 0 (no wrap) block2 has size 0 and the second CopyOp is a no-op.
+  void create1DCopies(Value block1, Value block2, Value dst, Location loc,
+                      ConversionPatternRewriter &rewriter) const {
+    auto zero =
+        arith::ConstantOp::create(rewriter, loc, rewriter.getIndexAttr(0));
+    auto one =
+        arith::ConstantOp::create(rewriter, loc, rewriter.getIndexAttr(1));
+
+    Value d1 = memref::DimOp::create(rewriter, loc, block1, 0);
+    Value d2 = memref::DimOp::create(rewriter, loc, block2, 0);
+
+    // dst[0 : d1]  <- block1
+    auto block1Dst = memref::SubViewOp::create(rewriter, loc, dst,
+                                               ValueRange{zero},
+                                               ValueRange{d1},
+                                               ValueRange{one});
+    // dst[d1 : d1 + d2]  <- block2
+    auto block2Dst = memref::SubViewOp::create(rewriter, loc, dst,
+                                               ValueRange{d1},
+                                               ValueRange{d2},
+                                               ValueRange{one});
+
+    memref::CopyOp::create(rewriter, loc, block1, block1Dst);
+    memref::CopyOp::create(rewriter, loc, block2, block2Dst);
+  }
+
   memref::SubViewOp createSubview(Value src, ArrayRef<OpFoldResult> offsets,
                                   ArrayRef<OpFoldResult> sizes,
                                   ArrayRef<OpFoldResult> strides, Location loc,
@@ -695,6 +799,25 @@ struct LoadConverter : public OpConversionPattern<tts::LoadOp> {
     return {sv1, sv2};
   }
 
+  // 1D masked subviews: each chunk is already sized to d1/d2 by the cast ops,
+  // so we read the actual dim and clip to the mask dimension.
+  std::pair<memref::SubViewOp, memref::SubViewOp>
+  get1DSubviews(ArrayRef<OpFoldResult> dims, Value block1, Value block2,
+                Location loc, ConversionPatternRewriter &rewriter) const {
+    // dims[0] is the mask size for the whole 1D block.
+    // chunk1 contributes min(d1, maskSize) elements; chunk2 the remainder.
+    // We use the actual sizes already baked into the ReinterpretCastOps (d1/d2)
+    // directly, clipped by the mask dimension via SubViewOp.
+    OpFoldResult d1 = memref::DimOp::create(rewriter, loc, block1, 0).getResult();
+    OpFoldResult d2 = memref::DimOp::create(rewriter, loc, block2, 0).getResult();
+
+    SmallVector<OpFoldResult> offsets{rewriter.getIndexAttr(0)};
+    SmallVector<OpFoldResult> strides{rewriter.getIndexAttr(1)};
+    auto sv1 = createSubview(block1, offsets, {d1}, strides, loc, rewriter);
+    auto sv2 = createSubview(block2, offsets, {d2}, strides, loc, rewriter);
+    return {sv1, sv2};
+  }
+
   LogicalResult
   rewriteStructuredLoad(tts::LoadOp op, OpAdaptor adaptor,
                         ConversionPatternRewriter &rewriter) const {
@@ -715,7 +838,8 @@ struct LoadConverter : public OpConversionPattern<tts::LoadOp> {
 
     auto ptrDefiningOp = ptr.getDefiningOp();
     if (ptrDefiningOp->hasAttr(WRAP_SIDE_BY_SIDE) ||
-        ptrDefiningOp->hasAttr(WRAP_STACKED)) {
+        ptrDefiningOp->hasAttr(WRAP_STACKED) ||
+        ptrDefiningOp->hasAttr(WRAP_1D)) {
 
       auto unrealizedCast = cast<UnrealizedConversionCastOp>(ptrDefiningOp);
       auto memrefs = unrealizedCast.getOperands();
@@ -727,6 +851,8 @@ struct LoadConverter : public OpConversionPattern<tts::LoadOp> {
         createSideBySideCopies(block1, block2, alloc, loc, rewriter);
       } else if (unrealizedCast->hasAttr(WRAP_STACKED)) {
         createStackedCopies(block1, block2, alloc, loc, rewriter);
+      } else if (unrealizedCast->hasAttr(WRAP_1D)) {
+        create1DCopies(block1, block2, alloc, loc, rewriter);
       } else {
         llvm_unreachable("unexpected wraparound type");
       }
@@ -765,7 +891,8 @@ struct LoadConverter : public OpConversionPattern<tts::LoadOp> {
 
     auto ptrDefiningOp = ptr.getDefiningOp();
     if (ptrDefiningOp->hasAttr(WRAP_SIDE_BY_SIDE) ||
-        ptrDefiningOp->hasAttr(WRAP_STACKED)) {
+        ptrDefiningOp->hasAttr(WRAP_STACKED) ||
+        ptrDefiningOp->hasAttr(WRAP_1D)) {
 
       auto unrealizedCast = cast<UnrealizedConversionCastOp>(ptrDefiningOp);
 
@@ -782,6 +909,10 @@ struct LoadConverter : public OpConversionPattern<tts::LoadOp> {
         auto [subview1, subview2] =
             getStackedSubviews(mixedDims, block1, block2, loc, rewriter);
         createStackedCopies(subview1, subview2, alloc, loc, rewriter);
+      } else if (unrealizedCast->hasAttr(WRAP_1D)) {
+        auto [subview1, subview2] =
+            get1DSubviews(mixedDims, block1, block2, loc, rewriter);
+        create1DCopies(subview1, subview2, alloc, loc, rewriter);
       } else {
         llvm_unreachable("unexpected wraparound type");
       }
@@ -1123,6 +1254,75 @@ struct StoreConverter : public OpConversionPattern<tts::StoreOp> {
     auto storeValue = op.getValue();
     auto rank = cast<RankedTensorType>(storeValue.getType()).getRank();
 
+    // Handle 1D circular-buffer wrap-around store: store into two contiguous
+    // segments that together span the block.
+    auto ptrDefiningOp = ptr.getDefiningOp();
+    if (ptrDefiningOp && ptrDefiningOp->hasAttr(WRAP_1D)) {
+      auto unrealizedCast = cast<UnrealizedConversionCastOp>(ptrDefiningOp);
+      auto memrefs = unrealizedCast.getOperands();
+      assert(memrefs.size() == 2);
+      Value block1 = memrefs[0];
+      Value block2 = memrefs[1];
+
+      Value d1 = memref::DimOp::create(rewriter, loc, block1, 0);
+      Value d2 = memref::DimOp::create(rewriter, loc, block2, 0);
+      Value zero =
+          arith::ConstantOp::create(rewriter, loc, rewriter.getIndexAttr(0));
+      Value one =
+          arith::ConstantOp::create(rewriter, loc, rewriter.getIndexAttr(1));
+
+      if (op.hasMask()) {
+        auto mixedDims = op.getMixedMaskDims();
+        // Clip store slices to the mask size (dims[0] for the 1D case).
+        // Use d1/d2 directly from the cast ops since they are already sized.
+        auto slice1 = tensor::ExtractSliceOp::create(
+            rewriter, loc, storeValue,
+            SmallVector<OpFoldResult>{rewriter.getIndexAttr(0)},
+            SmallVector<OpFoldResult>{OpFoldResult(d1)},
+            SmallVector<OpFoldResult>{rewriter.getIndexAttr(1)});
+        auto dst1Subview = memref::SubViewOp::create(
+            rewriter, loc, block1, ValueRange{zero}, ValueRange{d1},
+            ValueRange{one});
+        auto store1 = bufferization::MaterializeInDestinationOp::create(
+            rewriter, loc, slice1, dst1Subview);
+        store1.setWritable(true);
+
+        auto slice2 = tensor::ExtractSliceOp::create(
+            rewriter, loc, storeValue,
+            SmallVector<OpFoldResult>{OpFoldResult(d1)},
+            SmallVector<OpFoldResult>{OpFoldResult(d2)},
+            SmallVector<OpFoldResult>{rewriter.getIndexAttr(1)});
+        auto dst2Subview = memref::SubViewOp::create(
+            rewriter, loc, block2, ValueRange{zero}, ValueRange{d2},
+            ValueRange{one});
+        auto store2 = bufferization::MaterializeInDestinationOp::create(
+            rewriter, loc, slice2, dst2Subview);
+        store2.setWritable(true);
+      } else {
+        // Unmasked: store first d1 elements to chunk1, next d2 to chunk2.
+        auto slice1 = tensor::ExtractSliceOp::create(
+            rewriter, loc, storeValue,
+            SmallVector<OpFoldResult>{rewriter.getIndexAttr(0)},
+            SmallVector<OpFoldResult>{OpFoldResult(d1)},
+            SmallVector<OpFoldResult>{rewriter.getIndexAttr(1)});
+        auto store1 = bufferization::MaterializeInDestinationOp::create(
+            rewriter, loc, slice1, block1);
+        store1.setWritable(true);
+
+        auto slice2 = tensor::ExtractSliceOp::create(
+            rewriter, loc, storeValue,
+            SmallVector<OpFoldResult>{OpFoldResult(d1)},
+            SmallVector<OpFoldResult>{OpFoldResult(d2)},
+            SmallVector<OpFoldResult>{rewriter.getIndexAttr(1)});
+        auto store2 = bufferization::MaterializeInDestinationOp::create(
+            rewriter, loc, slice2, block2);
+        store2.setWritable(true);
+      }
+
+      rewriter.eraseOp(op);
+      return success();
+    }
+
     if (op.hasMask()) {
       auto mixedDims = op.getMixedMaskDims();