Fallback performance experiment

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

Author: whitneywhtsang

benchmarks/triton_kernels_benchmark/gemm_tensor_desc_benchmark.py

@@ -25,11 +25,11 @@ def matmul_kernel_with_tensor_descriptors(
         # Pointers to matrices
         a_ptr, b_ptr, c_ptr,
         # Matrix dimensions
-        M: tl.constexpr, N: tl.constexpr, K: tl.constexpr,
+        M, N, K,
         # Stride variables
-        stride_am: tl.constexpr, stride_ak: tl.constexpr,  #
-        stride_bk: tl.constexpr, stride_bn: tl.constexpr,  #
-        stride_cm: tl.constexpr, stride_cn: tl.constexpr,
+        stride_am, stride_ak,  #
+        stride_bk, stride_bn,  #
+        stride_cm, stride_cn,
         # Meta-parameters
         BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr, GROUP_SIZE_M: tl.constexpr):
     pid = tl.program_id(axis=0)
@@ -71,11 +71,11 @@ def matmul_kernel_with_tensor_descriptors_batched(
         # Pointers to matrices
         a_ptr, b_ptr, c_ptr,
         # Matrix dimensions
-        B: tl.constexpr, M: tl.constexpr, N: tl.constexpr, K: tl.constexpr,
+        B, M, N, K,
         # Stride variables
-        stride_az: tl.constexpr, stride_am: tl.constexpr, stride_ak: tl.constexpr,  #
-        stride_bz: tl.constexpr, stride_bk: tl.constexpr, stride_bn: tl.constexpr,  #
-        stride_cz: tl.constexpr, stride_cm: tl.constexpr, stride_cn: tl.constexpr,
+        stride_az, stride_am, stride_ak,  #
+        stride_bz, stride_bk, stride_bn,  #
+        stride_cz, stride_cm, stride_cn,
         # Meta-parameters
         BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr, GROUP_SIZE_M: tl.constexpr):
     bid = tl.program_id(axis=1)

benchmarks/triton_kernels_benchmark/gemm_tensor_of_ptr_benchmark.py

@@ -25,11 +25,11 @@ def matmul_kernel(
         # Pointers to matrices
         a_ptr, b_ptr, c_ptr,
         # Matrix dimensions
-        M: tl.constexpr, N: tl.constexpr, K: tl.constexpr,
+        M, N, K,
         # Stride variables
-        stride_am: tl.constexpr, stride_ak: tl.constexpr,  #
-        stride_bk: tl.constexpr, stride_bn: tl.constexpr,  #
-        stride_cm: tl.constexpr, stride_cn: tl.constexpr,
+        stride_am, stride_ak,  #
+        stride_bk, stride_bn,  #
+        stride_cm, stride_cn,
         # Meta-parameters
         BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr, GROUP_SIZE_M: tl.constexpr):
     pid = tl.program_id(axis=0)
@@ -76,11 +76,11 @@ def matmul_kernel_batched(
         # Pointers to matrices
         a_ptr, b_ptr, c_ptr,
         # Matrix dimensions
-        B: tl.constexpr, M: tl.constexpr, N: tl.constexpr, K: tl.constexpr,
+        B, M, N, K,
         # Stride variables
-        stride_az: tl.constexpr, stride_am: tl.constexpr, stride_ak: tl.constexpr,  #
-        stride_bz: tl.constexpr, stride_bk: tl.constexpr, stride_bn: tl.constexpr,  #
-        stride_cz: tl.constexpr, stride_cm: tl.constexpr, stride_cn: tl.constexpr,
+        stride_az, stride_am, stride_ak,  #
+        stride_bz, stride_bk, stride_bn,  #
+        stride_cz, stride_cm, stride_cn,
         # Meta-parameters
         BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr, GROUP_SIZE_M: tl.constexpr):
     bid = tl.program_id(axis=1)

third_party/intel/backend/compiler.py

@@ -207,7 +207,7 @@ def make_ttir(cls, mod, metadata, opt):
         pm = ir.pass_manager(mod.context)
         pm.enable_debug()
         passes.common.add_inliner(pm)
-        intel.passes.ttir.add_convert_tdesc_to_block_pointer(pm)
+        #intel.passes.ttir.add_convert_tdesc_to_block_pointer(pm)
         passes.ttir.add_rewrite_tensor_descriptor_to_pointer(pm)
         passes.common.add_cse(pm)
         passes.common.add_licm(pm)