class_docstr: Responsible for emitting a CUTLASS template definition.
class_docstr_tok: ["Responsible","for","emitting","a","CUTLASS","template","definition","."]
code: def __init__(self):
      self.epilogue_default = """
      ${epilogue_functor}<
        ${element_c},
        ${epilogue_vector_length},
        ${element_accumulator},
        ${element_epilogue}
      >"""
  
      self.epilogue_no_beta_scaling = """
      ${epilogue_functor}<
        ${element_c},
        ${epilogue_vector_length},
        ${element_accumulator},
        ${element_epilogue},
        cutlass::epilogue::thread::ScaleType::NoBetaScaling
      >"""
  
      self.epilogue_residual_block = """
      ${epilogue_functor}<
        ${element_c},
        ${element_accumulator},
        ${element_epilogue},
        ${element_c},
        ${epilogue_vector_length},
        ${activation},
        ${binary_op},
        ${unary_op}
      >"""
  
      self.epilogue_wgrad = """
      ${epilogue_functor}<
        ${element_c},
        4,
        float,
        float
      >"""
  
      self.template = """
    // Conv2d${conv_kind_name} ${iterator_algorithm_name} kernel instance "${operation_name}"
    using ${operation_name} =
    typename cutlass::conv::kernel::DefaultConv2d${conv_kind_name}${conv_kernel_postfix}<
      ${element_a},
      ${layout_a},
      ${element_b},
      ${layout_b},
      ${element_c},
      ${layout_c},
      ${element_accumulator},
      ${opcode_class},
      ${arch},
      cutlass::gemm::GemmShape<${threadblock_shape_m}, ${threadblock_shape_n}, ${threadblock_shape_k}>,
      cutlass::gemm::GemmShape<${warp_shape_m}, ${warp_shape_n}, ${warp_shape_k} >,
      cutlass::gemm::GemmShape<${instruction_shape_m}, ${instruction_shape_n}, ${instruction_shape_k}>,
      ${epilogue},
      ${swizzling_functor}, // cutlass::gemm::threadblock::GemmSplitKIdentityThreadblockSwizzle<>,
      ${stages},
      ${math_operator},
      ${iterator_algorithm},
      ${stride_support},
      ${align_a},
      ${align_b}
    >::Kernel;
  
    ${reduction}
  """
  
      self.reduction_template = """
  using EpilogueOutputOp = ${epilogue};
  using ReductionOp = cutlass::reduction::thread::ReduceAdd<
      ${element_accumulator},
      ${element_accumulator},
        EpilogueOutputOp::kCount
        >;
  
  using ReductionKernel = cutlass::reduction::kernel::ReduceSplitK<
      cutlass::MatrixShape<4, 32 * EpilogueOutputOp::kCount>,
        EpilogueOutputOp,
        ReductionOp
        >;
  
  using ReductionDevice = cutlass::reduction::device::ReduceSplitK<ReductionKernel>;
  using ReductionStrideIndex = typename ReductionDevice::StrideIndex;
  """
  
file_path: tvm/python/tvm/contrib/cutlass/conv2d_operation.py
imports: from .library import *
label_desc: Use image node_id 1 to create a new EmitConv2dInstance object  with example: obj = EmitConv2dInstance()