code: def schedule_pragmas(sch):
      """Add pragmas to the operators that require them.
  
      This adds the pragmas used for codegen to the NPU ops.
      They are taken directly from the TE compute op's attributes.
      Modifies the schedule in-place.
  
      Parameters
      ----------
      sch : tvm.te.Schedule
          The schedule.
  
      """
  
      def _add_pragmas(stage, ax):
          if stage.op.name == "T_concat":
              stage.pragma(ax, "op", "ethosu_concatenate")
          if "op" in [attr for attr, val in stage.op.attrs.items()]:
              stage.pragma(ax, "op", stage.op.attrs["op"])
              for attr, val in stage.op.attrs.items():
                  if attr not in ("op", "lut") and not isinstance(
                      val, Propagator
                  ):
                      stage.pragma(ax, str(attr), val)
          if stage.op.axis[0] in stage.iter_var_attrs:
              attrs = stage.iter_var_attrs[stage.op.axis[0]]
              if "block_config_height" in attrs.pragma_keys:
                  pragmas = dict(
                      zip(
                          [k.value for k in attrs.pragma_keys],
                          attrs.pragma_values,
                      )
                  )
                  stage.pragma(
                      ax,
                      "block_config_height",
                      pragmas["block_config_height"],
                  )
                  stage.pragma(
                      ax,
                      "block_config_width",
                      pragmas["block_config_width"],
                  )
                  stage.pragma(
                      ax,
                      "block_config_depth",
                      pragmas["block_config_depth"],
                  )
  
      for stage in sch.stages:
          if (
              isinstance(stage.op, tvm.te.ComputeOp)
              and len(stage.op.axis) + len(stage.op.reduce_axis) > 0
          ):
              # The logic ensures the pragmas are assigned to the inner tiling loops
              # rather than the outer ones (which end up getting unrolled).
              num_inner_loops = len(stage.op.axis) + len(
                  stage.op.reduce_axis
              )
              ax = stage.leaf_iter_vars[-num_inner_loops]
              _add_pragmas(stage, ax)
  
file_path: tvm/python/tvm/relay/backend/contrib/ethosu/tir/scheduler.py
imports: import tvm
  from tvm.contrib.ethosu.cascader import Propagator
label_desc: Use image node_id 6 for calling a global function with example usage: schedule_pragmas(sch) without return types