code: def __init__(
      self,
      manager: kp.Manager,
      local_size_x: int = -1,
      local_size_y: 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 local_size_x < 1:
          if local_size_y > 0:
              local_size_x = 1
              while (
                  2 * local_size_x * local_size_y
                  <= max_workgroup_invocation
                  and 2 * local_size_x <= max_workgroup_size[0]
              ):
                  local_size_x *= 2
          else:
              local_size_x = 1
              local_size_y = 1
              while (
                  2 * local_size_x * local_size_y
                  <= max_workgroup_invocation
              ):
                  if 2 * local_size_x <= max_workgroup_size[0]:
                      local_size_x *= 2
                  if 2 * local_size_y <= max_workgroup_size[1]:
                      local_size_y *= 2
                  elif (
                      2 * local_size_x > max_workgroup_size[0]
                  ):  # stop if neither x nor y can be double
                      break
      elif local_size_y < 0:
          local_size_y = 1
          while (
              2 * local_size_x * local_size_y
              <= max_workgroup_invocation
              and 2 * local_size_x <= max_workgroup_size[0]
          ):
              local_size_y *= 2
  
      assert local_size_x > 0
      assert local_size_y > 0
      assert local_size_x * local_size_y <= max_workgroup_invocation
      assert local_size_x <= max_workgroup_size[0]
      assert local_size_y <= max_workgroup_size[1]
      self.local_size_x = local_size_x
      self.local_size_y = local_size_y
  
      self.shader = """
  #version 450
  
  layout (local_size_x = {local_size_x}, 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;
  
  
  void main()
  {{
      uint globalRow = gl_GlobalInvocationID.x;
      uint globalCol = gl_GlobalInvocationID.y;
      uint tensor_size = uint(tensor_size_f);
      float acc = 0.0;
      for(uint k = 0u; k < tensor_size; k++)
          acc += in_tensor_1[(k * tensor_size) + globalRow] * in_tensor_2[(globalCol * tensor_size) + k];
      out_tensor[(globalCol * tensor_size) + globalRow] = acc;
  }}"""
      self.compiled_shader = kp.Shader.compile_source(
          self.shader.format(
              local_size_x=self.local_size_x,
              local_size_y=self.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, local_size_x, local_size_y)