code: def scatter_nd(data, indices, updates, mode):
      """Scatter elements from a n-dimension array.
  
      Given updates with shape (Y_0, ..., Y_{K-1}, X_M, ..., X_{N-1}), indices with shape
      (M, Y_0, ..., Y_{K-1}), and output copied from data with shape (X_0, X_1, ..., X_{N-1}),
      scatter_nd computes
  
      .. code-block::
  
          output[indices[0, y_0, ..., y_{K-1}],
                 ...,
                 indices[M-1, y_0, ..., y_{K-1}],
                 x_M,
                 ...,
                 x_{N-1}
                ] = f(output[...], updates[y_0, ..., y_{K-1}, x_M, ..., x_{N-1}])
  
      where the update function f is determinted by the mode.
  
      Parameters
      ----------
      data : tvm.te.Tensor
          The source array.
  
      indices : tvm.te.Tensor
          The indices of the values to extract.
  
      updates : tvm.te.Tensor
          The updates to apply at the Indices
  
      mode : string
          The update mode for the algorithm, either "update" or "add"
          If update, the update values will replace the input data
          If add, the update values will be added to the input data
  
      Returns
      -------
      ret : tvm.te.Tensor
      """
      _verify_scatter_nd_inputs(data, indices, updates)
  
      def gen_ir(data_ptr, indices_ptr, updates_ptr, out_ptr):
          # pylint: disable=invalid-name
          ib = tir.ir_builder.create()
  
          data = ib.buffer_ptr(data_ptr)
          indices = ib.buffer_ptr(indices_ptr)
          updates = ib.buffer_ptr(updates_ptr)
          out = ib.buffer_ptr(out_ptr)
  
          # We combine all the indices dimensions but the first one into a single
          # dimension so we can iterate it in single loop instead of an arbitrary
          # number of loops. We do the same thing for all the update dimensions.
          fused_indices_dimension = 1
          for i in indices_ptr.shape[1:]:
              fused_indices_dimension *= i
  
          fused_updates_dimension = 1
          for i in updates_ptr.shape[len(indices_ptr.shape) - 1 :]:
              fused_updates_dimension *= i
  
          fused_shape = 1
          for i in data_ptr.shape:
              fused_shape *= i
  
          with ib.for_range(0, fused_shape) as i:
              out[i] = data[i]
  
          with ib.for_range(0, fused_indices_dimension) as i:
              with ib.for_range(
                  0, fused_updates_dimension, kind="parallel"
              ) as j:
                  offset = fused_updates_dimension
                  index = j  # This is x_M, .. x_{N-1} part of the index into out.
                  # Build up the indices[0, y_0, .. y_{K-1}], .. indices[M-1, y_0, .. y_{K-1}] part
                  # of the index into out.
                  for l in reversed(range(indices_ptr.shape[0].value)):
                      # indices[i * l * fused_indices_dimension] = indices[l, y_0, ... y_{k-1}]
                      index += (
                          offset
                          * indices[i + l * fused_indices_dimension]
                      )
                      offset *= data_ptr.shape[l]
                  if mode == "update":
                      out[index] = updates[
                          i * fused_updates_dimension + j
                      ]
                  elif mode == "add":
                      out[index] += updates[
                          i * fused_updates_dimension + j
                      ]
                  elif mode == "mul":
                      out[index] *= updates[
                          i * fused_updates_dimension + j
                      ]
                  elif mode == "min":
                      out[index] = tir.min(
                          out[index],
                          updates[i * fused_updates_dimension + j],
                      )
                  elif mode == "max":
                      out[index] = tir.max(
                          out[index],
                          updates[i * fused_updates_dimension + j],
                      )
                  else:
                      raise NotImplementedError(
                          "scatter_nd mode not in [update, add, mul, min, max]:",
                          mode,
                      )
  
          return ib.get()
  
      out_buf = tir.decl_buffer(data.shape, data.dtype, "out_buf")
      return te.extern(
          [data.shape],
          [data, indices, updates],
          lambda ins, outs: gen_ir(ins[0], ins[1], ins[2], outs[0]),
          dtype=data.dtype,
          out_buffers=[out_buf],
          name="scatter_nd.generic",
          tag="scatter_nd.generic",
      )
  
imports: from tvm import te, tir
  from tvm.tir import expr
label_desc: Use image node_id 2 for calling a global function with example usage: scatter_nd(data, indices, updates, mode) and returns: te, ib