Use stride instead of order to determine block attr

Author: alexbaden

third_party/intel/lib/TritonIntelGPUTransforms/MaterializeBlockPointer.cpp

@@ -51,16 +51,28 @@ struct TritonIntelGPUMaterializeBlockPointerPass
       LDBG("Found make tensor ptr op: " << makeTensorPtrOp);
       auto ptrType = cast<tt::PointerType>(makeTensorPtrOp.getType());
       auto tensorType = cast<RankedTensorType>(ptrType.getPointeeType());
-      ArrayRef<int32_t> order = makeTensorPtrOp.getOrder();
-      unsigned rank = order.size();
+      Operation::operand_range shape = makeTensorPtrOp.getShape();
+      unsigned rank = shape.size();
       LDBG("Rank: " << rank);
       if (rank == 1)
         return;
 
-      unsigned fastChangeDim = order[0];
+      Operation::operand_range strides = makeTensorPtrOp.getStrides();
+      int fastChangeDim = -1;
+      for (size_t i = 0; i < strides.size(); i++) {
+        if (mlir::triton::gpu::intel::isConstant(strides[i], 1)) {
+          fastChangeDim = i;
+          break;
+        }
+      }
       LDBG("Fast change dim: " << fastChangeDim);
+      if (fastChangeDim < 0) {
+        return;
+      }
+      ArrayRef<int32_t> order = makeTensorPtrOp.getOrder();
+
+      // unsigned fastChangeDim = order[0];
       if (fastChangeDim >= (rank - 2)) {
-        Operation::operand_range strides = makeTensorPtrOp.getStrides();
 
         // HW 2D block read instruction only supports contiguous access.
         Value fastChangeStride = strides[fastChangeDim];