Minor changes

Author: mieshkiwrk

benchmarks/triton_kernels_benchmark/gemm_benchmark.py

@@ -201,7 +201,7 @@ def get_dpas_layout(num_warps: ttgl.constexpr, m_shape: ttgl.constexpr, n_shape:
     key=['M', 'N', 'K'],
 )
 @gluon.jit
-def gluon_matmul_kernel_dpas_tensor_desc(
+def gluon_matmul_kernel_with_tensor_descriptors(
         # Pointers to matrices
         a_ptr, b_ptr, c_ptr,
         # Matrix dimensions
@@ -268,7 +268,7 @@ def gluon_matmul_kernel_dpas_tensor_desc(
     key=['B', 'M', 'N', 'K'],
 )
 @gluon.jit
-def gluon_matmul_kernel_dpas_tensor_desc_batched(
+def gluon_matmul_kernel_with_tensor_descriptors_batched(
         # Pointers to matrices
         a_ptr, b_ptr, c_ptr,
         # Matrix dimensions
@@ -461,8 +461,10 @@ def get_benchmark(
     providers_filter: Optional[list[str]] = None,
     transpose_a=False,
     transpose_b=False,
-    matmul_kernel=matmul_kernel_with_block_pointers,
-    matmul_kernel_batched=matmul_kernel_with_block_pointers_batched,
+    triton_matmul_kernel=matmul_kernel_with_block_pointers,
+    triton_matmul_kernel_batched=matmul_kernel_with_block_pointers_batched,
+    gluon_matmul_kernel=gluon_matmul_kernel_with_tensor_descriptors,
+    gluon_matmul_kernel_batched=gluon_matmul_kernel_with_tensor_descriptors_batched,
     plot_name='matmul-performance',
 ):
     """
@@ -472,13 +474,13 @@ def get_benchmark(
     supported_providers = {
         'gluon': 'Gluon',
         'triton': 'Triton',
-        #'onednn': 'OneDNN',
+        'onednn': 'OneDNN',
     }
     # use_cutlass
-    # if not (transpose_a or transpose_b):
-    #     if torch.xpu.get_device_name() != 'Intel(R) Arc(TM) Graphics':
-    #         # FIXME: enable cutlass on LNL
-    #         supported_providers['cutlass'] = 'CUTLASS'
+    if not (transpose_a or transpose_b):
+        if torch.xpu.get_device_name() != 'Intel(R) Arc(TM) Graphics':
+            # FIXME: enable cutlass on LNL
+            supported_providers['cutlass'] = 'CUTLASS'
     providers = benchmark_suite.filter_providers(supported_providers, providers_filter)
 
     # Benchmark Performance
@@ -532,8 +534,8 @@ def benchmark(B, M, N, K, provider):
             else:
                 raise AssertionError(f'Unexpected shape of length {len(a.shape)}')
 
-            kernel = matmul_kernel if provider == 'triton' else gluon_matmul_kernel_dpas_tensor_desc
-            batched_kernel = matmul_kernel_batched if provider == 'triton' else gluon_matmul_kernel_dpas_tensor_desc_batched
+            kernel = triton_matmul_kernel if provider == 'triton' else gluon_matmul_kernel
+            batched_kernel = triton_matmul_kernel_batched if provider == 'triton' else gluon_matmul_kernel_batched
 
             matmul_fn = lambda: matmul(
                 a,

third_party/intel/include/Dialect/TritonIntelGPU/Transforms/Utility.h

@@ -71,7 +71,6 @@ calculateRepCluster(unsigned capRepeatCount, unsigned capSystolicDepth,
                     ArrayRef<int64_t> retShape, unsigned threadsPerWarp,
                     unsigned int a_bitwidth, bool is_a_FP8,
                     ArrayRef<int64_t> a_shape, ArrayRef<int64_t> b_shape,
-                    // RankedTensorType oldRetType,
                     SmallVector<unsigned> warpsPerTile);
 
 } // namespace mlir::triton::gpu::intel

third_party/intel/lib/Dialect/Triton/Transforms/TensorDescToBlockPointer.cpp

@@ -147,7 +147,7 @@ struct TritonIntelTensorDescToBlockPointer
           tensorType, pointerType.getAddressSpace());
 
       auto makeTensorPtr = builder.create<tt::MakeTensorPtrOp>(
-          loc, resultType, base, shape, strides, offsets,
+          builder, loc, resultType, base, shape, strides, offsets,
           builder.getDenseI32ArrayAttr({1, 0}));
       return makeTensorPtr;
     };

third_party/intel/lib/Dialect/TritonIntelGPU/IR/Dialect.cpp

@@ -211,7 +211,7 @@ SmallVector<int64_t> DpasEncodingAttr::calculateDPASRepetitions(
     ArrayRef<unsigned> repCluster, unsigned repeatCount, unsigned systolicDepth,
     unsigned executionSize, unsigned opsPerChannel) {
   // Always return a 3D shape repetitions for the ease of value handling, same
-  // to mma
+  // to mma.
   size_t rank = shape.size();
   SmallVector<int64_t> rep(3, 1);
 

third_party/intel/lib/TritonIntelGPUTransforms/AccelerateMatmul.cpp

@@ -212,17 +212,11 @@ class BlockedToDPAS : public OpRewritePattern<tt::DotOp> {
     size_t rank = retShape.size();
 
     SmallVector<unsigned> repCluster = ttgi::calculateRepCluster(
-        // dpasCap,
         dpasCap.repeatCount, dpasCap.systolicDepth, dpasCap.executionSize,
-        opsPerChan,
-        // rank,
-        retShape,
-        // mod,
-        threadsPerWarp, oldAType.getElementType().getIntOrFloatBitWidth(),
+        opsPerChan, retShape, threadsPerWarp,
+        oldAType.getElementType().getIntOrFloatBitWidth(),
         isa<Float8E5M2Type, Float8E4M3FNType>(oldAType.getElementType()),
-        oldAType.getShape(), oldBType.getShape(),
-        // oldRetType,
-        warpsPerTile);
+        oldAType.getShape(), oldBType.getShape(), warpsPerTile);
 
     unsigned repeatCount =
         std::min(dpasCap.repeatCount, (unsigned)retShape[rank - 2] /*M*/);
@@ -237,39 +231,6 @@ class BlockedToDPAS : public OpRewritePattern<tt::DotOp> {
         dpasCap.executionSize, opsPerChan, warpsPerTile, repCluster,
         threadsPerWarp);
 
-    // if (dpasCap.isPVC() || dpasCap.isFalconShore()) {
-    //   unsigned dpasElemBitWidths =
-    //       oldAType.getElementType().getIntOrFloatBitWidth();
-    //
-    //   // We are upcasting FP8 to FP16
-    //   if (isa<Float8E5M2Type, Float8E4M3FNType>(oldAType.getElementType()))
-    //     dpasElemBitWidths = 2 * dpasElemBitWidths;
-    //
-    //   // Enlarge the repCluster size to use the large 2D load for A and B
-    //   // operands.
-    //   unsigned maxRepClusterM =
-    //       PVC_2D_LOAD_MAXIMUM_NUMBER_OF_ROWS / dpasCap.repeatCount;
-    //   SmallVector<int64_t> repA =
-    //       dpasEnc.getDPASRepetitions(oldAType.getShape(), 0);
-    //   unsigned repClusterDimM =
-    //       std::min(maxRepClusterM, static_cast<unsigned>(repA[1]));
-    //
-    //   unsigned maxRepClusterN =
-    //       PVC_2D_LOAD_MAXIMUM_BYTES_OF_COLS /
-    //       ((dpasElemBitWidths / 8) * dpasCap.executionSize);
-    //   SmallVector<int64_t> repB =
-    //       dpasEnc.getDPASRepetitions(oldBType.getShape(), 1);
-    //   unsigned repClusterDimN =
-    //       std::min(maxRepClusterN, static_cast<unsigned>(repB[2]));
-    //   repCluster[rank - 2] = repClusterDimM;
-    //   repCluster[rank - 1] = repClusterDimN;
-    //
-    //   dpasEnc = ttgi::DpasEncodingAttr::get(
-    //       oldRetType.getContext(), repeatCount, dpasCap.systolicDepth,
-    //       dpasCap.executionSize, opsPerChan, warpsPerTile, repCluster,
-    //       threadsPerWarp);
-    // }
-
     RankedTensorType newRetType =
         RankedTensorType::get(retShape, oldRetType.getElementType(), dpasEnc);