code: def __init__(
      self,
      manager: kp.Manager,
      tile_size: int = -1,
      thread_work_ratio: int = 16,
  ):
      self.mgr = manager
  
      props = self.mgr.get_device_properties()
      max_workgroup_invocation = props["max_work_group_invocations"]
      max_workgroup_size = props["max_work_group_size"]
      if tile_size < 0:
          tile_size = 1
          local_size_y = tile_size // thread_work_ratio
          while (
              4 * tile_size * tile_size <= max_workgroup_invocation
              and 2 * tile_size <= max_workgroup_size[0]
              and 2 * tile_size <= max_workgroup_size[1]
          ):
              tile_size *= 2
              local_size_y = tile_size // thread_work_ratio
      else:
          local_size_y = tile_size // thread_work_ratio
  
      assert tile_size > 0
      assert thread_work_ratio > 0
      assert tile_size * local_size_y <= max_workgroup_invocation
      assert tile_size <= max_workgroup_size[0]
      assert local_size_y <= max_workgroup_size[1]
      self.tile_size = tile_size
      self.thread_work_ratio = thread_work_ratio
  
      self.local_size_x = tile_size
      self.local_size_y = tile_size // thread_work_ratio
      self.shader = """
  #version 450
  
  layout (local_size_x = {tile_size}, local_size_y = {local_size_y}) in;
  
  layout (set = 0, binding = 0) readonly buffer buf_in_tensor_1 {{ float in_tensor_1[]; }};
  layout (set = 0, binding = 1) readonly buffer buf_in_tensor_2 {{ float in_tensor_2[]; }};
  layout (set = 0, binding = 2) writeonly buffer buf_out_tensor {{ float out_tensor[]; }};
  
  layout (constant_id = 0) const float tensor_size_f = 0;
  
  shared float sub_tensor_1[{tile_size}][{tile_size}];
  shared float sub_tensor_2[{tile_size}][{tile_size}];
  
  void main()
  {{
      uint row = gl_LocalInvocationID.x;
      uint col = gl_LocalInvocationID.y;
      uint globalRow = {tile_size} * gl_WorkGroupID.x + row;
      uint globalCol = {tile_size} * gl_WorkGroupID.y + col;
  
      uint tensor_size = uint(tensor_size_f);
      float acc[{thread_work_ratio}];
      for(uint w = 0u; w < {thread_work_ratio}; w++)
          acc[w] = 0.0;
  
      uint numTiles = tensor_size / {tile_size};
      for(uint t = 0u; t < numTiles; t++)
      {{
          for(uint w = 0u; w < {thread_work_ratio}; w++)
          {{
              uint tiledRow = {tile_size} * t + row;
              uint tiledCol = {tile_size} * t + col;
              sub_tensor_1[col + w * {local_size_y}][row] = in_tensor_1[
                  (tiledCol + w * {local_size_y}) * tensor_size + globalRow];
              sub_tensor_2[col + w * {local_size_y}][row] = in_tensor_2[
                  (globalCol + w * {local_size_y})* tensor_size + tiledRow];
          }}
  
          memoryBarrierShared();
          barrier();
  
          for(uint k = 0u; k < {tile_size}; k++)
              for(uint w = 0u; w < {thread_work_ratio}; w++)
                  acc[w] += sub_tensor_1[k][row] * sub_tensor_2[col + w * {local_size_y}][k];
  
          barrier();
      }}
      for(uint w = 0u; w < {thread_work_ratio}; w++)
          out_tensor[(globalCol + w * {local_size_y}) * tensor_size + globalRow] = acc[w];
  }}"""
      self.compiled_shader = kp.Shader.compile_source(
          self.shader.format(
              tile_size=tile_size,
              thread_work_ratio=thread_work_ratio,
              local_size_y=local_size_y,
          )
      )
      self.tensor_shape: tuple[int, int] = (0, 0)
      self.params: list[kp.Tensor] = []
      self.algo = None
  
imports: import time
  import kp
  import numpy
label_desc: Use image node_id 1 to create a new MatMulOp object  with example: obj = MatMulOp(manager, tile_size, thread_work_ratio)