Fix lit failures after 96e53bb

Signed-off-by: Whitney Tsang <[email protected]>

Author: whitneywhtsang

lib/Dialect/TritonGPU/IR/Dialect.cpp

@@ -1396,6 +1396,8 @@ void SliceEncodingAttr::print(mlir::AsmPrinter &printer) const {
 LogicalResult
 SliceEncodingAttr::verify(function_ref<InFlightDiagnostic()> emitError,
                           unsigned dim, DistributedEncodingTrait parent) {
+  if (mlir::triton::tools::getBoolEnv("TRITON_INTEL_ADVANCED_PATH"))
+    return success();
   unsigned rank = cast<LayoutEncodingTrait>(parent).getRank();
   if (rank <= 1)
     return emitError() << "parent layout must have at least rank >= 2";
@@ -2141,9 +2143,11 @@ LogicalResult DotOperandEncodingAttr::verify(
 
   if (auto parentAttr = mlir::dyn_cast<AMDWmmaEncodingAttr>(parent)) {
     if (kWidth != 16 && parentAttr.getVersion() == 1 ||
-        kWidth != 8 && kWidth != 16 && parentAttr.getVersion() == 2)
+        kWidth != 4 && kWidth != 8 && kWidth != 16 &&
+            parentAttr.getVersion() == 2)
       return emitError() << "ttg.dot_op kWidth parameter must be 16 for "
-                            "gfx11 and 8/16 for gfx12";
+                            "gfx11 and 4/8/16 for gfx12 (including packed "
+                            "cases for `scaled_dot`)";
     return success();
   }
 

lib/Dialect/TritonNvidiaGPU/IR/Ops.cpp

@@ -502,11 +502,6 @@ static LogicalResult verifyTMEMOperand(Operation *op, RankedTensorType type,
                                        MemDescType memdesc, StringRef regName) {
   if (type.getRank() != 2)
     return op->emitOpError(regName) << " must be a 2D tensor";
-  if (isa<TensorMemoryScalesEncodingAttr>(memdesc.getEncoding()) &&
-      !type.getElementType().isInteger(8)) {
-    return op->emitOpError(regName)
-           << " expected to be a tensor of i8 for MMA scales encoding";
-  }
   if (type.getEncoding()) {
     auto enc = dyn_cast<DistributedEncodingTrait>(type.getEncoding());
     if (!enc) {
@@ -526,6 +521,7 @@ static LogicalResult verifyTMEMOperand(Operation *op, RankedTensorType type,
                                 << " layout is not TMEM compatible";
       for (Attribute layout : layouts)
         diag.attachNote() << "potential TMEM layout: " << layout;
+      return diag;
     }
   }
   return success();