code: def __init__(self, manager: kp.Manager, tile_size: int = -1):
      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
          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
  
      assert tile_size > 0
      assert tile_size * tile_size <= max_workgroup_invocation
      assert tile_size <= max_workgroup_size[0]
      assert tile_size <= max_workgroup_size[1]
      self.tile_size = tile_size
  
      self.shader = """
  #version 450
  
  layout (local_size_x = {tile_size}, local_size_y = {tile_size}) 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; // 0 .. tile_size
      uint col = gl_LocalInvocationID.y; // 0 .. tile_size
      // gl_WorkGroupID : 0 .. tensor_size / tile_size
      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 = 0.0;
      uint numTiles = tensor_size / {tile_size};
      for(uint t = 0u; t < numTiles; t++)
      {{
          uint tiledRow = ({tile_size} * t) + row;
          uint tiledCol = ({tile_size} * t) + col;
          sub_tensor_1[col][row] = in_tensor_1[(tiledCol * tensor_size) + globalRow];
          sub_tensor_2[col][row] = in_tensor_2[(globalCol * tensor_size) + tiledRow];
  
          memoryBarrierShared();
          barrier();
  
          for(uint k = 0u; k < {tile_size}; k++)
              acc += sub_tensor_1[k][row] * sub_tensor_2[col][k];
  
          barrier();
      }}
      out_tensor[tensor_size * globalCol + globalRow] = acc;
  }}"""
      self.compiled_shader = kp.Shader.compile_source(
          self.shader.format(tile_size=tile_size)
      )
      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)