class_docstr: Responsible for emitting a CUTLASS template definition.
docstr: Instantiate a GEMM kernel from given `operation`.
docstr_tok: ["Instantiate","a","GEMM","kernel","from","given","`","operation","`","."]
code: def emit(self, operation, no_beta_scaling=False, batched=False):
      """Instantiate a GEMM kernel from given `operation`."""
      warp_shape = [
          operation.tile_description.threadblock_shape[idx]
          // operation.tile_description.warp_count[idx]
          for idx in range(3)
      ]
      epilogue_vector_length = (
          min(
              operation.C.alignment * DataTypeSize[operation.C.element],
              128,
          )
          // DataTypeSize[operation.C.element]
      )
      values = {
          "operation_name": operation.procedural_name(),
          "element_a": DataTypeTag[operation.A.element],
          "layout_a": LayoutTag[operation.A.layout],
          "element_b": DataTypeTag[operation.B.element],
          "layout_b": LayoutTag[operation.B.layout],
          "element_c": DataTypeTag[operation.C.element],
          "layout_c": LayoutTag[operation.C.layout],
          "element_accumulator": DataTypeTag[
              operation.accumulator_type()
          ],
          "opcode_class": OpcodeClassTag[
              operation.tile_description.math_instruction.opcode_class
          ],
          "arch": f"cutlass::arch::Sm{operation.arch}",
          "threadblock_shape_m": str(
              operation.tile_description.threadblock_shape[0]
          ),
          "threadblock_shape_n": str(
              operation.tile_description.threadblock_shape[1]
          ),
          "threadblock_shape_k": str(
              operation.tile_description.threadblock_shape[2]
          ),
          "warp_shape_m": str(warp_shape[0]),
          "warp_shape_n": str(warp_shape[1]),
          "warp_shape_k": str(warp_shape[2]),
          "instruction_shape_m": str(
              operation.tile_description.math_instruction.instruction_shape[
                  0
              ]
          ),
          "instruction_shape_n": str(
              operation.tile_description.math_instruction.instruction_shape[
                  1
              ]
          ),
          "instruction_shape_k": str(
              operation.tile_description.math_instruction.instruction_shape[
                  2
              ]
          ),
          "epilogue_vector_length": str(epilogue_vector_length),
          "element_epilogue": str(
              DataTypeTag[operation.element_epilogue]
          ),
          "epilogue_functor": EpilogueFunctorTag[
              operation.epilogue_functor
          ],
          "swizzling_functor": SwizzlingFunctorTag[
              operation.swizzling_functor
          ],
          "stages": str(operation.tile_description.stages),
          "align_a": str(operation.A.alignment),
          "align_b": str(operation.B.alignment),
          "math_operation": MathOperationTag[
              operation.tile_description.math_instruction.math_operation
          ],
      }
  
      values["kernel_name"] = "GemmBatched" if batched else "Gemm"
      values["split_k_serial"] = "" if batched else "false,"
  
      gemm_template = substitute_template(
          self.gemm_template,
          {
              "epilogue": self.epilogue_no_beta_scaling
              if no_beta_scaling
              else self.epilogue_default
          },
      )
      return substitute_template(gemm_template, values)
  
file_path: tvm/python/tvm/contrib/cutlass/gemm_operation.py
imports: from .library import *
label_desc: Use image node_id 2 for calling the EmitGemmInstance obj's underlying member method code with example usage: obj.emit(operation, no_beta_scaling, batched) and returns: substitute_template