Workaround test failures for the swizzling path

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

Author: whitneywhtsang

include/triton/Conversion/TritonGPUToLLVM/TargetInfoBase.h

@@ -97,6 +97,7 @@ class TargetInfoBase {
   virtual bool supportLdMatrix() const { return false; }
   virtual bool supportStMatrix() const { return false; }
   virtual bool isCuda() const { return false; }
+  virtual bool isXpu() const { return false; }
 
   // Annotate target specific information to local load operations during
   // lowering to LLVM. `llLoadOp` is the generated LLVM load op.

lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp

@@ -279,8 +279,9 @@ struct ConvertLayoutOpUsingLinearLayoutsConversion
     assert(cvtNeedsSharedMemory(op.getSrc().getType(), op.getType()));
 
     // Try to use swizzling to implement the conversion
-    if (succeeded(
-            transferWithinBlockSwizzling(op, adaptor.getSrc(), rewriter))) {
+    // HACK Remove once XPU tests pass for the swizzling path
+    if (!targetInfo.isXpu() && succeeded(transferWithinBlockSwizzling(
+                                   op, adaptor.getSrc(), rewriter))) {
       return success();
     }
 

third_party/intel/lib/TritonIntelGPUToLLVM/TargetInfo.h

@@ -76,6 +76,8 @@ class TargetInfo : public mlir::triton::TargetInfoBase {
                              StringRef name, StringRef value,
                              unsigned addressSpace) const;
 
+  bool isXpu() const override { return true; }
+
 protected:
   virtual bool isSupportedWarpReduceOp(Operation *op, unsigned numLanesToReduce,
                                        unsigned warpSize) const = 0;