Revert "Fix failures from llvm/llvm-project@1f20eee6dc36"

This reverts commit a876742.
intel · Dec 5, 2024 · 5c0e236 · 5c0e236
1 parent e7fb607
commit 5c0e236
Show file tree

Hide file tree

Showing 2 changed files with 7 additions and 7 deletions.
diff --git a/test/TritonIntelGPU/prefetch-block.mlir b/test/TritonIntelGPU/prefetch-block.mlir
@@ -33,7 +33,7 @@ module attributes {"ttg.num-warps" = 32 : i32, "ttg.threads-per-warp" = 1 : i32}
     // CHECK-NEXT: [[B3:%.*]] = tt.advance [[B2]], {{.*}} : <tensor<32x256xf16, #blocked2>>
     // CHECK-NEXT: [[B4:%.*]] = tt.make_tensor_ptr %arg1, {{.*}} : <tensor<32x256xf16, #ttg.dot_op<{opIdx = 1, parent = #blocked}>>>
 
-    // CHECK:      spirv.INTEL.ControlBarrierArrive <Workgroup> <Workgroup> <None>
+    // CHECK:      spirv.INTEL.ControlBarrierArrive <Workgroup>, <Workgroup>, <None>
     // CHECK-NEXT: scf.for [[IV:%.*]] = [[CST_ZERO]] to [[CST_4096]] step [[CST_32]]
     // CHECK-SAME:      iter_args([[CST:%.*]] = {{.*}}, [[A6:%.*]] = [[A4]], [[B6:%.*]] = [[B4]], [[A5:%.*]] = [[A3]], [[B5:%.*]] = [[B3]])
     // CHECK-NEXT:   [[LD_A:%.*]] = tt.load [[A6]]
@@ -45,11 +45,11 @@ module attributes {"ttg.num-warps" = 32 : i32, "ttg.threads-per-warp" = 1 : i32}
     // CHECK-DAG:    tt.advance [[A6]], {{.*}} : <tensor<256x32xf16, #ttg.dot_op<{opIdx = 0, parent = #blocked}>>>
     // CHECK-NEXT:   tt.advance [[B5]], {{.*}} : <tensor<32x256xf16, #blocked2>>
     // CHECK-DAG:    tt.advance [[B6]], {{.*}} : <tensor<32x256xf16, #ttg.dot_op<{opIdx = 1, parent = #blocked}>>>
-    // CHECK:        spirv.INTEL.ControlBarrierWait <Workgroup> <Workgroup> <None>
-    // CHECK-NEXT:   spirv.INTEL.ControlBarrierArrive <Workgroup> <Workgroup> <None>
+    // CHECK:        spirv.INTEL.ControlBarrierWait <Workgroup>, <Workgroup>, <None>
+    // CHECK-NEXT:   spirv.INTEL.ControlBarrierArrive <Workgroup>, <Workgroup>, <None>
     // CHECK-NEXT:   scf.yield {{.*}}
     // CHECK-NEXT: }
-    // CHECK-NEXT: spirv.INTEL.ControlBarrierWait <Workgroup> <Workgroup> <None>
+    // CHECK-NEXT: spirv.INTEL.ControlBarrierWait <Workgroup>, <Workgroup>, <None>
 
     %c64_i32 = arith.constant 64 : i32
     %c16_i32 = arith.constant 16 : i32

diff --git a/third_party/intel/include/Dialect/TritonGEN/IR/TritonGENOps.td b/third_party/intel/include/Dialect/TritonGEN/IR/TritonGENOps.td
@@ -35,7 +35,7 @@ class TritonGEN_Op<string mnemonic, list<Trait> traits = []> :
 def TritonGEN_MatrixElemType : AnyTypeOf<[AnyI8, AnyI16, AnyI32, F32, F16, BF16]>;
 
 def TritonGEN_MatrixDPASOp : TritonGEN_Op<"dpas">,
-  Results<(outs FixedVectorOfAnyRank<[TritonGEN_MatrixElemType]>:$d)>,
+  Results<(outs FixedVectorOf<[TritonGEN_MatrixElemType]>:$d)>,
   Arguments<(ins
     FixedVectorOfRankAndType<[1], [TritonGEN_MatrixElemType]>:$c,
     FixedVectorOfRankAndType<[1], [TritonGEN_MatrixElemType]>:$a,
@@ -82,7 +82,7 @@ def TritonGEN_MatrixDPASOp : TritonGEN_Op<"dpas">,
 }
 
 def TritonGEN_Matrix2DBlockLoadOp : TritonGEN_Op<"2Dblockload">,
-  Results<(outs FixedVectorOfAnyRank<[TritonGEN_MatrixElemType]>:$res)>,
+  Results<(outs FixedVectorOf<[TritonGEN_MatrixElemType]>:$res)>,
   Arguments<(ins
     Arg<LLVM_AnyPointer, "", [MemRead]>:$ptr,
     I32:$base_width,
@@ -145,7 +145,7 @@ def TritonGEN_Matrix2DBlockStoreOp : TritonGEN_Op<"2Dblockstore">,
     I32Attr:$tile_width,
     I32Attr:$tile_height,
     I32Attr:$v_blocks,
-    FixedVectorOfAnyRank<[TritonGEN_MatrixElemType]>:$stored_val,
+    FixedVectorOf<[TritonGEN_MatrixElemType]>:$stored_val,
     DefaultValuedAttr<TritonGEN_StoreCacheControl, "::mlir::triton::TritonGEN::StoreCacheControl::DEFAULT">:$cache_control
   )> {