code: def generate_sm80_tensor_op_16816(
      out_dtype,
      arg0_dtype,
      arg1_dtype,
      op_creator,
      check_align,
      use_3xtf32=True,
      profile_all_alignments=False,
      accumlator_dtype="float32",
  ):
      """Generate GEMM or Conv2D kernels for Ampere."""
      min_cc = 80
      max_cc = 1024
      max_cc_smem_limited = 80
  
      def get_default_tile_descriptions(block_k_factor):
          return [
              (
                  [128, 256, int(32 * block_k_factor)],
                  3,
                  [2, 4, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [256, 128, int(32 * block_k_factor)],
                  3,
                  [4, 2, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [256, 64, int(32 * block_k_factor)],
                  3,
                  [4, 1, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [256, 64, int(32 * block_k_factor)],
                  4,
                  [4, 1, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [64, 256, int(32 * block_k_factor)],
                  4,
                  [1, 4, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [128, 128, int(32 * block_k_factor)],
                  3,
                  [2, 2, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [128, 128, int(32 * block_k_factor)],
                  4,
                  [2, 2, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [128, 128, int(32 * block_k_factor)],
                  5,
                  [2, 2, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [128, 64, int(32 * block_k_factor)],
                  6,
                  [2, 2, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [64, 128, int(32 * block_k_factor)],
                  6,
                  [2, 2, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [64, 64, int(32 * block_k_factor)],
                  10,
                  [2, 2, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [256, 128, int(64 * block_k_factor)],
                  3,
                  [4, 2, 1],
                  min_cc,
                  max_cc_smem_limited,
              ),
              (
                  [128, 256, int(64 * block_k_factor)],
                  3,
                  [2, 4, 1],
                  min_cc,
                  max_cc_smem_limited,
              ),
              (
                  [256, 64, int(64 * block_k_factor)],
                  4,
                  [4, 1, 1],
                  min_cc,
                  max_cc_smem_limited,
              ),
              (
                  [64, 256, int(64 * block_k_factor)],
                  4,
                  [1, 4, 1],
                  min_cc,
                  max_cc_smem_limited,
              ),
              (
                  [128, 128, int(64 * block_k_factor)],
                  4,
                  [2, 2, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [256, 64, int(64 * block_k_factor)],
                  3,
                  [4, 1, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [64, 256, int(64 * block_k_factor)],
                  3,
                  [1, 4, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [128, 128, int(64 * block_k_factor)],
                  3,
                  [2, 2, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [128, 64, int(64 * block_k_factor)],
                  3,
                  [2, 2, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [64, 128, int(64 * block_k_factor)],
                  3,
                  [2, 2, 1],
                  min_cc,
                  max_cc,
              ),
              (
                  [64, 64, int(64 * block_k_factor)],
                  5,
                  [2, 2, 1],
                  min_cc,
                  max_cc,
              ),
          ]
  
      if arg0_dtype == "float16" and arg1_dtype == "float16":
          math_instructions = [
              MathInstruction(
                  [16, 8, 16],
                  DataType.f16,
                  DataType.f16,
                  dtype_map[out_dtype],
                  dtype_map[accumlator_dtype],
                  OpcodeClass.TensorOp,
                  MathOperation.multiply_add,
              )
          ]
          alignment_constraints = [8, 4, 2]
          tile_descriptions = get_default_tile_descriptions(1)
      elif arg0_dtype == "float32" and arg1_dtype == "float32":
          math_instructions = [
              MathInstruction(
                  [16, 8, 8],
                  DataType.f32,
                  DataType.f32,
                  DataType.f32,
                  DataType.f32,
                  OpcodeClass.TensorOp,
                  MathOperation.multiply_add_fast_f32
                  if use_3xtf32
                  else MathOperation.multiply_add,
              )
          ]
          alignment_constraints = [4, 2, 1]
  
          if use_3xtf32:
              # tf32
              tile_descriptions = [
                  ([128, 128, 16], 4, [4, 2, 1], min_cc, max_cc),
                  ([128, 128, 16], 3, [4, 2, 1], min_cc, max_cc),
                  ([256, 64, 16], 3, [4, 2, 1], min_cc, max_cc),
                  ([64, 256, 16], 3, [2, 4, 1], min_cc, max_cc),
                  ([128, 64, 16], 4, [2, 2, 1], min_cc, max_cc),
                  ([64, 128, 16], 4, [2, 2, 1], min_cc, max_cc),
                  ([64, 64, 16], 3, [2, 2, 1], min_cc, max_cc),
                  ([128, 128, 32], 3, [4, 2, 1], min_cc, max_cc),
                  (
                      [256, 64, 32],
                      3,
                      [4, 2, 1],
                      min_cc,
                      max_cc_smem_limited,
                  ),
                  (
                      [64, 256, 32],
                      3,
                      [2, 4, 1],
                      min_cc,
                      max_cc_smem_limited,
                  ),
                  ([128, 64, 32], 3, [2, 2, 1], min_cc, max_cc),
                  ([64, 128, 32], 3, [2, 2, 1], min_cc, max_cc),
                  ([64, 64, 32], 3, [2, 2, 1], min_cc, max_cc),
              ]
          else:
              tile_descriptions = get_default_tile_descriptions(0.5)
      else:
          assert out_dtype == "int32"
          math_instructions = [
              MathInstruction(
                  [16, 8, 32],
                  dtype_map[arg0_dtype],
                  dtype_map[arg1_dtype],
                  DataType.s32,
                  DataType.s32,
                  OpcodeClass.TensorOp,
                  MathOperation.multiply_add_saturate,
              )
          ]
          alignment_constraints = [16, 8, 4]
          tile_descriptions = get_default_tile_descriptions(2)
  
      def get_tile_descriptions(math_inst):
          return [
              TileDescription(
                  threadblock_shape,
                  stages,
                  warp_count,
                  math_inst,
                  min_cc,
                  max_cc,
              )
              for threadblock_shape, stages, warp_count, min_cc, max_cc in tile_descriptions
          ]
  
      alignment_constraints = [
          align for align in alignment_constraints if check_align(align)
      ]
  
      if len(alignment_constraints) > 0 and not profile_all_alignments:
          alignment_constraints = [alignment_constraints[0]]
  
      if arg0_dtype != "float32" and arg1_dtype != "float32":
          sm75_kernels = generate_sm75_tensor_op_1688(
              out_dtype,
              arg0_dtype,
              arg1_dtype,
              op_creator,
              check_align,
              False,
              profile_all_alignments,
              accumlator_dtype=accumlator_dtype,
          )
      else:
          # TF32 (float32 + float32 case) is only supported on sm80
          sm75_kernels = []
  
      if len(alignment_constraints) > 0:
          sm80_kernels = generate_tensor_op_common(
              math_instructions,
              alignment_constraints,
              get_tile_descriptions,
              op_creator,
          )
      else:
          sm80_kernels = []
  
      # TODO(masahi): For int8 kernels, The CUTLASS generator modifies the output tensor alignment
      # after ops are created. Revisit how important this modification is.
      # for op in operations:
      #     if op.tile_description.threadblock_shape[1] >= 128:
      #         op.C.alignment = 16
      #     else:
      #         op.C.alignment = 8
  
      return sm75_kernels + sm80_kernels
  
imports: import logging
  import os
  import re
  import tempfile
  import subprocess
  import multiprocessing
  import tvm._ffi
  from tvm.tir import IntImm
  from tvm.runtime import Object
  from .None import _ffi_api
  from .library import MathInstruction, DataType, DataTypeTag, OpcodeClass, MathOperation, TileDescription, EpilogueFunctor
  from .gemm_operation import instantiate_gemm_template
  from .conv2d_operation import instantiate_conv2d_template
label_desc: Use image node_id 4 for calling a global function with example usage: generate_sm80_tensor_op_16816(out_dtype, arg0_dtype, arg1_dtype, op_creator, check_align, use_3xtf32, profile_all_alignments, accumlator_dtype) and returns: unknown, list, unknown